From 113c90f42373a8d7bc7d717d3f3316b98e3142f8 Mon Sep 17 00:00:00 2001 From: Eric Niebler Date: Mon, 4 Nov 2024 12:03:33 -0800 Subject: [PATCH] avoid race in `nvexec::launch` test; fixes #1430 --- test/nvexec/launch.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/test/nvexec/launch.cpp b/test/nvexec/launch.cpp index 5a1fd6940..16a497d44 100644 --- a/test/nvexec/launch.cpp +++ b/test/nvexec/launch.cpp @@ -61,14 +61,14 @@ namespace { namespace { | nvexec::launch( // {NUM_BLOCKS, THREAD_BLOCK_SIZE}, // [flags](cudaStream_t stm, int* first, int* last) -> void { + const int32_t idx = blockIdx.x * blockDim.x + threadIdx.x; + const ptrdiff_t size = last - first; + // this should be executing on the GPU - if (nvexec::is_on_gpu()) { + if (idx == 0 && nvexec::is_on_gpu()) { flags.set(0); } - const int32_t idx = blockIdx.x * blockDim.x + threadIdx.x; - const ptrdiff_t size = last - first; - if (idx < size) { first[idx] *= scaling; }