CUDA内存传输与核函数调用并非总是重叠的原因咨询

阿华AIGC实验室

2026-4-16

嘿，我来帮你拆解这个问题～你用--default-stream per-thread编译代码，想用OpenMP开两个线程分别跑异步内存拷贝和核函数，结果发现它们并不是每次都能并发执行，对吧？先把你没写完的代码补全（核函数调用部分），方便咱们分析：

__global__ void zero(int *Y)
{
    for (int i = 0; i < 100; i++) {
        Y[(blockIdx.x + i * gridDim.x) * blockDim.x + threadIdx.x] = 0;
    }
}

int main(int argc, char *argv[])
{
    int n = 32 * 100000;
    int *h_X;
    cudaHostAlloc((void **)&h_X, sizeof(int) * n, cudaHostAllocWriteCombined);
    for (int i = 0; i < n; i++) {
        h_X[i] = 0;
    }

    int *X;
    int *Y;
    cudaMalloc((void **)&X, sizeof(int) * n);
    cudaMalloc((void **)&Y, sizeof(int) * n);

#pragma omp parallel sections num_threads(2) firstprivate(n, h_X, X, Y)
    {
#pragma omp section
        {
            for (int i = 0; i < 20; i++) {
                cudaMemcpyAsync(X, h_X, sizeof(int) * n, cudaMemcpyHostToDevice);
            }
        }
#pragma omp section
        {
            for (int i = 0; i < 20; i++) {
                zero<<<n/256, 256>>>(Y); // 补全核函数调用逻辑
            }
        }
    }

    // 省略后续资源清理代码
    return 0;
}

下面咱们聊聊为啥会出现“有时候重叠、有时候不重叠”的情况：

主机线程的启动延迟：OpenMP的两个section线程并不是绝对同时启动的，比如负责内存拷贝的线程可能先跑了1、2次拷贝任务，核函数线程才开始调度。这时候前几次拷贝自然没法和核函数重叠，后面的任务才会进入并发状态，看起来就像“部分时候才并发”。
任务粒度不匹配：
你每次内存拷贝的大小是128MB左右（321000004字节），这个量级的拷贝在GPU上耗时不算太长；而你的zero核函数，每个线程要执行100次赋值，总计算量大概是3200万次操作，在现代GPU上可能几毫秒就跑完了。如果拷贝和核函数的执行时间都很短，很容易出现“一个任务已经跑完，另一个才刚启动”的情况，自然看不到重叠。
CUDA流的串行特性：
虽然你用了per-thread默认流，每个OMP线程有独立的默认流，但同一个流里的任务是串行执行的。也就是拷贝线程的20次cudaMemcpyAsync是按顺序排队执行的，核函数线程的20次核调用也是串行的。如果其中一个流的任务队列被快速填满，或者GPU的拷贝/计算引擎暂时被占满，后续任务就会排队，导致重叠中断。
GPU硬件资源限制：
如果你的GPU只有一个拷贝引擎（比如一些入门级NVIDIA GPU），当有连续的拷贝任务时，新的拷贝会排队等待；同理，如果计算引擎被核任务占满，新的核函数也会延迟启动。这种硬件资源的瓶颈会导致任务无法持续重叠。