cudaErrorIllegalInstruction on recursive function call

Question

I am writing a path tracer for GPU using CUDA 10.2. The entire program ran fine until i added a recursive call to the trace function. nvcc still compiles it, although with the warning: "Severity Code Description Project File Line Suppression State Warning Stack size for entry function '' cannot be statically determined". When the GPU reaches the point it stops and the next time CPU gets an cudaError from an API call it is cuda error 715, which is cudaErrorIllegalInstruction. I tried recreating the issue by writing another recursive kernel/function pair, and the compiler gave the same warning, but it executed expectedly. Unfortunately this means i have to dump my entire function here (if there are any questions to the functions and types used i will happily answer them):

__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;

}

Has anyone else had this issue and in that case, how was it fixed? I could probably redesign to a non-recursive design, although it wouldn't be an optimal solution. I don't even know where to start with debugging this issue, so any ideas are greatly appreciated.

Answer 1

The problem is that CUDA usually selects a fitting max stack size for a kernel call, but it is unable to because nvcc cannot predict the necessary size for a recursive functions.

#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;

}

This program will run, but if 2048 is replaced with 1024, it will output the cudaErrorIllegalInstruction.