[libclc] r346083 - amdgcn-amdhsa: Convert get_{global, local}_size to clc for all llvm versions
Jan Vesely via cfe-commits
cfe-commits at lists.llvm.org
Sat Nov 3 17:39:30 PDT 2018
Author: jvesely
Date: Sat Nov 3 17:39:30 2018
New Revision: 346083
URL: http://llvm.org/viewvc/llvm-project?rev=346083&view=rev
Log:
amdgcn-amdhsa: Convert get_{global,local}_size to clc for all llvm versions
Signed-off-by: Jan Vesely <jan.vesely at rutgers.edu>
Acked-By: Aaron Watry <awatry at gmail.com>
Removed:
libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES
libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES_3.9
libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES_4.0
libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES_5.0
libclc/trunk/amdgcn-amdhsa/lib/SOURCES_3.9
libclc/trunk/amdgcn-amdhsa/lib/SOURCES_4.0
libclc/trunk/amdgcn-amdhsa/lib/SOURCES_5.0
libclc/trunk/amdgcn-amdhsa/lib/workitem/get_global_size.39.ll
libclc/trunk/amdgcn-amdhsa/lib/workitem/get_global_size.50.ll
libclc/trunk/amdgcn-amdhsa/lib/workitem/get_local_size.39.ll
libclc/trunk/amdgcn-amdhsa/lib/workitem/get_local_size.50.ll
Modified:
libclc/trunk/amdgcn-amdhsa/lib/workitem/get_global_size.cl
libclc/trunk/amdgcn-amdhsa/lib/workitem/get_local_size.cl
Removed: libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES?rev=346082&view=auto
==============================================================================
--- libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES (original)
+++ libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES (removed)
@@ -1,6 +0,0 @@
-workitem/get_num_groups.ll
-workitem/get_global_size.ll
-workitem/get_local_size.ll
-workitem/get_num_groups.40.ll
-workitem/get_global_size.40.ll
-workitem/get_local_size.40.ll
Removed: libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES_3.9
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES_3.9?rev=346082&view=auto
==============================================================================
--- libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES_3.9 (original)
+++ libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES_3.9 (removed)
@@ -1,3 +0,0 @@
-workitem/get_global_size.cl
-workitem/get_local_size.cl
-workitem/get_num_groups.39.ll
Removed: libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES_4.0
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES_4.0?rev=346082&view=auto
==============================================================================
--- libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES_4.0 (original)
+++ libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES_4.0 (removed)
@@ -1,2 +0,0 @@
-workitem/get_global_size.cl
-workitem/get_local_size.cl
Removed: libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES_5.0
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES_5.0?rev=346082&view=auto
==============================================================================
--- libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES_5.0 (original)
+++ libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES_5.0 (removed)
@@ -1,2 +0,0 @@
-workitem/get_global_size.cl
-workitem/get_local_size.cl
Removed: libclc/trunk/amdgcn-amdhsa/lib/SOURCES_3.9
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn-amdhsa/lib/SOURCES_3.9?rev=346082&view=auto
==============================================================================
--- libclc/trunk/amdgcn-amdhsa/lib/SOURCES_3.9 (original)
+++ libclc/trunk/amdgcn-amdhsa/lib/SOURCES_3.9 (removed)
@@ -1,2 +0,0 @@
-workitem/get_global_size.39.ll
-workitem/get_local_size.39.ll
Removed: libclc/trunk/amdgcn-amdhsa/lib/SOURCES_4.0
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn-amdhsa/lib/SOURCES_4.0?rev=346082&view=auto
==============================================================================
--- libclc/trunk/amdgcn-amdhsa/lib/SOURCES_4.0 (original)
+++ libclc/trunk/amdgcn-amdhsa/lib/SOURCES_4.0 (removed)
@@ -1,2 +0,0 @@
-workitem/get_global_size.50.ll
-workitem/get_local_size.50.ll
Removed: libclc/trunk/amdgcn-amdhsa/lib/SOURCES_5.0
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn-amdhsa/lib/SOURCES_5.0?rev=346082&view=auto
==============================================================================
--- libclc/trunk/amdgcn-amdhsa/lib/SOURCES_5.0 (original)
+++ libclc/trunk/amdgcn-amdhsa/lib/SOURCES_5.0 (removed)
@@ -1,2 +0,0 @@
-workitem/get_global_size.50.ll
-workitem/get_local_size.50.ll
Removed: libclc/trunk/amdgcn-amdhsa/lib/workitem/get_global_size.39.ll
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn-amdhsa/lib/workitem/get_global_size.39.ll?rev=346082&view=auto
==============================================================================
--- libclc/trunk/amdgcn-amdhsa/lib/workitem/get_global_size.39.ll (original)
+++ libclc/trunk/amdgcn-amdhsa/lib/workitem/get_global_size.39.ll (removed)
@@ -1,36 +0,0 @@
-declare i8 addrspace(2)* @llvm.amdgcn.dispatch.ptr() #0
-
-define i32 @get_global_size(i32 %dim) #1 {
- %dispatch_ptr = call noalias nonnull dereferenceable(64) i8 addrspace(2)* @llvm.amdgcn.dispatch.ptr()
- switch i32 %dim, label %default [
- i32 0, label %x
- i32 1, label %y
- i32 2, label %z
- ]
-
-x:
- %ptr_x = getelementptr inbounds i8, i8 addrspace(2)* %dispatch_ptr, i32 12
- %ptr_x32 = bitcast i8 addrspace(2)* %ptr_x to i32 addrspace(2)*
- %x32 = load i32, i32 addrspace(2)* %ptr_x32, align 4, !invariant.load !0
- ret i32 %x32
-
-y:
- %ptr_y = getelementptr inbounds i8, i8 addrspace(2)* %dispatch_ptr, i32 16
- %ptr_y32 = bitcast i8 addrspace(2)* %ptr_y to i32 addrspace(2)*
- %y32 = load i32, i32 addrspace(2)* %ptr_y32, align 4, !invariant.load !0
- ret i32 %y32
-
-z:
- %ptr_z = getelementptr inbounds i8, i8 addrspace(2)* %dispatch_ptr, i32 20
- %ptr_z32 = bitcast i8 addrspace(2)* %ptr_z to i32 addrspace(2)*
- %z32 = load i32, i32 addrspace(2)* %ptr_z32, align 4, !invariant.load !0
- ret i32 %z32
-
-default:
- ret i32 1
-}
-
-attributes #0 = { nounwind readnone }
-attributes #1 = { alwaysinline norecurse nounwind readonly }
-
-!0 = !{}
Removed: libclc/trunk/amdgcn-amdhsa/lib/workitem/get_global_size.50.ll
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn-amdhsa/lib/workitem/get_global_size.50.ll?rev=346082&view=auto
==============================================================================
--- libclc/trunk/amdgcn-amdhsa/lib/workitem/get_global_size.50.ll (original)
+++ libclc/trunk/amdgcn-amdhsa/lib/workitem/get_global_size.50.ll (removed)
@@ -1,39 +0,0 @@
-declare i8 addrspace(2)* @llvm.amdgcn.dispatch.ptr() #0
-
-define i64 @get_global_size(i32 %dim) #1 {
- %dispatch_ptr = call noalias nonnull dereferenceable(64) i8 addrspace(2)* @llvm.amdgcn.dispatch.ptr()
- switch i32 %dim, label %default [
- i32 0, label %x
- i32 1, label %y
- i32 2, label %z
- ]
-
-x:
- %ptr_x = getelementptr inbounds i8, i8 addrspace(2)* %dispatch_ptr, i64 12
- %ptr_x32 = bitcast i8 addrspace(2)* %ptr_x to i32 addrspace(2)*
- %x32 = load i32, i32 addrspace(2)* %ptr_x32, align 4, !invariant.load !0
- %size_x = zext i32 %x32 to i64
- ret i64 %size_x
-
-y:
- %ptr_y = getelementptr inbounds i8, i8 addrspace(2)* %dispatch_ptr, i64 16
- %ptr_y32 = bitcast i8 addrspace(2)* %ptr_y to i32 addrspace(2)*
- %y32 = load i32, i32 addrspace(2)* %ptr_y32, align 4, !invariant.load !0
- %size_y = zext i32 %y32 to i64
- ret i64 %size_y
-
-z:
- %ptr_z = getelementptr inbounds i8, i8 addrspace(2)* %dispatch_ptr, i64 20
- %ptr_z32 = bitcast i8 addrspace(2)* %ptr_z to i32 addrspace(2)*
- %z32 = load i32, i32 addrspace(2)* %ptr_z32, align 4, !invariant.load !0
- %size_z = zext i32 %z32 to i64
- ret i64 %size_z
-
-default:
- ret i64 1
-}
-
-attributes #0 = { nounwind readnone }
-attributes #1 = { alwaysinline norecurse nounwind readonly }
-
-!0 = !{}
Modified: libclc/trunk/amdgcn-amdhsa/lib/workitem/get_global_size.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn-amdhsa/lib/workitem/get_global_size.cl?rev=346083&r1=346082&r2=346083&view=diff
==============================================================================
--- libclc/trunk/amdgcn-amdhsa/lib/workitem/get_global_size.cl (original)
+++ libclc/trunk/amdgcn-amdhsa/lib/workitem/get_global_size.cl Sat Nov 3 17:39:30 2018
@@ -8,10 +8,16 @@
#define CONST_AS __attribute__((address_space(2)))
#endif
+#if __clang_major__ >= 6
+#define __dispatch_ptr __builtin_amdgcn_dispatch_ptr
+#else
+#define __dispatch_ptr __clc_amdgcn_dispatch_ptr
+CONST_AS uchar * __clc_amdgcn_dispatch_ptr(void) __asm("llvm.amdgcn.dispatch.ptr");
+#endif
+
_CLC_DEF size_t get_global_size(uint dim)
{
- CONST_AS uint * ptr =
- (CONST_AS uint *) __builtin_amdgcn_dispatch_ptr();
+ CONST_AS uint * ptr = (CONST_AS uint *) __dispatch_ptr();
if (dim < 3)
return ptr[3 + dim];
return 1;
Removed: libclc/trunk/amdgcn-amdhsa/lib/workitem/get_local_size.39.ll
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn-amdhsa/lib/workitem/get_local_size.39.ll?rev=346082&view=auto
==============================================================================
--- libclc/trunk/amdgcn-amdhsa/lib/workitem/get_local_size.39.ll (original)
+++ libclc/trunk/amdgcn-amdhsa/lib/workitem/get_local_size.39.ll (removed)
@@ -1,35 +0,0 @@
-declare i8 addrspace(2)* @llvm.amdgcn.dispatch.ptr() #0
-
-define i32 @get_local_size(i32 %dim) #1 {
- %dispatch_ptr = call noalias nonnull dereferenceable(64) i8 addrspace(2)* @llvm.amdgcn.dispatch.ptr()
- %dispatch_ptr_i32 = bitcast i8 addrspace(2)* %dispatch_ptr to i32 addrspace(2)*
- %xy_size_ptr = getelementptr inbounds i32, i32 addrspace(2)* %dispatch_ptr_i32, i32 1
- %xy_size = load i32, i32 addrspace(2)* %xy_size_ptr, align 4, !invariant.load !0
- switch i32 %dim, label %default [
- i32 0, label %x_dim
- i32 1, label %y_dim
- i32 2, label %z_dim
- ]
-
-x_dim:
- %x_size = and i32 %xy_size, 65535
- ret i32 %x_size
-
-y_dim:
- %y_size = lshr i32 %xy_size, 16
- ret i32 %y_size
-
-z_dim:
- %z_size_ptr = getelementptr inbounds i32, i32 addrspace(2)* %dispatch_ptr_i32, i32 2
- %z_size = load i32, i32 addrspace(2)* %z_size_ptr, align 4, !invariant.load !0, !range !1
- ret i32 %z_size
-
-default:
- ret i32 1
-}
-
-attributes #0 = { nounwind readnone }
-attributes #1 = { alwaysinline norecurse nounwind readonly }
-
-!0 = !{}
-!1 = !{ i32 0, i32 257 }
Removed: libclc/trunk/amdgcn-amdhsa/lib/workitem/get_local_size.50.ll
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn-amdhsa/lib/workitem/get_local_size.50.ll?rev=346082&view=auto
==============================================================================
--- libclc/trunk/amdgcn-amdhsa/lib/workitem/get_local_size.50.ll (original)
+++ libclc/trunk/amdgcn-amdhsa/lib/workitem/get_local_size.50.ll (removed)
@@ -1,38 +0,0 @@
-declare i8 addrspace(2)* @llvm.amdgcn.dispatch.ptr() #0
-
-define i64 @get_local_size(i32 %dim) #1 {
- %dispatch_ptr = call noalias nonnull dereferenceable(64) i8 addrspace(2)* @llvm.amdgcn.dispatch.ptr()
- %dispatch_ptr_i32 = bitcast i8 addrspace(2)* %dispatch_ptr to i32 addrspace(2)*
- %xy_size_ptr = getelementptr inbounds i32, i32 addrspace(2)* %dispatch_ptr_i32, i64 1
- %xy_size = load i32, i32 addrspace(2)* %xy_size_ptr, align 4, !invariant.load !0
- switch i32 %dim, label %default [
- i32 0, label %x_dim
- i32 1, label %y_dim
- i32 2, label %z_dim
- ]
-
-x_dim:
- %x_size = and i32 %xy_size, 65535
- %x_size.ext = zext i32 %x_size to i64
- ret i64 %x_size.ext
-
-y_dim:
- %y_size = lshr i32 %xy_size, 16
- %y_size.ext = zext i32 %y_size to i64
- ret i64 %y_size.ext
-
-z_dim:
- %z_size_ptr = getelementptr inbounds i32, i32 addrspace(2)* %dispatch_ptr_i32, i64 2
- %z_size = load i32, i32 addrspace(2)* %z_size_ptr, align 4, !invariant.load !0, !range !1
- %z_size.ext = zext i32 %z_size to i64
- ret i64 %z_size.ext
-
-default:
- ret i64 1
-}
-
-attributes #0 = { nounwind readnone }
-attributes #1 = { alwaysinline norecurse nounwind readonly }
-
-!0 = !{}
-!1 = !{ i32 0, i32 257 }
Modified: libclc/trunk/amdgcn-amdhsa/lib/workitem/get_local_size.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn-amdhsa/lib/workitem/get_local_size.cl?rev=346083&r1=346082&r2=346083&view=diff
==============================================================================
--- libclc/trunk/amdgcn-amdhsa/lib/workitem/get_local_size.cl (original)
+++ libclc/trunk/amdgcn-amdhsa/lib/workitem/get_local_size.cl Sat Nov 3 17:39:30 2018
@@ -8,10 +8,16 @@
#define CONST_AS __attribute__((address_space(2)))
#endif
+#if __clang_major__ >= 6
+#define __dispatch_ptr __builtin_amdgcn_dispatch_ptr
+#else
+#define __dispatch_ptr __clc_amdgcn_dispatch_ptr
+CONST_AS char * __clc_amdgcn_dispatch_ptr(void) __asm("llvm.amdgcn.dispatch.ptr");
+#endif
+
_CLC_DEF size_t get_local_size(uint dim)
{
- CONST_AS uint * ptr =
- (CONST_AS uint *) __builtin_amdgcn_dispatch_ptr();
+ CONST_AS uint * ptr = (CONST_AS uint *) __dispatch_ptr();
switch (dim) {
case 0:
return ptr[1] & 0xffffu;
More information about the cfe-commits
mailing list