cudaGraph: захват многопоточных потоков вызывает ошибки только при запуске в cuda-memcheck

Question

У меня есть программа, в которой несколько потоков хоста пытаются захватить граф cuda и выполнить его. Он выдает правильные результаты, но не может быть запущен с помощью cuda-memcheck.

При запуске с помощью cuda-memcheck появляется следующая ошибка.

Хит программы cudaErrorStreamCaptureInvalidated (ошибка 901) из-за «операции не удалось из-за предыдущей ошибки во время захвата» при вызове API CUDA к cudaLaunchKernel.

Когда используется только один хост-поток, cuda-memcheck не показывает ошибку.

Здесь это пример кода, который можно скомпилировать с помощью nv cc 10.2: nv cc -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;
}

Почему ошибка появляется только при использовании нескольких потоков и почему именно захват признан недействительным?

Answer 1

Графики Cuda не являются потокобезопасными. Если вы читаете документацию, она говорит, что:

Графические объекты (cudaGraph_t, CUgraph) не синхронизированы внутри и не должны быть доступны одновременно из нескольких потоков. Вызовы API, обращающиеся к одному и тому же графическому объекту, должны быть сериализованы извне.

Вам необходим доступ к графическому объекту в критической секции.