diff --git a/numba_cuda/numba/cuda/cudadrv/driver.py b/numba_cuda/numba/cuda/cudadrv/driver.py index 7246efcf8..7b9245248 100644 --- a/numba_cuda/numba/cuda/cudadrv/driver.py +++ b/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): """ diff --git a/numba_cuda/numba/cuda/tests/cudadrv/test_events.py b/numba_cuda/numba/cuda/tests/cudadrv/test_events.py index 2320acbbf..9fc4ce912 100644 --- a/numba_cuda/numba/cuda/tests/cudadrv/test_events.py +++ b/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()