CUDA 并发内核在使用事件时序列化

Question

我在需要并发执行的 CUDA 内核中遇到了序列化问题。我使用 cudaEvents 作为标记来跟踪内核执行。

在我对具有多个流的并发内核的实验中，我们观察到在它们各自的流上使用事件会导致并发内核被序列化。

下面的代码演示了这个问题。我在两个不同的设备上对此进行了测试，这些设备具有下面列出的并发内核执行功能:

Tesla C2070，驱动程序版本 4.10，运行时版本 4.10，CUDA 功能 2.0

Tesla M2090，驱动程序版本 4.10，运行时版本 4.10，CUDA 功能 2.0

您可以通过更改 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;

}

关于为什么会发生这种情况以及如何规避这种序列化的任何建议都将很有用。

Answer 1

上面的示例问题按以下顺序工作:

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 之前等待操作 2 完成，并在发出操作 5 之前等待操作 4 完成，...如果事件记录被删除，则操作是

1 launch on stream A
2 launch on stream B

操作 1 和 2 位于不同的流上，因此 GPU 可以同时执行这两个操作。

Parallel Nsight 和 CUDA 命令行分析器 (v4.2) 可用于计时并发操作。命令行分析器选项是“conckerneltrace”。此功能应出现在 NVIDIA Visual Profiler 的 future 版本中。