Обработка ошибок в CUDA: cudaError_t, cudaGetLastError и отладка ядер

Введение

CUDA API возвращает коды ошибок практически из каждого вызова, но в отличие от исключений в C++, эти ошибки не прерывают выполнение программы автоматически — если не проверять возвращаемое значение явно, программа может продолжить работу с неверными данными или незаметно «проглотить» сбой. Особенно коварны ошибки внутри самих ядер: запуск ядра (kernel<<<...>>>) асинхронен и не возвращает код ошибки напрямую. В статье разберём правильные практики обработки ошибок CUDA.

Концепция

Большинство функций CUDA Runtime API возвращают cudaError_t. Для проверки ошибок асинхронного запуска ядра используется cudaGetLastError() (возвращает ошибку запуска, например, неверную конфигурацию сетки) в сочетании с cudaDeviceSynchronize() (дожидается завершения ядра и возвращает ошибку, произошедшую уже во время его выполнения, например, выход за границы памяти). Поскольку ручная проверка каждого вызова многословна, стандартная практика — оборачивать вызовы CUDA API в макрос, который автоматически проверяет код ошибки и завершает программу с понятным сообщением при сбое.

Пример кода

#include <cuda_runtime.h>
#include <iostream>
#include <cstdlib>

// Макрос для проверки ошибок CUDA API
#define CUDA_CHECK(call)                                                       
    do {                                                                        
        cudaError_t error = call;                                               
        if (error != cudaSuccess) {                                             
            std::cerr << "Ошибка CUDA в " << __FILE__ << ":" << __LINE__         
                      << " - " << cudaGetErrorString(error) << std::endl;        
            std::exit(EXIT_FAILURE);                                            
        }                                                                        
    } while (0)

__global__ void faultyKernel(float *data, int n)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    // Намеренная ошибка: отсутствует проверка idx < n
    data[idx] = data[idx] * 2.0f;
}

int main()
{
    const int n = 1000;
    float *deviceData;

    CUDA_CHECK(cudaMalloc(&deviceData, n * sizeof(float)));

    const int threadsPerBlock = 256;
    // Намеренно запускаем больше потоков, чем элементов в массиве
    const int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;

    faultyKernel<<<blocksPerGrid, threadsPerBlock>>>(deviceData, n);

    // Проверка ошибки ЗАПУСКА ядра (неверная конфигурация, ресурсы и т.п.)
    CUDA_CHECK(cudaGetLastError());

    // Дожидаемся завершения ядра и проверяем ошибки ВЫПОЛНЕНИЯ
    // (например, выход за границы памяти внутри faultyKernel)
    CUDA_CHECK(cudaDeviceSynchronize());

    std::cout << "Ядро выполнено успешно" << std::endl;

    CUDA_CHECK(cudaFree(deviceData));
    return 0;
}

Пояснения к коду

Макрос CUDA_CHECK оборачивает любой вызов CUDA API, сравнивает результат с cudaSuccess и при ошибке выводит читаемое сообщение через cudaGetErrorString() с указанием файла и строки, после чего завершает программу — этот паттерн стандартен почти во всех серьёзных CUDA-проектах и сильно упрощает диагностику по сравнению с молчаливым игнорированием ошибок.

Ядро faultyKernel намеренно не проверяет границы массива — поскольку число запущенных потоков (blocksPerGrid * threadsPerBlock = 1024) больше числа элементов (n = 1000), потоки с индексом от 1000 до 1023 обратятся за пределы выделенного буфера. Этот баг не проявится на этапе cudaGetLastError() (конфигурация запуска валидна), но будет обнаружен cudaDeviceSynchronize(), который вернёт ошибку вроде an illegal memory access was encountered.

Подводные камни

• Ошибки выполнения ядра обнаруживаются только при следующей синхронизирующей операции, которая может находиться далеко от места реального возникновения ошибки в коде — это усложняет локализацию проблемы без использования специализированных инструментов отладки (cuda-memcheck, Nsight Compute, или их современная замена compute-sanitizer).
• «Залипшая» ошибка влияет на все последующие вызовы. После возникновения ошибки в контексте CUDA все последующие вызовы API будут возвращать ту же или связанные ошибки, пока контекст не будет сброшен (обычно перезапуском процесса) — это значит, что важно проверять ошибки сразу после операции, а не «накопленно» в конце программы.
• cudaGetLastError() сбрасывает сохранённый код ошибки после вызова. Если в коде есть несколько мест проверки ошибок подряд без промежуточных операций, повторный вызов cudaGetLastError() вернёт cudaSuccess, даже если фактическая ошибка была обнаружена и обработана первым вызовом — важно не путать это с «ошибка исчезла сама по себе».
• Отсутствие compute-sanitizer/cuda-memcheck в процессе разработки. Эти инструменты обнаруживают выходы за границы памяти, гонки данных между потоками и неинициализированную память — без них многие ошибки проявляются только спорадически на специфичных конфигурациях GPU и входных данных, что делает их крайне сложными для воспроизведения и исправления постфактум.