对于上一段的内容，可以用一句幽默又准确的话来概括，即“知道得越多越危险”。关于CUDA的好消息是它提供了一种自然的方式将您作为程序员的思路转换到大量平行的程序。坏消息是要使这些程序健壮又高效需要更高的理解力。
放开胆量，开始试验，动手做一下吧！CUDA提供了用来创建优秀软件的编程工具和结构，要想真正学会它，就得做试验。实际上，这些专栏文章通过用简短的示例重点介绍CUDA功能并为您提供Internet上的优秀信息资源来补充您的试验和学习过程。记住，CUDA Zone是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]）。
// includes, system
#include
#include
// Simple utility function to check for CUDA runtime errors
void checkCUDAError(const char* msg);
// Part3: implement the kernel
__global__ void reverseArrayBlock(int *d_out, int *d_in)
{
int inOffset = blockDim.x * blockIdx.x;
int utOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x);
int in = inOffset + threadIdx.x;
int ut = outOffset + (blockDim.x - 1 - threadIdx.x);
d_out[out] = d_in[in];
}
/////////////////////////////////////////////////////////////////////
// Program main
/////////////////////////////////////////////////////////////////////
int main( int argc, char** argv)
{
// pointer for host memory and size
int *h_a;
int dimA = 256 * 1024; // 256K elements (1MB total)
// pointer for device memory
int *d_b, *d_a;
// define grid and block size
int numThreadsPerBlock = 256;
// Part 1: compute number of blocks needed based on
// array size and desired block size
int numBlocks = dimA / numThreadsPerBlock;
// allocate host and device memory
size_t memSize = numBlocks * numThreadsPerBlock * sizeof(int);
h_a = (int *) malloc(memSize);
cudaMalloc( (void **) &d_a, memSize );
cudaMalloc( (void **) &d_b, memSize );
// Initialize input array on host
for (int i = 0; i < dimA; ++i)
{
h_a[i] = i;
}
// Copy host array to device array
cudaMemcpy( d_a, h_a, memSize, cudaMemcpyHostToDevice );
// launch kernel
dim3 dimGrid(numBlocks);
dim3 dimBlock(numThreadsPerBlock);
reverseArrayBlock<<< dimGrid,
dimBlock >>>( d_b, d_a );
// block until the device has completed
cudaThreadSynchronize();
// check if kernel execution generated an error
// Check for any CUDA errors
checkCUDAError("kernel invocation");
// device to host copy
cudaMemcpy( h_a, d_b, memSize, cudaMemcpyDeviceToHost );
// Check for any CUDA errors
checkCUDAError("memcpy");
// verify the data returned to the host is correct
for (int i = 0; i < dimA; i++)
{
assert(h_a[i] == dimA - 1 - i );
}
// free device memory
cudaFree(d_a);
cudaFree(d_b);
// free host memory
free(h_a);
// If the program makes it this far, then the results are
// correct and there are no run-time errors. Good work!
printf("Correct!\n");
return 0;
}
void checkCUDAError(const char *msg)
{
cudaError_t err = cudaGetLastError();
if( cudaSuccess != err)
{
fprintf(stderr, "Cuda error: %s: %s.\n", msg,
cudaGetErrorString( err) );
exit(EXIT_FAILURE);
}
}
该程序一个重要的设计特性是两个数组d_a和d_b都驻留在设备上的全局内存中。CUDA SDK提供了示例程序bandwidthTest，它提供了关于设备特征的一些信息。在我的系统中，全局内存带宽刚刚超过60 GB/s。如果要为128个硬件线程提供服务，这将是很有用的——每个线程都能提供大量的浮点操作。一个32位浮点值占据4个字节，该硬件上全局内存带宽受限的应用程序将只能提供大概15 GF/s——或者很少的可用性能百分比（假设应用程序只读取全局内存，不向它写入东西）。很显然，性能较高的应用程序必须以某种方式重用数据。这是共享和寄存器内存的作用。我们程序员的工作就是获得这些内存类型的最大效益。要想更好的了解浮点能力与内存带宽之间的机器平衡法则（和其他的机器特征），请阅读我的文章HPC Balance and Common Sense。
共享内存版
以下资源列表是关于arrayReversal_multiblock_fast.cu的，我会在下一部分中介绍。我现在提供它是为了方便您了解如何在这个问题上使用共享内存。
// includes, system
#include
#include
// Simple utility function to check for CUDA runtime errors
void checkCUDAError(const char* msg);
// Part 2 of 2: implement the fast kernel using shared memory
__global__ void reverseArrayBlock(int *d_out, int *d_in)
{
extern __shared__ int s_data[];
int inOffset = blockDim.x * blockIdx.x;
int in = inOffset + threadIdx.x;
// Load one element per thread from device memory and store it
// *in reversed order* into temporary shared memory
s_data[blockDim.x - 1 - threadIdx.x] = d_in[in];
// Block until all threads in the block have
// written their data to shared mem
__syncthreads();
// write the data from shared memory in forward order,
// but to the reversed block offset as before
int utOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x);
int ut = outOffset + threadIdx.x;
d_out[out] = s_data[threadIdx.x];
}
/////////////////////////////////////////////////////////////////////
// Program main
/////////////////////////////////////////////////////////////////////
int main( int argc, char** argv)
{
// pointer for host memory and size
int *h_a;
int dimA = 256 * 1024; // 256K elements (1MB total)
// pointer for device memory
int *d_b, *d_a;
// define grid and block size
int numThreadsPerBlock = 256;
// Compute number of blocks needed based on array size
// and desired block size
int numBlocks = dimA / numThreadsPerBlock;
// Part 1 of 2: Compute number of bytes of shared memory needed
// This is used in the kernel invocation below
int sharedMemSize = numThreadsPerBlock * sizeof(int);
// allocate host and device memory
size_t memSize = numBlocks * numThreadsPerBlock * sizeof(int);
h_a = (int *) malloc(memSize);
cudaMalloc( (void **) &d_a, memSize );
cudaMalloc( (void **) &d_b, memSize );
// Initialize input array on host
for (int i = 0; i < dimA; ++i)
{
h_a[i] = i;
}
// Copy host array to device array
cudaMemcpy( d_a, h_a, memSize, cudaMemcpyHostToDevice );
// launch kernel
dim3 dimGrid(numBlocks);
dim3 dimBlock(numThreadsPerBlock);
reverseArrayBlock<<< dimGrid, dimBlock,
sharedMemSize >>>( d_b, d_a );
// block until the device has completed
cudaThreadSynchronize();
// check if kernel execution generated an error
// Check for any CUDA errors
checkCUDAError("kernel invocation");
// device to host copy
cudaMemcpy( h_a, d_b, memSize, cudaMemcpyDeviceToHost );
// Check for any CUDA errors
checkCUDAError("memcpy");
// verify the data returned to the host is correct
for (int i = 0; i < dimA; i++)
{
assert(h_a[i] == dimA - 1 - i );
}
// free device memory
cudaFree(d_a);
cudaFree(d_b);
// free host memory
free(h_a);
// If the program makes it this far, then results are correct and
// there are no run-time errors. Good work!
printf("Correct!\n");
return 0;
}
void checkCUDAError(const char *msg)
{
cudaError_t err = cudaGetLastError();
if( cudaSuccess != err)
{
fprintf(stderr, "Cuda error: %s: %s.\n", msg,
cudaGetErrorString( err) );
exit(EXIT_FAILURE);
}
}
在下一专栏文章中，我将介绍共享内存的使用以提高性能。那时，我会深入介绍CUDA内存类型——特别是 __shared__、__constant__和register memory。
来自 “ ITPUB博客 ” ，链接：https://blog.itpub.net/14741601/viewspace-374350/ ，如需转载，请注明出处，否则将追究法律责任。
