Skip to content

GitLab

"...composable_kernel.git" did not exist on "1f3870ca19e6392c2453e63ed187f0cca6872295"
Commit b188c0d2 authored by Jing Zhang's avatar Jing Zhang
Browse files

fixed wait

parent 0de4286a
Show whitespace changes
Inline Side-by-side
Showing with 8 additions and 13 deletions
src/include/inline_asm.hpp
View file @ b188c0d2
...@@ -5,41 +5,36 @@ typedef float Float4 __attribute__((ext_vector_type(4))); ...@@ -5,41 +5,36 @@ typedef float Float4 __attribute__((ext_vector_type(4)));
extern "C" __attribute__((address_space(3))) void* __to_local(void* p)[[hc]]; extern "C" __attribute__((address_space(3))) void* __to_local(void* p)[[hc]];
inline __device__ void lgkmcnt(int cnt){ inline __device__ void lgkmcnt(int cnt){
#if 1
if(cnt == 0) { if(cnt == 0) {
asm volatile("\n \ asm volatile("\n \
s_waitcnt lgkmcnt(0) \n \ s_waitcnt lgkmcnt(0) \n \
"::); "::);
} }
if(cnt == 1) { else if(cnt == 1) {
asm volatile("\n \ asm volatile("\n \
s_waitcnt lgkmcnt(1) \n \ s_waitcnt lgkmcnt(1) \n \
"::); "::);
} }
if(cnt == 2) { else if(cnt == 2) {
asm volatile("\n \ asm volatile("\n \
s_waitcnt lgkmcnt(2) \n \ s_waitcnt lgkmcnt(2) \n \
"::); "::);
} }
if(cnt == 3) { else if(cnt == 3) {
asm volatile("\n \ asm volatile("\n \
s_waitcnt lgkmcnt(3) \n \ s_waitcnt lgkmcnt(3) \n \
"::); "::);
} }
if(cnt == 4) { else if(cnt == 4) {
asm volatile("\n \ asm volatile("\n \
s_waitcnt lgkmcnt(4) \n \ s_waitcnt lgkmcnt(4) \n \
"::); "::);
} }
if(cnt == 5) { else {
asm volatile("\n \ assert(0);
s_waitcnt lgkmcnt(5) \n \
"::);
}
if(cnt == 6) {
asm volatile("\n \
s_waitcnt lgkmcnt(6) \n \
"::);
} }
#endif
} }
inline __device__ void outerProduct1x4(const float *a, const float *b, float *c) { inline __device__ void outerProduct1x4(const float *a, const float *b, float *c) {
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment