如何让高优先级cudaMemcpyAsync抢占低优先级GPU数据拷贝任务？

核心问题原因 #

你的代码存在两个关键问题导致高优先级拷贝无法抢占：

1. 任务提交顺序不合理：你先一次性将所有10次低优先级大拷贝任务全部提交到低优先级流，之后才提交高优先级任务。此时低优先级的拷贝任务已经填满GPU的任务队列，高优先级任务只能排队等待，无法插入执行。
2. 计时逻辑误导：高优先级流的起始事件start是在所有低优先级任务提交完成后才记录的，这导致你测量的是高优先级任务从提交到完成的时间，而非它与低优先级任务并行/抢占的真实耗时。

实现抢占的关键调整 #

要让高优先级拷贝能够抢占低优先级任务，需要满足两个条件：

• 确保高优先级任务在低优先级任务执行过程中提交，而非等所有低优先级任务都提交完毕。
• 利用CUDA流的优先级调度机制，结合A30的多拷贝DMA引擎实现任务抢占。

修改后的示例代码 #

#include <cuda_runtime.h>
#include <stdint.h>
#include <stdio.h>
#include <unistd.h>

int main()
{
    cudaStream_t high, low;
    int low_priority_range, high_priority_range;
    cudaDeviceGetStreamPriorityRange(&low_priority_range, &high_priority_range);
    printf("low_priority_range: %d, high_priority_range: %d\n",
           low_priority_range,
           high_priority_range);

    // 创建非阻塞的优先级流
    cudaStreamCreateWithPriority(&low, cudaStreamNonBlocking, low_priority_range);
    cudaStreamCreateWithPriority(&high, cudaStreamNonBlocking, high_priority_range);

    cudaEvent_t low_start, low_stop, high_start, high_stop;
    cudaEventCreate(&low_start);
    cudaEventCreate(&low_stop);
    cudaEventCreate(&high_start);
    cudaEventCreate(&high_stop);

    // 分配内存
    void *host1, *host2;
    cudaMallocHost(&host1, 1024LL * 1024 * 1024 * 2);
    cudaMallocHost(&host2, 1024);
    void *device1, *device2;
    cudaMalloc(&device1, 1024LL * 1024 * 1024 * 2);
    cudaMalloc(&device2, 1024);

    // 启动低优先级拷贝，先提交一次大拷贝（避免一次性填满队列）
    cudaEventRecord(low_start, low);
    cudaMemcpyAsync(device1, host1, 1024LL * 1024 * 1024 * 2, cudaMemcpyHostToDevice, low);

    // 主机端短暂等待，让低优先级拷贝开始执行
    usleep(1000);

    // 此时提交高优先级拷贝，低优先级拷贝正在执行中
    cudaEventRecord(high_start, high);
    for (int i = 0; i < 10; i++)
    {
        cudaMemcpyAsync(device2, host2, 1024, cudaMemcpyHostToDevice, high);
    }
    cudaEventRecord(high_stop, high);

    // 补充提交剩余低优先级拷贝，等待全部完成
    cudaMemcpyAsync(device1, host1, 1024LL * 1024 * 1024 * 2, cudaMemcpyHostToDevice, low);
    cudaEventRecord(low_stop, low);

    cudaDeviceSynchronize();

    // 计算耗时
    float high_time, low_time;
    cudaEventElapsedTime(&high_time, high_start, high_stop);
    cudaEventElapsedTime(&low_time, low_start, low_stop);
    printf("high priority stream time: %f ms\n", high_time);
    printf("low priority stream time: %f ms\n", low_time);

    // 释放资源
    cudaEventDestroy(low_start);
    cudaEventDestroy(low_stop);
    cudaEventDestroy(high_start);
    cudaEventDestroy(high_stop);
    cudaFreeHost(host1);
    cudaFreeHost(host2);
    cudaFree(device1);
    cudaFree(device2);
    cudaStreamDestroy(low);
    cudaStreamDestroy(high);

    return 0;
}

额外说明 #

1. A30硬件特性：A30属于Ampere架构，支持多DMA拷贝引擎和流优先级调度，高优先级拷贝任务可占用空闲拷贝引擎，或抢占低优先级任务的时间片。
2. 任务提交策略：避免一次性提交大量低优先级任务，分批提交才能给高优先级任务留出插入执行的空间。
3. 非阻塞流属性：必须使用cudaStreamNonBlocking创建流，确保不同流的任务可以并行调度，不受默认流的同步限制。

内容的提问来源于stack exchange，提问作者GoodGame