[libclc] r346082 - amdgcn: Move __clc_amdgcn_s_waitcnt definition to clc file

[Jan Vesely via cfe-commits]

Author: jvesely
Date: Sat Nov  3 17:39:27 2018
New Revision: 346082

URL: http://llvm.org/viewvc/llvm-project?rev=346082&view=rev
Log:
amdgcn: Move __clc_amdgcn_s_waitcnt definition to clc file

Signed-off-by: Jan Vesely <jan.vesely at rutgers.edu>
Reviewed-By: Aaron Watry <awatry at gmail.com>

Removed:
    libclc/trunk/amdgcn/lib/mem_fence/waitcnt.ll
Modified:
    libclc/trunk/amdgcn/lib/SOURCES_3.9
    libclc/trunk/amdgcn/lib/SOURCES_4.0
    libclc/trunk/amdgcn/lib/mem_fence/fence.cl

Modified: libclc/trunk/amdgcn/lib/SOURCES_3.9
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn/lib/SOURCES_3.9?rev=346082&r1=346081&r2=346082&view=diff
==============================================================================
--- libclc/trunk/amdgcn/lib/SOURCES_3.9 (original)
+++ libclc/trunk/amdgcn/lib/SOURCES_3.9 Sat Nov  3 17:39:27 2018
@@ -1,2 +1 @@
 cl_khr_int64_extended_atomics/minmax_helpers.39.ll
-mem_fence/waitcnt.ll

Modified: libclc/trunk/amdgcn/lib/SOURCES_4.0
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn/lib/SOURCES_4.0?rev=346082&r1=346081&r2=346082&view=diff
==============================================================================
--- libclc/trunk/amdgcn/lib/SOURCES_4.0 (original)
+++ libclc/trunk/amdgcn/lib/SOURCES_4.0 Sat Nov  3 17:39:27 2018
@@ -1,2 +1 @@
 cl_khr_int64_extended_atomics/minmax_helpers.39.ll
-mem_fence/waitcnt.ll

Modified: libclc/trunk/amdgcn/lib/mem_fence/fence.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn/lib/mem_fence/fence.cl?rev=346082&r1=346081&r2=346082&view=diff
==============================================================================
--- libclc/trunk/amdgcn/lib/mem_fence/fence.cl (original)
+++ libclc/trunk/amdgcn/lib/mem_fence/fence.cl Sat Nov  3 17:39:27 2018
@@ -14,6 +14,7 @@ void __clc_amdgcn_s_waitcnt(unsigned fla
 #  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)

Removed: libclc/trunk/amdgcn/lib/mem_fence/waitcnt.ll
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn/lib/mem_fence/waitcnt.ll?rev=346081&view=auto
==============================================================================
--- libclc/trunk/amdgcn/lib/mem_fence/waitcnt.ll (original)
+++ libclc/trunk/amdgcn/lib/mem_fence/waitcnt.ll (removed)
@@ -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 }