CUDA has lots of documentation and guides all over the place, but one I haven't been able to find has been any form of instruction in how to diagnose kernels that compile but get cryptic, vague error messages such as 'unspecified launch failure' beyond the normal "Do these block/grid structures make sense?" etc. Can I intercept the cubin file somehow and do some static analysis on the memory structures etc? Forgive my noobness but I can't find any definitive, idiots guide, anywhere. Have a good weekend everyone. What I'm looking for How to separate out the cubin intermediate file What to do with it afterwards to work out what's going on, specifically register and memory configuration to see if my code is violating any hardware requirements, or if I'm just missing an off-by-one error somewhere. For anyone coming across this later (I seem to have a habit of creating SO questions that keep showing up in my own queries months later...) CUDA-Memcheck gives much more interesting responses that the 'check error' handlers. eg ========= Error: process didn't terminate successfully ========= Invalid __global__ write of size 4 ========= at 0x00000040 in decomp ========= by thread (1,0,0) in block (0,0,0) ========= Address 0x00101024 is out of bounds ========= ========= ERROR SUMMARY: 1 error I don't even have to explain that error message...

Reputation: 7926

Diagnosing CUDA Kernel problems

CUDA has lots of documentation and guides all over the place, but one I haven't been able to find has been any form of instruction in how to diagnose kernels that compile but get cryptic, vague error messages such as 'unspecified launch failure' beyond the normal "Do these block/grid structures make sense?" etc.

Can I intercept the cubin file somehow and do some static analysis on the memory structures etc? Forgive my noobness but I can't find any definitive, idiots guide, anywhere.

Have a good weekend everyone.

What I'm looking for

How to separate out the cubin intermediate file
What to do with it afterwards to work out what's going on, specifically register and memory configuration to see if my code is violating any hardware requirements, or if I'm just missing an off-by-one error somewhere.

For anyone coming across this later (I seem to have a habit of creating SO questions that keep showing up in my own queries months later...) CUDA-Memcheck gives much more interesting responses that the 'check error' handlers. eg

========= Error: process didn't terminate successfully
========= Invalid __global__ write of size 4
=========     at 0x00000040 in decomp
=========     by thread (1,0,0) in block (0,0,0)
=========     Address 0x00101024 is out of bounds
=========
========= ERROR SUMMARY: 1 error

I don't even have to explain that error message...

Upvotes: 3

Answers (3)

talonmies

Reputation: 72372

In CUDA, "unspecified launch failure" is the equivalent of a segfault.

Recent toolkit versions ship with a utility called cuda-memcheck. It performs valgrind like analysis of memory transactions inside an executing kernel, and will report buffer overflows or any illegal pointer usage in a kernel. You can use that as a launching point for further analysis. If you are using a Fermi card, there is also in-kernel printf support, it isn't hard to generate your own assert function to test and report for error conditions inside a kernel.

CUDA also ships with a source level debugger, but you need a dedicated GPU to use it. If you are on linux and only have a single GPU, quit out of X11 and run it from a console TTY.

Upvotes: 3

Adam S.

Reputation: 1271

Are you using cudaGetLastError()? That could help give more information if it's not already used to give 'unspecified launch failure'.

Upvotes: 2

Ade Miller

Reputation: 13753

If you set the Keep Preprocessed Files flag --keep this will leave the CUBIN files and a host of others lying around for you to take a look at. But I'm not sure this will help that much.

Upvotes: 2

Diagnosing CUDA Kernel problems

Answers (3)

Related Questions