最大化GPU利用率

使用异步函数和流，最大化主机端（host）任务、设备端（device）任务和主机设备通信任务的并行性。

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在内存传输时可以继续工作。这里使用了CUDA流技术，通过将数据拷贝与内核计算操作提交到不同的流中，实现异步并发执行。例如，一个流可同时负责将主机数据异步传输至GPU内存，而另一个流则在GPU上并行执行计算任务。当计算完成后，结果可异步回传至主机内存。这种多流并发调度机制，有效实现了数据传输与计算的重叠，显著提升了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的性能。