Repository: numba/numba
Pull request: CUDA: Fix #7806, Division by zero stops the kernel
The expected behaviour is that execution proceeds after division by zero, as the CUDA target doesn't normally have a way to signal the exception. When debug is enabled, it can signal the exception so an exception should be raised and execution should stop. This is implemented by setting CUDA kernels to use the NumPy error model in general, and only switching back to the Python error model when debug is on.

This PR also adds support for raising on divide by zero when debug is on in the simulator - this helps align the behaviour of the simulator with that of the real device. The error behaviour in NumPy is configured on a per-thread basis, so it's essential to use `np.seterr()` when the target thread starts running. Since threads are ephemeral for the life of the kernel launch, no attempt is made to save and restore the initial state of the error behaviour.

Finally, there were some illegal uses of debug with optimization in the `test_exception.py` file I was working on, so I fixed that while I was working in this area.

Fixes #7806.

Problem statement:
CUDA: Division by zero stops the kernel
The following:

```python
from numba import cuda
import numpy as np

r = np.zeros(1)
x = np.zeros(1)
y = np.ones(1)


@cuda.jit
def f(r, x, y):
    r[0] = y[0] / x[0]
    print("Here")


f[1, 1](r, x, y)
cuda.synchronize()

print(r)
```

prints:

```
[0.]
```

instead of the expected:

```
Here
[inf]
```

Hints from issue discussion:
It seems like this has existed for quite some time - even with Numba 0.48.0 the same result is observed. The PTX contains:

```assembly
       ; %rd5 points to the value of x[0]

	ld.global.f64 	%fd1, [%rd5];
	setp.eq.f64	%p1, %fd1, 0d0000000000000000;
	@%p1 bra 	BB0_2;

       ; Do division, print("Here")

BB0_2:
	ret;

```
Think the issue is in code gen. The above PTX probably comes from LLVM IR:
```llvm
  %".87" = load double, double* %".86"
  %".88" = fcmp oeq double %".87",              0x0
  br i1 %".88", label %"B0.if", label %"B0.endif", !prof !2
B0.if:
  ret i32 1
B0.endif:
  %".91" = fdiv double %".58", %".87"
```
which is from:
https://github.com/numba/numba/blob/62df0f056c70981c47c32faa32e3ad4651618129/numba/cpython/numbers.py#L599-L600
If those lines are commented, then the MWR behaves as expected.

Commit message: Fix flake8 errors

Patch to review:
diff --git a/numba/cuda/compiler.py b/numba/cuda/compiler.py
index 19b0c7e1c90..29dfcc76d40 100644
--- a/numba/cuda/compiler.py
+++ b/numba/cuda/compiler.py
@@ -196,6 +196,9 @@ def compile_cuda(pyfunc, return_type, args, debug=False, lineinfo=False,
         # later want to overload some other behavior on the debug flag.
         # In particular, -opt=3 is not supported with -g.
         flags.debuginfo = True
+        flags.error_model = 'python'
+    else:
+        flags.error_model = 'numpy'
     if inline:
         flags.forceinline = True
     if fastmath:
diff --git a/numba/cuda/simulator/api.py b/numba/cuda/simulator/api.py
index 15e233547cc..974cd83800c 100644
--- a/numba/cuda/simulator/api.py
+++ b/numba/cuda/simulator/api.py
@@ -77,7 +77,7 @@ def elapsed_time(self, event):
 
 def jit(func_or_sig=None, device=False, debug=False, argtypes=None,
         inline=False, restype=None, fastmath=False, link=None,
-        boundscheck=None,
+        boundscheck=None, opt=True
         ):
     # Here for API compatibility
     if boundscheck:
@@ -91,9 +91,10 @@ def jit(func_or_sig=None, device=False, debug=False, argtypes=None,
         def jitwrapper(fn):
             return FakeCUDAKernel(fn,
                                   device=device,
-                                  fastmath=fastmath)
+                                  fastmath=fastmath,
+                                  debug=debug)
         return jitwrapper
-    return FakeCUDAKernel(func_or_sig, device=device)
+    return FakeCUDAKernel(func_or_sig, device=device, debug=debug)
 
 
 @contextmanager
diff --git a/numba/cuda/simulator/kernel.py b/numba/cuda/simulator/kernel.py
index 2499d513910..8297680de2b 100644
--- a/numba/cuda/simulator/kernel.py
+++ b/numba/cuda/simulator/kernel.py
@@ -62,10 +62,11 @@ class FakeCUDAKernel(object):
     Wraps a @cuda.jit-ed function.
     '''
 
-    def __init__(self, fn, device, fastmath=False, extensions=[]):
+    def __init__(self, fn, device, fastmath=False, extensions=[], debug=False):
         self.fn = fn
         self._device = device
         self._fastmath = fastmath
+        self._debug = debug
         self.extensions = list(extensions) # defensive copy
         # Initial configuration: grid unconfigured, stream 0, no dynamic shared
         # memory.
@@ -117,7 +118,7 @@ def fake_arg(arg):
             with swapped_cuda_module(self.fn, fake_cuda_module):
                 # Execute one block at a time
                 for grid_point in np.ndindex(*grid_dim):
-                    bm = BlockManager(self.fn, grid_dim, block_dim)
+                    bm = BlockManager(self.fn, grid_dim, block_dim, self._debug)
                     bm.run(grid_point, *fake_args)
 
             for wb in retr:
@@ -159,8 +160,16 @@ class BlockThread(threading.Thread):
     '''
     Manages the execution of a function for a single CUDA thread.
     '''
-    def __init__(self, f, manager, blockIdx, threadIdx):
-        super(BlockThread, self).__init__(target=f)
+    def __init__(self, f, manager, blockIdx, threadIdx, debug):
+        if debug:
+            def debug_wrapper(*args, **kwargs):
+                np.seterr(divide='raise')
+                f(*args, **kwargs)
+            target = debug_wrapper
+        else:
+            target = f
+
+        super(BlockThread, self).__init__(target=target)
         self.syncthreads_event = threading.Event()
         self.syncthreads_blocked = False
         self._manager = manager
@@ -169,6 +178,7 @@ def __init__(self, f, manager, blockIdx, threadIdx):
         self.exception = None
         self.daemon = True
         self.abort = False
+        self.debug = debug
         blockDim = Dim3(*self._manager._block_dim)
         self.thread_id = self.threadIdx.x + (blockDim.x * (self.threadIdx.y +
                                                            blockDim.y *
@@ -247,10 +257,11 @@ class BlockManager(object):
     The polling continues until no threads are alive, when execution is
     complete.
     '''
-    def __init__(self, f, grid_dim, block_dim):
+    def __init__(self, f, grid_dim, block_dim, debug):
         self._grid_dim = grid_dim
         self._block_dim = block_dim
         self._f = f
+        self._debug = debug
         self.block_state = np.zeros(block_dim, dtype=np.bool_)
 
     def run(self, grid_point, *args):
@@ -261,7 +272,7 @@ def run(self, grid_point, *args):
         for block_point in np.ndindex(*self._block_dim):
             def target():
                 self._f(*args)
-            t = BlockThread(target, self, grid_point, block_point)
+            t = BlockThread(target, self, grid_point, block_point, self._debug)
             t.start()
             threads.add(t)
             livethreads.add(t)
diff --git a/numba/cuda/tests/cudapy/test_exception.py b/numba/cuda/tests/cudapy/test_exception.py
index cd5e0308470..db248ca3aec 100644
--- a/numba/cuda/tests/cudapy/test_exception.py
+++ b/numba/cuda/tests/cudapy/test_exception.py
@@ -16,7 +16,7 @@ def foo(ary):
                 ary.shape[-x]
 
         unsafe_foo = cuda.jit(foo)
-        safe_foo = cuda.jit(debug=True)(foo)
+        safe_foo = cuda.jit(debug=True, opt=False)(foo)
 
         if not config.ENABLE_CUDASIM:
             # Simulator throws exceptions regardless of debug
@@ -28,7 +28,7 @@ def foo(ary):
         self.assertIn("tuple index out of range", str(cm.exception))
 
     def test_user_raise(self):
-        @cuda.jit(debug=True)
+        @cuda.jit(debug=True, opt=False)
         def foo(do_raise):
             if do_raise:
                 raise ValueError
@@ -44,7 +44,9 @@ def case_raise_causing_warp_diverge(self, with_debug_mode):
         of unifying branch target and resulting in unexpected warp
         divergence.
         """
-        @cuda.jit(debug=with_debug_mode)
+        with_opt_mode = not with_debug_mode
+
+        @cuda.jit(debug=with_debug_mode, opt=with_opt_mode)
         def problematic(x, y):
             tid = cuda.threadIdx.x
             ntid = cuda.blockDim.x
@@ -91,6 +93,56 @@ def test_raise_causing_warp_diverge(self):
         """
         self.case_raise_causing_warp_diverge(with_debug_mode=False)
 
+    # The following two cases relate to Issue #7806: Division by zero stops the
+    # kernel. https://github.com/numba/numba/issues/7806.
+
+    def test_no_zero_division_error(self):
+        # When debug is False:
+        # - Zero by division raises no exception
+        # - Execution proceeds after a divide by zero
+        @cuda.jit
+        def f(r, x, y):
+            r[0] = y[0] / x[0]
+            r[1] = y[0]
+
+        r = np.zeros(2)
+        x = np.zeros(1)
+        y = np.ones(1)
+
+        f[1, 1](r, x, y)
+
+        self.assertTrue(np.isinf(r[0]), 'Expected inf from div by zero')
+        self.assertEqual(r[1], y[0], 'Expected execution to continue')
+
+    def test_zero_division_error_in_debug(self):
+        # When debug is True:
+        # - Zero by division raises an exception
+        # - Execution halts at the point of division by zero
+        @cuda.jit(debug=True, opt=False)
+        def f(r, x, y):
+            r[0] = y[0] / x[0]
+            r[1] = y[0]
+
+        r = np.zeros(2)
+        x = np.zeros(1)
+        y = np.ones(1)
+
+        # Simulator and device behaviour differs slightly in the exception
+        # raised - in debug mode, the CUDA target uses the Python error model,
+        # which gives a ZeroDivision error. The simulator uses NumPy with the
+        # error mode for division by zero set to raise, which results in a
+        # FloatingPointError instead.
+        if config.ENABLE_CUDASIM:
+            exc = FloatingPointError
+        else:
+            exc = ZeroDivisionError
+
+        with self.assertRaises(exc):
+            f[1, 1](r, x, y)
+
+        self.assertEqual(r[0], 0, 'Expected result to be left unset')
+        self.assertEqual(r[1], 0, 'Expected execution to stop')
+
 
 if __name__ == '__main__':
     unittest.main()


Code review comments:
diff_hunk: @@ -91,6 +93,56 @@ def test_raise_causing_warp_diverge(self):
         """
         self.case_raise_causing_warp_diverge(with_debug_mode=False)
 
+    # The following two cases relate to Issue #7806: Division by zero stops the
+    # kernel. https://github.com/numba/numba/issues/7806.
+
+    def test_no_zero_division_error(self):
+        # When debug is False:
+        # - Zero by division raises no exception
path: numba/cuda/tests/cudapy/test_exception.py
text: @user1:
```suggestion
        # - Division by zero raises no exception
```


@author:
Thanks, not quite sure how I managed to write that!