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

numba_cuda/numba/cuda/cudadrv/driver.py

@@ -2076,14 +2076,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 == enums.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 numba.cuda._compat import Device
 from numba.cuda.testing import skip_on_cudasim
@@ -48,6 +48,41 @@ def event_elapsed_inner(self, stream):
         # Exercise the code path
         evtstart.elapsed_time(evtend)
 
+    def test_event_query(self):
+        from time import perf_counter
+
+        @cuda.jit
+        def spin(ms):
+            # Sleep for ms
+            for i in range(ms):
+                cuda.nanosleep(int32(1_000_000))  # 1 ms
+
+        stream = cuda.stream()
+        evt = cuda.event()
+
+        # Run once to compile
+        spin[1, 1, stream](1)
+
+        t0 = perf_counter()
+        spin_ms = 250
+        spin[1, 1, stream](250)
+        evt.record(stream)
+
+        # Query immediately.
+        event_time = perf_counter() - t0
+        while not evt.query():
+            event_time = perf_counter() - t0
+
+        # Synchronize and capture stream-finished time.
+        evt.synchronize()
+        sync_time = perf_counter() - t0
+
+        # If this assertion fails, it was nanosleep inaccuracy that caused it
+        assert sync_time * 1000 > spin_ms * 0.9
+
+        # If this assertion fails, the event query returned early
+        assert event_time * 1000 > spin_ms * 0.9
+
 
 if __name__ == "__main__":
     unittest.main()