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

Posted 2023-02-16

【中文标题】cudaGraph：多线程流捕获仅在 cuda-memcheck 中运行时才会导致错误

【英文标题】：cudaGraph: Multi-threaded stream capturing causes errors only when run in 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> deviceIds0;

    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;

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

【参考方案1】：

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

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

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

【讨论】：

我认为这不是图形对象的竞争条件。每个图以及每个流只能由一个线程访问。指南中引用的部分不包括这种情况。

你说得对，我之前没有彻底检查过你的代码。我用 cuda-memcheck 和不用 cuda-memcheck 运行了你的代码，我无法重现任何错误。

【参考方案2】：

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

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