第五节：了解和使用共享内存(2)
Rob Farber 是西北太平洋国家实验室（Pacific Northwest National Laboratory）的高级科研人员。他在多个国家级的实验室进行大型并行运算的研究，并且是几个新创企业的合伙人。大家可以发邮件到rmfarber@gmail.com与他沟通和交流。
在CUDA系列文章的第四节里，我探讨了执行模型和内核启动执行配置是如何影响寄存器的数量和本地多处理器资源如共享内存的数量的。在本小节，我会继续探讨内存性能，和在reverseArray_multiblock_fast.cu.内使用共享内存。
CUDA 内存性能
本地内存空间和全局内存空间不会缓存，这意味着每次对全局内存（或本地内存）进行访问都将导致一次实际的内存访问。那么访问（例如读取或写入）各种类型的内存的“开销”是多少？
多处理器需要四（4）个时钟周期才能为每次“warp（交换）”发出一条内存指令。读取本地或全局内存。访问本地或全局内存将导致400到600个时钟周期的延迟。举例说明，在以下代码片段中的赋值运算符需要4个时钟周期才能从全局内存中进行一次读取，4个时钟周期从共享内存进行一次写入，需要400到600个时钟周期从全局内存读取一个浮点值。注意：4个时钟周期从共享内存进行一次写入，需要400到600个时钟周期从全局内存读取一个浮点值。注意：使用__device__变量类型限定符表示全局内存中的变量（有关其他变量特征，请参见CUDA Programming Guide第4.2.2.1节）。主代码不能访问变量类型__device__。
__shared__ float shared[32];
__device__ float device[32];
shared[threadIdx.x] = device[threadIdx.x];
当访问时间出现100-150倍差别时，难怪程序开发人员会需要最小化对全局内存的访问，并且在本地多处理器内存内重新使用数据。CUDA设计者对线程调度程序的设计十分巧妙，大量的全局内存延迟都可以透明地隐藏起来：只需在执行配置中指定大量数据块，并尽可能在内核中使用寄存器、__shared__和__constant__存储器类型处理变量即可。
因为共享内存在芯片上，因此访问速度要比读取全局内存快很多，并且主要的优化是为了避免存储器组冲突。共享内存速度较快（有些文章认为它和寄存器访问一样快捷）。
然而，最近CUBLAS和CUFFT性能获得极大的改进：通过尽量使用共享内存而非寄存器－因此如可能的话，尽量使用寄存器吧。CUDA共享内存被分为大小相同的内存模块，称为存储器组（memorybank）。每个存储器组都保存有一个连续的32位值（如int和float），因此通过连续线程对连续数组访问非常的快捷。当向同一个存储数组（可能是同一个地址或映射到同一个存储数组的多个地址）发出数据请求时，存储器组发生冲突。如果发生这种情况，硬件会有效地序列化内存运算，强迫所有线程等待直到内存请求得以完成。如果所有的线程从同一个共享内存地址读取，那就会自动调用广播机制，避免序列化。共享内存广播是一个同时向多个线程提供数据的有效方法。使用共享内存时，您完全可以尝试下利用这一特点。
我将在下一个专栏文章中更为详细地讨论存储器组冲突。现在，只需要知道reverseArray_multiblock_fast.cu 没有存储器冲突，因为连续线程方位连续值。
具有读取/写入功能的多处理器本地存储器类型总结如下：
寄存器
- 多处理器上最快的内存形式；
- 仅可通过线程访问；
- 有线程的生命周期
共享内存
- 在没有存储器组冲突（从同一个地址读取）时与寄存器一样快。
- 可从创建线程的任何块访问；
- 有线程的生命周期
全局内存:
- 可能比寄存器或共享内存慢150倍，注意非联合读取和写入（将在下一专栏中讨论）。
- 可从主机或设备访问；
- 有应用程序的生命周期
本地内存：
- 潜在的性能缺陷，位于全局内存中，可能比寄存器或共享内存慢150倍。
- 仅可通过线程访问；
- 有线程的生命周期
共享内存注意事项
小心共享内存存储器冲突，可能会导致性能降低；
所有在内核里的动态分配的共享变量在同一个内存地址开始。使用至少两个动态分配共享内存数组要求手动生成偏移量。例如，如果你想动态分配共享内存，以包含两个数组, a 和b, 你需要进行如下操作：
1 __global__ void kernel(int aSize)
2
3 {
4 extern __shared__ float sData[];
5 float *a, *b;
6
7 a = sData;
8 b = &a[aSize]; __global__ void kernel(int aSize)
9 {
10 extern __shared__ float sData[];
11 float *a, *b;
12
13 a = sData;
14 b = &a[aSize];
寄存器/本地内存注意事项
寄存器内存可以透明地存入本地内存。这可能会造成性能欠佳。检查ptx 汇编代码或在通过nvcc在输出中查找lmem，命令如下：“—ptxas-options=-v”。
编译时已经知道的通过常量索引的数组通常存于寄存器中，但是如果使用变量索引，它们就不能存于寄存器中。这就为开发人员弄出了一个难题，因为可能需要循环展开以在寄存器内存，而非较慢的全局内存中，保存数组元素。然而，展开循环会急剧增加寄存器使用量，因而就可能会使得变量被保存在本地内存中――从而抵消了循环展开的诸多好处。可以使用使用nvcc选项，“—maxrregcount=value”告诉编译器使用更多寄存器（注意：可以指定的最大寄存器数量为128）。这需要在“使用更多的寄存器”和“创建更少的线程”之间权衡利弊，有可能会妨碍隐藏存储器延迟。在某些架构中，使用该选项可能造成资源不足，从而导致内核无法启动。
共享内存内核
程序reverseArray_multiblock.cu和revereseArray_multiblock_fast.cu 执行同样的任务。它们创建一个一维整数数组h_a，数组包含整数值[0 ..dimA-1]。数组可以通过cudaMemcpy移动到设备，然后主机启动reverseArrayBlock内核就地逆转数组内容的顺序。再次使用cudaMemcpy将数据从设备传回主机，并执行一个检查，检查设备生成了正确的结果（例如[dimA-1 .. 0]）。
区别就是：reverseArray_multiblock_fast.cu使用共享内存以改进内核的性能，而reverseArray_multiblock.cu在全局内存里执行。尝试下给这两个程序计时，验证下它们的性能区别。另外，reverseArray_multiblock.cu 访问全局内存的效率不高。在以后的专栏文章里，我们将使用CUDAPROFILER协助诊断和纠正这个性能，并且展示最新的10个系列架构中的改进如何减少了许多情况下对这些优化的需求。
1 #include <stdio.h>
2 #include <assert.h>
3 #include "cuda.h"
4 #include "cuda_runtime.h"
5 #include "device_launch_parameters.h"
6 #include <device_functions.h>
7 //检查CUDA运行时是否有错误
8 void checkCUDAError(const char* msg);
9 // Part 2 of 2: 使用共享内存执行内核
10 __global__ void reverseArrayBlock(int *d_out, int *d_in)
11 {
12 extern __shared__ int s_data[];
13 int inOffset = blockDim.x * blockIdx.x;
14 int in = inOffset + threadIdx.x;
15 // Load one element per thread from device memory and store it
16 // *in reversed order* into temporary shared memory
17 /*
18 每个线程从设备内存加载一个数据元素并按逆序存储在共享存储器上
19 */
20 s_data[blockDim.x - 1 - threadIdx.x] = d_in[in];
21 /*
22 阻塞，一直到所有线程将他们的数据都写入到共享内存中
23 */
24 __syncthreads();
25 // write the data from shared memory in forward order,
26 // but to the reversed block offset as before
27 /*
28 将共享内存中的数据s_data写入到d_out中，按照前序
29 */
30 int outOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x);
31 int out = outOffset + threadIdx.x;
32 d_out[out] = s_data[threadIdx.x];
33 }
34 ////////////////////////////////////////////////////////////////////
35 //主函数
36 ////////////////////////////////////////////////////////////////////
37 int main(int argc, char** argv)
38 {
39 //指向主机的内存空间和大小
40 int *h_a;
41 int dimA = 256 * 1024; // 256K elements (1MB total)
42 // pointer for device memory
43 int *d_b, *d_a;
44 //指向设备的指针和大小
45 int numThreadsPerBlock = 256;
46
47 /*
48 根据数组大小和预设的块大小来计算需要的块数
49 */
50 int numBlocks = dimA / numThreadsPerBlock;
51 /*
52 Part 1 of 2:
53 计算共享内存所需的内存空间大小，这在下面的内核调用时被使用
54 */
55 int sharedMemSize = numThreadsPerBlock * sizeof(int);
56 //申请主机及设备上的存储空间
57 size_t memSize = numBlocks * numThreadsPerBlock * sizeof(int);
58 //主机上的大小
59 h_a = (int *)malloc(memSize);
60 //设备上的大小
61 cudaMalloc((void **)&d_a, memSize);
62 cudaMalloc((void **)&d_b, memSize);
63 //在主机上初始化输入数组
64 for (int i = 0; i < dimA; ++i)
65 {
66 h_a[i] = i;
67 }
68 //将主机数组拷贝到设备上，h_a-->d_a
69 cudaMemcpy(d_a, h_a, memSize, cudaMemcpyHostToDevice);
70 //启动内核
71 dim3 dimGrid(numBlocks);
72 dim3 dimBlock(numThreadsPerBlock);
73 reverseArrayBlock << < dimGrid, dimBlock, sharedMemSize >> >(d_b, d_a);
74 //阻塞，一直到设备完成计算
75 cudaThreadSynchronize();
76 //检查是否设备产生了错误
77 //检查任何CUDA错误
78 checkCUDAError("kernel invocation");
79 //将结果从设备拷贝到主机，d_b-->h_a
80 cudaMemcpy(h_a, d_b, memSize, cudaMemcpyDeviceToHost);
81 //检查任何CUDA错误
82 checkCUDAError("memcpy");
83 //核对返回到主机上的结果是否正确
84 for (int i = 0; i < dimA; i++)
85 {
86 assert(h_a[i] == dimA - 1 - i);
87 }
88 //释放设备内存
89 cudaFree(d_a);
90 cudaFree(d_b);
91 //释放主机内存
92 free(h_a);
93 printf("Correct!
");
94 return 0;
95 }
96
97 void checkCUDAError(const char *msg)
98 {
99 cudaError_t err = cudaGetLastError();
100 if (cudaSuccess != err)
101 {
102 fprintf(stderr, "Cuda error: %s: %s.
", msg, cudaGetErrorString(err));
103 exit(EXIT_FAILURE);
104 }
105 }
确定运行时共享内存的量需要在主机和设备代码里进行设置。在本示例中，内核中每个块的共享内存的量（字节为单位）在主机上的执行配置中予以明确，作为可选的第三个参数。（仅当共享内存的容量在内核启动时指定后，才设置主机端。如果在编译时修复，则主机端不需要任何设置。）例如，在主机代码arrayReversal_multiblock_fast.cu中，以下代码片段为一个整数数组（包含的元素数等于数据块中的线程数）分配共享内存：
// Part 1 of 2: Compute the number of bytes of share memory needed
// This is used in the kernel invocation below
int sharedMemSize = numThreadsPerBlock * sizeof(int);
检测reverseArrayBlock 内核, 共享内存使用以下代码声明：
extern __shared__ int s_data[]; extern __shared__ int s_data[];
注意：内核中没有指定大小－大小通过执行配置从主机获取。
在下一个关于配置的专栏前，我建议研究一下reverseArray_multiblock.cu。您认为访问全局内存时存在性能问题吗？如果您认为存在，请尝试解决这一问题。