10.12 最优线程块大小选择以及性能

当以“显而易见”的方式生成纹理坐标，如tex2d_addressing.cu中的方式

row = blockIdx.y*blockDim.y + threadIdx.y;  
col = blockIdx.x*blockDim.x + threadIdx.x;  
... tex2D(tex, (float) col, (float) row);

则纹理操作的性能取决于线程块的大小。

为了找到线程块的最优大小，程序tex2D_shmoo.cu与surf2Dmemset_shmoo.cu通过计时的方式记录了当线程块的宽度和高度分别设置从 $4\sim 64$ 之间的任意一个值时的性能。其中部分线程块的宽高组合是无效的，因为线程的总数目超出限制。

对于该程序，为了最大程度地体现硬件纹理单元的性能，纹理操作内核将设计为“执行尽可能少的任务”，但编译器仍将生成可执行代码。内核中每个线程将计算读出的所有值的总和，并在输出参数为非空时将总和写入输出中。然而，在执行的时候我们将不会传递一个非空指针给内核。这样设计内核的原因主要是：若直接去掉这段将结果写回输出内存的代码，编译器将发现内核没有做任何工作，因此将生成没有任何纹理操作的可执行代码。

extern "C" _global _ void   
TexSums(float \*out,size_t Width,size_t Height)   
{ float sum  $= 0$  .0f; for ( int row  $=$ (blockIdx.y\*blockDim.y+threadIdx.y; row  $<$  Height; row  $+ =$  blockDim.y\*gridDim.y) { for (int col  $=$  blockIdx.x\*blockDim.x  $^+$  threadIdx.x; col  $<$ Width; col  $+ =$  blockDim.x\*gridDim.x } sum  $+ =$  tex2D( tex,(float)col,(float)row); } if (out){ out[blocks.x\*blockDim.x+threadIdx.x]  $=$  sum;

尽管我们使用了“技巧”，编译器仍可能会生成“先检查参数out的值，并在看到out为NULL时则直接退出内核的”可执行代码。因此，我们需要合成一些不会太影响性能的输出（例如，每个线程块在共享内存上进行和的归约计算，并将结果写到输出参数out中）。但在编译程序时选择--keep选项，并使用cuobjdump---dump-sass查看微码（microcode），可以看到在执行了双重嵌套的for循环之后，编译器才检查参数out。

结果

在 GeForce GTX 280（GT200）上，线程块的最优大小为128个线程，得到的带宽为35.7GB/s。线程块配置为 $32 \times 4$ 的速度与 $16 \times 8$ 或 $8 \times 16$ 相同，均在1.88ms内传输了4K $\times$ 4K的浮点型纹理。在特斯拉M2050上，线程块的最优大小为192个线程，带宽为35.4GB/s。与GT200一样，宽高不同但整体线程数相同的线程块的执行速度相同，宽和高分别为 $6 \times 32$ 、 $16 \times 12$ 与 $8 \times 24$ 的线程块在传输数据上性能相同。

二维表面内存填充程序的测试结果说服力较小。线程块的线程数必须超过128个线程才会有较好的性能，并且需要线程数目能够被线程束的大小32整除。在cg1.4xlarge上不开启ECC，二维表面内存上的内存填充的性能最快能达到48GB/s。

对于这两个GPU卡均使用浮点值数据进行测试，纹理操作和表面内存写操作的峰值带宽分别是“全局加载及存储”峰值带宽的1/4和1/2。