From 21d8301c837af05ed6eb2f8f4da2a09d12859cba Mon Sep 17 00:00:00 2001 From: Brandon Potter Date: Tue, 30 Jul 2024 14:53:19 -0500 Subject: [PATCH] Move inline asm into separate file [ROCm/rocshmem commit: 862ef5713f0aebbba5292fc6bd23d6ab33eb1b14] --- projects/rocshmem/src/assembly.hpp | 236 ++++++++++++++++++ .../src/reverse_offload/context_ro_device.cpp | 22 -- projects/rocshmem/src/util.cpp | 11 - projects/rocshmem/src/util.hpp | 110 +------- 4 files changed, 237 insertions(+), 142 deletions(-) create mode 100644 projects/rocshmem/src/assembly.hpp diff --git a/projects/rocshmem/src/assembly.hpp b/projects/rocshmem/src/assembly.hpp new file mode 100644 index 0000000000..8d69249561 --- /dev/null +++ b/projects/rocshmem/src/assembly.hpp @@ -0,0 +1,236 @@ +/****************************************************************************** + * 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 __gfx906__ +#endif +#if __gfx908__ +#endif +#if __gfx90a__ + asm volatile( + "global_load_ubyte %0 %1 off glc slc \n" + "s_waitcnt vmcnt(0)" + : "=v"(ret) + : "v"(src)); +#endif +#if __gfx940__ +#endif + return ret; +} + +__device__ __forceinline__ void refresh_volatile_sbyte(volatile int *assigned_value, + volatile char *read_value) { +#if __gfx906__ +#endif +#if __gfx908__ +#endif +#if __gfx90a__ + asm volatile( + "global_load_sbyte %0 %1 off glc slc\n " + "s_waitcnt vmcnt(0)" + : "=v"(*assigned_value) + : "v"(read_value)); +#endif +#if __gfx940__ +#endif +} + +__device__ __forceinline__ void refresh_volatile_dwordx2(volatile uint64_t *assigned_value, + volatile uint64_t *read_value) { +#if __gfx906__ +#endif +#if __gfx908__ +#endif +#if __gfx90a__ + asm volatile( + "global_load_dwordx2 %0 %1 off glc slc\n " + "s_waitcnt vmcnt(0)" + : "=v"(*assigned_value) + : "v"(read_value)); +#endif +#if __gfx940__ +#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 __gfx906__ +#endif +#if __gfx908__ +#endif +#if __gfx90a__ + asm volatile( + "global_load_dword %0 %1 off glc slc \n" + "s_waitcnt vmcnt(0)" + : "=v"(ret) + : "v"(src)); +#endif +#if __gfx940__ +#endif + break; + case 8: +#if __gfx906__ +#endif +#if __gfx908__ +#endif +#if __gfx90a__ + asm volatile( + "global_load_dwordx2 %0 %1 off glc slc \n" + "s_waitcnt vmcnt(0)" + : "=v"(ret) + : "v"(src)); +#endif +#if __gfx940__ +#endif + break; + default: + break; + } + return ret; + } +) +// clang-format on + +__device__ __forceinline__ void __roc_inv() { +#if defined USE_COHERENT_HEAP +#if __gfx906__ +#endif +#if __gfx908__ +#endif +#if __gfx90a__ + asm volatile("buffer_wbinvl1;"); +#endif +#if __gfx940__ +#endif +#endif +} + +__device__ __forceinline__ void __roc_flush() { +#if defined USE_COHERENT_HEAP +#if __gfx906__ +#endif +#if __gfx908__ +#endif +#if __gfx90a__ + asm volatile("s_dcache_wb;"); + asm volatile("buffer_wbl2;"); +#endif +#if __gfx940__ +#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 __gfx906__ +#endif +#if __gfx908__ +#endif +#if __gfx90a__ + asm volatile("flat_store_short %0 %1 glc slc" : : "v"(dst), "v"(val16)); +#endif +#if __gfx940__ +#endif + break; + } + case 4: { + int32_t val32{*(reinterpret_cast(val))}; +#if __gfx906__ +#endif +#if __gfx908__ +#endif +#if __gfx90a__ + asm volatile("flat_store_dword %0 %1 glc slc" : : "v"(dst), "v"(val32)); +#endif +#if __gfx940__ +#endif + break; + } + case 8: { + int64_t val64{*(reinterpret_cast(val))}; +#if __gfx906__ +#endif +#if __gfx908__ +#endif +#if __gfx90a__ + asm volatile("flat_store_dwordx2 %0 %1 glc slc" : : "v"(dst), "v"(val64)); +#endif +#if __gfx940__ +#endif + break; + } + default: + break; + } +} + +__device__ __forceinline__ uint64_t __read_clock() { + uint64_t clock{}; +#if __gfx906__ +#endif +#if __gfx908__ +#endif +#if __gfx90a__ + asm volatile( + "s_memrealtime %0\n" + "s_waitcnt lgkmcnt(0)\n" + : "=s"(clock)); +#endif +#if __gfx940__ +#endif + return clock; +} + +} // namespace rocshmem + +#endif // LIBRARY_SRC_ASSEMBLY_HPP_ diff --git a/projects/rocshmem/src/reverse_offload/context_ro_device.cpp b/projects/rocshmem/src/reverse_offload/context_ro_device.cpp index 6c6ea8f809..8e933e9777 100644 --- a/projects/rocshmem/src/reverse_offload/context_ro_device.cpp +++ b/projects/rocshmem/src/reverse_offload/context_ro_device.cpp @@ -389,28 +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) { -#if __gfx90a__ - asm volatile( - "global_load_dwordx2 %0 %1 off glc slc\n " - "s_waitcnt vmcnt(0)" - : "=v"(*assigned_value) - : "v"(read_value)); -#endif -} - -__device__ void refresh_volatile_sbyte(volatile int *assigned_value, - volatile char *read_value) { -#if __gfx90a__ - asm volatile( - "global_load_sbyte %0 %1 off glc slc\n " - "s_waitcnt vmcnt(0)" - : "=v"(*assigned_value) - : "v"(read_value)); -#endif -} - __device__ void acquire_lock(BlockHandle *handle) { while(atomicCAS((uint64_t *)&handle->lock, 0, 1) == 1) ; } diff --git a/projects/rocshmem/src/util.cpp b/projects/rocshmem/src/util.cpp index c68bfc55c1..7318b14208 100644 --- a/projects/rocshmem/src/util.cpp +++ b/projects/rocshmem/src/util.cpp @@ -40,17 +40,6 @@ typedef struct device_agent { std::vector gpu_agents; std::vector cpu_agents; -__device__ uint64_t __read_clock() { - uint64_t clock{}; -#if __gfx90a__ - asm volatile( - "s_memrealtime %0\n" - "s_waitcnt lgkmcnt(0)\n" - : "=s"(clock)); -#endif - 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/projects/rocshmem/src/util.hpp b/projects/rocshmem/src/util.hpp index 1fa9330509..c967ec1a9c 100644 --- a/projects/rocshmem/src/util.hpp +++ b/projects/rocshmem/src/util.hpp @@ -29,66 +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; -#if __gfx90a__ - asm volatile( - "global_load_ubyte %0 %1 off glc slc \n" - "s_waitcnt vmcnt(0)" - : "=v"(ret) - : "v"(src)); -#endif - 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: -#if __gfx90a__ - asm volatile( - "global_load_dword %0 %1 off glc slc \n" - "s_waitcnt vmcnt(0)" - : "=v"(ret) - : "v"(src)); -#endif - break; - case 8: -#if __gfx90a__ - asm volatile( - "global_load_dwordx2 %0 %1 off glc slc \n" - "s_waitcnt vmcnt(0)" - : "=v"(ret) - : "v"(src)); -#endif - 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) @@ -102,8 +48,6 @@ NOWARN(-Wdeprecated-volatile, } \ } -#define SFENCE() asm volatile("sfence" ::: "memory") - #ifdef DEBUG #define DPRINTF(...) \ do { \ @@ -129,29 +73,6 @@ NOWARN(-Wdeprecated-volatile, extern const int gpu_clock_freq_mhz; /* Device-side internal functions */ -__device__ __forceinline__ void __roc_inv() { -#if defined USE_COHERENT_HEAP -#if __gfx90a__ - asm volatile("buffer_wbinvl1;"); -#endif -#endif -} - -__device__ __forceinline__ void __roc_flush() { -#if defined USE_COHERENT_HEAP -#if __gfx90a__ - asm volatile("s_dcache_wb;"); - asm volatile("buffer_wbl2;"); -#endif -#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; } @@ -238,35 +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))}; -#if __gfx90a__ - asm volatile("flat_store_short %0 %1 glc slc" : : "v"(dst), "v"(val16)); -#endif - break; - } - case 4: { - int32_t val32{*(reinterpret_cast(val))}; -#if __gfx90a__ - asm volatile("flat_store_dword %0 %1 glc slc" : : "v"(dst), "v"(val32)); -#endif - break; - } - case 8: { - int64_t val64{*(reinterpret_cast(val))}; -#if __gfx90a__ - asm volatile("flat_store_dwordx2 %0 %1 glc slc" : : "v"(dst), "v"(val64)); -#endif - 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)};