avoid race in nvexec::launch test; fixes #1430

NVIDIA · Nov 4, 2024 · 113c90f · 113c90f
1 parent 35e8941
commit 113c90f
Showing 1 changed file with 4 additions and 4 deletions.
diff --git 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;
                    }