[BUG] Raising an exception or asserting inside a kernel is silently ignored. #388
Description
@cuda.jit
def throwing_kernel():
assert False, "TODO: You need to implement this kernel!"
raise NotImplementedError("TODO: You need to implement this kernel!")
throwing_kernel[1024, 256]()
cp.cuda.stream.get_current_stream().synchronize()
Today, this compiles and runs with no error.
It would be better to either:
- Trap on the GPU when an exception is raised (this is the optimal solution).
- Fail to compile this code.
This is slightly annoying for the CUDA Python DLI and other Numba CUDA teaching material. We want to be able to create notebooks that have placeholders that students need to fill in. It needs to be very obvious where they should fill in code, e.g. the cell should fail if they haven't written what they need to. Raising the exception in host code would be annoying - they're not necessarily writing the entire kernel for the exercise, just part of it, and I don't want to put the exception raising someplace else in the cell that they have to track down. So I need some way to trap within a Numba CUDA kernel.