在cuda设备端的内存包括,全局内存(global memory),共享内存(shared memory),纹理内存(texture memory),常量内存(constant memory)等。
在我是小将的博文中,详细画出了内存的分布情况,很清晰,一目了然。
全局内存,每个block中的thread都可以访问,可以用来存储和host端交互数据。
这里主要来学习共享内存。
共享内存
共享内存在block内的线程都可以访问,但是不同的block不能互相访问。
共享内存的存取速度非常块,极大程度上解决出现内存延迟的问题。像cpu上的多线程一样,同时操作一片内存空间也会出现问题,发生竞态条件。所以需要对读写进行同步。
点乘实现
利用共享内存实现点乘操作,依然先看代码。
void CUDA_Base::matrix_dot_plus(float *a, float *b, int length) {
cudaSetDeviceFlags(cudaDeviceMapHost);
float *dev_src1, *dev_src2, *dev_dst, *dst, *src1, *src2;
int byte_size = length * sizeof(float);
cudaHostAlloc((void**)&src1, byte_size, cudaHostAllocWriteCombined | cudaHostAllocMapped);
cudaHostAlloc((void**)&src2, byte_size, cudaHostAllocWriteCombined | cudaHostAllocMapped);
memcpy((void*)src1, (void*)a, byte_size);
memcpy((void*)src2, (void*)b, byte_size);
std::cout << src1[2] <<std::endl;
cudaHostGetDevicePointer(&dev_src1, src1, 0);
cudaHostGetDevicePointer(&dev_src2, src2, 0);
dim3 block_size(256);
dim3 grid_size((length / block_size.x) + 1);
cudaHostAlloc((void**)&dst, grid_size.x * sizeof(float), cudaHostAllocWriteCombined | cudaHostAllocMapped);
cudaHostGetDevicePointer(&dev_dst, dst, 0);
dot_plus_1D<<<grid_size, block_size>>>(dev_src1, dev_src2, dev_dst, length);
// cudaThreadSynchronize();
cudaDeviceSynchronize();
float num = 0;
for (int i = 0; i < grid_size.x; ++i) {
num += dst[i];
}
std::cout << num << std::endl;
cudaFreeHost(src1);
cudaFreeHost(src2);
cudaFreeHost(dst);
}
cpu端代码依然和之前的类似,主要是kernel的实现。
__global__ void dot_plus_1D(float* a, float* b, float* c, float length){
__shared__ float cache[256];
int x = threadIdx.x + blockIdx.x * blockDim.x;
// int y = threadIdx.y + blockIdx.y * blockDim.y;
int curr_index = threadIdx.x;
float temp = 0;
if (x < length){
temp += a[x] * b[x];
}
cache[curr_index] = temp;
__syncthreads();
int i = blockDim.x / 2;
while (i != 0){
if (curr_index < i){
cache[curr_index] += cache[curr_index + i];
}
__syncthreads();
i /= 2;
}
if (curr_index == 0){
c[blockIdx.x] = cache[0];
}
}
这里使用的依然是1D的方式,下标索引计算,和前文相同。
首先__shared__ float cache[256];
申请一个shared memory,用于缓存中间结果。需要注意的是,在申请的时候不能使用动态申请,大小需要时已知常量。 256是block内thread的大小,保证了每个线程的临时变量都能顺利存储。
int x = threadIdx.x + blockIdx.x * blockDim.x;
int curr_index = threadIdx.x;
float temp = 0;
if (x < length){
temp += a[x] * b[x];
}
cache[curr_index] = temp;
这段代码和之前的也是一样,找到指定下标的元素,相乘后存入对应线程编号的缓存内,由于shared memory只在block内,所以不需要计算便宜了。
下面就是要把所有计算结果累加起来,这里采用的方式是一种归约的计算方式,折半相加。缓存数组大小是256,第一步先把1和128,2和129,…,127和256相加。前128维数据就是相加后的结果,再在前128上重复此操作,以此类推,最终得到第1个元素就是最终的结果。此时需要注意的是原始缓存数组长度需要是2的指数,不然结果就会有问题。
在执行这个步骤之前,需要保证所有的乘法运算已经执行完成,不然求和是没有意义的,此时就需要__syncthreads();
进行线程同步。其意义就是,在这个block内所有的线程都执行完这条语句之前的代码段后,才会执行下面的代码。
int i = blockDim.x / 2;
while (i != 0){
if (curr_index < i){
cache[curr_index] += cache[curr_index + i];
}
__syncthreads(); // 归约的每一步,也需要等前面一步全部执行完成
i /= 2;
}
// 最后在0号线程上,把当前block的结果返回。
if (curr_index == 0){
c[blockIdx.x] = cache[0];
}
因为shared memory只能在block内共享,之间无法互相通信。所以我们申请一个和我们所申请的block的数目一样大小的数组保存,每个block的结果。这个数组的长度已经很小,只需要在cpu上完成计算即可。
注意:
__syncthreads();
必须在if分支外面,放在if分支里面,将会出现有的线程永远不会执行这句代码,处理器将会永远挂起。所以使用的时候要谨慎。
Reference
GPU高性能变成CUDA实战(桑德斯)
https://blog.csdn.net/xiaohu2022/article/details/79599947