diff --git a/CMakeLists.txt b/CMakeLists.txt index ee7f42fb12..8603737308 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -60,7 +60,6 @@ option(USE_IPC "Enable IPC support (using HIP)" OFF) option(USE_THREADS "Enable workgroup threads to share network queues" OFF) option(USE_WF_COAL "Enable wavefront message coalescing" OFF) option(USE_COHERENT_HEAP "Enable support for coherent systems" OFF) -option(USE_CACHED_HEAP "Enable support for cached systems" OFF) option(USE_MANAGED_HEAP "Enable managed memory" OFF) option(USE_HOST_HEAP "Enable host memory using malloc/free" OFF) option(USE_HIP_HOST_HEAP "Enable host memory using hip api" OFF) diff --git a/README.md b/README.md index c974fd7052..ba671ff51a 100644 --- a/README.md +++ b/README.md @@ -45,7 +45,7 @@ ROC_SHMEM base requirements: * May work with other versions, but not tested * AMD GFX9 GPUs (e.g.: MI25, Vega 56, Vega 64, MI50, MI60, MI100, Radeon VII) * AMD MI200 GPUs: To enable the support on MI200, please configure the library - with USE_CACHED_HEAP + with USE_COHERENT_HEAP * ROCm-aware MPI as described in [Building the Dependencies](#building-the-dependencies) * InfiniBand adaptor compatable with ROCm RDMA technology diff --git a/cmake/config.h.in b/cmake/config.h.in index 6ec07c2853..f409c10ba3 100644 --- a/cmake/config.h.in +++ b/cmake/config.h.in @@ -7,10 +7,9 @@ #cmakedefine USE_SHARED_CTX #cmakedefine USE_WF_COAL #cmakedefine USE_COHERENT_HEAP -#cmakedefine USE_CACHED_HEAP #cmakedefine USE_MANAGED_HEAP #cmakedefine USE_HOST_HEAP #cmakedefine USE_HIP_HOST_HEAP #cmakedefine USE_FUNC_CALL #cmakedefine USE_SINGLE_NODE -#cmakedefine USE_HOST_SIDE_HDP_FLUSH \ No newline at end of file +#cmakedefine USE_HOST_SIDE_HDP_FLUSH diff --git a/scripts/build_configs/rc_single b/scripts/build_configs/rc_single index 31eb7283c6..aceca7dbed 100755 --- a/scripts/build_configs/rc_single +++ b/scripts/build_configs/rc_single @@ -20,7 +20,6 @@ cmake \ -DUSE_DC=OFF \ -DUSE_IPC=OFF \ -DUSE_COHERENT_HEAP=OFF \ - -DUSE_CACHED_HEAP=OFF \ -DUSE_THREADS=OFF \ -DUSE_WF_COAL=OFF \ $src_path diff --git a/scripts/build_configs/rc_single_single_node b/scripts/build_configs/rc_single_single_node index 8bcf9fcb0d..a5223c379f 100755 --- a/scripts/build_configs/rc_single_single_node +++ b/scripts/build_configs/rc_single_single_node @@ -20,12 +20,10 @@ cmake \ -DUSE_DC=OFF \ -DUSE_IPC=ON \ -DUSE_COHERENT_HEAP=OFF \ - -DUSE_CACHED_HEAP=OFF \ -DUSE_THREADS=OFF \ -DUSE_WF_COAL=OFF \ -DUSE_SINGLE_NODE=ON \ -DUSE_HOST_SIDE_HDP_FLUSH=ON\ - -DROCM_PATH="/opt/rocm-5.4.2/"\ $src_path cmake --build . --parallel 8 cmake --install . diff --git a/scripts/build_configs/rc_single_single_node_debug b/scripts/build_configs/rc_single_single_node_debug index b535f1b853..c4e2477da9 100755 --- a/scripts/build_configs/rc_single_single_node_debug +++ b/scripts/build_configs/rc_single_single_node_debug @@ -20,7 +20,6 @@ cmake \ -DUSE_DC=OFF \ -DUSE_IPC=ON \ -DUSE_COHERENT_HEAP=OFF \ - -DUSE_CACHED_HEAP=OFF \ -DUSE_THREADS=OFF \ -DUSE_WF_COAL=OFF \ -DUSE_SINGLE_NODE=ON \ diff --git a/scripts/build_configs/ro_ipc b/scripts/build_configs/ro_ipc new file mode 100755 index 0000000000..d57404b88c --- /dev/null +++ b/scripts/build_configs/ro_ipc @@ -0,0 +1,29 @@ +#!/bin/bash +# Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. + +if [ -z $1 ] +then + install_path=~/rocshmem +else + install_path=$1 +fi + +src_path=$(dirname "$(realpath $0)")/../../ + +cmake \ + -DCMAKE_BUILD_TYPE=Release \ + -DCMAKE_INSTALL_PREFIX=$install_path \ + -DCMAKE_VERBOSE_MAKEFILE=OFF \ + -DDEBUG=OFF \ + -DPROFILE=OFF \ + -DUSE_GPU_IB=OFF \ + -DUSE_DC=OFF \ + -DUSE_IPC=ON \ + -DUSE_COHERENT_HEAP=ON \ + -DUSE_THREADS=OFF \ + -DUSE_WF_COAL=OFF \ + -DUSE_SINGLE_NODE=ON \ + -DUSE_HOST_SIDE_HDP_FLUSH=OFF\ + $src_path +cmake --build . --parallel 8 +cmake --install . diff --git a/src/assembly.hpp b/src/assembly.hpp new file mode 100644 index 0000000000..076600067e --- /dev/null +++ b/src/assembly.hpp @@ -0,0 +1,271 @@ +/****************************************************************************** + * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS + * IN THE SOFTWARE. + *****************************************************************************/ + +#ifndef LIBRARY_SRC_ASSEMBLY_HPP_ +#define LIBRARY_SRC_ASSEMBLY_HPP_ + +#include +#include +#include + +namespace rocshmem { + +#define DO_PRAGMA(x) _Pragma(#x) +#define NOWARN(warnoption, ...) \ + DO_PRAGMA(GCC diagnostic push) \ + DO_PRAGMA(GCC diagnostic ignored #warnoption) \ + __VA_ARGS__ \ + DO_PRAGMA(GCC diagnostic pop) + +#define SFENCE() asm volatile("sfence" ::: "memory") + +__device__ __forceinline__ int uncached_load_ubyte(uint8_t* src) { + int ret; +#if defined(__gfx906__) +#endif +#if defined(__gfx908__) +#endif +#if defined(__gfx90a__) + asm volatile( + "global_load_ubyte %0 %1 off glc slc \n" + "s_waitcnt vmcnt(0)" + : "=v"(ret) + : "v"(src)); +#endif +#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) + asm volatile( + "global_load_ubyte %0 %1 off sc0 sc1 \n" + "s_waitcnt vmcnt(0)" + : "=v"(ret) + : "v"(src)); +#endif + return ret; +} + +__device__ __forceinline__ void refresh_volatile_sbyte(volatile int *assigned_value, + volatile char *read_value) { +#if defined(__gfx906__) +#endif +#if defined(__gfx908__) +#endif +#if defined(__gfx90a__) + asm volatile( + "global_load_sbyte %0 %1 off glc slc\n " + "s_waitcnt vmcnt(0)" + : "=v"(*assigned_value) + : "v"(read_value)); +#endif +#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) + asm volatile( + "global_load_sbyte %0 %1 off sc0 sc1\n " + "s_waitcnt vmcnt(0)" + : "=v"(*assigned_value) + : "v"(read_value)); +#endif +} + +__device__ __forceinline__ void refresh_volatile_dwordx2(volatile uint64_t *assigned_value, + volatile uint64_t *read_value) { +#if defined(__gfx906__) +#endif +#if defined(__gfx908__) +#endif +#if defined(__gfx90a__) + asm volatile( + "global_load_dwordx2 %0 %1 off glc slc\n " + "s_waitcnt vmcnt(0)" + : "=v"(*assigned_value) + : "v"(read_value)); +#endif +#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) + asm volatile( + "global_load_dwordx2 %0 %1 off sc0 sc1\n " + "s_waitcnt vmcnt(0)" + : "=v"(*assigned_value) + : "v"(read_value)); +#endif +} + +/* Ignore the warning about deprecated volatile. + * The only usage of volatile is to force the compiler to generate + * the assembly instruction. If volatile is omitted, the compiler + * will NOT generate the non-temporal load or the waitcnt. + */ +// clang-format off +NOWARN(-Wdeprecated-volatile, + template __device__ __forceinline__ T uncached_load(T* src) { + T ret; + switch (sizeof(T)) { + case 4: +#if defined(__gfx906__) +#endif +#if defined(__gfx908__) +#endif +#if defined(__gfx90a__) + asm volatile( + "global_load_dword %0 %1 off glc slc \n" + "s_waitcnt vmcnt(0)" + : "=v"(ret) + : "v"(src)); +#endif +#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) + asm volatile( + "global_load_dword %0 %1 off sc0 sc1 \n" + "s_waitcnt vmcnt(0)" + : "=v"(ret) + : "v"(src)); +#endif + break; +#if defined(__gfx906__) +#endif +#if defined(__gfx908__) +#endif +#if defined(__gfx90a__) + case 8: + asm volatile( + "global_load_dwordx2 %0 %1 off glc slc \n" + "s_waitcnt vmcnt(0)" + : "=v"(ret) + : "v"(src)); +#endif +#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) + asm volatile( + "global_load_dwordx2 %0 %1 off sc0 sc1 \n" + "s_waitcnt vmcnt(0)" + : "=v"(ret) + : "v"(src)); +#endif + break; + default: + break; + } + return ret; + } +) +// clang-format on + +__device__ __forceinline__ void __roc_inv() { +#if defined USE_COHERENT_HEAP +#if defined(__gfx906__) +#endif +#if defined(__gfx908__) +#endif +#if defined(__gfx90a__) +// asm volatile("buffer_wbinvl1;"); +#endif +#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) +// asm volatile("buffer_inv sc0 sc1;"); +#endif +#endif +} + +__device__ __forceinline__ void __roc_flush() { +#if defined USE_COHERENT_HEAP +#if defined(__gfx906__) +#endif +#if defined(__gfx908__) +#endif +#if defined(__gfx90a__) +// asm volatile("s_dcache_wb;"); +// asm volatile("buffer_wbl2;"); +#endif +#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) +// asm volatile("s_dcache_wb;"); +// asm volatile("buffer_wbl2;"); +#endif +#endif +} + +__device__ __forceinline__ void store_asm(uint8_t* val, uint8_t* dst, + int size) { + switch (size) { + case 2: { + int16_t val16{*(reinterpret_cast(val))}; +#if defined(__gfx906__) +#endif +#if defined(__gfx908__) +#endif +#if defined(__gfx90a__) + asm volatile("flat_store_short %0 %1 glc slc" : : "v"(dst), "v"(val16)); +#endif +#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) + asm volatile("flat_store_short %0 %1 sc0 sc1" : : "v"(dst), "v"(val16)); +#endif + break; + } + case 4: { + int32_t val32{*(reinterpret_cast(val))}; +#if defined(__gfx906__) +#endif +#if defined(__gfx908__) +#endif +#if defined(__gfx90a__) + asm volatile("flat_store_dword %0 %1 glc slc" : : "v"(dst), "v"(val32)); +#endif +#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) + asm volatile("flat_store_dword %0 %1 sc0 sc1" : : "v"(dst), "v"(val32)); +#endif + break; + } + case 8: { + int64_t val64{*(reinterpret_cast(val))}; +#if defined(__gfx906__) +#endif +#if defined(__gfx908__) +#endif +#if defined(__gfx90a__) + asm volatile("flat_store_dwordx2 %0 %1 glc slc" : : "v"(dst), "v"(val64)); +#endif +#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) + asm volatile("flat_store_dwordx2 %0 %1 sc0 sc1" : : "v"(dst), "v"(val64)); +#endif + break; + } + default: + break; + } +} + +__device__ __forceinline__ uint64_t __read_clock() { + uint64_t clock{}; +#if defined(__gfx906__) +#endif +#if defined(__gfx908__) +#endif +#if defined(__gfx90a__) + asm volatile( + "s_memrealtime %0\n" + "s_waitcnt lgkmcnt(0)\n" + : "=s"(clock)); +#endif +#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) + asm volatile( + "s_memrealtime %0\n" + "s_waitcnt lgkmcnt(0)\n" + : "=s"(clock)); +#endif + return clock; +} + +} // namespace rocshmem + +#endif // LIBRARY_SRC_ASSEMBLY_HPP_ diff --git a/src/gpu_ib/connection.cpp b/src/gpu_ib/connection.cpp index 631f72f324..b6b2ae4dcd 100644 --- a/src/gpu_ib/connection.cpp +++ b/src/gpu_ib/connection.cpp @@ -256,7 +256,7 @@ void* Connection::buf_alloc([[maybe_unused]] struct ibv_pd* pd, if (use_gpu_mem) { void* dev_ptr; if (coherent_cq == 1) { -#if defined USE_COHERENT_HEAP || defined USE_CACHED_HEAP +#if defined USE_COHERENT_HEAP CHECK_HIP(hipMalloc(reinterpret_cast(&dev_ptr), size)); #else #ifdef HIP_SUPPORTS_MALLOC_UNCACHED diff --git a/src/memory/heap_type.hpp b/src/memory/heap_type.hpp index 0008317498..3f48edd11e 100644 --- a/src/memory/heap_type.hpp +++ b/src/memory/heap_type.hpp @@ -40,7 +40,7 @@ namespace rocshmem { #if defined USE_MANAGED_HEAP using HEAP_T = HeapMemory; -#elif defined USE_COHERENT_HEAP || defined USE_CACHED_HEAP +#elif defined USE_COHERENT_HEAP using HEAP_T = HeapMemory; #elif defined USE_HOST_HEAP using HEAP_T = HeapMemory; diff --git a/src/reverse_offload/context_ro_device.cpp b/src/reverse_offload/context_ro_device.cpp index 6182ce0917..8e933e9777 100644 --- a/src/reverse_offload/context_ro_device.cpp +++ b/src/reverse_offload/context_ro_device.cpp @@ -389,24 +389,6 @@ __device__ bool enough_space(BlockHandle *h, uint64_t required) { return (h->queue_size - (h->write_index - h->read_index)) >= required; } -__device__ void refresh_volatile_dwordx2(volatile uint64_t *assigned_value, - volatile uint64_t *read_value) { - __asm__ volatile( - "global_load_dwordx2 %0 %1 off glc slc\n " - "s_waitcnt vmcnt(0)" - : "=v"(*assigned_value) - : "v"(read_value)); -} - -__device__ void refresh_volatile_sbyte(volatile int *assigned_value, - volatile char *read_value) { - __asm__ volatile( - "global_load_sbyte %0 %1 off glc slc\n " - "s_waitcnt vmcnt(0)" - : "=v"(*assigned_value) - : "v"(read_value)); -} - __device__ void acquire_lock(BlockHandle *handle) { while(atomicCAS((uint64_t *)&handle->lock, 0, 1) == 1) ; } diff --git a/src/util.cpp b/src/util.cpp index 11122d7b16..7318b14208 100644 --- a/src/util.cpp +++ b/src/util.cpp @@ -40,15 +40,6 @@ typedef struct device_agent { std::vector gpu_agents; std::vector cpu_agents; -__device__ uint64_t __read_clock() { - uint64_t clock{}; - asm volatile( - "s_memrealtime %0\n" - "s_waitcnt lgkmcnt(0)\n" - : "=s"(clock)); - return clock; -} - hsa_status_t rocm_hsa_amd_memory_pool_callback( hsa_amd_memory_pool_t memory_pool, void* data) { hsa_amd_memory_pool_global_flag_t pool_flag{}; diff --git a/src/util.hpp b/src/util.hpp index 0d4dad806a..c967ec1a9c 100644 --- a/src/util.hpp +++ b/src/util.hpp @@ -29,60 +29,12 @@ #include +#include "assembly.hpp" #include "config.h" // NOLINT(build/include_subdir) #include "constants.hpp" namespace rocshmem { -#define DO_PRAGMA(x) _Pragma(#x) -#define NOWARN(warnoption, ...) \ - DO_PRAGMA(GCC diagnostic push) \ - DO_PRAGMA(GCC diagnostic ignored #warnoption) \ - __VA_ARGS__ \ - DO_PRAGMA(GCC diagnostic pop) - -__device__ __forceinline__ int uncached_load_ubyte(uint8_t* src) { - int ret; - __asm__ volatile( - "global_load_ubyte %0 %1 off glc slc \n" - "s_waitcnt vmcnt(0)" - : "=v"(ret) - : "v"(src)); - return ret; -} - -/* Ignore the warning about deprecated volatile. - * The only usage of volatile is to force the compiler to generate - * the assembly instruction. If volatile is omitted, the compiler - * will NOT generate the non-temporal load or the waitcnt. - */ -// clang-format off -NOWARN(-Wdeprecated-volatile, - template __device__ __forceinline__ T uncached_load(T* src) { - T ret; - switch (sizeof(T)) { - case 4: - __asm__ volatile( - "global_load_dword %0 %1 off glc slc \n" - "s_waitcnt vmcnt(0)" - : "=v"(ret) - : "v"(src)); - break; - case 8: - __asm__ volatile( - "global_load_dwordx2 %0 %1 off glc slc \n" - "s_waitcnt vmcnt(0)" - : "=v"(ret) - : "v"(src)); - break; - default: - break; - } - return ret; - } -) -// clang-format on - #define LOAD(VAR) __atomic_load_n((VAR), __ATOMIC_SEQ_CST) #define STORE(DST, SRC) __atomic_store_n((DST), (SRC), __ATOMIC_SEQ_CST) @@ -96,8 +48,6 @@ NOWARN(-Wdeprecated-volatile, } \ } -#define SFENCE() asm volatile("sfence" ::: "memory") - #ifdef DEBUG #define DPRINTF(...) \ do { \ @@ -123,19 +73,6 @@ NOWARN(-Wdeprecated-volatile, extern const int gpu_clock_freq_mhz; /* Device-side internal functions */ -__device__ __forceinline__ void __roc_inv() { asm volatile("buffer_wbinvl1;"); } - -__device__ __forceinline__ void __roc_flush() { -#if defined USE_COHERENT_HEAP || defined USE_CACHED_HEAP -#if __gfx90a__ - asm volatile("s_dcache_wb;"); - asm volatile("buffer_wbl2;"); -#endif -#endif -} - -__device__ uint64_t __read_clock(); - __device__ __forceinline__ uint32_t lowerID() { return __ffsll(__ballot(1)) - 1; } @@ -222,29 +159,6 @@ __device__ void gpu_dprintf(const char* fmt, const Args&... args) { } } -__device__ __forceinline__ void store_asm(uint8_t* val, uint8_t* dst, - int size) { - switch (size) { - case 2: { - int16_t val16{*(reinterpret_cast(val))}; - asm volatile("flat_store_short %0 %1 glc slc" : : "v"(dst), "v"(val16)); - break; - } - case 4: { - int32_t val32{*(reinterpret_cast(val))}; - asm volatile("flat_store_dword %0 %1 glc slc" : : "v"(dst), "v"(val32)); - break; - } - case 8: { - int64_t val64{*(reinterpret_cast(val))}; - asm volatile("flat_store_dwordx2 %0 %1 glc slc" : : "v"(dst), "v"(val64)); - break; - } - default: - break; - } -} - __device__ __forceinline__ void memcpy(void* dst, void* src, size_t size) { uint8_t* dst_bytes{static_cast(dst)}; uint8_t* src_bytes{static_cast(src)};