[libclc] 61/79: amdgcn: Fix build after GDS/const AS swap in r325030
Andreas Boll
aboll-guest at moszumanska.debian.org
Mon Mar 19 16:51:01 UTC 2018
This is an automated email from the git hooks/post-receive script.
aboll-guest pushed a commit to branch master
in repository libclc.
commit fb39e8206d1eee316691e311311b9f1262c63e39
Author: Jan Vesely <jan.vesely at rutgers.edu>
Date: Fri Feb 23 07:37:01 2018 +0000
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>
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@325866 91177308-0d34-0410-b5e6-96231b3b80d8
---
amdgcn-amdhsa/lib/OVERRIDES | 2 ++
amdgcn-amdhsa/lib/OVERRIDES_3.9 | 4 ++--
amdgcn-amdhsa/lib/OVERRIDES_4.0 | 2 ++
amdgcn-amdhsa/lib/OVERRIDES_5.0 | 2 ++
amdgcn-amdhsa/lib/SOURCES | 4 ++--
amdgcn-amdhsa/lib/SOURCES_4.0 | 2 ++
amdgcn-amdhsa/lib/SOURCES_5.0 | 2 ++
.../{get_global_size.ll => get_global_size.50.ll} | 0
amdgcn-amdhsa/lib/workitem/get_global_size.cl | 16 ++++++++++++++++
.../{get_local_size.ll => get_local_size.50.ll} | 0
amdgcn-amdhsa/lib/workitem/get_local_size.cl | 22 ++++++++++++++++++++++
.../minmax_helpers.ll | 2 +-
amdgcn/lib/workitem/get_global_offset.cl | 11 ++++++++---
amdgcn/lib/workitem/get_global_size.ll | 2 +-
amdgcn/lib/workitem/get_local_size.ll | 2 +-
amdgcn/lib/workitem/get_num_groups.ll | 2 +-
amdgcn/lib/workitem/get_work_dim.cl | 11 ++++++++---
17 files changed, 72 insertions(+), 14 deletions(-)
diff --git a/amdgcn-amdhsa/lib/OVERRIDES b/amdgcn-amdhsa/lib/OVERRIDES
index 18ce3f8..a7a694a 100644
--- a/amdgcn-amdhsa/lib/OVERRIDES
+++ b/amdgcn-amdhsa/lib/OVERRIDES
@@ -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
diff --git a/amdgcn-amdhsa/lib/OVERRIDES_3.9 b/amdgcn-amdhsa/lib/OVERRIDES_3.9
index 4c5c851..dfe9c8e 100644
--- a/amdgcn-amdhsa/lib/OVERRIDES_3.9
+++ b/amdgcn-amdhsa/lib/OVERRIDES_3.9
@@ -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
diff --git a/amdgcn-amdhsa/lib/OVERRIDES_4.0 b/amdgcn-amdhsa/lib/OVERRIDES_4.0
new file mode 100644
index 0000000..ee3a48c
--- /dev/null
+++ b/amdgcn-amdhsa/lib/OVERRIDES_4.0
@@ -0,0 +1,2 @@
+workitem/get_global_size.cl
+workitem/get_local_size.cl
diff --git a/amdgcn-amdhsa/lib/OVERRIDES_5.0 b/amdgcn-amdhsa/lib/OVERRIDES_5.0
new file mode 100644
index 0000000..ee3a48c
--- /dev/null
+++ b/amdgcn-amdhsa/lib/OVERRIDES_5.0
@@ -0,0 +1,2 @@
+workitem/get_global_size.cl
+workitem/get_local_size.cl
diff --git a/amdgcn-amdhsa/lib/SOURCES b/amdgcn-amdhsa/lib/SOURCES
index 7b7ea4a..8224b77 100644
--- a/amdgcn-amdhsa/lib/SOURCES
+++ b/amdgcn-amdhsa/lib/SOURCES
@@ -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
diff --git a/amdgcn-amdhsa/lib/SOURCES_4.0 b/amdgcn-amdhsa/lib/SOURCES_4.0
new file mode 100644
index 0000000..2b957ed
--- /dev/null
+++ b/amdgcn-amdhsa/lib/SOURCES_4.0
@@ -0,0 +1,2 @@
+workitem/get_global_size.50.ll
+workitem/get_local_size.50.ll
diff --git a/amdgcn-amdhsa/lib/SOURCES_5.0 b/amdgcn-amdhsa/lib/SOURCES_5.0
new file mode 100644
index 0000000..2b957ed
--- /dev/null
+++ b/amdgcn-amdhsa/lib/SOURCES_5.0
@@ -0,0 +1,2 @@
+workitem/get_global_size.50.ll
+workitem/get_local_size.50.ll
diff --git a/amdgcn-amdhsa/lib/workitem/get_global_size.ll b/amdgcn-amdhsa/lib/workitem/get_global_size.50.ll
similarity index 100%
rename from amdgcn-amdhsa/lib/workitem/get_global_size.ll
rename to amdgcn-amdhsa/lib/workitem/get_global_size.50.ll
diff --git a/amdgcn-amdhsa/lib/workitem/get_global_size.cl b/amdgcn-amdhsa/lib/workitem/get_global_size.cl
new file mode 100644
index 0000000..2289615
--- /dev/null
+++ b/amdgcn-amdhsa/lib/workitem/get_global_size.cl
@@ -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;
+}
diff --git a/amdgcn-amdhsa/lib/workitem/get_local_size.ll b/amdgcn-amdhsa/lib/workitem/get_local_size.50.ll
similarity index 100%
rename from amdgcn-amdhsa/lib/workitem/get_local_size.ll
rename to amdgcn-amdhsa/lib/workitem/get_local_size.50.ll
diff --git a/amdgcn-amdhsa/lib/workitem/get_local_size.cl b/amdgcn-amdhsa/lib/workitem/get_local_size.cl
new file mode 100644
index 0000000..034c6d9
--- /dev/null
+++ b/amdgcn-amdhsa/lib/workitem/get_local_size.cl
@@ -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;
+}
diff --git a/amdgcn/lib/cl_khr_int64_extended_atomics/minmax_helpers.ll b/amdgcn/lib/cl_khr_int64_extended_atomics/minmax_helpers.ll
index 8ab771c..e678434 100644
--- a/amdgcn/lib/cl_khr_int64_extended_atomics/minmax_helpers.ll
+++ b/amdgcn/lib/cl_khr_int64_extended_atomics/minmax_helpers.ll
@@ -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:
diff --git a/amdgcn/lib/workitem/get_global_offset.cl b/amdgcn/lib/workitem/get_global_offset.cl
index 3b359ae..0c2b948 100644
--- a/amdgcn/lib/workitem/get_global_offset.cl
+++ b/amdgcn/lib/workitem/get_global_offset.cl
@@ -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;
diff --git a/amdgcn/lib/workitem/get_global_size.ll b/amdgcn/lib/workitem/get_global_size.ll
index c87cd0f..40636a1 100644
--- a/amdgcn/lib/workitem/get_global_size.ll
+++ b/amdgcn/lib/workitem/get_global_size.ll
@@ -2,7 +2,7 @@ declare i32 @llvm.r600.read.global.size.x() nounwind readnone
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]
diff --git a/amdgcn/lib/workitem/get_local_size.ll b/amdgcn/lib/workitem/get_local_size.ll
index 77f5670..8ed4925 100644
--- a/amdgcn/lib/workitem/get_local_size.ll
+++ b/amdgcn/lib/workitem/get_local_size.ll
@@ -2,7 +2,7 @@ declare i32 @llvm.r600.read.local.size.x() nounwind readnone
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]
diff --git a/amdgcn/lib/workitem/get_num_groups.ll b/amdgcn/lib/workitem/get_num_groups.ll
index 92c7553..4283197 100644
--- a/amdgcn/lib/workitem/get_num_groups.ll
+++ b/amdgcn/lib/workitem/get_num_groups.ll
@@ -2,7 +2,7 @@ declare i32 @llvm.r600.read.ngroups.x() nounwind readnone
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]
diff --git a/amdgcn/lib/workitem/get_work_dim.cl b/amdgcn/lib/workitem/get_work_dim.cl
index 8d8a30c..8eb81ba 100644
--- a/amdgcn/lib/workitem/get_work_dim.cl
+++ b/amdgcn/lib/workitem/get_work_dim.cl
@@ -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];
}
--
Alioth's /usr/local/bin/git-commit-notice on /srv/git.debian.org/git/pkg-opencl/libclc.git
More information about the Pkg-opencl-commits
mailing list