cudaGraph：多线程流捕获仅在 cuda-memcheck 中运行时才会导致错误答案

【问题标题】：cudaGraph: Multi-threaded stream capturing causes errors only when run in cuda-memcheckcudaGraph：多线程流捕获仅在 cuda-memcheck 中运行时才会导致错误
【发布时间】：2020-04-25 09:16:11
【问题描述】：

我有一个程序，其中多个主机线程尝试捕获 cuda 图并执行它。它会产生正确的结果，但不能使用 cuda-memcheck 运行。

使用 cuda-memcheck 运行时，出现以下错误。

程序命中 cudaErrorStreamCaptureInvalidated（错误 901），原因是在对 cudaLaunchKernel 的 CUDA API 调用中“由于先前的错误导致操作失败”。

当只使用一个主机线程时，cuda-memcheck 不会显示错误。

这是可以使用 nvcc 10.2 编译的示例代码：nvcc -arch=sm_61 -O3 main.cu -o main

#include <iostream>
#include <memory>
#include <algorithm>
#include <cassert>
#include <vector>
#include <thread>
#include <iterator>


#ifndef CUERR

    #define CUERR {                                                            \
        cudaError_t err;                                                       \
        if ((err = cudaGetLastError()) != cudaSuccess) {                       \
            std::cout << "CUDA error: " << cudaGetErrorString(err) << " : "    \
                      << __FILE__ << ", line " << __LINE__ << std::endl;       \
            exit(1);                                                           \
        }                                                                      \
    }

#endif


__global__
void kernel(int id, int num){
    printf("kernel %d, id %d\n", num, id);
}

struct Data{
    bool isValidGraph = false;
    int id = 0;
    int deviceId = 0;
    cudaGraphExec_t execGraph = nullptr;
    cudaStream_t stream = nullptr;
};

void buildGraphViaCapture(Data& data){
    cudaSetDevice(data.deviceId); CUERR;

    if(!data.isValidGraph){
        std::cerr << "rebuild graph\n";

        if(data.execGraph != nullptr){
            cudaGraphExecDestroy(data.execGraph); CUERR;
        }

        assert(data.stream != cudaStreamLegacy);

        cudaStreamCaptureStatus captureStatus;
        cudaStreamIsCapturing(data.stream, &captureStatus); CUERR;

        assert(captureStatus == cudaStreamCaptureStatusNone);

        cudaStreamBeginCapture(data.stream, cudaStreamCaptureModeRelaxed); CUERR;

        for(int i = 0; i < 64; i++){
            kernel<<<1,1,0,data.stream>>>(data.id, i);
        }

        cudaGraph_t graph;
        cudaStreamEndCapture(data.stream, &graph); CUERR;

        cudaGraphExec_t execGraph;
        cudaGraphNode_t errorNode;
        auto logBuffer = std::make_unique<char[]>(1025);
        std::fill_n(logBuffer.get(), 1025, 0);
        cudaError_t status = cudaGraphInstantiate(&execGraph, graph, &errorNode, logBuffer.get(), 1025);
        if(status != cudaSuccess){
            if(logBuffer[1024] != '\0'){
                std::cerr << "cudaGraphInstantiate: truncated error message: ";
                std::copy_n(logBuffer.get(), 1025, std::ostream_iterator<char>(std::cerr, ""));
                std::cerr << "\n";
            }else{
                std::cerr << "cudaGraphInstantiate: error message: ";
                std::cerr << logBuffer.get();
                std::cerr << "\n";
            }
            CUERR;
        }            

        cudaGraphDestroy(graph); CUERR;

        data.execGraph = execGraph;

        data.isValidGraph = true;
    }
}

void execute(Data& data){
    buildGraphViaCapture(data);

    assert(data.isValidGraph);

    cudaGraphLaunch(data.execGraph, data.stream); CUERR;
}


void initData(Data& data, int id, int deviceId){
    data.id = id;
    data.deviceId = deviceId;
    cudaStreamCreate(&data.stream); CUERR;
}

void destroyData(Data& data){
    if(data.execGraph != nullptr){
        cudaGraphExecDestroy(data.execGraph); CUERR;
    }
    cudaStreamDestroy(data.stream); CUERR; 
}

int main(){

    std::vector<int> deviceIds{0};

    std::vector<std::thread> threads;

    for(int deviceId : deviceIds){
        for(int k = 0; k < 4; k++){
            threads.emplace_back([&,deviceId](){

                std::vector<Data> vec(3);

                initData(vec[0], deviceId * 10 + 4*k + 0, deviceId);
                initData(vec[1], deviceId * 10 + 4*k + 1, deviceId);

                int cur = 0;

                for(int iter = 0; iter < 10; iter++){
                    cudaStreamSynchronize(vec[cur].stream); CUERR;
                    execute(vec[cur]); CUERR;
                    cur = 1 - cur;
                }

                cudaStreamSynchronize(vec[0].stream); CUERR;
                cudaStreamSynchronize(vec[1].stream); CUERR;

                destroyData(vec[0]);
                destroyData(vec[1]);

            });
        }
    }

    for(auto& t : threads){
        t.join();
    }



    cudaDeviceReset();
    return 0;
}

为什么只有在使用多个线程时才会出现错误，为什么捕获无效？

【问题讨论】：

标签： c++ cuda

【解决方案1】：

Cuda 图不是线程安全的。如果您阅读文档，它会说：

图形对象（cudaGraph_t、CUgraph）在内部不是同步的，并且不能从多个线程同时访问。访问同一图形对象的 API 调用必须在外部进行序列化。

您需要在关键部分访问图形对象。

【讨论】：

我认为这不是图形对象的竞争条件。每个图以及每个流只能由一个线程访问。指南中引用的部分不包括这种情况。
你说得对，我之前没有彻底检查过你的代码。我用 cuda-memcheck 和不用 cuda-memcheck 运行了你的代码，我无法重现任何错误。

【解决方案2】：

我们也遇到了这个问题 - 即使我们正在处理不同的 CUDA 图形对象，我们仍然会遇到错误。我们的（ugly）解决方案是将cudaStreamBeginCapture 和cudaStreamEndCapture 包装在带有静态互斥锁的RAII 结构中。

它现在解决了这个问题，但我将在 CUDA 开发者论坛上进一步询问。

【讨论】：