amdgcn-amdhsa: Convert get_{global,local}_size to clc for all llvm versions

Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Acked-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 346083

diff --git a/libclc/amdgcn-amdhsa/lib/OVERRIDES b/libclc/amdgcn-amdhsa/lib/OVERRIDES
deleted file mode 100644 (file)
index a7a694a..0000000
--- a/libclc/amdgcn-amdhsa/lib/OVERRIDES
+++ /dev/null
@@ -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

diff --git a/libclc/amdgcn-amdhsa/lib/OVERRIDES_3.9 b/libclc/amdgcn-amdhsa/lib/OVERRIDES_3.9
deleted file mode 100644 (file)
index dfe9c8e..0000000
--- a/libclc/amdgcn-amdhsa/lib/OVERRIDES_3.9
+++ /dev/null
@@ -1,3 +0,0 @@
-workitem/get_global_size.cl
-workitem/get_local_size.cl
-workitem/get_num_groups.39.ll

diff --git a/libclc/amdgcn-amdhsa/lib/OVERRIDES_4.0 b/libclc/amdgcn-amdhsa/lib/OVERRIDES_4.0
deleted file mode 100644 (file)
index ee3a48c..0000000
--- a/libclc/amdgcn-amdhsa/lib/OVERRIDES_4.0
+++ /dev/null
@@ -1,2 +0,0 @@
-workitem/get_global_size.cl
-workitem/get_local_size.cl

diff --git a/libclc/amdgcn-amdhsa/lib/OVERRIDES_5.0 b/libclc/amdgcn-amdhsa/lib/OVERRIDES_5.0
deleted file mode 100644 (file)
index ee3a48c..0000000
--- a/libclc/amdgcn-amdhsa/lib/OVERRIDES_5.0
+++ /dev/null
@@ -1,2 +0,0 @@
-workitem/get_global_size.cl
-workitem/get_local_size.cl

diff --git a/libclc/amdgcn-amdhsa/lib/SOURCES_3.9 b/libclc/amdgcn-amdhsa/lib/SOURCES_3.9
deleted file mode 100644 (file)
index a6a08af..0000000
--- a/libclc/amdgcn-amdhsa/lib/SOURCES_3.9
+++ /dev/null
@@ -1,2 +0,0 @@
-workitem/get_global_size.39.ll
-workitem/get_local_size.39.ll

diff --git a/libclc/amdgcn-amdhsa/lib/SOURCES_4.0 b/libclc/amdgcn-amdhsa/lib/SOURCES_4.0
deleted file mode 100644 (file)
index 2b957ed..0000000
--- a/libclc/amdgcn-amdhsa/lib/SOURCES_4.0
+++ /dev/null
@@ -1,2 +0,0 @@
-workitem/get_global_size.50.ll
-workitem/get_local_size.50.ll

diff --git a/libclc/amdgcn-amdhsa/lib/SOURCES_5.0 b/libclc/amdgcn-amdhsa/lib/SOURCES_5.0
deleted file mode 100644 (file)
index 2b957ed..0000000
--- a/libclc/amdgcn-amdhsa/lib/SOURCES_5.0
+++ /dev/null
@@ -1,2 +0,0 @@
-workitem/get_global_size.50.ll
-workitem/get_local_size.50.ll

diff --git a/libclc/amdgcn-amdhsa/lib/workitem/get_global_size.39.ll b/libclc/amdgcn-amdhsa/lib/workitem/get_global_size.39.ll
deleted file mode 100644 (file)
index b5e7db2..0000000
--- a/libclc/amdgcn-amdhsa/lib/workitem/get_global_size.39.ll
+++ /dev/null
@@ -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 = !{}

diff --git a/libclc/amdgcn-amdhsa/lib/workitem/get_global_size.50.ll b/libclc/amdgcn-amdhsa/lib/workitem/get_global_size.50.ll
deleted file mode 100644 (file)
index af0f2ea..0000000
--- a/libclc/amdgcn-amdhsa/lib/workitem/get_global_size.50.ll
+++ /dev/null
@@ -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 = !{}

diff --git a/libclc/amdgcn-amdhsa/lib/workitem/get_global_size.cl b/libclc/amdgcn-amdhsa/lib/workitem/get_global_size.cl
index 392cd08..2f95f99 100644 (file)
--- a/libclc/amdgcn-amdhsa/lib/workitem/get_global_size.cl
+++ b/libclc/amdgcn-amdhsa/lib/workitem/get_global_size.cl
@@ -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;

diff --git a/libclc/amdgcn-amdhsa/lib/workitem/get_local_size.39.ll b/libclc/amdgcn-amdhsa/lib/workitem/get_local_size.39.ll
deleted file mode 100644 (file)
index ecb5e8f..0000000
--- a/libclc/amdgcn-amdhsa/lib/workitem/get_local_size.39.ll
+++ /dev/null
@@ -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 }

diff --git a/libclc/amdgcn-amdhsa/lib/workitem/get_local_size.50.ll b/libclc/amdgcn-amdhsa/lib/workitem/get_local_size.50.ll
deleted file mode 100644 (file)
index ff4b811..0000000
--- a/libclc/amdgcn-amdhsa/lib/workitem/get_local_size.50.ll
+++ /dev/null
@@ -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 }

diff --git a/libclc/amdgcn-amdhsa/lib/workitem/get_local_size.cl b/libclc/amdgcn-amdhsa/lib/workitem/get_local_size.cl
index 64d1cf4..9f208d8 100644 (file)
--- a/libclc/amdgcn-amdhsa/lib/workitem/get_local_size.cl
+++ b/libclc/amdgcn-amdhsa/lib/workitem/get_local_size.cl
@@ -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;