cudaErrorIllegalInstruction при рекурсивном вызове функции

Question

Я пишу трассировщик пути для GPU, используя CUDA 10.2. Вся программа работала нормально, пока я не добавил рекурсивный вызов функции трассировки. nv cc все еще компилирует его, хотя и с предупреждением: «Код серьезности Описание Предупреждение о состоянии подавления строки файла проекта Размер стека для функции входа« не может быть определен статически ». Когда графический процессор достигает точки, он останавливается, и в следующий раз, когда центральный процессор получает cudaError от вызова API, это ошибка cuda 715, которая является cudaErrorIllegalInstruction. Я попытался воссоздать проблему, написав другую рекурсивную пару ядро ​​/ функция, и компилятор выдал то же самое предупреждение, но оно выполнялось ожидаемым образом. К сожалению, это означает, что я должен сбросить всю свою функцию здесь (если есть какие-либо вопросы к используемым функциям и типам, я с радостью на них отвечу):

__device__ Vec3 trace(
    const Settings& settings,
    const Ray& r,
    const Shape* shapes,
    const size_t nshapes,
    uint8_t bounces,
    curandState& randState) {

    if (bounces >= settings.maxBounces) {
        return Vec3(0.0f);
    }

    const Shape* shape = nullptr;
    float t = inf;
    bool flipNormal;

    float dist;

    for (size_t i = 0; i < nshapes; i++) {
        if (shapes[i].intersect(r, dist, flipNormal) && dist < t) {
            shape = shapes + i;
            t = dist;
        }
    }

    if (shape == nullptr) 
        return settings.background;

    const Vec3 hitPos = r.ori + t * r.dir;
    const Vec3 normal = flipNormal ? -shape->normal(hitPos) : shape->normal(hitPos);

    const Vec3 hemiDir = cosineSample(normal, randState);

    const Vec3 traceCol = trace(
        settings,
        Ray(hitPos + normal * settings.bias, hemiDir),
        shapes,
        nshapes,
        bounces + 1,
        randState
        );

    return shape->surface.emittance + shape->surface.color * traceCol;

}

У кого-нибудь еще была эта проблема, и в этом случае, как это было исправлено? Я мог бы, вероятно, изменить дизайн на нерекурсивный дизайн, хотя это не было бы оптимальным решением. Я даже не знаю, с чего начать с отладки этой проблемы, поэтому любые идеи очень ценятся.

Answer 1

Проблема в том, что CUDA обычно выбирает подходящий максимальный размер стека для вызова ядра, но не может этого сделать, потому что nv cc не может предсказать необходимый размер для рекурсивных функций.

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <iostream>

#include <stdint.h>

__device__ int recurse(uint64_t n, uint64_t max) {
    if (n < max)
        return recurse(n + 1, max);
    else 
        return n;
}

__global__ void start(uint64_t max) {
    uint32_t idx = threadIdx.x + (blockIdx.x * blockDim.x);

    if(idx == 256 * 256 - 1)
        printf("%i: %i\n", idx, recurse(0, max));

    return;
}

int main() {

    cudaError_t status;

    status = cudaSetDevice(0);
    if (status != cudaSuccess) {
        std::cerr << "failed: " << cudaGetErrorString(status) << std::endl;
        return status;
    }

    cudaThreadSetLimit(cudaLimitStackSize, 2048);

    start<<<256, 256>>>(126);

    status = cudaDeviceSynchronize();
    if (status != cudaSuccess) {
        std::cerr << "failed: " << cudaGetErrorString(status) << std::endl;
        return status;
    }

    return 0;

}

Это Программа запустится, но если 2048 заменить на 1024, она выведет инструкцию cudaErrorIllegalInstruction.