[libclc] r325866 - amdgcn: Fix build after GDS/const AS swap in r325030

Jan Vesely via cfe-commits cfe-commits at lists.llvm.org
Thu Feb 22 23:37:01 PST 2018


Author: jvesely
Date: Thu Feb 22 23:37:01 2018
New Revision: 325866

URL: http://llvm.org/viewvc/llvm-project?rev=325866&view=rev
Log:
amdgcn: Fix build after GDS/const AS swap in r325030

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

Added:
    libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES_4.0
    libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES_5.0
    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.50.ll
      - copied, changed from r325865, libclc/trunk/amdgcn-amdhsa/lib/workitem/get_global_size.ll
    libclc/trunk/amdgcn-amdhsa/lib/workitem/get_global_size.cl
    libclc/trunk/amdgcn-amdhsa/lib/workitem/get_local_size.50.ll
      - copied, changed from r325865, libclc/trunk/amdgcn-amdhsa/lib/workitem/get_local_size.ll
    libclc/trunk/amdgcn-amdhsa/lib/workitem/get_local_size.cl
Removed:
    libclc/trunk/amdgcn-amdhsa/lib/workitem/get_global_size.ll
    libclc/trunk/amdgcn-amdhsa/lib/workitem/get_local_size.ll
Modified:
    libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES
    libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES_3.9
    libclc/trunk/amdgcn-amdhsa/lib/SOURCES
    libclc/trunk/amdgcn/lib/cl_khr_int64_extended_atomics/minmax_helpers.ll
    libclc/trunk/amdgcn/lib/workitem/get_global_offset.cl
    libclc/trunk/amdgcn/lib/workitem/get_global_size.ll
    libclc/trunk/amdgcn/lib/workitem/get_local_size.ll
    libclc/trunk/amdgcn/lib/workitem/get_num_groups.ll
    libclc/trunk/amdgcn/lib/workitem/get_work_dim.cl

Modified: libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES?rev=325866&r1=325865&r2=325866&view=diff
==============================================================================
--- libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES (original)
+++ libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES Thu Feb 22 23:37:01 2018
@@ -1,4 +1,6 @@
 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

Modified: libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES_3.9
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES_3.9?rev=325866&r1=325865&r2=325866&view=diff
==============================================================================
--- libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES_3.9 (original)
+++ libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES_3.9 Thu Feb 22 23:37:01 2018
@@ -1,3 +1,3 @@
-workitem/get_global_size.ll
-workitem/get_local_size.ll
+workitem/get_global_size.cl
+workitem/get_local_size.cl
 workitem/get_num_groups.39.ll

Added: libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES_4.0
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES_4.0?rev=325866&view=auto
==============================================================================
--- libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES_4.0 (added)
+++ libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES_4.0 Thu Feb 22 23:37:01 2018
@@ -0,0 +1,2 @@
+workitem/get_global_size.cl
+workitem/get_local_size.cl

Added: libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES_5.0
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES_5.0?rev=325866&view=auto
==============================================================================
--- libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES_5.0 (added)
+++ libclc/trunk/amdgcn-amdhsa/lib/OVERRIDES_5.0 Thu Feb 22 23:37:01 2018
@@ -0,0 +1,2 @@
+workitem/get_global_size.cl
+workitem/get_local_size.cl

Modified: libclc/trunk/amdgcn-amdhsa/lib/SOURCES
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn-amdhsa/lib/SOURCES?rev=325866&r1=325865&r2=325866&view=diff
==============================================================================
--- libclc/trunk/amdgcn-amdhsa/lib/SOURCES (original)
+++ libclc/trunk/amdgcn-amdhsa/lib/SOURCES Thu Feb 22 23:37:01 2018
@@ -1,3 +1,3 @@
-workitem/get_global_size.ll
-workitem/get_local_size.ll
+workitem/get_global_size.cl
+workitem/get_local_size.cl
 workitem/get_num_groups.cl

Added: libclc/trunk/amdgcn-amdhsa/lib/SOURCES_4.0
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn-amdhsa/lib/SOURCES_4.0?rev=325866&view=auto
==============================================================================
--- libclc/trunk/amdgcn-amdhsa/lib/SOURCES_4.0 (added)
+++ libclc/trunk/amdgcn-amdhsa/lib/SOURCES_4.0 Thu Feb 22 23:37:01 2018
@@ -0,0 +1,2 @@
+workitem/get_global_size.50.ll
+workitem/get_local_size.50.ll

Added: libclc/trunk/amdgcn-amdhsa/lib/SOURCES_5.0
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn-amdhsa/lib/SOURCES_5.0?rev=325866&view=auto
==============================================================================
--- libclc/trunk/amdgcn-amdhsa/lib/SOURCES_5.0 (added)
+++ libclc/trunk/amdgcn-amdhsa/lib/SOURCES_5.0 Thu Feb 22 23:37:01 2018
@@ -0,0 +1,2 @@
+workitem/get_global_size.50.ll
+workitem/get_local_size.50.ll

Copied: libclc/trunk/amdgcn-amdhsa/lib/workitem/get_global_size.50.ll (from r325865, libclc/trunk/amdgcn-amdhsa/lib/workitem/get_global_size.ll)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn-amdhsa/lib/workitem/get_global_size.50.ll?p2=libclc/trunk/amdgcn-amdhsa/lib/workitem/get_global_size.50.ll&p1=libclc/trunk/amdgcn-amdhsa/lib/workitem/get_global_size.ll&r1=325865&r2=325866&rev=325866&view=diff
==============================================================================
    (empty)

Added: 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=325866&view=auto
==============================================================================
--- libclc/trunk/amdgcn-amdhsa/lib/workitem/get_global_size.cl (added)
+++ libclc/trunk/amdgcn-amdhsa/lib/workitem/get_global_size.cl Thu Feb 22 23:37:01 2018
@@ -0,0 +1,16 @@
+#include <clc/clc.h>
+
+#if __clang_major__ >= 7
+#define CONST_AS __attribute__((address_space(4)))
+#else
+#define CONST_AS __attribute__((address_space(2)))
+#endif
+
+_CLC_DEF size_t get_global_size(uint dim)
+{
+	CONST_AS uint * ptr =
+		(CONST_AS uint *) __builtin_amdgcn_dispatch_ptr();
+	if (dim < 3)
+		return ptr[3 + dim];
+	return 1;
+}

Removed: libclc/trunk/amdgcn-amdhsa/lib/workitem/get_global_size.ll
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn-amdhsa/lib/workitem/get_global_size.ll?rev=325865&view=auto
==============================================================================
--- libclc/trunk/amdgcn-amdhsa/lib/workitem/get_global_size.ll (original)
+++ libclc/trunk/amdgcn-amdhsa/lib/workitem/get_global_size.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 = !{}

Copied: libclc/trunk/amdgcn-amdhsa/lib/workitem/get_local_size.50.ll (from r325865, libclc/trunk/amdgcn-amdhsa/lib/workitem/get_local_size.ll)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn-amdhsa/lib/workitem/get_local_size.50.ll?p2=libclc/trunk/amdgcn-amdhsa/lib/workitem/get_local_size.50.ll&p1=libclc/trunk/amdgcn-amdhsa/lib/workitem/get_local_size.ll&r1=325865&r2=325866&rev=325866&view=diff
==============================================================================
    (empty)

Added: 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=325866&view=auto
==============================================================================
--- libclc/trunk/amdgcn-amdhsa/lib/workitem/get_local_size.cl (added)
+++ libclc/trunk/amdgcn-amdhsa/lib/workitem/get_local_size.cl Thu Feb 22 23:37:01 2018
@@ -0,0 +1,22 @@
+#include <clc/clc.h>
+
+#if __clang_major__ >= 7
+#define CONST_AS __attribute__((address_space(4)))
+#else
+#define CONST_AS __attribute__((address_space(2)))
+#endif
+
+_CLC_DEF size_t get_local_size(uint dim)
+{
+	CONST_AS uint * ptr =
+		(CONST_AS uint *) __builtin_amdgcn_dispatch_ptr();
+	switch (dim) {
+	case 0:
+		return ptr[1] & 0xffffu;
+	case 1:
+		return ptr[1] >> 16;
+	case 2:
+		return ptr[2] & 0xffffu;
+	}
+	return 1;
+}

Removed: libclc/trunk/amdgcn-amdhsa/lib/workitem/get_local_size.ll
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn-amdhsa/lib/workitem/get_local_size.ll?rev=325865&view=auto
==============================================================================
--- libclc/trunk/amdgcn-amdhsa/lib/workitem/get_local_size.ll (original)
+++ libclc/trunk/amdgcn-amdhsa/lib/workitem/get_local_size.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/lib/cl_khr_int64_extended_atomics/minmax_helpers.ll
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn/lib/cl_khr_int64_extended_atomics/minmax_helpers.ll?rev=325866&r1=325865&r2=325866&view=diff
==============================================================================
--- libclc/trunk/amdgcn/lib/cl_khr_int64_extended_atomics/minmax_helpers.ll (original)
+++ libclc/trunk/amdgcn/lib/cl_khr_int64_extended_atomics/minmax_helpers.ll Thu Feb 22 23:37:01 2018
@@ -1,4 +1,4 @@
-target datalayout = "e-p:64:64-p1:64:64-p2:64:64-p3:32:32-p4:32:32-p5:32:32-p6: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-A5"
+target datalayout = "e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6: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-A5"
 
 define i64 @__clc__sync_fetch_and_min_global_8(i64 addrspace(1)* nocapture %ptr, i64 %value) nounwind alwaysinline {
 entry:

Modified: libclc/trunk/amdgcn/lib/workitem/get_global_offset.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn/lib/workitem/get_global_offset.cl?rev=325866&r1=325865&r2=325866&view=diff
==============================================================================
--- libclc/trunk/amdgcn/lib/workitem/get_global_offset.cl (original)
+++ libclc/trunk/amdgcn/lib/workitem/get_global_offset.cl Thu Feb 22 23:37:01 2018
@@ -1,10 +1,15 @@
 #include <clc/clc.h>
 
+#if __clang_major__ >= 7
+#define CONST_AS __attribute__((address_space(4)))
+#else
+#define CONST_AS __attribute__((address_space(2)))
+#endif
+
 _CLC_DEF size_t get_global_offset(uint dim)
 {
-	__attribute__((address_space(2))) uint * ptr =
-		(__attribute__((address_space(2))) uint *)
-		__builtin_amdgcn_implicitarg_ptr();
+	CONST_AS uint * ptr =
+		(CONST_AS uint *) __builtin_amdgcn_implicitarg_ptr();
 	if (dim < 3)
 		return ptr[dim + 1];
 	return 0;

Modified: libclc/trunk/amdgcn/lib/workitem/get_global_size.ll
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn/lib/workitem/get_global_size.ll?rev=325866&r1=325865&r2=325866&view=diff
==============================================================================
--- libclc/trunk/amdgcn/lib/workitem/get_global_size.ll (original)
+++ libclc/trunk/amdgcn/lib/workitem/get_global_size.ll Thu Feb 22 23:37:01 2018
@@ -2,7 +2,7 @@ declare i32 @llvm.r600.read.global.size.
 declare i32 @llvm.r600.read.global.size.y() nounwind readnone
 declare i32 @llvm.r600.read.global.size.z() nounwind readnone
 
-target datalayout = "e-p:64:64-p1:64:64-p2:64:64-p3:32:32-p4:32:32-p5:32:32-p6: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-A5"
+target datalayout = "e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6: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-A5"
 
 define i64 @get_global_size(i32 %dim) nounwind readnone alwaysinline {
   switch i32 %dim, label %default [i32 0, label %x_dim i32 1, label %y_dim i32 2, label %z_dim]

Modified: libclc/trunk/amdgcn/lib/workitem/get_local_size.ll
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn/lib/workitem/get_local_size.ll?rev=325866&r1=325865&r2=325866&view=diff
==============================================================================
--- libclc/trunk/amdgcn/lib/workitem/get_local_size.ll (original)
+++ libclc/trunk/amdgcn/lib/workitem/get_local_size.ll Thu Feb 22 23:37:01 2018
@@ -2,7 +2,7 @@ declare i32 @llvm.r600.read.local.size.x
 declare i32 @llvm.r600.read.local.size.y() nounwind readnone
 declare i32 @llvm.r600.read.local.size.z() nounwind readnone
 
-target datalayout = "e-p:64:64-p1:64:64-p2:64:64-p3:32:32-p4:32:32-p5:32:32-p6: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-A5"
+target datalayout = "e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6: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-A5"
 
 define i64 @get_local_size(i32 %dim) nounwind readnone alwaysinline {
   switch i32 %dim, label %default [i32 0, label %x_dim i32 1, label %y_dim i32 2, label %z_dim]

Modified: libclc/trunk/amdgcn/lib/workitem/get_num_groups.ll
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn/lib/workitem/get_num_groups.ll?rev=325866&r1=325865&r2=325866&view=diff
==============================================================================
--- libclc/trunk/amdgcn/lib/workitem/get_num_groups.ll (original)
+++ libclc/trunk/amdgcn/lib/workitem/get_num_groups.ll Thu Feb 22 23:37:01 2018
@@ -2,7 +2,7 @@ declare i32 @llvm.r600.read.ngroups.x()
 declare i32 @llvm.r600.read.ngroups.y() nounwind readnone
 declare i32 @llvm.r600.read.ngroups.z() nounwind readnone
 
-target datalayout = "e-p:64:64-p1:64:64-p2:64:64-p3:32:32-p4:32:32-p5:32:32-p6: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-A5"
+target datalayout = "e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6: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-A5"
 
 define i64 @get_num_groups(i32 %dim) nounwind readnone alwaysinline {
   switch i32 %dim, label %default [i32 0, label %x_dim i32 1, label %y_dim i32 2, label %z_dim]

Modified: libclc/trunk/amdgcn/lib/workitem/get_work_dim.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn/lib/workitem/get_work_dim.cl?rev=325866&r1=325865&r2=325866&view=diff
==============================================================================
--- libclc/trunk/amdgcn/lib/workitem/get_work_dim.cl (original)
+++ libclc/trunk/amdgcn/lib/workitem/get_work_dim.cl Thu Feb 22 23:37:01 2018
@@ -1,9 +1,14 @@
 #include <clc/clc.h>
 
+#if __clang_major__ >= 7
+#define CONST_AS __attribute__((address_space(4)))
+#else
+#define CONST_AS __attribute__((address_space(2)))
+#endif
+
 _CLC_DEF uint get_work_dim(void)
 {
-	__attribute__((address_space(2))) uint * ptr =
-		(__attribute__((address_space(2))) uint *)
-		__builtin_amdgcn_implicitarg_ptr();
+	CONST_AS uint * ptr =
+		(CONST_AS uint *) __builtin_amdgcn_implicitarg_ptr();
 	return ptr[0];
 }




More information about the cfe-commits mailing list