AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Move inline assembly into arch defines blocks

Browse Source 
[ROCm/rocshmem commit: 73303ca2d2]
This commit is contained in:
Brandon Potter
2024-07-30 12:56:32 -05:00
parent 13ba4a1894
commit 61140c9a10
3 changed files with 34 additions and 6 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rocshmem/src/reverse_offload/context_ro_device.cpp
+6 -2
View File
@@ -391,20 +391,24 @@ __device__ bool enough_space(BlockHandle *h, uint64_t required) {
__device__ void refresh_volatile_dwordx2(volatile uint64_t *assigned_value,
volatile uint64_t *read_value) {
__asm__ volatile(
#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) {
__asm__ volatile(
#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) {