[BUG] Integer overflow in device array when numpy version > 2.0 for T4 GPU:s

[lcniel]

Describe the bug
It seems like the MemoryPointer object relies on the auto-casting behaviour of old numpy versions to upgrade int32 to int64 on demand, which results in semi-deterministic errors with large GPU:s on newer numpy versions, at least in some cases. I can reproduce the error on T4 GPU:s with driver version 535.161.07, which is what Gitlab runners use.

numba-cuda/numba_cuda/numba/cuda/cudadrv/driver.py

Line 1839 in 1dd1e7a

def view(self, start, stop=None):

Here with numpy==2.2.6 and numba_cuda==0.21.3.
E.g. pytest traceback:

self = <numba.cuda.cudadrv.driver.AutoFreePointer object at 0x7f4590b57f70>
start = np.int32(0), stop = np.int32(1152)
    def view(self, start, stop=None):
        if stop is None:
            size = self.size - start
        else:
            size = stop - start
    
        # Handle NULL/empty memory buffer
        if not self.device_pointer_value:
            if size != 0:
                raise RuntimeError("non-empty slice into empty slice")
            view = self  # new view is just a reference to self
        # Handle normal case
        else:
>           base = self.device_pointer_value + start
E           OverflowError: Python integer 139939146685440 out of bounds for int32
local_installation_linux/numba_cuda/numba/cuda/cudadrv/driver.py:1852: OverflowError

Steps/Code to reproduce bug
It happens in my Gitlab CI when something much like the following is done:

import numba.cuda as cuda
import numpy as np
arr = np.arange(50 * 50 * 100, dtype=np.float32).reshape(50, 50, 100)
arr2 = cuda.to_device(arr)

The problem goes away when I downgrade numpy below 2.0. I can't reproduce it on e.g. an L40S with 580.95.05 drivers.

Expected behavior
There should be no reliance on auto-casting of typed integers in the code. Wrapping start in int64 would probably fix the issue?

Environment details (please complete the following information):

• Environment location: Gitlab enterprise GPU runner (VM with T4 GPU)
• Method of numba-cuda install: pip

Additional context
It's a bit hard to deterministically reproduce since I don't know where device_pointer_value comes from, but it should be easy enough to make sure this doesn't happen at all.

Here's an example of a recent CI failure of mine:

https://gitlab.com/liebi-group/software/mumott/-/jobs/12288859623

[gmarkall]

Thanks for the report. I'll look into this and see if I can work out a reproducer and potential fix.

[gmarkall]

The main thing confounding me at the moment is working out how start and stop got to be np.int32. I'm observing them as Python ints in my experiments so far.

[lcniel]

The main thing confounding me at the moment is working out how start and stop got to be np.int32. I'm observing them as Python ints in my experiments so far.

Hmm, for sure there is a check in dummyarray.py:

    def from_desc(cls, offset, shape, strides, itemsize):
        dims = []
        for ashape, astride in zip(shape, strides):
            if not isinstance(ashape, (int, np.integer)):
                raise TypeError("all elements of shape must be ints")
            dim = Dim(
                offset, offset + ashape * astride, ashape, astride, single=False
            )
            dims.append(dim)
            offset = 0  # offset only applies to first dimension
        return cls(dims, itemsize)

isinstance(shape, (int, np.integer)) will at least permit np.int32 but I don't know why it expects that possibility. Everywhere I can see it uses some .shape or .stride attribute of an array which I'm pretty sure should be an int object.

I also double checked if my own code was doing anything funky with those arrays, but no, they're going to be numpy arrays typically created with np.zeros and then filled with values or otherwise initialized. No room for anything strange to happen, really.

[gmarkall]

Also if I change that check to only permit int and not np.integer, running the mumott tests locally still seems to succeed for me.

[gmarkall]

Whoops, I was completely wrong in my previous comment, I get

FAILED tests/gpu_tests/test_async_pipelines.py::test_motr - TypeError: all elements of shape must be ints

[gmarkall]

I wonder if this is a regression introduced by #542.

[gmarkall]

For reference, I'm ran mumott by:

git clone https://gitlab.com/liebi-group/software/mumott.git
cd mumott
pip install -e .
pytest -k test_motr

after applying

diff --git a/numba_cuda/numba/cuda/cudadrv/dummyarray.py b/numba_cuda/numba/cuda/cudadrv/dummyarray.py
index 16545954..5746586f 100644
--- a/numba_cuda/numba/cuda/cudadrv/dummyarray.py
+++ b/numba_cuda/numba/cuda/cudadrv/dummyarray.py
@@ -250,7 +250,7 @@ class Array(object):
     def from_desc(cls, offset, shape, strides, itemsize):
         dims = []
         for ashape, astride in zip(shape, strides):
-            if not isinstance(ashape, (int, np.integer)):
+            if not isinstance(ashape, int):
                 raise TypeError("all elements of shape must be ints")
             dim = Dim(
                 offset, offset + ashape * astride, ashape, astride, single=False

[lcniel]

For reference, I'm ran mumott by:

git clone https://gitlab.com/liebi-group/software/mumott.git
cd mumott
pip install -e .
pytest -k test_motr

after applying

[code]

I tested this on 0.20.1 now (just adding the isinstance(ashape, int) condition rather than replacing an existing one) and it's reproduced prior to #542.

cls = <class 'numba.cuda.cudadrv.dummyarray.Array'>, offset = 0, shape = (1, np.int32(4), np.int32(4), 18)
strides = (np.int32(1152), np.int32(288), 72, 4), itemsize = 4

    @classmethod
    def from_desc(cls, offset, shape, strides, itemsize):
        dims = []
        for ashape, astride in zip(shape, strides):
            if not isinstance(ashape, int):
>               raise TypeError("all elements of shape must be ints")
E               TypeError: all elements of shape must be ints

../../py313/lib64/python3.13/site-packages/numba_cuda/numba/cuda/cudadrv/dummyarray.py:252: TypeError

In this case I do actually set the shape in part using an array which has dtype np.int32 (which I do internally for consistency in state representation when loading and saving data). It seems like numba.cuda.device_array does not care to convert np.int32 to int (unlike numpy arrays). I could patch this particular case in my code by just exposing that as int, but it seems to pop up in some versions and some cases where I use cuda.to_device as well... probably device array should just always use int or int64 as an integer representation internally?

I can't recreate the error that occurs in my CI when cuda.to_device is used with a numpy.ndarray runs, just the one that happens when numba.devicearray is used. But it seems like the device array gets int-type shape parameters but the underlying dummyarray gets int32 - maybe that is due to the aforementioned caching.

[gmarkall]

@kaeun97 This is more of a debugging / investigation-type issue. Would you be interested in digging into this and seeing if you can work out the root cause for how we're running into this error? (no worries if not - just let me know 🙂 )

[gmarkall]

@kaeun97 A quick reminder of this issue - do let me know if it's not interesting to you 🙂

[kaeun97]

Apologies - this must've slipped through.. Would love to investigate further. Let me jump into it. Thanks for the reminder! Appreciate it.