CUDA多流构建流水线失败原因：跨设备系统差异分析

我尝试使用CUDA多流机制构建流水线，编写代码后通过Nsight Systems分析发现初始代码呈串行执行；但将Device-to-Host（D2H）复制操作单独提取到独立循环中后，流水线可正常工作。

后续在4种不同设备环境下开展测试，结果如下：

• 3060 GPU（Windows系统）：无法构建流水线
• 2080TI GPU（Linux系统）：流水线构建成功
• 3090 GPU（Linux系统）：流水线构建成功
• A100 GPU（Linux系统）：流水线构建成功

恳请分析仅3060-Windows环境无法构建流水线的原因。

初始代码 #

using namespace std;

__global__ void vecAdd(float *c, const float *a, const float *b);

void initBuffer(float *data, int size);

int main() {
    int size = 1 << 22;
    int bufsize = size * sizeof(float);
    int nStream = 4;

    float* ha[4];
    float* hb[4];
    float* hc[4];
    float* da[4];
    float* db[4];
    float* dc[4];

    srand(2019);
    for (int i = 0; i < 4; ++i) {
        cudaMallocHost((void **) &ha[i], bufsize);
        cudaMallocHost((void **) &hb[i], bufsize);
        cudaMallocHost((void **) &hc[i], bufsize);
        initBuffer(ha[i], size);
        initBuffer(hb[i], size);

        cudaMalloc((void **) &da[i], bufsize);
        cudaMalloc((void **) &db[i], bufsize);
        cudaMalloc((void **) &dc[i], bufsize);
    }

    auto *streams = new cudaStream_t[nStream];
    for (int i = 0; i < nStream; i++) {
        cudaStreamCreate(&streams[i]);
    }

    for (int i = 0; i < nStream; i++) {
        cudaMemcpyAsync(da[i], ha[i], bufsize, cudaMemcpyHostToDevice, streams[i]);
        cudaMemcpyAsync(db[i], hb[i], bufsize, cudaMemcpyHostToDevice, streams[i]);

        dim3 dimBlock(256);
        dim3 dimGrid(size / dimBlock.x);
        vecAdd<<< dimGrid, dimBlock, 0, streams[i] >>>(dc[i], da[i], db[i]);
        cudaMemcpyAsync(hc[i], dc[i], bufsize, cudaMemcpyDeviceToHost, streams[i]);
    }

    cudaDeviceSynchronize();

    // terminate operators
    delete[] streams;

    for (int i = 0; i < 4; ++i) {

        // terminate device memories
        cudaFree(da[i]);
        cudaFree(db[i]);
        cudaFree(dc[i]);

        // terminate host memories
        cudaFreeHost(ha[i]);
        cudaFreeHost(hb[i]);
        cudaFreeHost(hc[i]);
    }

    return 0;
}

void initBuffer(float *data, const int size) {
    for (int i = 0; i < size; i++)
        data[i] = rand() / (float) RAND_MAX;
}

__global__ void vecAdd(float *c, const float *a, const float *b) {
    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
    for (int i = 0; i < 200; i++)
        c[idx] = a[idx] + b[idx];
}

修改后代码片段 #

auto *streams = new cudaStream_t[nStream];
    for (int i = 0; i < nStream; i++) {
        cudaStreamCreate(&streams[i]);
    }

    for (int i = 0; i < nStream; i++) {
        cudaMemcpyAsync(da[i], ha[i], bufsize, cudaMemcpyHostToDevice, streams[i]);
        cudaMemcpyAsync(db[i], hb[i], bufsize, cudaMemcpyHostToDevice, streams[i]);

        dim3 dimBlock(256);
        dim3 dimGrid(size / dimBlock.x);
        vecAdd<<< dimGrid, dimBlock, 0, streams[i] >>>(dc[i], da[i], db[i]);
//        cudaMemcpyAsync(hc[i], dc[i], bufsize, cudaMemcpyDeviceToHost, streams[i]);
    }

    for (int i = 0; i < nStream; ++i)
        cudaMemcpyAsync(hc[i], dc[i], bufsize, cudaMemcpyDeviceToHost, streams[i]);

    cudaDeviceSynchronize();

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