fix: empty array type mismatch between host and device (#612)

Fixes #483. 

As mentioned
[here](numba/numba#9654 (comment)),
the bug was due to:
- empty arrays incorrectly identified as broadcast arrays
- incorrect contiguity flags for empty arrays

When running this:
```
from numba import cuda, typeof
import numpy as np

h_values = np.random.randint(low=0, high=2, size=(10, 0))
d_values = cuda.to_device(h_values)

print(typeof(h_values))
print(typeof(d_values))

assert typeof(h_values) == typeof(d_values)
```

Below is the output:
```
array(int64, 2d, C)
array(int64, 2d, C)
```

Author: kaeun97

numba_cuda/numba/cuda/cudadrv/devicearray.py

@@ -178,7 +178,8 @@ def _numba_type_(self):
         # of which will be 0, will not match those hardcoded in for 'C' or 'F'
         # layouts.
 
-        broadcast = 0 in self.strides
+        broadcast = 0 in self.strides and (self.size != 0)
+
         if self.flags["C_CONTIGUOUS"] and not broadcast:
             layout = "C"
         elif self.flags["F_CONTIGUOUS"] and not broadcast:

numba_cuda/numba/cuda/cudadrv/dummyarray.py

@@ -279,6 +279,10 @@ def _compute_layout(self):
         if not self.dims:
             return {"C_CONTIGUOUS": True, "F_CONTIGUOUS": True}
 
+        # All 0-size arrays are considered contiguous, even if they are multidimensional
+        if self.size == 0:
+            return {"C_CONTIGUOUS": True, "F_CONTIGUOUS": True}
+
         # If this is a broadcast array then it is not contiguous
         if any([dim.stride == 0 for dim in self.dims]):
             return {"C_CONTIGUOUS": False, "F_CONTIGUOUS": False}

numba_cuda/numba/cuda/tests/nocuda/test_dummyarray.py

@@ -387,5 +387,66 @@ def test_for_loop(self):
             x = val  # noqa: F841
 
 
+@skip_on_cudasim("Tests internals of the CUDA driver device array")
+class TestEmptyArrays(unittest.TestCase):
+    def test_empty_array_flags(self):
+        test_shapes = [
+            (0,),
+            (10, 0),
+            (0, 10),
+            (0, 0),
+            (5, 0, 3),
+            (0, 5, 3),
+            (5, 3, 0),
+            (0, 0, 0),
+        ]
+        for shape in test_shapes:
+            with self.subTest(shape=shape):
+                nparr = np.empty(shape)
+                arr = Array.from_desc(
+                    0, nparr.shape, nparr.strides, nparr.dtype.itemsize
+                )
+                # Empty arrays should be both C and F contiguous
+                self.assertEqual(
+                    arr.flags["C_CONTIGUOUS"],
+                    nparr.flags["C_CONTIGUOUS"],
+                    f"C_CONTIGUOUS mismatch for shape {shape}",
+                )
+                self.assertEqual(
+                    arr.flags["F_CONTIGUOUS"],
+                    nparr.flags["F_CONTIGUOUS"],
+                    f"F_CONTIGUOUS mismatch for shape {shape}",
+                )
+                self.assertTrue(arr.flags["C_CONTIGUOUS"])
+                self.assertTrue(arr.flags["F_CONTIGUOUS"])
+
+
+@skip_on_cudasim("Tests CUDA device array type inference")
+class TestEmptyArrayTypeInference(unittest.TestCase):
+    def test_empty_array_typeof(self):
+        from numba import cuda, typeof
+
+        test_cases = [
+            ((0,), np.int64),
+            ((10, 0), np.int64),
+            ((0, 10), np.int64),
+            ((0, 0), np.float32),
+            ((5, 0, 3), np.float32),
+            ((0, 5, 3), np.int32),
+            ((5, 3, 0), np.float64),
+        ]
+
+        for shape, dtype in test_cases:
+            with self.subTest(shape=shape, dtype=dtype):
+                h_values = np.empty(shape, dtype=dtype)
+                d_values = cuda.to_device(h_values)
+                self.assertEqual(
+                    typeof(h_values),
+                    typeof(d_values),
+                    f"Type mismatch for shape {shape}, dtype {dtype}: "
+                    f"host={typeof(h_values)}, device={typeof(d_values)}",
+                )
+
+
 if __name__ == "__main__":
     unittest.main()