[Libclc-dev] [PATCH 4/5] amdgcn: Move __clc_amdgcn_s_waitcnt definition to clc file

[Jan Vesely via Libclc-dev]

Signed-off-by: Jan Vesely <jan.vesely at rutgers.edu>
---
 amdgcn/lib/SOURCES_3.9          |  1 -
 amdgcn/lib/SOURCES_4.0          |  1 -
 amdgcn/lib/mem_fence/fence.cl   |  1 +
 amdgcn/lib/mem_fence/waitcnt.ll | 13 -------------
 4 files changed, 1 insertion(+), 15 deletions(-)
 delete mode 100644 amdgcn/lib/mem_fence/waitcnt.ll

diff --git a/amdgcn/lib/SOURCES_3.9 b/amdgcn/lib/SOURCES_3.9
index 86a222e..c97d406 100644
--- a/amdgcn/lib/SOURCES_3.9
+++ b/amdgcn/lib/SOURCES_3.9
@@ -1,2 +1 @@
 cl_khr_int64_extended_atomics/minmax_helpers.39.ll
-mem_fence/waitcnt.ll
diff --git a/amdgcn/lib/SOURCES_4.0 b/amdgcn/lib/SOURCES_4.0
index 86a222e..c97d406 100644
--- a/amdgcn/lib/SOURCES_4.0
+++ b/amdgcn/lib/SOURCES_4.0
@@ -1,2 +1 @@
 cl_khr_int64_extended_atomics/minmax_helpers.39.ll
-mem_fence/waitcnt.ll
diff --git a/amdgcn/lib/mem_fence/fence.cl b/amdgcn/lib/mem_fence/fence.cl
index 408ffc3..b85baf7 100644
--- a/amdgcn/lib/mem_fence/fence.cl
+++ b/amdgcn/lib/mem_fence/fence.cl
@@ -14,6 +14,7 @@ void __clc_amdgcn_s_waitcnt(unsigned flags);
 #  define __waitcnt(x) __builtin_amdgcn_s_waitcnt(x)
 #else
 #  define __waitcnt(x) __clc_amdgcn_s_waitcnt(x)
+_CLC_DEF void __clc_amdgcn_s_waitcnt(unsigned)  __asm("llvm.amdgcn.s.waitcnt");
 #endif
 
 _CLC_DEF void mem_fence(cl_mem_fence_flags flags)
diff --git a/amdgcn/lib/mem_fence/waitcnt.ll b/amdgcn/lib/mem_fence/waitcnt.ll
deleted file mode 100644
index ccf016a..0000000
--- a/amdgcn/lib/mem_fence/waitcnt.ll
+++ /dev/null
@@ -1,13 +0,0 @@
-declare void @llvm.amdgcn.s.waitcnt(i32) #0
-
-target datalayout = "e-p:32:32-p1:64:64-p2:64:64-p3:32:32-p4:64:64-p5:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64"
-
-; Export waitcnt intrinsic for clang < 5
-define void @__clc_amdgcn_s_waitcnt(i32 %flags) #1 {
-entry:
-  tail call void @llvm.amdgcn.s.waitcnt(i32 %flags)
-  ret void
-}
-
-attributes #0 = { nounwind }
-attributes #1 = { nounwind alwaysinline }
-- 
2.18.1