ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-21 13:29:20 +00:00
Code Issues Packages Projects Releases Wiki Activity

optimize global_load

Browse Source 
[ROCm/composable_kernel commit: 155d78594b]
This commit is contained in:
Jing Zhang
2019-04-03 21:47:58 -05:00
parent 6bf787f524
commit 5fce583246
2 changed files with 61 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

46
src/include/inline_asm.hpp
View File

@@ -4,6 +4,32 @@ typedef float Float4 __attribute__((ext_vector_type(4)));
extern "C" __attribute__((address_space(3))) void* __to_local(void* p)[[hc]];
inline __device__ void vmcnt(int cnt) {
if(cnt == 0) {
asm volatile ("\n \
s_waitcnt vmcnt(0) \n \
"::);
}
else if(cnt == 1) {
asm volatile ("\n \
s_waitcnt vmcnt(1) \n \
"::);
}
else if(cnt == 2) {
asm volatile ("\n \
s_waitcnt vmcnt(2) \n \
"::);
}
else if(cnt == 4) {
asm volatile ("\n \
s_waitcnt vmcnt(2) \n \
"::);
}
else {
assert(0);
}
}
inline __device__ void lgkmcnt(int cnt)
{
#if 1
@@ -370,3 +396,23 @@ inline __device__ void ds_read_b128(Float4& r, void* lds, int offset = 0)
assert(0);
}
}
inline __device__ void global_load(Float4 &r, Float4* ptr) {
asm volatile("\n \
global_load_dwordx4 %0, %1, off \n \
"
:"=v"(r)
:"v"(ptr)
);
}
inline __device__ void ds_write_b128(Float4& r, void* lds, int offset = 0)
{
asm volatile("\n \
ds_write_b128 %0, %1 \n \
"
:
: "v"(__to_local(lds)), "v"(r)
);
}