第六节：全局内存和CUDA RPOFILER
Rob Farber 是西北太平洋国家实验室（Pacific Northwest National Laboratory）的高级科研人员。他在多个国家级的实验室进行大型并行运算的研究，并且是几个新创企业的合伙人。大家可以发邮件到rmfarber@gmail.com与他沟通和交流。
在关于CUDA（Compute Unified DeviceArchitecture，即计算统一设备架构的简称）的系列文章的第二节，我探讨了内存性能和在reverseArray_multiblock_fast.cu.内使用共享内存。在本节，我探讨使用CUDA PROFILER检测全局内存。
本系列文章的细心读者已经了解了第四节和第五节里讨论的两个反向数组样例，但是仍然困扰他们的是，共享内存版本为什么比全局内存版的速度要快一些。回忆下共享内存版本吧，reverseArray_multiblock_fast.cu,内核将数组数据从全局内存复制到共享内存，然后回到全局内存，而较慢的内核reverseArray_multiblock.cu,仅将数据从全局内存复制到全局内存。因为全局内存性能比共享内存要慢100～150倍，慢得多的全局存储器性能占据了两个示例的绝大部分运行时。为什么共享存储器版本更快？回答这个问题需要先了解更多有关全局内存的信息，还需要使用来自CUDA开发环境的附加工具-－特别是CUDAPROFILER。CUDA软件的配置简单快捷，因为文本和可视化版本的profiler都在CUDA启动的设备上读取硬件配置计数器。启动文本配置非常简易：设置开始和控制profiler的环境变量。使用可视化profiler同样很简单：启动cudaprof并开始在GUI中进行单击操作。通过配置可以了解许多有价值的信息。配置事件集合完全由支持CUDA的设备内部的硬件来处理。然而，经过配置的内核不再具有异步特征。只有在每个内核完成之后，才将结果报告给主机，这样可以最小化所有通信带来的影响。
全局内存
了解如何有效使用全局内存是成为一名CUDA编程高手的基本要求。下面对全局内存进行了简要介绍，应该可以有助于你了解reverseArray_multiblock.cu和reverseArray_multiblock_fast.cu之间的性能区别。如有需要，以后的专栏文章会继续探索如何有效利用全局内存。同时，我们会采用图示的方式详细探讨全局内存（见CUDA Programming Guide的第5.1.2.1节）。
只有当全局存储器访问能够合并到一个half-warp时，硬件才能以最少的事务量获取（或存储）数据，全局存储器才能交付最高的存储器带宽。CUDAComputeCapability设备（1.0和1.1）能够在单个64字节或128字节事务中获取数据。如果无法合并存储器事务，那么将会为half-warp中的每个线程发出一个独立的存储器事务，这不是期望的结果。未合并的存储器操作的性能损失取决于数据类型的大小。CUDA文档对各种数据类型大小决定的预期性能降低给出了一些简单指南：
- 32位数据类型将减慢大约10x
- 64位数据类型将减慢大约4x
- 128位数据类型将减慢大约2x
当满足下列条件时，数据块的half-warp中的所有线程执行的全局存储器访问可以被合并到G80架构上一个有效的存储器事务中：
线程访问32、64或128位数据类型。
事务的所有16个字所在的分段的大小必须和内存事务大小一致（当为128位字时，为内存事务大小的2倍）。这就意味着起始地址和校准非常的重要了。
线程必须依次访问这些字：half-warp中的第k个线程必须访问第k个字。注意：不是warp中的所有线程都需要访问某个线程所访问的存储器才能进行合并。这称为发散warp。
较新的架构（比如GT200系列设备）的合并要求比刚才讨论的架构更宽松。我们将在未来的专栏中更深入地讨论它们之间的架构差异。从本专栏的主题看，可以肯定，如果经过调优的代码能够在支持CUDA的G80设备上进行很好的合并，那么它将能够在GT200设备上进行很好地合并。
启动和控制文本配置
控制CUDA profiler的文本版本的环境变量是：
- CUDA_PROFILE – 设置为1（或0）可以启用（或禁用）profiler
- CUDA_PROFILE_LOG – 设置为日志文件的名称（默认设置为./cuda_profile.log）
- CUDA_PROFILE_CSV – 设置为1（或0）可以启用（或禁用）使用逗号分隔的日志版本。
- CUDA_PROFILE_CONFIG – 指定最多带有4个信号的配置文件
最后一点非常重要，因为一次只能配置四个信号。通过在名为CUDA_PROFILE_CONFIG的文件中的单独行上指定名称，开发人员可以用profiler收集以下任何事件：gld_incoherent：未合并的全局存储器负载单元的数量
- gld_coherent：已合并的全局存储器负载单元的数量
- gst_incoherent：未合并的全局存储器存储单元的数量
- gst_coherent：已合并的全局存储器存储单元的数量
- local_load：局部存储器负载单元的数量
- local_store：局部存储器存储单元的数量
- branch：线程执行的分支事件的数量
- divergent_branch：warp中发散分支的数量
- instructions：指令计数
- warp_serialize：warp中基于与共享或常量存储器的地址冲突进行序列化的线程数量
- cta_launched：执行的线程块
Profiler 计数器注意问题：
注意，性能计数器值与单个的线程活动无联系。实际上，这些值代表了线程warp内的事件。例如，一个线程warp中的一个不连贯的存储将会递增gst_incoherent一次。因此，最终的计数器值存储的是关于所有warp中的所有不连贯存储的信息。
此外，profiler仅以GPU中的一个多处理器为目标，因此计数器值与为特定内核启动的warp的总数不会相关。因此，当使用profiler内的性能计数器选项时，用户应该总是启动足够的线程块以确保为目标多处理器分配固定百分比的工作。实际上，NVIDIA建议最好启动至少100个块，以获得一致的结果。
结果就是，用户不应该期待计数器值与通过检查内核代码所确定的数值一致。计数器值最好用于确定未优化和已优化代码之间的相对性能差异。例如，如果profiler报告软件的初始部分有一定数量的未合并全局负载，那么很容易确定更精细的代码版本是否会利用更少数量的未合并负载。在大多数情形下，我们的目标是将未合并的全局负载数量减少为0，因此，计数器值对于跟踪此目标的实现进度非常有用。
配置结果
我们来使用profiler.看下reverseArray_multiblock.cu 和reverseArray_multiblock_fast.cu。在本样例中，我们会在Linux下的bash shell中对环境变量和配置文件进行如下设置：
1 export CUDA_PROFILE=1
2 export CUDA_PROFILE_CONFIG=$HOME/.cuda_profile_configexport CUDA_PROFILE=1
3 export CUDA_PROFILE_CONFIG=$HOME/.cuda_profile_config
在Linux下使用bash比较Profile配置和环境变量
1 gld_coherent
2 gld_incoherent
3 gst_coherent
4 gst_incoherent
5 [code]
6
7 CUDA_PROFILE_CONFIG文件内容
8 运行reverseArray_multiblock.cu可执行文件，在./cuda_profile.log中生成以下配置报告：
9 [code]
10 method,gputime,cputime,occupancy,gld_incoherent,gld_coherent,gst_incoherent,gst_coherent
11 method=[ memcopy ]
12 gputime=[ 438.432 ]
13 method=[ _Z17reverseArrayBlockPiS_ ]
14 gputime=[ 267.520 ]
15 cputime=[ 297.000 ]
16 occupancy=[ 1.000 ]
17 gld_incoherent=[ 0 ]
18 gld_coherent=[ 1952 ]
19 gst_incoherent=[ 62464 ]
20 gst_coherent=[ 0 ]
21 method=[ memcopy ]
22 gputime=[ 349.344 ]
reverseArray_multiblock.cu配置报告
类似地，运行reverseArray_multiblock_fast.cu可执行文件生成以下输出，这些输出会覆盖.cuda_profile.log中以前的输出
1 method,gputime,cputime,occupancy,gld_incoherent,gld_coherent,gst_incoherent,gst_coherent
2 method=[ memcopy ]
3 gputime=[ 449.600 ]
4 method=[ _Z17reverseArrayBlockPiS_ ]
5 gputime=[ 50.464 ]
6 cputime=[ 108.000 ]
7 occupancy=[ 1.000 ]
8 gld_incoherent=[ 0 ]
9 gld_coherent=[ 2032 ]
10 gst_incoherent=[ 0 ]
11 gst_coherent=[ 8128 ]
12 method=[ memcopy ]
13 gputime=[ 509.984 ]
reverseArray_multiblock_fast.cu配置报告
比较这两个profiler结果，可看到reverseArray_multiblock_fast.cu内没有不连贯的存储，而 reverseArray_multiblock.cu却相反，它包含很多不连贯存储。看一下reverseArray_multiblock.cu的源，并看一下您是否可以修复不连贯存储的性能问题。修复之后，测量一下这两个程序彼此的相对速度。
为方便起见，列表1显示了reverseArray_multiblock.cu的情况，列表2显示了reverseArray_multiblock_fast.cu的情况。
1 // includes, system
2 #include <stdio.h>
3 #include <assert.h>
4 // Simple utility function to check for CUDA runtime errors
5 void checkCUDAError(const char* msg);
6 // Part3: implement the kernel
7 __global__ void reverseArrayBlock(int *d_out, int *d_in)
8 {
9 int inOffset = blockDim.x * blockIdx.x;
10 int outOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x);
11 int in = inOffset + threadIdx.x;
12 int out = outOffset + (blockDim.x - 1 - threadIdx.x);
13 d_out[out] = d_in[in];
14 }
15 ////////////////////////////////////////////////////////////////////////////////
16 // Program main
17 ////////////////////////////////////////////////////////////////////////////////
18 int main( int argc, char** argv)
19 {
20 // pointer for host memory and size
21 int *h_a;
22 int dimA = 256 * 1024; // 256K elements (1MB total)
23 // pointer for device memory
24 int *d_b, *d_a;
25 // define grid and block size
26 int numThreadsPerBlock = 256;
27 // Part 1: compute number of blocks needed based on array size and desired block size
28 int numBlocks = dimA / numThreadsPerBlock;
29 // allocate host and device memory
30 size_t memSize = numBlocks * numThreadsPerBlock * sizeof(int);
31 h_a = (int *) malloc(memSize);
32 cudaMalloc( (void **) &d_a, memSize );
33 cudaMalloc( (void **) &d_b, memSize );
34 // Initialize input array on host
35 for (int i = 0; i < dimA; ++i)
36 {
37 h_a[i] = i;
38 }
39 // Copy host array to device array
40 cudaMemcpy( d_a, h_a, memSize, cudaMemcpyHostToDevice );
41 // launch kernel
42 dim3 dimGrid(numBlocks);
43 dim3 dimBlock(numThreadsPerBlock);
44 reverseArrayBlock<<< dimGrid, dimBlock >>>( d_b, d_a );
45 // block until the device has completed
46 cudaThreadSynchronize();
47 // check if kernel execution generated an error
48 // Check for any CUDA errors
49 checkCUDAError("kernel invocation");
50 // device to host copy
51 cudaMemcpy( h_a, d_b, memSize, cudaMemcpyDeviceToHost );
52 // Check for any CUDA errors
53 checkCUDAError("memcpy");
54 // verify the data returned to the host is correct
55 for (int i = 0; i < dimA; i++)
56 {
57 assert(h_a[i] == dimA - 1 - i );
58 }
59 // free device memory
60 cudaFree(d_a);
61 cudaFree(d_b);
62 // free host memory
63 free(h_a);
64 // If the program makes it this far, then the results are correct and
65 // there are no run-time errors. Good work!
66 printf("Correct!
");
67
68 return 0;
69 }
70 void checkCUDAError(const char *msg)
71 {
72 cudaError_t err = cudaGetLastError();
73 if( cudaSuccess != err)
74 {
75 fprintf(stderr, "Cuda error: %s: %s.
", msg, cudaGetErrorString( err) );
76 exit(EXIT_FAILURE);
77 }
78 }
reverseArray_multiblock.cu
1 // includes, system
2 #include <stdio.h>
3 #include <assert.h>
4 // Simple utility function to check for CUDA runtime errors
5 void checkCUDAError(const char* msg);
6 // Part 2 of 2: implement the fast kernel using shared memory
7 __global__ void reverseArrayBlock(int *d_out, int *d_in)
8 {
9 extern __shared__ int s_data[];
10 int inOffset = blockDim.x * blockIdx.x;
11 int in = inOffset + threadIdx.x;
12 // Load one element per thread from device memory and store it
13 // *in reversed order* into temporary shared memory
14 s_data[blockDim.x - 1 - threadIdx.x] = d_in[in];
15 // Block until all threads in the block have written their data to shared mem
16 __syncthreads();
17 // write the data from shared memory in forward order,
18 // but to the reversed block offset as before
19 int outOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x);
20 int out = outOffset + threadIdx.x;
21 d_out[out] = s_data[threadIdx.x];
22 }
23 ////////////////////////////////////////////////////////////////////////////////
24 // Program main
25 ////////////////////////////////////////////////////////////////////////////////
26 int main( int argc, char** argv)
27 {
28 // pointer for host memory and size
29 int *h_a;
30 int dimA = 256 * 1024; // 256K elements (1MB total)
31 // pointer for device memory
32 int *d_b, *d_a;
33 // define grid and block size
34 int numThreadsPerBlock = 256;
35 // Compute number of blocks needed based on array size and desired block size
36 int numBlocks = dimA / numThreadsPerBlock;
37 // Part 1 of 2: Compute the number of bytes of shared memory needed
38 // This is used in the kernel invocation below
39 int sharedMemSize = numThreadsPerBlock * sizeof(int);
40 // allocate host and device memory
41 size_t memSize = numBlocks * numThreadsPerBlock * sizeof(int);
42 h_a = (int *) malloc(memSize);
43 cudaMalloc( (void **) &d_a, memSize );
44 cudaMalloc( (void **) &d_b, memSize );
45 // Initialize input array on host
46 for (int i = 0; i < dimA; ++i)
47 {
48 h_a[i] = i;
49 }
50 // Copy host array to device array
51 cudaMemcpy( d_a, h_a, memSize, cudaMemcpyHostToDevice );
52 // launch kernel
53 dim3 dimGrid(numBlocks);
54 dim3 dimBlock(numThreadsPerBlock);
55 reverseArrayBlock<<< dimGrid, dimBlock, sharedMemSize >>>( d_b, d_a );
56 // block until the device has completed
57 cudaThreadSynchronize();
58 // check if kernel execution generated an error
59 // Check for any CUDA errors
60 checkCUDAError("kernel invocation");
61 // device to host copy
62 cudaMemcpy( h_a, d_b, memSize, cudaMemcpyDeviceToHost );
63 // Check for any CUDA errors
64 checkCUDAError("memcpy");
65 // verify the data returned to the host is correct
66 for (int i = 0; i < dimA; i++)
67 {
68 assert(h_a[i] == dimA - 1 - i );
69 }
70 // free device memory
71 cudaFree(d_a);
72 cudaFree(d_b);
73 // free host memory
74 free(h_a);
75 // If the program makes it this far, then the results are correct and
76 // there are no run-time errors. Good work!
77 printf("Correct!
");
78 return 0;
79 }
80
81 void checkCUDAError(const char *msg)
82 {
83 cudaError_t err = cudaGetLastError();
84 if( cudaSuccess != err)
85 {
86 fprintf(stderr, "Cuda error: %s: %s.
", msg, cudaGetErrorString( err) );
87 exit(EXIT_FAILURE);
88 }
89 }
reverseArray_multiblock_fast .cu