-1

So, I'm trying to port this raytracer to CUDA, and I'm running into some trouble when trying to adapt the code to test for collision with a list of hittable objects.

More specifically, I have the following code to test for collisions:

__device__ bool hittable_list_hit(const ray& r, double t_min, double t_max, hit_record& rec, uint8_t* data, int& offset) {
    hit_record temp_rec = { point3(0,0,0), point3(0,0,0), 0.0, false };
    bool hit_anything = false;
    auto closest_so_far = t_max;

    //printf("%f", temp_rec.t); <-- UNCOMMENTING THIS LINE MAKES CUDA CRASH

    uint8_t size = data[offset++];

    for (uint8_t i = 0; i < size; i++) {
        if (top_level_hit(r, t_min, closest_so_far, temp_rec, data, offset)) {
            hit_anything = true;
            //closest_so_far = temp_rec.t;
            //rec = temp_rec;
        }
    }

    return hit_anything;
}

If I keep the highlighted line commented, the output of the program is a nice sky background (as in the raytracing tutorial). But if I uncomment it, the output is a black screen, and I suspect it might be a CUDA crash. It seems like reading this variable makes CUDA crash for some reason.

I'm relatively new to CUDA, so I have no clue why this might be happening. Does anybody have any idea on why reading a normal variable can cause this?

If it helps, this function is 4 levels deep after the kernel launch. Perhaps this nesting can cause problems?

Yxuer
  • 44
  • 7
  • 1
    https://developer.nvidia.com/blog/accelerated-ray-tracing-cuda/ – Robert Crovella Jul 19 '23 at 19:36
  • Oh, I didn't know somebody had made that already. Thanks a lot! However, I would like to give it a try at doing this project myself, so my original question still stands – Yxuer Jul 19 '23 at 20:19
  • 2
    compile the code with `-lineinfo` and use compute-sanitizer to find the line that causes the crash. – Abator Abetor Jul 20 '23 at 04:40

1 Answers1

-1

Alright, I managed to figure out what was happening! Thanks to Abator Abetor's answer, I took a look at what compute-sanitizer outputted, and I was getting a "Too Many Resources Requested for Launch" error.

Also thanks to this Stack Overflow question, I saw that a potential solution was reducing the number of threads per block, and that worked for me! I also changed all doubles in my code to floats, just to be sure.

In short: whenever you are faced with a weird CUDA crash, compute-sanitizer is your friend :-)

Yxuer
  • 44
  • 7
  • 1
    You don't even need `compute-sanitizer` for this kind of issue, just use proper CUDA runtime error checking: [What is the canonical way to check for errors using the CUDA runtime API?](https://stackoverflow.com/questions/14038589/what-is-the-canonical-way-to-check-for-errors-using-the-cuda-runtime-api) – paleonix Jul 20 '23 at 19:39
  • Oh, I tried that one too, and it gave me the same "too many resources" error, but there was not even a black screen, so I thought this way of checking was causing some other problem with the code. Thanks for the tip! – Yxuer Jul 20 '23 at 22:01