From b188c0d2437448cf7a17f904c9ea721b8a2fdd89 Mon Sep 17 00:00:00 2001 From: Jing Zhang Date: Wed, 3 Apr 2019 15:30:19 -0500 Subject: [PATCH] fixed wait --- src/include/inline_asm.hpp | 21 ++++++++------------- 1 file changed, 8 insertions(+), 13 deletions(-) diff --git a/src/include/inline_asm.hpp b/src/include/inline_asm.hpp index 53f58e3539..7cec3dc5ec 100644 --- a/src/include/inline_asm.hpp +++ b/src/include/inline_asm.hpp @@ -5,41 +5,36 @@ typedef float Float4 __attribute__((ext_vector_type(4))); extern "C" __attribute__((address_space(3))) void* __to_local(void* p)[[hc]]; inline __device__ void lgkmcnt(int cnt){ +#if 1 if(cnt == 0) { asm volatile("\n \ s_waitcnt lgkmcnt(0) \n \ "::); } - if(cnt == 1) { + else if(cnt == 1) { asm volatile("\n \ s_waitcnt lgkmcnt(1) \n \ "::); } - if(cnt == 2) { + else if(cnt == 2) { asm volatile("\n \ s_waitcnt lgkmcnt(2) \n \ "::); } - if(cnt == 3) { + else if(cnt == 3) { asm volatile("\n \ s_waitcnt lgkmcnt(3) \n \ "::); } - if(cnt == 4) { + else if(cnt == 4) { asm volatile("\n \ s_waitcnt lgkmcnt(4) \n \ "::); } - if(cnt == 5) { - asm volatile("\n \ - s_waitcnt lgkmcnt(5) \n \ - "::); - } - if(cnt == 6) { - asm volatile("\n \ - s_waitcnt lgkmcnt(6) \n \ - "::); + else { + assert(0); } +#endif } inline __device__ void outerProduct1x4(const float *a, const float *b, float *c) {