diff --git a/projects/rocshmem/src/memory/notifier.hpp b/projects/rocshmem/src/memory/notifier.hpp index e6364a26cf..b2211f0cc5 100644 --- a/projects/rocshmem/src/memory/notifier.hpp +++ b/projects/rocshmem/src/memory/notifier.hpp @@ -32,11 +32,11 @@ namespace rocshmem { template class Notifier { public: - __device__ uint64_t read() { + __device__ uint64_t load() { return detail::atomic::load(&value_, orders); } - __device__ void write(uint64_t val) { + __device__ void store(uint64_t val) { detail::atomic::store(&value_, val, orders); } diff --git a/projects/rocshmem/src/memory/slab_heap.cpp b/projects/rocshmem/src/memory/slab_heap.cpp index 89067b78da..80b8a8cef9 100644 --- a/projects/rocshmem/src/memory/slab_heap.cpp +++ b/projects/rocshmem/src/memory/slab_heap.cpp @@ -75,8 +75,8 @@ __device__ void SlabHeap::malloc(void** ptr, size_t size) { * Notify other threads in block about the allocation result. */ auto notifier{notifier_.get()}; - notifier->write(ptr_deref_u64); - uint64_t notification_u64{notifier->read()}; + notifier->store(ptr_deref_u64); + uint64_t notification_u64{notifier->load()}; notifier->done(); /* diff --git a/projects/rocshmem/tests/unit_tests/notifier_gtest.hpp b/projects/rocshmem/tests/unit_tests/notifier_gtest.hpp index 4a72212f5d..b914ca4f22 100644 --- a/projects/rocshmem/tests/unit_tests/notifier_gtest.hpp +++ b/projects/rocshmem/tests/unit_tests/notifier_gtest.hpp @@ -53,8 +53,8 @@ __global__ void all_threads_once(uint8_t* raw_memory, Notifier * notifier) { - notifier->write(NOTIFIER_OFFSET); - uint64_t offset_u64 {notifier->read()}; + notifier->store(NOTIFIER_OFFSET); + uint64_t offset_u64 {notifier->load()}; notifier->done(); uint64_t raw_memory_u64 {reinterpret_cast(raw_memory)};