最大化GPU利用率

cudaStream_t stream[2];
for (int i = 0; i < 2; ++i)
cudaStreamCreate(&stream[i]);
float *hostPtr;
cudaMallocHost(&hostPtr, 2 * size);
for (int i = 0; i < 2; ++i) {
cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size, size, cudaMemcpyHostToDevice, stream[i]);
MyKernel<<<100, 512, 0, stream[i]>>>(outputDevPtr + i * size, inputDevPtr + i * size, size);
cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size, size, cudaMemcpyDeviceToHost, stream[i]);
}
for (int i = 0; i < 2; ++i)
cudaStreamDestroy(stream[i]);

在上面的例子中如果使用cudaMemcpy，CPU会阻塞直到主机内存拷贝完毕，但是使用cudaMemcpyAsync，程序会继续执行，也就是说CPU在内存传输时可以继续工作。这里使用到了流的技术，它会把GPU的内存分为多个部分，一部分用来接收主机的内存并且立即开始计算，当计算完成后，流会自动复制回主机内存，另外一半会复制另外一部分主机内存，达到重叠数据传输和计算的效果。如果kernel比较小，那这种模式可以充分利用GPU，如果不指定流，默认会安排在0号流中排队执行，如果指定了流，就可以提高并行效率。在以下的例子中，如果所有操作都是1s内就可以完成，那么运行时间都是一致的。

现在大部分系统CPU和GPU都是通过PCIe进行通信的，少部分会使用NVLink，通信的代价通常比较高，如果能够异步地进行传输可以提升效率。

以下代码示例创建了两个流，每个流都用来作为主机内存拷贝到设备以及设备内存拷贝到主机的并行操作。

cudastream_t s1，s2;
cudastreamCreate(&s1); cudaStreamCreate(&s2):
# 执行3s
cudaMemcpy(&d_arr, &h._arr,numbytes,cudaH2D);
A<<<1,128>>>(d_arr);
cudaMemcpy(&h_arr,&d_arr，numbytes，cudaD2H);
# 执行3s
cudaMemcpyAsync(&d_arr，&h_arr，numbytes，cudaH2D, s1);
A<<<1,128, s1>>>(d_arr);
cudaMemcpyAsync(&h_arr,&d_arr, numbytes,cudaD2H,s1);

# 执行3s
cudaMemcpyAsync(&d_arr1,&h_arr1, numbytes，cudaH2D，s1);
A<<<1, 128, s1>>>(d_arr1);
cudaMemcpyAsync(&h_arr1,&d_arr1, numbytes, cudaD2H, s1);
cudaMemcpyAsync(&d_arr2,&h_arr2, numbytes，cudaH2D，s2);
B<<<1,192, s2>>>(d_arr2);
cudaMemcpyAsync(&h-arr2,d_arr2，numbytes，cudaD2H, cudaH2D, s1);

# 执行3s
cudaMemcpyAsync(&d_arr1,&h_arr1，numbytes，cudaH2D,s1);  
cudaMemcpyAsync(&d_arr2,&h_arr2, numbytes , cudaH2D,s2);  
A<<<1,128,s1>>>(d_arr1);
B<<<1,192,s2>>>(d_arr2];
cudaMemcpyAsync(&h_arr1,&d_arr1,numbytes，cudaD2H，s1);
cudaMemcpyAsync(&h_arr2,&d_arr2，numbytes, cudaD2H，s2); 

将没有数据依赖，可并行的kernel使用流加速。

编写kernel代码时谨慎使用寄存器和共享内存，防止影响占有率（occupancy）。

A100每个SM上面的占有率有如下受限因素：

• 线程块数量2个
• 每个线程块中最多运行的线程数量1024
• 所有线程块共享的寄存器个数65536
• 共享内存大小20KB

• block大小应该为warp的倍数

对于GPU而言，如果一个线程被分配更多的work时，可能会更好地覆盖延时。如果线程有更多的work时，对于编译器而言，就可能有更多的机会对相关指令进行重排，从而去覆盖访存时的巨大延时。因此block大小设置为warp的倍数，可以充分发挥GPU的性能。