Merge pull request #1431 from NVIDIA/fix-nvexec-launch-test

avoid race in `nvexec::launch` test; fixes #1430
NVIDIA · Nov 4, 2024 · c1508b1 · c1508b1
2 parents 35e8941 + 113c90f
commit c1508b1
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;
                    }