AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

Move inline asm into separate file

Просмотр исходного кода 
[ROCm/rocshmem commit: 862ef5713f]
Этот коммит содержится в:
Brandon Potter
2024-07-30 14:53:19 -05:00
родитель 61140c9a10
Коммит 21d8301c83
4 изменённых файлов: 237 добавлений и 142 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
projects/rocshmem/src/assembly.hpp
+236
Просмотреть файл
@@ -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 <hip/hip_runtime.h>
#include <hsa/hsa.h>
#include <hsa/hsa_ext_amd.h>
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 <typename T> __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<int16_t*>(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<int32_t*>(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<int64_t*>(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_
projects/rocshmem/src/reverse_offload/context_ro_device.cpp
-22
Просмотреть файл
@@ -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) ;
}
projects/rocshmem/src/util.cpp
-11
Просмотреть файл
@@ -40,17 +40,6 @@ typedef struct device_agent {
std::vector<device_agent_t> gpu_agents;
std::vector<device_agent_t> 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{};
projects/rocshmem/src/util.hpp
+1 -109
Просмотреть файл
@@ -29,66 +29,12 @@
#include <cstdio>
#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 <typename T> __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<int16_t*>(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<int32_t*>(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<int64_t*>(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<uint8_t*>(dst)};
uint8_t* src_bytes{static_cast<uint8_t*>(src)};