cuda 共享存储器

CUDA 共享存储器（Shared Memory）是 GPU 编程中一种重要的片上高速缓存，它允许同一线程块内的所有线程进行高效的数据交换和通信，以下是关于 CUDA 共享存储器的详细内容：

1、定义与特性

定义方式：在 CUDA 核函数中，共享存储器通常有两种定义方式，一种是静态方式，即在定义时同时指定大小，例如__shared__ type shared[SIZE];；另一种是动态方式，通过内核函数的第三个参数设置大小，如extern __shared__ type shared[];。

访问权限：共享存储器可以被同一块中的所有线程访问，生存期是块的生命期。

2、组织形式

共享存储器被组织为多个 bank，每个 bank 拥有一定的宽度（如 32bit），当多个线程同时访问不同的 bank 时，可以实现并行访问，提高访存效率；但如果多个线程同时访问同一个 bank，则会产生 bank conflict，导致只能顺序处理，降低访存效率。

3、应用场景

线程间通信：在同一个线程块内的线程可以通过共享存储器高效地交换数据，而无需通过全局内存进行通信，从而减少了数据传输的延迟和带宽消耗。

数据重用：可以将经常使用的数据存储在共享存储器中，以便线程可以快速访问，提高程序的性能。

4、性能优化

减少 bank conflicts：通过合理地安排线程对共享存储器的访问方式，可以减少 bank conflicts，提高访存效率，将线程按照访问的地址进行分组，使得每组线程访问不同的 bank。

利用广播机制：当 half-warp 的线程访问同一地址时，会产生一次广播，不会产生 bank conflict，可以利用这一特性来优化程序的性能。

5、示例代码

 // 静态共享存储器示例
   __global__ void static_shared_memory_kernel(int *input, int *output) {
       __shared__ int shared[256];
       int idx = threadIdx.x + blockIdx.x * blockDim.x;
       if (idx < 256) {
           shared[threadIdx.x] = input[idx];
       }
       __syncthreads();
       output[threadIdx.x] = shared[threadIdx.x] * 2;
   }
   // 动态共享存储器示例
   __global__ void dynamic_shared_memory_kernel(int *input, int *output) {
       extern __shared__ int shared[];
       int idx = threadIdx.x + blockIdx.x * blockDim.x;
       if (idx < 256) {
           shared[threadIdx.x] = input[idx];
       }
       __syncthreads();
       output[threadIdx.x] = shared[threadIdx.x] * 2;
   }

6、常见问题解答

Q: 如何查看 CUDA 设备的共享内存信息？

A: 可以使用 CUDA 提供的cudaGetDeviceCount 和cudaGetDeviceProperties 函数来获取 CUDA 设备的信息，包括每个线程块的共享内存大小等。

 void show_GPU_info() {
         int deviceCount;
         cudaGetDeviceCount(&deviceCount);
         for (int i = 0; i < deviceCount; i++) {
             cudaDeviceProp devProp;
             cudaGetDeviceProperties(&devProp, i);
             std::cout << "使用 GPU device " << i << ": " << devProp.name << std::endl;
             std::cout << "每个线程块的共享内存大小：" << devProp.sharedMemPerBlock / 1024.0 << " KB" << std::endl;
         }
     }

Q: 如何避免共享存储器的 bank conflicts？

A: 可以通过以下几种方法来避免或减少 bank conflicts：

尽量使线程按照访问的地址进行分组，使得每组线程访问不同的 bank。

使用循环展开、偏移等技术来打散线程对共享存储器的访问。

利用共享存储器的广播机制，当 half-warp 的线程访问同一地址时，不会产生 bank conflict。

7、小编有话说

CUDA 共享存储器是一种非常强大的工具，可以为 GPU 程序带来显著的性能提升，要充分发挥其优势，需要深入理解其原理和特性，并在实际编程中灵活运用，通过合理地使用共享存储器，可以有效地减少线程间的通信开销，提高程序的执行效率。