现在的位置: 首页 > 综合 > 正文

CUDA, 用于大量数据的超级运算:第十二节

2014年01月22日 ⁄ 综合 ⁄ 共 6338字 ⁄ 字号 评论关闭

 http://www.ddj.com/architect/217500110 

转自:http://blog.csdn.net/gemin/article/details/4807927
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。

[c-sharp] view
plain
copy

  1. // incrementMappedArrayInPlace.cu  
  2. #include <stdio.h>  
  3. #include <assert.h>  
  4. #include <cuda.h>  
  5. // define the problem and block size  
  6. #define NUMBER_OF_ARRAY_ELEMENTS 100000  
  7. #define N_THREADS_PER_BLOCK 256  
  8. void incrementArrayOnHost(float *a, int N)  
  9. {  
  10.   int i;  
  11.   for (i=0; i < N; i++) a[i] = a[i]+1.f;  
  12. }  
  13. __global__ void incrementArrayOnDevice(float *a, int N)  
  14. {  
  15.   int idx = blockIdx.x*blockDim.x + threadIdx.x;  
  16.   if (idx < N) a[idx] = a[idx]+1.f;  
  17. }  
  18. void checkCUDAError(const char *msg)  
  19. {  
  20.   cudaError_t err = cudaGetLastError();  
  21.   if( cudaSuccess != err) {  
  22.     fprintf(stderr, "Cuda error: %s: %s./n", msg, cudaGetErrorString( err) );  
  23.     exit(EXIT_FAILURE);  
  24.   }                           
  25. }  
  26. int main(void)  
  27. {  
  28.   float *a_m; // pointer to host memory  
  29.   float *a_d; // pointer to mapped device memory  
  30.   float *check_h;   // pointer to host memory used to check results  
  31.   int i, N = NUMBER_OF_ARRAY_ELEMENTS;  
  32.   size_t size = N*sizeof(float);  
  33.   cudaDeviceProp deviceProp;  
  34. #if CUDART_VERSION < 2020  
  35. #error "This CUDART version does not support mapped memory!/n"  
  36. #endif  
  37.   // Get properties and verify device 0 supports mapped memory  
  38.   cudaGetDeviceProperties(&deviceProp, 0);  
  39.   checkCUDAError("cudaGetDeviceProperties");  
  40.   if(!deviceProp.canMapHostMemory) {  
  41.     fprintf(stderr, "Device %d cannot map host memory!/n", 0);  
  42.     exit(EXIT_FAILURE);  
  43.   }  
  44.   // set the device flags for mapping host memory  
  45.   cudaSetDeviceFlags(cudaDeviceMapHost);  
  46.   checkCUDAError("cudaSetDeviceFlags");  
  47.   // allocate mapped arrays   
  48.   cudaHostAlloc((void **)&a_m, size, cudaHostAllocMapped);  
  49.   checkCUDAError("cudaHostAllocMapped");  
  50.   // Get the device pointers to the mapped memory  
  51.   cudaHostGetDevicePointer((void **)&a_d, (void *)a_m, 0);  
  52.   checkCUDAError("cudaHostGetDevicePointer");  
  53.   // initialization of host data  
  54.   for (i=0; i<N; i++) a_m[i] = (float)i;  
  55.   // do calculation on device:  
  56.   // Part 1 of 2. Compute execution configuration  
  57.   int blockSize = N_THREADS_PER_BLOCK;  
  58.   int nBlocks = N/blockSize + (N%blockSize > 0?1:0);  
  59.   // Part 2 of 2. Call incrementArrayOnDevice kernel   
  60.   incrementArrayOnDevice <<< nBlocks, blockSize >>> (a_d, N);  
  61.   checkCUDAError("incrementArrayOnDevice");  
  62.   /* Note the allocation, initialization and call to incrementArrayOnHost 
  63.      occurs asynchronously to the GPU */  
  64.   check_h = (float *)malloc(size);  
  65.   for (i=0; i<N; i++) check_h[i] = (float)i;  
  66.   incrementArrayOnHost(check_h, N);  
  67.   // Make certain that all threads are idle before proceeding  
  68.   cudaThreadSynchronize();  
  69.   checkCUDAError("cudaThreadSynchronize");  
  70.   // check results  
  71.   for (i=0; i<N; i++) assert(check_h[i] == a_m[i]);  
  72.   // cleanup  
  73.   free(check_h); // free host memory  
  74.   cudaFreeHost(a_m); // free mapped memory (and device pointers)  
  75. }  

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

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

[c-sharp] view
plain
copy

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

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

[c-sharp] view
plain

抱歉!评论已关闭.