11.1 设备内存

如果输入和输出数据都在设备内存上，优化如SAXPY这样低计算密度的任务就变成了优化全局内存访问的问题。除了对齐和合并读取约束外，CUDA内核的性能也对线程块数目和每块的线程数目比较敏感。globalRead、globalWrite、globalCopy和globalCopy2应用程序（在本书配套源代码的memory/子目录）生成针对多种操作数大小、块大小和循环展开因子的带宽报告。由globalCopy2（它遵循类似于SAXPY的内存访问模式：每个循环迭代包含两次读操作和一次写操作）生成的一个示例报告如代码清单11-3所示。

如果我们引用第5章的globalCopy2.cu应用程序（参见代码清单5-8），在GK104运行它，得到4个字节操作数时的输出，如代码清单11-3所示。最上面一行（展开因子为1）对应于简单实现（类似于代码清单11-2）。从中可以看到，循环展开具有少量的性能优势。当使用的展开因子为4时，得到约 $10\%$ 的加速，提供的带宽为128GB/s，而不是简单实现时的116GB/s。

有趣的是，使用#pragma unroll编译器指令仅提升性能到约118GB/s，而修改带模板的globalCopy2.cu内核来执行SAXPY，性能则可提升至135GB/s。代码清单11-4给出了最终的内核，它在

stream1Device.cu应用程序中实现（见本书配套源代码codahandbook/Streaming/目录）。

对于大多数应用程序，这些小的性能差异并不能证明这种重写内核方式的价值。但是，如果内核被编写为“跟线程块无关”（即能在任何网格或线程块大小下正确工作），那么最优的设置可以凭经验确定，而无须花费太多精力。

代码清单11-3 globalCopy2输出（GK104）

代码清单11-4 saxpyGPU内核函数（带模板的展开）

template<const int n>
    __device__void
saxpy_unrolled(
        float *out,
        const float *px,
        const float *py,

size_t N, float alpha)   
{ float  $\mathbf{x}[\mathbf{n}]$  ,y[n]; size_t i; for  $\mathrm{i} = \mathrm{n}^{\star}$  blockIdx.x\*blockDim.x+threadIdx.x; i  $<$  N-n\*blockDim.x\*gridDim.x; i  $^{**}$  n\*blockDim.x\*gridDim.x} { for (int j  $= 0$  ;j<n;j++) { size_t index  $=$  i+j\*blockDim.x; x[j]  $=$  px[index]; y[j]  $=$  py[index]; } for (int j  $= 0$  ;j<n;j++) { size_t index  $=$  i+j\*blockDim.x; out[index]  $=$  alpha\*x[j]+y[j]; } //to avoid the (index<N) conditional in the inner loop, //we left off some work at the end for (int j  $= 0$  ;j<n;j++) { for (int j  $= 0$  ;j<n;j++) { size_t index  $=$  i+j\*blockDim.x; if (index<N）{ x[j]  $=$  px[index]; y[j]  $=$  py(index]; } } for (int j  $= 0$  ;j<n;j++) { size_t index  $=$  i+j\*blockDim.x; if (index<N）out[index]  $=$  alpha\*x[j]+y[j]; } } _global_void saxpyGPU(float \*out,const float \*px,const float \*py,size_t N, float alpha) { saxpy_unrolled4>(out,px,py,N,alpha); }

该stream1Device.cu应用程序报告了从基于分页的系统内存复制数据到设备内存，在数据上运行代码清单11-4所示的内核，并将数据传输回去所需的系统时钟时间（wall clock time）。测试系统采用 GeForce GTX680，运行在Intel i7和Windows 7环境下，这个应用程序的输出如下所示。

Measuring times with 128M floats (use --N to specify number of Mfloats)  
Memcpy(host->device): 365.95 ms (2934.15 MB/s)  
Kernel processing: 11.94 ms (134920.75 MB/s)  
Memcpy (device->host): 188.72 ms (2844.73 MB/s)  
Total time (wall clock): 570.22 ms (2815.30 MB/s)

内核的执行仅占用整体执行时间的一小部分，大约 $2\%$ 的系统时钟时间。其余 $98\%$ 的时间花费在把数据传入和传出GPU！对于像这样的传输受限型负载，如果待处理的数据的一部分或者全部是在主机内存中，那么优化应用程序的最好方法是提高CPU/GPU执行的重叠程度和传输性能。

11.1_设备内存

11.1 设备内存