CUDA, 用于大量数据的超级运算：第12节

 http://www.ddj.com/architect/217500110
Rob Farber
CUDA 2.2改变数据移动样例
Rob Farber 是西北太平洋国家实验室（Pacific Northwest National Laboratory）的高级科研人员。他在多个国家级的实验室进行大型并行运算的研究，并且是几个新创企业的合伙人。大家可以发邮件到rmfarber@gmail.com与他沟通和交流。

在关于CUDA的系列文章第11节CUDA,用于大量数据的超级运算: 里，我重新讨论了CUDA内存空间，介绍了“纹理内存”的概念。在本小节，我讨论了新发布的CUDA2.2版的一些主要特点改变－即，介绍了“映射”固定系统内存。有了固定系统内存，计算机内核可共享主机系统内存，当在诸多CUDA启动的图形处理器上运行时，为对主机系统内存的直接读取提供零拷贝支持。本系列文章的下一节将继续讨论纹理内存，并且还会涉及有关CUDA2.2版的一些信息，如向在GPU（有纹理与之绑定）上的全局内存吸入的能力。（点击
这里，以了解更多有关CUDA 2.2的信息.)。
在CUDA2.2之前，CUDA内核不能直接访问主机系统内存。因此，CUDA程序员使用了第一节和第二节里介绍的设计模式：
将数据移到GPU；
在GPU上执行计算；
将结果从GPU移动到主机
这个样例现在有所调整，因为CUDA2.2引入了新的API，这样主机内存通过一个新的函数cudaHostAlloc (或CUDA驱动API里的cuMemHostAlloc)被映射到设备内存里。这个新的内存类型支持如下特点：
所有GPU都可获得的“便携”固定缓冲区：
在以后的文章里会讨论对多个GPU的使用
"映射"固定缓冲区，将主机内存映射到CUDA地址空间，无需清晰的程序员初始化复制，就可异步透明读取数据。
整合的GPU通过主机处理器共享物理内存（相对于分离的GPU自带的快速全局内存）。对于较新的来说，映射固定缓存区可做为零拷贝缓冲区（特别是整合图形处理器），因为它们避免了过多拷贝。当为整合GPU开发代码时，使用映射固定内存真的可以发挥作用。
对于分离的GPU, 映射固定内存在特定的情况下，仅仅是一个性能优势。因为内存不是由GPU缓冲的：
它应该被读取或写入一次；
读取或写入内存的全局下载和存储必须被聚合以避免，2x-7x PCIe性能惩罚；
在最好的情况下，仅执行PCIe带宽性能，但是可以比cudaMemcpy 提高两倍速度，因为映射内存能通过同时读取和写入，充分利用PCIe bus的全双工能力。对cudaMemcpy 的调用一次仅在一个方向移动数据（即，半双工）。
而且，现在CUDA2.2版的一个缺陷是：所有固定分配被映射到GPU的32位线性地址空间，不管是否需要设备指针。(NVIDIA显示在随后的版本里，这将被改变每个分配基础)。

"WC" (合并式写入)内存可提供更高的内存：
因为WC内存没有被缓冲，或者缓冲不一致。可以获得更高的PCIe性能，因为在PCI Express bus转移时，内存没有被探听。NVIDIA 在它们的”CUDA2.2固定内存API“文件里提到，WC内存的性能可能比在特定的PCI Express 2.0上要快。
It may increase the host processor(s) write performance to host memory because individual writes are first combined (via an internal processor write-buffer) so that only a single burst write containing many aggregated individual writes need be issued. (Intel claims they have observed actual performance increases of over 10x but this is not typical). For more information, please see the Intel publication 合并式写入内存执行指导
主机端计算和应用程序可能会运行得更快些，因为合并式写入内存不污染内部处理器如L1 和L2 缓存。这是因为WC不执行缓存一致性，从而在执行缓存一致性时，通过减少缓存不中率和避免产生运行时间，可以增加主机处理器的效率。合并式写入通过利用单独的专用内部写缓存缓冲也可以避免缓存污染，因为它绕过并且留下了其它未动的内部处理器缓冲。
WC内存确实有缺陷，CUDA程序员不应该考虑把WC内存区当作通用目的的内存，因为它是弱有序的。换言之，从WC内存定位读取可能会返回未期待的－并且是不正确的－数据，因为之前对内存定位的写入可能会被延迟，以与其它写入合并。尽管没有程序员通过“栅栏”操作强制执行一致，对WC内存的读取还是有可能真正的“读取”旧的，或甚至是初始化的数据。
不幸的是，从WC内存强制执行的一致读取可能会导致对一些主机处理器构架的性能惩罚。不过，配有SSE4指导集的处理器提供流加载指令(MOVNTDQA) ，可从WC内存有效读取。（检查CPUID 指令是否使用EAX==1执行，ECX19位，看看SSE4.1是否可得）。请看INTEL文章， Increasing Memory Throughput With Intel Streaming SIMD Extensions 4 (Intel SSE4) Streaming Load（ 采用INTEL流SIMD拓展4（INTEL SSE4）流加载增加内存通量）
尚不清楚是否CUDA程序员需要采取行动（如使用内存栅栏)以确保WC内存到位，准备由主机或图形处理器使用）。INTEL文件声明，"[a] '内存栅栏'指令被用来确保数据生产者和数据消费者之间的一致”. CUDA驱动器内部使用WC内存，并且一旦它向GPU发出指令，就要发出一个存储栅栏指令。因此，NVIDIA文件注明，“应用程序可能根本不必要使用存储栅栏”（添加重点）。一个看上去行得通的法则就是: 在引用WC内存前，参考CUDA指令，并假定它们发出栅栏指令。或者，利用你的编译器本身的运算发布一个存储栅栏指令，确保每个之前的存储全局可视。这依赖于编译器。Linux编译器可能需要了解mm_sfence而Windows 编译器可能会使用_WriteBarrier.
这些内存特点可以单独使用或合并使用－你可以分配一个便携式，合并式写入缓存，一个便携式固定缓存，一个既不便携也不固定的合并式写入缓存，或任何其它标识启动的任一序列。
总而言之，这些特点给我们提供了便捷和改善的性能，但是也使得操作更为复杂，在CUDA驱动器，CUDA硬件和主机处理器上就要视具体版本而定了。然而，诸多应用程序获益于这些新特点。

incrementMappedArrayInPlace.cu的以下源码表是第二节中incrementArrays.cu样例的修订版，使用了新的被映射的，固定运行时间API。

// incrementMappedArrayInPlace.cu #include <stdio.h> #include <assert.h> #include <cuda.h> // define the problem and block size #define NUMBER_OF_ARRAY_ELEMENTS 100000 #define N_THREADS_PER_BLOCK 256 void incrementArrayOnHost(float *a, int N) { int i; for (i=0; i < N; i++) a[i] = a[i]+1.f; } __global__ void incrementArrayOnDevice(float *a, int N) { int idx = blockIdx.x*blockDim.x + threadIdx.x; if (idx < N) a[idx] = a[idx]+1.f; } 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); } } int main(void) { float *a_m; // pointer to host memory float *a_d; // pointer to mapped device memory float *check_h; // pointer to host memory used to check results int i, N = NUMBER_OF_ARRAY_ELEMENTS; size_t size = N*sizeof(float); cudaDeviceProp deviceProp; #if CUDART_VERSION < 2020 #error "This CUDART version does not support mapped memory!/n" #endif // Get properties and verify device 0 supports mapped memory cudaGetDeviceProperties(&deviceProp, 0); checkCUDAError("cudaGetDeviceProperties"); if(!deviceProp.canMapHostMemory) { fprintf(stderr, "Device %d cannot map host memory!/n", 0); exit(EXIT_FAILURE); } // set the device flags for mapping host memory cudaSetDeviceFlags(cudaDeviceMapHost); checkCUDAError("cudaSetDeviceFlags"); // allocate mapped arrays cudaHostAlloc((void **)&a_m, size, cudaHostAllocMapped); checkCUDAError("cudaHostAllocMapped"); // Get the device pointers to the mapped memory cudaHostGetDevicePointer((void **)&a_d, (void *)a_m, 0); checkCUDAError("cudaHostGetDevicePointer"); // initialization of host data for (i=0; i<N; i++) a_m[i] = (float)i; // do calculation on device: // Part 1 of 2. Compute execution configuration int blockSize = N_THREADS_PER_BLOCK; int nBlocks = N/blockSize + (N%blockSize > 0?1:0); // Part 2 of 2. Call incrementArrayOnDevice kernel incrementArrayOnDevice <<< nBlocks, blockSize >>> (a_d, N); checkCUDAError("incrementArrayOnDevice"); /* Note the allocation, initialization and call to incrementArrayOnHost occurs asynchronously to the GPU */ check_h = (float *)malloc(size); for (i=0; i<N; i++) check_h[i] = (float)i; incrementArrayOnHost(check_h, N); // Make certain that all threads are idle before proceeding cudaThreadSynchronize(); checkCUDAError("cudaThreadSynchronize"); // check results for (i=0; i<N; i++) assert(check_h[i] == a_m[i]); // cleanup free(check_h); // free host memory cudaFreeHost(a_m); // free mapped memory (and device pointers) }

CUDA 2.2 向cudaGetDeviceProperties检索的 cudaDeviceProp 结构添加了以下两个设备属性，这样你可以决定设备是否支持新的映像内存API（检查GPU是否是整合的图形处理器）：

下面的代码块使用前处理器检查以确定正在使用的是CUDA的有效版本，以编译映射代码。此外，函数cudaGetDeviceProperties 被调用，以进行运行时间检查确保CUDA设备支持映像内存：

#if CUDART_VERSION < 2020 #error "This CUDART version does not support mapped memory!/n" #endif // Get properties and verify device 0 supports mapped memory cudaGetDeviceProperties(&deviceProp, 0); checkCUDAError("cudaGetDeviceProperties"); if(!deviceProp.canMapHostMemory) { fprintf(stderr, "Device %d cannot map host memory!/n", 0); exit(EXIT_FAILURE); }

在设备上启动主机内存映射：

// set the device flags for mapping host memory cudaSetDeviceFlags(cudaDeviceMapHost); checkCUDAError("cudaSetDeviceFlags");

被映射数组 a_m, 然后在主机上被分配。（注：内存在这个点上被映射，但是没有设备指针。让设备指针遵循以下步骤）

// allocate host mapped arrays cudaHostAlloc((void **)&a_m, size, cudaHostAllocMapped); checkCUDAError("cudaHostAllocMapped");// allocate host mapped arrays cudaHostAlloc((void **)&a_m, size, cudaHostAllocMapped); checkCUDAError("cudaHostAllocMapped");

向映像内存添加设备指针：

// Get the device pointers to the mapped memory cudaHostGetDevicePointer((void **)&a_d, (void *)a_m, 0); checkCUDAError("cudaHostGetDevicePointer")

数据初始化，内核在GPU上被执行。与起初的incrementArrays.cu样例不同，cudaMemcpy没有明确的程序员初始化数据移动。注意数据移动和内核执行与主机运算异步。结果就是， 验证数组check_h的主机创建和计算执行，而GPU同时运行incrementArrayOnDevice 内核以通过设备内存指针a_d更新主机数组a_m。
通过调用cudaThreadSynchronize ，同步发生，之后GPU结果根据主机产生的结果进行验证。假定主机和GPU内核结果一致，程序然后进行清理。函数cudaFreeHost 被用来释放主机上的被映射数组和GPU上的指针。
在Linux下，程序使用命令行编译：

nvcc "o incrementMappedArrayInPlace incrementMappedArrayInPlace.cu

对映像内存执行安装升级的意义尚不清楚。为了确保PCIe运算为最小数，似乎在分离的数组之间进行流数据是谨慎之道。换言之，当有一个为专用读取运算，而另外一个为专用写入运算时，使用分离的数组。
演示合并式写入
下面的程序incrementMappedArrayWC.cu显示了使用分离的合并式写入，映射固定内存以一次性增加数组元素。这要求改变incrementArrayOnHost 和incrementArrayOnDevice 以从数组a 读取，并写入数组b。这样，可以避免一致性问题，并获得高流性能。cudaHostAllocWriteCombined 标识被添加到 cudaHostAlloc调用。我们依靠CUDA对驱动器的调用发布恰当的栅栏运算以确保写入全局可视。

// incrementMappedArrayWC.cu #include <stdio.h> #include <assert.h> #include <cuda.h> // define the problem and block size #define NUMBER_OF_ARRAY_ELEMENTS 100000 #define N_THREADS_PER_BLOCK 256 void incrementArrayOnHost(float *b, float *a, int N) { int i; for (i=0; i < N; i++) b[i] = a[i]+1.f; } __global__ void incrementArrayOnDevice(float *b, float *a, int N) { int idx = blockIdx.x*blockDim.x + threadIdx.x; if (idx < N) b[idx] = a[idx]+1.f; } 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); } } int main(void) { float *a_m, *b_m; // pointers to mapped host memory float *a_d, *b_d; // pointers to mapped device memory float *check_h; // pointer to host memory used to check results int i, N = NUMBER_OF_ARRAY_ELEMENTS; size_t size = N*sizeof(float); cudaDeviceProp deviceProp; #if CUDART_VERSION < 2020 #error "This CUDART version does not support mapped memory!/n" #endif // Get properties and verify device 0 supports mapped memory cudaGetDeviceProperties(&deviceProp, 0); checkCUDAError("cudaGetDeviceProperties"); if(!deviceProp.canMapHostMemory) { fprintf(stderr, "Device %d cannot map host memory!/n", 0); exit(EXIT_FAILURE); } // set the device flags for mapping host memory cudaSetDeviceFlags(cudaDeviceMapHost); checkCUDAError("cudaSetDeviceFlags"); // allocate host mapped arrays int flags = cudaHostAllocMapped|cudaHostAllocWriteCombined; cudaHostAlloc((void **)&a_m, size, flags); cudaHostAlloc((void **)&b_m, size, flags); checkCUDAError("cudaHostAllocMapped"); // Get the device pointers to memory mapped cudaHostGetDevicePointer((void **)&a_d, (void *)a_m, 0); cudaHostGetDevicePointer((void **)&b_d, (void *)b_m, 0); checkCUDAError("cudaHostGetDevicePointer"); /* initialization of the mapped data. Since a_m is write-combined, it is not guaranteed to be initialized until a fence operation is called. In this case that should happen when the kernel is invoked on the GPU */ for (i=0; i<N; i++) a_m[i] = (float)i; // do calculation on device: // Part 1 of 2. Compute execution configuration int blockSize = N_THREADS_PER_BLOCK; int nBlocks = N/blockSize + (N%blockSize > 0?1:0); // Part 2 of 2. Call incrementArrayOnDevice kernel incrementArrayOnDevice <<< nBlocks, blockSize >>> (b_d, a_d, N); checkCUDAError("incrementArrayOnDevice"); // Note the allocation and call to incrementArrayOnHost occurs // asynchronously to the GPU check_h = (float *)malloc(size); incrementArrayOnHost(check_h, a_m,N); // Make certain that all threads are idle before proceeding cudaThreadSynchronize(); checkCUDAError("cudaThreadSynchronize"); // cudaThreadSynchronize() should have caused an sfence // to be issued, which will guarantee that all writes are done // check results. Note: the updated array is in b_m, not b_d for (i=0; i<N; i++) assert(check_h[i] == b_m[i]); // cleanup free(check_h); // free mapped memory (and device pointers) cudaFreeHost(a_m); cudaFreeHost(b_m); }

结论
CUDA 2.2 通过为主机和GPU之间的被映射的，透明数据转移提供APIs，改变数据运动范式。这些APIs允许CUDA程序员在主机和图形处理器之间更为有效地共享数据。程序员们可通过使用合并式写入内存，利用异步运算和全双工PCIe数据转移。程序员还能够使用多个GPU共享固定内存。
就我个人而言，当移植现有的科学代码到GPU上时，我把这些APIs做为一个便捷条件在使用，因为被映射的内存让我可以保持主机和设备数据的同步，并且可以逐步尽可能地将计算移到GPU上。这样就可以让我在每个改变之后验证结果，以确保没有问题。实际上，当使用互相依赖的复杂代码时，这样可以实时节省很多精力。此外，我还使用这些APIs，通过利用异步主机和多个GPU计算，还有全双工PCIe转移和其它一些CUDA2.2的特点提高效率。
我还看到新的CUDA2.2 APIs促进了应用程序如操作系统和实时系统的整个全新的类型的发展。
有个样例就是Alabama 大学和Sandia国家实验室进行的RAID 研究，它将CUDA启动的GPU变成了高性能的RAID加速器，可以为高通量硬盘子系统实时计算Reed-Solomon代码（参见Matthew Curry, Lee Ward, Tony Skjellum, Ron Brightwell写的Accelerating Reed-Solomon Coding in RAID Systems with GPUs)。他们的摘要声称，“性能结果显示，在这个问题上，GPU的表现能超过现代的CPU一个数量级，并且确定了GPU可被用来支持至少有三个同位碟的系统，而不会产生性能惩罚。”
我的猜想是，我们会看到在不久的将来看到由CUDA提高的Linuxmd (多个设备或软件RAID)驱动器。想象下，无需锁入专用的RAID控制器中。如果有问题，将你的RAID数组和另外一个Linux盒连接起来，读取数据。如果计算机没有NVIDIA GPU，那么就使用标准的Linux 软件md驱动器读取数据。
不要忘记CUDA启动的设备可以加速和同时运行多个应用程序。下面的文章会讨论如何将合并图形和CUDA。该文章也将利用这一能力。试着运行下单独的图形应用程序，同时也运行你的一个CUDA应用程序吧。我想，你会对这两种应用程序的性能感到惊喜的。