在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，第一步先把１和128，2和129，…，127和256相加。前128维数据就是相加后的结果，再在前128上重复此操作，以此类推，最终得到第１个元素就是最终的结果。此时需要注意的是原始缓存数组长度需要是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上完成计算即可。

---

Reference

GPU高性能变成CUDA实战（桑德斯）
https://blog.csdn.net/xiaohu2022/article/details/79599947