Use Kokkos atomic_inc in place of atomic_increment

[CKegel]

Approximately a month ago, the development branch of kokkos deprecated calls to atomic_increment in favor of atomic_inc (see kokkos/kokkos@e61165e). atomic_inc is present in the master and release branches, so this branch replaces calls to Kokkos:atomic_increment with Kokkos:atomic_inc.

[cwsmith]

Thank you. I think we should merge this, but we'll want to add a version check in cmakelists.txt to be sure the kokkos being used against supports atomic_[inc|dec]. Do we know what version of Kokkos added support for it?

[CKegel]

I should be able to find it rather quickly.

[CKegel]

It looks like the atomic_[inc/dec] api was added 2 years ago (kokkos/kokkos#4690) in the same PR as atomic_increment/decrement. However I can not find it in any version branch earlier than Kokkos 4.5 which does not seem correct.

[cwsmith]

Running git log -S ' atomic_inc(' --source --all [1] indicates that it first showed up in its current form in 3.5.00:

(ins)cws@turtle: /home/cwsmith/develop/kokkos (develop)$ git grep ' atomic_inc('
core/src/Kokkos_Atomics_Desul_Wrapper.hpp:template<class T> KOKKOS_FUNCTION Impl::enable_if_atomic_t<T, void> atomic_inc(T* ptr) { desul::atomic_inc(const_cast<std::remove_volatile_t<T>*>(ptr), desul::MemoryOrderRelaxed(), KOKKOS_DESUL_MEM_SCOPE); }
core/src/Kokkos_Atomics_Desul_Wrapper.hpp:template<class T> KOKKOS_DEPRECATED_WITH_COMMENT("Use atomic_inc() instead!") KOKKOS_FUNCTION Impl::enable_if_atomic_t<T, void> atomic_increment(T* ptr) { atomic_inc(ptr); }
core/src/Kokkos_Crs.hpp:  void operator()(index_type i) const { atomic_inc(&out[in.entries(i)]); }
core/src/Kokkos_WorkGraphPolicy.hpp:        atomic_inc(begin_hint);
core/src/Kokkos_WorkGraphPolicy.hpp:    atomic_inc(count_queue + m_graph.entries[i]);
tpls/desul/include/desul/atomics/Generic.hpp:DESUL_INLINE_FUNCTION void atomic_inc(T* const dest,


(ins)cws@turtle: /home/cwsmith/develop/kokkos (develop)$ git checkout 3.5.00
(cmd)cws@turtle: /home/cwsmith/develop/kokkos ((3.5.00))$ git grep ' atomic_inc('
core/src/Kokkos_Atomics_Desul_Volatile_Wrapper.hpp:void atomic_inc(volatile T* const dest) { return desul::atomic_inc(const_cast<T*>(dest),desul::MemoryOrderRelaxed(), desul::MemoryScopeDevice()); }
core/src/Kokkos_Atomics_Desul_Wrapper.hpp:void atomic_inc(T* const dest) { return desul::atomic_inc(dest, desul::MemoryOrderRelaxed(), desul::MemoryScopeDevice()); }
core/src/desul/atomics/CUDA.hpp:    inline void atomic_inc(TYPE* const dest, ORDER order, SCOPE scope) { \
core/src/desul/atomics/Generic.hpp:DESUL_INLINE_FUNCTION void atomic_inc(T* const dest,
core/src/desul/atomics/cuda/cuda_cc7_asm_atomic_op.inc_forceglobal:inline __device__ void atomic_inc(type* dest, __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER, __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE) { \
core/src/desul/atomics/cuda/cuda_cc7_asm_atomic_op.inc_generic:inline __device__ void atomic_inc(type* dest, __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER, __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE) { \
core/src/desul/atomics/cuda/cuda_cc7_asm_atomic_op.inc_isglobal:inline __device__ void atomic_inc(type* dest, __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER, __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE) { \
core/src/desul/atomics/cuda/cuda_cc7_asm_atomic_op.inc_predicate:inline __device__ void atomic_inc(type* dest, __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER, __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE) { \

[cwsmith]

Since we already require Kokkos 3.7 we are good to go.

omega_h/CMakeLists.txt

Line 60 in e64bba9

set(Kokkos_REQUIRED_VERSION 3.7)