Add warp sync for Cuda parallel reduce to avoid race condition

[tcclevenger]

Potential race condition discovered by compute-sanitizer --tool=racecheck on a simple reduce using Cuda:

Real result;
Kokkos::parallel_reduce("simple_reduce", N, KOKKOS_LAMBDA(int, Real&) {
  // Do nothing;
}, result);

Basically, the value of shared[i] below __syncwarp() was being updated on a single thread within a warp inKokkos_Cuda_ReduceScan.hpp:349, but no synchronization is guarenteed before reading shared[i] on (potentially) different threads.

Closes #6217.

[dalg24]

Please comment about the performance implications of that fence

[dalg24]

Typo. Add the missing white space if you somehow make any change or retrigger CI. Otherwise not worth fixing.

[masterleinad]

I would at least mention that final is also only run by a single thread and that the result is not necessarily available to the whole warp.

[crtrott]

Please comment about the performance implications of that fence

We can certainly check, but that thing semantically needs to be there - even though I doubt anyone can trigger a bug right now for it missing.

[tcclevenger]

@dalg24 I'm not seeing any significant performance impact. Launching parallel reduce kernels and the relative difference in timings is under 1% compared to develop.

[dalg24]

Is that the only place it is needed? Thre are many instances of

      if (CudaTraits::WarpSize < word_count.value) {
        __syncthreads();
      }

[tcclevenger]

@dalg24 This is the only place I've had compute-sanitizer racecheck complain. I could do a more detailed investigation in to other places.

[tcclevenger]

@crtrott The cherry pick PR into 4.2.01 (#6631) got merged, but this PR into develop did not. This does include an extra commit (changing the comment wording).

[dalg24]

The cherry-pick should not have been merged before that one. That was a mistake.

[dalg24]

Specifically thinking

kokkos/core/src/Cuda/Kokkos_Cuda_Parallel_Team.hpp

Lines 765 to 767 in dcf93fc

	if (CudaTraits::WarpSize < word_count.value) {
	__syncthreads();
	}

and

kokkos/core/src/Cuda/Kokkos_Cuda_Parallel_MDRange.hpp

Lines 309 to 311 in dcf93fc

	if (CudaTraits::WarpSize < word_count.value) {
	__syncthreads();
	}

[crtrott]

Yes those two places also need it.