4

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

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

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

  1. Tesla C2070，驱动程序版本 4.10，运行时版本 4.10，CUDA 功能 2.0
  2. 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;

}

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

4

2 回答 2

3

上述示例问题按以下顺序工作：

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 的未来版本中。

于 2012-05-09T03:19:37.160 回答
1

我正在调试基本相同的问题。格雷格的回答非常有帮助，尽管解释似乎并不完整。真正的问题是，当 4 发出时，操作 3 正在等待 2。即使 4 在不同的流中，如果已经有内核/事件在问题队列中等待，它不能被发出。这类似于每个流连续发布多个内核的情况。这可以通过延迟流结束事件来解决，如下所示：

  1. 流 A 上的事件记录（启动计时器）
  2. 在流 A 上启动
  3. 流 B 上的事件记录（启动计时器）
  4. 在流 B 上启动
  5. 流 A 上的事件记录（结束计时器）
  6. 流 B 上的事件记录（结束计时器）

由于启动是异步的，因此流结束事件将等到该流中的先前内核启动完成，并且所有其他流的内核问题都已启动。显然，如果给定硬件上的流数多于可以同时发出的流，这将导致结束计时器发出得太晚。

于 2012-06-14T13:20:54.263 回答