CUDA：Supercomputing for the Masses (用于大量数据的超级计算)-第三节
第三节：错误处理和全局内存性能局限
恭喜！通过对CUDA（Compute Unified DeviceArchitecture，即计算统一设备架构的首字母缩写）系列文章第一节和第二节，您现在已经是能够使用CUDA的程序员了，您可以创建和运行在支持CUDA的设备上使用成百上千同步线程的程序。在第二节的incrementArrays.cu中，我提供了一个常见的CUDA应用程序模式的工作样例——将数据移动到设备，运行一个或多个内核以进行计算并获得结果。本质上，只需使用您自己的内核并加载自己的数据（我在本篇专栏文章的示例中就是这样做的）就可以将incrementArrays.cu变形到任何您需要的应用程序中。以后的专栏文章将介绍CUDA异步I/O和流。
“你知道的越多，就越危险！“是对上一段内容的幽默而又准确的总结。CUDA的好消息就是，它提供了一种自然的方式将您作为程序员的思路转换到大量平行的程序。不过坏消息是， 要让这些程序更“结实”，更为有效，需要更高的理解力。
不要小心翼翼。开始试验吧，动手起来！CUDA提供了创建优秀软件的编程工具和结构。想要真正学会它，就要动手试验。实际上，这些专栏文章通过简短的样例重点介绍CUDA功能，并为您提供网路上的优秀信息资源，以补充您的经验和学习过程。记住，CUDAZone是CUDA资源的集散地，而论坛上会提供很多你提出的问题的答案。另外，它们的互动优势能让你提出问题，得到答案。
本专栏文章和以下几个专栏文章将利用一个简单的数组反向应用程序来扩充您的知识，将重点介绍共享内存的性能影响。我将和CUDA剖析工具一起介绍错误检查和性能行为。此外，还包含了下一专栏文章的资源列表，这样您可以看到如何通过共享内存实现数组反向。程序reverseArray_multiblock.cu采用明显但性能较低的方式实现了在CUDA设备上反向全局内存中的数组。不要将它用作应用程序的模型，因为对于此类应用程序来说全局内存并不是最佳的内存类型——而且此版本还要进行不结合的内存访问，这会对全局内存性能产生不利影响。只有当同时的内存访问能够结合成单个的内存事务时，我们才能获得最佳的全局内存带宽。在后续的专栏文章中，我将介绍全局内存和共享内存之间的不同，以及根据设备的计算能力对要结合的内存访问的各种要求。
CUDA 错误处理
创建健壮和实用的软件，发现和处理错误非常重要。当用户的应用程序出现故障，或产生错误的结果时，他们通常会非常恼怒。对于程序开发人员来说，添加错误处理代码是个令人恼火和繁琐的工作。它会使原本整洁的代码变得散乱，为了处理每个能想到的错误，开发的进程被延缓。是的，错误处理是个不讨好的工作，但是记住，你不是为自己做这个工作（尽管良好的错误检查机制已经挽救了我无数次）――您是为了将来使用这些程序的人做这项工作。如果出现故障，用户需要了解为什么会这样，更重要的是，他们能够做些什么来解决问题。有效的错误处理和回复的确可以让你的应用程序受到用户的欢迎。商业开发人员尤其需要注意到这一点。
CUDA设计人员明了有效错误处理的重要性。为了更好的处理错误，每个CUDA调用（包括内核启动异常）都会返回一个类型为cudaError_t.的错误代码。一旦成功完成，返回cudaSuccess。否则，返回错误代码。
char *cudaGetErrorString(cudaError_t code);
C语言程序员会发现此方法和C库之间的相似处，C库使用变量errno来表示错误，使用perror和strerror来获得适合阅读的错误消息报告。C库样例已经很好地为几百万行C代码服务了，无疑它将来也会很好地为CUDA软件服务。
CUDA还提供了一个方法，cudaGetLastError, 它报告主机线程里的任何之前的运行调用的最后一次错误。它有几个作用：
内核启动的不同步本质排除了通过cudaGetLastError显示检查错误的可能。相反，使用cudaThreadSynchronize会阻止错误检查直到设备完成所有之前的调用，包括内核调用，并且如果前面的某个任务失败它会返回错误。多个内核启动排队则意味着只有在所有内核都完成以后才能进行错误检查——除非程序员在内核内进行了明显的错误检查并向主机报告。
错误被报告给正确的主机线程。如果主机正在运行多线程，很有可能当应用程序正在使用多个CUDA设备时，错误被报告给正确的主机线程。
当有多个错误在对cudaGetLastError的调用之间发生时，仅最后一个错误会被报告。这意味着程序员必须注意将错误与生成该错误的运行时调用相连或冒险给用户发送一个不正确的错误报告。
查看源代码
查看reverseArray_multiblock.cu的源代码，你会注意到该程序的结构非常类似于第二节的moveArrays.cu的结构。提供了一个错误例程checkCUDAError，这样主机可打印出可阅读的消息，并当错误被报告时（通过cudaGetLastError），退出。您看到了，在整个程序中，我们巧妙地利用了checkCUDAError来检查错误。
程序reverseArray_multiblock.cu实质上创建了一个1D整数数组，h_a, 包含整数值 [0 ..dimA-1]。数组h_a通过cudaMemcpy移动到数组d_a，后者位于设备的全局内存里。主机然后启动reverseArrayBlock内核，以反向顺序从d_a到 d_b拷贝数组内容（这是另外一个全局内存数组）。使用cudaMemcpy 来传输数据－这次是从d_b到主机。然后进行主机检查，以确认设备给出了正确的结果（比如，[dimA-1 .. 0]）。
1 #include <stdio.h> 2 #include <assert.h> 3 #include "cuda.h" 4 #include "cuda_runtime.h" 5 #include "device_launch_parameters.h" 6 //检查CUDA运行时是否有错误 7 void checkCUDAError(const char* msg); 8 // Part3: 在全局内存执行内核 9 /* 10 blockDim块内的线程数 11 blockIdx网格内的块索引 12 gridDim网格内块个数 13 threadIdx块内线程索引 14 */ 15 __global__ void reverseArrayBlock(int *d_out, int *d_in) 16 { 17 int inOffset = blockDim.x * blockIdx.x; 18 int outOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x); 19 int in = inOffset + threadIdx.x; 20 int out = outOffset + (blockDim.x - 1 - threadIdx.x); 21 d_out[out] = d_in[in]; 22 } 23 ///////////////////////////////////////////////////////////////////// 24 //主函数 25 ///////////////////////////////////////////////////////////////////// 26 int main(int argc, char** argv) 27 { 28 //指向主机的内存空间和大小 29 int *h_a; 30 int dimA = 256 * 1024; // 256K elements (1MB total) 31 //指向设备的指针和大小 32 int *d_b, *d_a; 33 //定义网格和块大小，每个块的线程数量 34 int numThreadsPerBlock = 256; 35 36 /* 37 根据数组大小和预设的块大小来计算需要的块数 38 */ 39 int numBlocks = dimA / numThreadsPerBlock; 40 //申请主机及设备上的存储空间 41 size_t memSize = numBlocks * numThreadsPerBlock * sizeof(int); 42 //主机上的大小 43 h_a = (int *)malloc(memSize); 44 //设备上的大小 45 cudaMalloc((void **)&d_a, memSize); 46 cudaMalloc((void **)&d_b, memSize); 47 //在主机上初始化输入数组 48 for (int i = 0; i < dimA; ++i) 49 { 50 h_a[i] = i; 51 } 52 //将主机数组拷贝到设备上，h_a-->d_a 53 cudaMemcpy(d_a, h_a, memSize, cudaMemcpyHostToDevice); 54 //启动内核 55 dim3 dimGrid(numBlocks); 56 dim3 dimBlock(numThreadsPerBlock); 57 reverseArrayBlock <<< dimGrid, dimBlock >>>(d_b, d_a); 58 //阻塞，一直到设备完成计算 59 cudaThreadSynchronize(); 60 //检查是否设备产生了错误 61 //检查任何CUDA错误 62 checkCUDAError("kernel invocation"); 63 //将结果从设备拷贝到主机，d_b-->h_a 64 cudaMemcpy(h_a, d_b, memSize, cudaMemcpyDeviceToHost); 65 //检查任何CUDA错误 66 checkCUDAError("memcpy"); 67 //核对返回到主机上的结果是否正确 68 for (int i = 0; i < dimA; i++) 69 { 70 assert(h_a[i] == dimA - 1 - i); 71 } 72 //释放设备内存 73 cudaFree(d_a); 74 cudaFree(d_b); 75 //释放主机内存 76 free(h_a); 77 printf("Correct!\n"); 78 return 0; 79 } 80 void checkCUDAError(const char *msg) 81 { 82 cudaError_t err = cudaGetLastError(); 83 if (cudaSuccess != err) 84 { 85 fprintf(stderr, "Cuda error: %s: %s.\n", msg,cudaGetErrorString(err)); 86 exit(EXIT_FAILURE); 87 } 88 }
该程序的一个关键设计特征是，两个数组d_a 和d_b 位于设备上的全局内存里。CUDA SDK提供样例程序bandwidthTest，它提供了一些关于设备特点的信息。在我的系统中，全局内存带宽刚刚超过60GB/s。如果要为128个硬件线程提供服务，这将是很有用的—－每个线程都能提供大量的浮点操作。因为一个32位浮点值占据四个字节，此设备上的全局内存带宽受限应用程序只能提供大约1515 GF/s–或可用性能很小一部分的百分比。（假定应用程序仅从全局内存读取，并且不向全局内存写入东西）。显然，性能较高的应用程序必须重新以某种方式使用数据。这是共享和寄存器内存的功能。我们程序员的工作就是获得这些内存类型的最大效益。要想更好的了解浮点能力与内存带宽之间的机器平衡法则（和其他的机器特征），请阅读我的文章HPC Balance and Common Sense。
共享内存版
以下资源列表是关于arrayReversal_multiblock_fast.cu, 我会在下一部分里介绍。
我现在提供它是为了方便您了解如何在这个问题上使用共享内存。
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!\n"); 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.\n", msg, cudaGetErrorString(err)); 103 exit(EXIT_FAILURE); 104 } 105 }
在下一专栏文章中，我将介绍共享内存的使用以提高性能。那时，我会深入介绍CUDA内存类型——特别是 __shared__、__constant__和register memory。
