Exclude execution fence from memory fence

janciesko · May 6, 2024 · eb5acd5 · eb5acd5
1 parent 645b2f9
commit eb5acd5
Show file tree

Hide file tree

Showing 3 changed files with 3 additions and 12 deletions.
diff --git a/src/impl/nvshmemspace/Kokkos_NVSHMEMSpace.cpp b/src/impl/nvshmemspace/Kokkos_NVSHMEMSpace.cpp
@@ -130,10 +130,7 @@ void NVSHMEMSpace::impl_deallocate(
   }
 }
 
-void NVSHMEMSpace::fence() {
-  Kokkos::fence();
-  nvshmem_barrier_all();
-}
+void NVSHMEMSpace::fence() { nvshmem_barrier_all(); }
 
 KOKKOS_FUNCTION
 int get_num_pes() { return nvshmem_n_pes(); }

diff --git a/src/impl/rocshmemspace/Kokkos_ROCSHMEMSpace.cpp b/src/impl/rocshmemspace/Kokkos_ROCSHMEMSpace.cpp
@@ -128,10 +128,7 @@ void ROCSHMEMSpace::impl_deallocate(
   }
 }
 
-void ROCSHMEMSpace::fence() {
-  Kokkos::fence();
-  roc_shmem_barrier_all();
-}
+void ROCSHMEMSpace::fence() { roc_shmem_barrier_all(); }
 
 KOKKOS_FUNCTION
 size_t get_num_pes() { return roc_shmem_n_pes(); }

diff --git a/src/impl/shmemspace/Kokkos_SHMEMSpace.cpp b/src/impl/shmemspace/Kokkos_SHMEMSpace.cpp
@@ -141,10 +141,7 @@ void SHMEMSpace::impl_deallocate(
   }
 }
 
-void SHMEMSpace::fence() {
-  Kokkos::fence();
-  shmem_barrier_all();
-}
+void SHMEMSpace::fence() { shmem_barrier_all(); }
 
 size_t get_num_pes() { return shmem_n_pes(); }
 size_t get_my_pe() { return shmem_my_pe(); }