CUDA并发内核在使用事件时序列化
我在遇到并行执行的CUDA内核中遇到了序列化问题。 我使用cudaEvents作为标记来跟踪内核执行情况。
在我使用多个流的并发内核的实验中，我们观察到在各自的流上使用事件会导致并发内核序列化。
下面的代码演示了这个问题。 我在两个具有以下列出的并发内核执行功能的不同设备上进行了测试：
您可以通过更改USE_EVENTS宏来运行带有和不带有事件的程序，并且您将观察到由于并发执行与串行执行的差异。
#include<cuda.h>
#include<pthread.h>
#include<stdio.h>
#include<stdlib.h>
#include<stdint.h>
#define CUDA_SAFE_CALL( call) do {
cudaError_t err = call;
if( cudaSuccess != err) {
fprintf(stderr, "Cuda error in call at file '%s' in line %i : %s.n",
__FILE__, __LINE__, cudaGetErrorString( err) );
exit(-1);
} } while (0)
// Device code
__global__ void VecAdd(uint64_t len)
{
volatile int a;
for(uint64_t n = 0 ; n < len ; n ++)
a++;
return ;
}
#define USE_EVENTS
int
main(int argc, char *argv[])
{
cudaStream_t stream[2];
for(int i = 0 ; i < 2 ; i++)
CUDA_SAFE_CALL(cudaStreamCreate(&stream[i]));
#ifdef USE_EVENTS
cudaEvent_t e[4];
CUDA_SAFE_CALL(cudaEventCreate(&e[0]));
CUDA_SAFE_CALL(cudaEventCreate(&e[1]));
CUDA_SAFE_CALL(cudaEventRecord(e[0],stream[0]));
#endif
VecAdd<<<1, 32, 0, stream[0]>>>(0xfffffff);
#ifdef USE_EVENTS
CUDA_SAFE_CALL(cudaEventRecord(e[1],stream[0]));
#endif
#ifdef USE_EVENTS
CUDA_SAFE_CALL(cudaEventCreate(&e[2]));
CUDA_SAFE_CALL(cudaEventCreate(&e[3]));
CUDA_SAFE_CALL(cudaEventRecord(e[2],stream[1]));
#endif
VecAdd<<<1, 32, 0, stream[1]>>>(0xfffffff);
#ifdef USE_EVENTS
CUDA_SAFE_CALL(cudaEventRecord(e[3],stream[1]));
#endif
CUDA_SAFE_CALL(cudaDeviceSynchronize());
for(int i = 0 ; i < 2 ; i++)
CUDA_SAFE_CALL(cudaStreamDestroy(stream[i]));
return 0;
}
为什么会出现这种情况以及如何规避这种序列化的任何建议将会很有用。
上面的示例问题按以下顺序工作：
1 event record on stream A
2 launch on stream A
3 event record on Stream A
4 event record on stream B
5 launch on stream B
6 event record on stream B
同一个流上的CUDA操作按问题顺序执行。 不同流中的CUDA操作可以同时运行。
通过编程模型定义，应该有并发性。 但是，在目前的设备上，这项工作是通过一个单独的缓冲区发布给GPU的。 这导致GPU在发布操作3和操作4完成之前等待操作2完成，然后发布5，...如果事件记录被删除，则操作是
1 launch on stream A
2 launch on stream B
操作1和2位于不同的流上，因此GPU有可能同时执行这两个操作。
Parallel Nsight和CUDA命令行分析器（v4.2）可用于定时并发操作。 命令行分析器选项是“conckerneltrace”。 此功能应出现在NVIDIA Visual Profiler的未来版本中。
我从根本上调试了同样的问题。 格雷格的回答非常有帮助，尽管解释似乎并不完整。 真正的问题是当4发布时，op 3正在等待2。 即使4处于不同的流中，如果已经有内核/事件在发布队列中等待，它也不能发布。 这与每个流连续发出多个内核的情况类似。 这可以通过延迟结束事件来解决，如下所示：
由于启动是异步的，因此流结束事件将一直等到该流中先前的内核启动完成，并且所有其他流的内核问题都已启动。 显然，如果在给定硬件上有更多的流可以同时发布，这将导致终端计算机发行太迟。
链接地址: http://www.djcxy.com/p/80101.html