CUDA学习（二十九）
占用率计算器：
有几个API函数可以帮助程序员根据寄存器和共享内存的要求来选择线程块大小：
-
占用计算器API cudaOccupancyMaxActiveBlocksPerMultiprocessor可以根据内核的块大小和共享内存使用情况提供占用率预测。 该函数根据每个多处理器的并发线程块数来报告占用情况。
- 请注意，此值可以转换为其他指标。 乘以每块的warp数量会得到每个multiprocessor的并发warp数量; 通过每个多处理器的最大warp进一步划分并发warp给出占用率的百分比。
- 基于占用的启动配置程序API cudaOccupancyMaxPotentialBlockSize和cudaOccupancyMaxPotentialBlockSizeVariableSMem启发式地计算实现最大多处理器级占用率的执行配置。
以下代码示例计算MyKernel的占用情况。 然后它会报告占用率级别，其中每个多处理器的并发warp与最大warp之间的比率。
// Device code
__global__ void MyKernel(int *d, int *a, int *b)
{
int idx = threadIdx.x + blockIdx.x * blockDim.x;
d[idx] = a[idx] * b[idx];
}
// Host code
int main()
{
int numBlocks; // Occupancy in terms of active blocks
int blockSize = 32;
// These variables are used to convert occupancy to warps
int device;
cudaDeviceProp prop;
int activeWarps;
int maxWarps;
cudaGetDevice(&device);
cudaGetDeviceProperties(&prop, device);
cudaOccupancyMaxActiveBlocksPerMultiprocessor(
&numBlocks,
MyKernel,
blockSize,
0);
activeWarps = numBlocks * blockSize / prop.warpSize;
maxWarps = prop.maxThreadsPerMultiProcessor / prop.warpSize;
std::cout << "Occupancy: " << (double)activeWarps / maxWarps * 100 << "%" <<
std::endl;
return 0;
}以下代码示例根据用户输入配置MyKernel的基于占位的内核启动。
// Device code
__global__ void MyKernel(int *array, int arrayCount)
{
int idx = threadIdx.x + blockIdx.x * blockDim.x;
if (idx < arrayCount) {
array[idx] *= array[idx];
}
}
// Host code
int launchMyKernel(int *array, int arrayCount)
{
int blockSize; // The launch configurator returned block size
int minGridSize; // The minimum grid size needed to achieve the
// maximum occupancy for a full device
// launch
int gridSize; // The actual grid size needed, based on input
cudaOccupancyMaxActiveBlocksPerMultiprocessor // size
cudaOccupancyMaxPotentialBlockSize(
&minGridSize,
&blockSize,
(void*)MyKernel,
0,
arrayCount);
// Round up according to array size
gridSize = (arrayCount + blockSize - 1) / blockSize;
MyKernel << <gridSize, blockSize >> >(array, arrayCount);
cudaDeviceSynchronize();
// If interested, the occupancy can be calculated with
// cudaOccupancyMaxActiveBlocksPerMultiprocessor
return 0;
}CUDA工具包还为 / include / cuda_occupancy.h中的任何不依赖于CUDA软件堆栈的使用情况提供自记录，独立占用计算器和启动配置器实现。 还提供了电子表格版本的占用率计算器。 电子表格版本作为一种学习工具特别有用，可以将影响占用率的参数（块大小，每个线程的寄存器和每个线程的共享内存）
最大化内存吞吐量：
为应用程序最大化整体内存吞吐量的第一步是使带宽较低的数据传输最小化。
这意味着最大限度地减少主机和设备之间的数据传输，详见主机和设备之间的数据传输，因为它们比全局内存和设备之间的数据传输具有更低的带宽。
这也意味着通过最大限度地利用片上内存来最大限度地减少全局内存和设备之间的数据传输：共享内存和高速缓存（例如，计算能力2.x或更高的设备上可用的L1高速缓存和L2高速缓存，纹理高速缓存和常量高速缓存 适用于所有设备）。
共享内存相当于用户管理的缓存：应用程序显式分配并访问它。 如CUDA C运行时所示，典型的编程模式是将来自设备内存的数据放入共享内存; 换句话说，要让块的每个线程：
- 将数据从设备内存加载到共享内存
- 与块的所有其他线程同步，以便每个线程可以安全地读取由不同线程填充的共享内存位置，
- 处理共享内存中的数据，
- 如果需要再次同步以确保共享内存已更新结果，
- 将结果写回设备内存
对于某些应用程序（例如，全局内存访问模式依赖数据），传统的硬件管理缓存更适合利用数据局部性。 如计算能力3.x和计算能力7.x中所述，对于计算能力3.x和7.x的设备，L1和共享内存使用相同的片上存储器，并且其中有多少是专用的 到L1与共享内存是可配置的每个内核调用。
内核对内存访问的吞吐量可以根据每种内存的访问模式而变化一个数量级。 因此，最大化内存吞吐量的下一步是基于设备内存访问中描述的最佳内存访问模式尽可能最佳地组织内存访问。 这种优化对于全局内存访问尤为重要，因为全局内存带宽较低，所以非最佳全局内存访问对性能影响较大。