amdgcn: Fix build after GDS/const AS swap in r325030

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

11 files changed, 52 insertions, 4 deletions

diff --git a/libclc/amdgcn-amdhsa/lib/OVERRIDES b/libclc/amdgcn-amdhsa/lib/OVERRIDES
index 18ce3f8ea9c2..a7a694ade6a5 100644
--- a/libclc/amdgcn-amdhsa/lib/OVERRIDES
+++ b/libclc/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/libclc/amdgcn-amdhsa/lib/OVERRIDES_3.9 b/libclc/amdgcn-amdhsa/lib/OVERRIDES_3.9
index 4c5c851aea3c..dfe9c8e9049b 100644
--- a/libclc/amdgcn-amdhsa/lib/OVERRIDES_3.9
+++ b/libclc/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/libclc/amdgcn-amdhsa/lib/OVERRIDES_4.0 b/libclc/amdgcn-amdhsa/lib/OVERRIDES_4.0
new file mode 100644
index 000000000000..ee3a48ce2c47
--- /dev/null
+++ b/libclc/amdgcn-amdhsa/lib/OVERRIDES_4.0
@@ -0,0 +1,2 @@
+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
new file mode 100644
index 000000000000..ee3a48ce2c47
--- /dev/null
+++ b/libclc/amdgcn-amdhsa/lib/OVERRIDES_5.0
@@ -0,0 +1,2 @@
+workitem/get_global_size.cl
+workitem/get_local_size.cl
diff --git a/libclc/amdgcn-amdhsa/lib/SOURCES b/libclc/amdgcn-amdhsa/lib/SOURCES
index 7b7ea4a1a149..8224b7721b2c 100644
--- a/libclc/amdgcn-amdhsa/lib/SOURCES
+++ b/libclc/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/libclc/amdgcn-amdhsa/lib/SOURCES_4.0 b/libclc/amdgcn-amdhsa/lib/SOURCES_4.0
new file mode 100644
index 000000000000..2b957ed6f89a
--- /dev/null
+++ b/libclc/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/libclc/amdgcn-amdhsa/lib/SOURCES_5.0 b/libclc/amdgcn-amdhsa/lib/SOURCES_5.0
new file mode 100644
index 000000000000..2b957ed6f89a
--- /dev/null
+++ b/libclc/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/libclc/amdgcn-amdhsa/lib/workitem/get_global_size.ll b/libclc/amdgcn-amdhsa/lib/workitem/get_global_size.50.ll
index af0f2ea60f3d..af0f2ea60f3d 100644
--- a/libclc/amdgcn-amdhsa/lib/workitem/get_global_size.ll
+++ b/libclc/amdgcn-amdhsa/lib/workitem/get_global_size.50.ll
diff --git a/libclc/amdgcn-amdhsa/lib/workitem/get_global_size.cl b/libclc/amdgcn-amdhsa/lib/workitem/get_global_size.cl
new file mode 100644
index 000000000000..2289615a6d46
--- /dev/null
+++ b/libclc/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/libclc/amdgcn-amdhsa/lib/workitem/get_local_size.ll b/libclc/amdgcn-amdhsa/lib/workitem/get_local_size.50.ll
index ff4b81118476..ff4b81118476 100644
--- a/libclc/amdgcn-amdhsa/lib/workitem/get_local_size.ll
+++ b/libclc/amdgcn-amdhsa/lib/workitem/get_local_size.50.ll
diff --git a/libclc/amdgcn-amdhsa/lib/workitem/get_local_size.cl b/libclc/amdgcn-amdhsa/lib/workitem/get_local_size.cl
new file mode 100644
index 000000000000..034c6d95cf3a
--- /dev/null
+++ b/libclc/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;
+}