fix: extract handle value in Event.query() call to match other driver calls

numba_cuda/numba/cuda/cudadrv/driver.py

@@ -2028,14 +2028,14 @@ def query(self):
         otherwise, returns False.
         """
         try:
-            driver.cuEventQuery(self.handle)
+            handle = self.handle.value
+            driver.cuEventQuery(handle)
+            return True
         except CudaAPIError as e:
             if e.code == binding.CUresult.CUDA_ERROR_NOT_READY:
                 return False
             else:
                 raise
-        else:
-            return True
 
     def record(self, stream=0):
         """

numba_cuda/numba/cuda/tests/cudadrv/test_events.py

@@ -2,7 +2,7 @@
 # SPDX-License-Identifier: BSD-2-Clause
 
 import numpy as np
-from numba import cuda
+from numba import cuda, int32
 from numba.cuda.testing import unittest, CUDATestCase
 from cuda.core import Device
 from numba.cuda.testing import skip_on_cudasim
@@ -48,6 +48,53 @@ def event_elapsed_inner(self, stream):
         # Exercise the code path
         evtstart.elapsed_time(evtend)
 
+    @skip_on_cudasim("Testing cuda.core events requires driver")
+    def test_event_query(self):
+        stream = cuda.stream()
+        evt = cuda.event()
+
+        # Mapped arrays: host-side edits visible to device and vice versa.
+        started = cuda.mapped_array(1, dtype=np.int32)
+        release = cuda.mapped_array(1, dtype=np.int32)
+
+        @cuda.jit
+        def gated_kernel(started_flag, release_flag):
+            # Signal that kernel has started
+            started_flag[0] = 1
+            # Spin until host releases us
+            while release_flag[0] == 0:
+                cuda.nanosleep(int32(1_000))
+
+        # Compile first
+        started[0] = 0
+        release[0] = 1  # Don't block during warmup
+        gated_kernel[1, 1, stream](started, release)
+        stream.synchronize()
+
+        # Reset for actual test
+        started[0] = 0
+        release[0] = 0
+
+        # Launch - kernel will spin until we release it
+        gated_kernel[1, 1, stream](started, release)
+        evt.record(stream)
+
+        # Wait until kernel confirms it's running
+        while started[0] == 0:
+            pass
+
+        # Kernel is running until we release it - if query returns True, fail.
+        immediate_query = evt.query()
+        assert immediate_query is False, "Query returned True prematurely"
+
+        # Release the kernel and synchronize
+        release[0] = 1
+        evt.synchronize()
+
+        # If query returns False after synchronize, fail.
+        synced_query = evt.query()
+        assert synced_query is True, "Query returned False after sync"
+
 
 if __name__ == "__main__":
     unittest.main()