CUDA - 共享内存及核函数

staticReverse:

__global__ void staticReverse (int *d, int n)  
{  
  __shared__ int s [64]; // 声明一个大小为64的共享存储器数组s
  int t = threadIdx.x; // 获取当前线程在线程块内的索引t
  int tr = n-t-1; // 计算当前线程对应元素在反转后数组中的索引tr
  s [t] = d [t]; // 将全局存储器中d[t]处的元素复制到共享存储器中s[t]处
  __syncthreads (); // 等待所有线程完成复制操作
  d [t] = s [tr]; // 将共享存储器中s[tr]处的元素复制到全局存储器中d[t]处
}

dynamicReverse:

__global__ void dynamicReverse (int *d, int n)  
{  
  extern __shared__ int s []; // 声明一个动态大小的共享存储器数组s
  int t = threadIdx.x; // 获取当前线程在线程块内的索引t
  int tr = n-t-1; // 计算当前线程对应元素在反转后数组中的索引tr
  s [t] = d [t]; // 将全局存储器中d[t]处的元素复制到共享存储器中s[t]处
  __syncthreads (); // 等待所有线程完成复制操作
  d [t] = s [tr]; // 将共享存储器中s[tr]处的元素复制到全局存储器中d[t]处
}

这段代码的目的是实现一个数组的反转，即将数组中的元素按照逆序重新排列。例如，如果数组是[1, 2, 3, 4]，那么反转后的数组是[4, 3, 2, 1]。

这段代码中有两个核函数（kernel functions），分别是staticReverse和dynamicReverse。核函数是在GPU上执行的函数，它们可以被CPU上的主机函数（host functions）调用。

线程块是一组在同一个流多处理器（streaming multiprocessor）上执行的线程，它们可以共享一块片上存储器（on-chip memory），即共享存储器（shared memory）。共享存储器是一种高速缓存（cache），它比全局存储器（global memory）要快得多，但是容量要小得多。共享存储器可以用来存储线程块内的线程之间需要共享的数据，例如从全局存储器加载的数据。

在这段代码中，staticReverse和dynamicReverse都使用了共享存储器来实现数组的反转。它们的区别在于共享存储器的分配方式不同。staticReverse使用了一个固定大小的共享存储器数组，即__shared__ int s[64]，其中64是数组元素的个数，也是每个线程块内线程的个数。dynamicReverse使用了一个动态大小的共享存储器数组，即extern shared int s[]，其中s的大小由核函数调用时指定，即n*sizeof(int)，其中n是数组元素的个数。

核函数的调用

核函数的调用方式是使用<<<>>>符号，其中包含了执行配置（execution configuration），即指定了线程块（thread blocks）和线程（threads）的数量和大小，以及可选的共享存储器（shared memory）的大小和流（stream）的编号。例如：

histogram<<<2500, numBins, numBins * sizeof (unsigned int)>>>(...);

这个核函数调用指定了2500个线程块，每个线程块有numBins个线程，每个线程块还有numBins * sizeof (unsigned int)字节的共享存储器可用。这个核函数没有指定流编号，所以默认为0。

执行配置的选取准则和方法主要取决于以下几个因素：

• 核函数的功能和逻辑：不同的核函数可能有不同的并行度和数据依赖性，需要根据具体的算法设计合适的执行配置。例如，如果核函数需要进行归约操作，那么可能需要使用较小的线程块，并利用共享存储器和同步机制来加速计算。
• GPU的硬件特性：不同的GPU可能有不同的计算能力（compute capability），流多处理器（streaming multiprocessor）数量，寄存器（register）数量，共享存储器容量等等，需要根据具体的硬件特性来优化执行配置。例如，如果核函数使用了较多的寄存器，那么可能需要使用较小的线程块，以避免寄存器溢出和内存交换。
• GPU的占用率（occupancy）：占用率是指GPU上活跃的线程数占总线程数的比例，它反映了GPU的利用率和效率。一般来说，高占用率可以提高GPU的性能，因为它可以隐藏内存访问延迟和指令执行延迟。但是，并不是占用率越高越好，因为过高的占用率可能导致资源竞争和冲突。因此，需要根据实际情况来平衡占用率和资源利用率。

为了选择合适的执行配置，可以使用以下的工具和方法：

• CUDA占用率计算器（CUDA Occupancy Calculator）：这是一个Excel表格，可以根据你输入的核函数参数和GPU硬件特性，计算出不同执行配置下的占用率，并给出建议。
• CUDA Visual Profiler：这是一个可视化工具，可以对你编写的CUDA程序进行分析和优化，包括执行配置，内存访问模式，指令执行效率等等。
• CUDA动态并行（CUDA Dynamic Parallelism）：这是一种在GPU上动态生成和调度新核函数的技术，它可以让你在核函数中调用另一个核函数，并根据运行时数据来确定执行配置。这样可以避免在CPU和GPU之间传输数据和控制信息，提高程序灵活性和效率。