Выполнение графика CUDA из ядра CUDA ⇐ C++

[Anonymous]

Я пытаюсь запустить захваченный график CUDA из обычного ядра CUDA (т. е. запуск графа на стороне устройства).
Из блога NVIDIA, посвященного запуску графа устройства, кажется, что это должно поддерживаться на новых графических процессорах (таких как H100) с последними версиями CUDA.
Чтобы проверить это, я написал минимальный пример, который фиксирует простое добавление ядро в граф CUDA, создает его экземпляр с помощью cudaGraphInstantiateFlagDeviceLaunch, а затем пытается запустить его из другого ядра.
Вот код:

Код: Выделить всё

#include 
#include 
#include 

__global__ void add(int* a, int* b, int* out){
if(threadIdx.x==0 && blockIdx.x==0) *out = *a + *b;
}

__global__ void regularLauncher(cudaGraphExec_t gexec, int* deviceErr){
if(threadIdx.x==0 && blockIdx.x==0){
*deviceErr = 999; // mark kernel ran
cudaError_t e = cudaGraphLaunch(gexec, cudaStreamGraphFireAndForget);
*deviceErr = (e == cudaSuccess) ? 512 /*test value*/ : (int)e;
}
}

int main(){
int device;
cudaGetDevice(&device);
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, device);
printf("Device: %s, Compute Capability: %d.%d\n", prop.name, prop.major, prop.minor);

int *d_a, *d_b, *d_out, *d_err;
cudaMalloc(&d_a, sizeof(int));
cudaMalloc(&d_b, sizeof(int));
cudaMalloc(&d_out, sizeof(int));
cudaMalloc(&d_err, sizeof(int));

int ha=2, hb=3;
cudaMemcpy(d_a, &ha, sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(d_b, &hb, sizeof(int), cudaMemcpyHostToDevice);
cudaMemset(d_out, 0, sizeof(int));
cudaMemset(d_err, 0, sizeof(int));

cudaStream_t s;
cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking);
cudaGraph_t g;
cudaStreamBeginCapture(s, cudaStreamCaptureModeGlobal);
add(d_a, d_b, d_out);
cudaStreamEndCapture(s, &g);

cudaGraphExec_t exec;
cudaGraphInstantiate(&exec, g, cudaGraphInstantiateFlagDeviceLaunch);
cudaGraphUpload(exec, s);
cudaStreamSynchronize(s);

// Verify host launch works
cudaGraphLaunch(exec, s);
cudaStreamSynchronize(s);
printf("Host cudaGraphLaunch error: %s\n", cudaGetErrorString(cudaGetLastError()));

int hostOut = 0;
cudaMemcpy(&hostOut, d_out, sizeof(int), cudaMemcpyDeviceToHost);
printf("Host add kernel result = %d\n", hostOut);

// Reset and test device-side launch
cudaMemset(d_out, 0, sizeof(int));
regularLauncher(exec, d_err);
cudaStreamSynchronize(s);
printf("Kernel launch error: %s\n", cudaGetErrorString(cudaGetLastError()));

int err=0, out=0;
cudaMemcpy(&err, d_err, sizeof(int), cudaMemcpyDeviceToHost);
cudaMemcpy(&out, d_out, sizeof(int), cudaMemcpyDeviceToHost);
printf("device cudaGraphLaunch -> %d\n", err);
printf("add kernel result = %d\n", out);

cudaGraphExecDestroy(exec);
cudaGraphDestroy(g);
cudaStreamDestroy(s);
cudaFree(d_a); cudaFree(d_b); cudaFree(d_out); cudaFree(d_err);
return 0;
}

Команда сборки:

Код: Выделить всё

nvcc -std=c++17 -O3 -rdc=true -arch=sm_90 device_graph_test.cu -o device_graph_test -lcudadevrt

Выход:

Код: Выделить всё

Device: NVIDIA H100 NVL, Compute Capability: 9.0
Host cudaGraphLaunch error: no error
Host add kernel result = 5
Kernel launch error: operation not supported
device cudaGraphLaunch -> 0
add kernel result = 0

Вопрос:

• Почему запуск графа на стороне хоста завершается успешно, но cudaGraphLaunch на стороне устройства завершается с ошибкой, операция не поддерживается?
• Существует ли дополнительное требование или флаг API, необходимый для включения запуска графа устройства на H100 с помощью CUDA 12.9?
• Я что-то упустил в том, как я создаю или запускаю граф изнутри ядра?

Любые разъяснения или рабочий пример запуска CUDA Graph из обычного ядра будут очень признательны.

Подробнее здесь: https://stackoverflow.com/questions/797 ... uda-kernel