cudaErrorIllegalInstruction 递归函数调用

Posted 2023-02-22

【中文标题】cudaErrorIllegalInstruction 递归函数调用

【英文标题】：cudaErrorIllegalInstruction on recursive function call

【发布时间】：2020-03-26 16:16:26

【问题描述】：

我正在使用 CUDA 10.2 为 GPU 编写路径跟踪器。整个程序运行良好，直到我添加了对跟踪函数的递归调用。 nvcc 仍然编译它，虽然有警告：“严重性代码描述项目文件行抑制状态 警告“无法静态确定入口函数的堆栈大小”。当 GPU 到达它停止的点并且下一次 CPU 从 API 调用获取 cudaError 时，它是 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;

有没有其他人遇到过这个问题，在这种情况下，它是如何解决的？我可能会重新设计为非递归设计，尽管它不是最佳解决方案。 我什至不知道从哪里开始调试这个问题，所以非常感谢任何想法。

【问题讨论】：

您确定不只是运行时堆空间不足？

有可能，但是错误不会说这样的话吗？此外，我尝试将bounces + 1 与settings.maxBounces 交换，所以函数只会运行两次，我仍然得到同样的错误。

递归会有限制。大多数递归算法都可以非递归地实现。要继续调试非法指令错误，您可以尝试here 描述的方法。这可能有助于定位错误。但是，如果您的堆栈空间不足（递归的可能性），则可能很难“调试”。您可以尝试增加the stack size

寻求调试帮助的问题应包括minimal reproducible example

感谢您的建议。我没有包含 MRE 的原因是我无法在除此之外的其他情况下重现错误，从而有效地使包含的代码成为我的 MRE。

【参考方案1】：

问题在于 CUDA 通常会为内核调用选择合适的最大堆栈大小，但由于 nvcc 无法预测递归函数所需的大小，因此无法选择。

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

【讨论】：

cudaThreadSetLimit 是 deprecated。如果您关注 the link I gave you 到当前文档，您会发现更新后的 API 调用。