[Piglit] [PATCH RESEND 1/1] cl: Extend program scope arrays with more types

[Jan Vesely]

Signed-off-by: Jan Vesely <jan.vesely at rutgers.edu>
---

This used to trigger failure when LLVM put different types in different .rodata sections. Now there are no .rodata sections for radeon, but I think we should still test it.

Jan

 tests/cl/program/execute/program-scope-arrays.cl | 102 ++++++++++++++++++++++-
 1 file changed, 98 insertions(+), 4 deletions(-)

diff --git a/tests/cl/program/execute/program-scope-arrays.cl b/tests/cl/program/execute/program-scope-arrays.cl
index b9a962c..a930da9 100644
--- a/tests/cl/program/execute/program-scope-arrays.cl
+++ b/tests/cl/program/execute/program-scope-arrays.cl
@@ -6,6 +6,28 @@ global_size: 4 0 0
 
 
 [test]
+name: simple-constant-char
+kernel_name: simple_constant_char
+arg_out: 0 buffer char[4] 3 3 3 3
+
+[test]
+name: given-constant-char
+kernel_name: given_constant_char
+arg_in: 1 int 2
+arg_out: 0 buffer char[4] 2 2 2 2
+
+[test]
+name: simple-gid-char
+kernel_name: simple_gid_char
+arg_out: 0 buffer char[4] 4 3 2 1
+
+[test]
+name: indirection-char
+kernel_name: indirection_char
+arg_in: 1 buffer uchar[4] 3 2 1 0
+arg_out: 0 buffer char[4] 1 2 3 4
+
+[test]
 name: simple-constant
 kernel_name: simple_constant
 arg_out: 0 buffer float[4] 3.0 3.0 3.0 3.0
@@ -13,8 +35,8 @@ arg_out: 0 buffer float[4] 3.0 3.0 3.0 3.0
 [test]
 name: given-constant
 kernel_name: given_constant
-arg_in: 1 int 1
-arg_out: 0 buffer float[4] 3.0 3.0 3.0 3.0
+arg_in: 1 int 2
+arg_out: 0 buffer float[4] 2.0 2.0 2.0 2.0
 
 [test]
 name: simple-gid
@@ -24,8 +46,31 @@ arg_out: 0 buffer float[4] 4.0 3.0 2.0 1.0
 [test]
 name: indirection
 kernel_name: indirection
-arg_in: 1 buffer uchar[4] 0 1 2 3
-arg_out: 0 buffer float[4] 4.0 3.0 2.0 1.0
+arg_in: 1 buffer uchar[4] 3 2 1 0
+arg_out: 0 buffer float[4] 1.0 2.0 3.0 4.0
+
+
+[test]
+name: simple-constant-2
+kernel_name: simple_constant_2
+arg_out: 0 buffer float[8] 3.0 5.0 3.0 5.0 3.0 5.0 3.0 5.0
+
+[test]
+name: given-constant-2
+kernel_name: given_constant_2
+arg_in: 1 int 2
+arg_out: 0 buffer float[8] 2.0 5.0 2.0 5.0 2.0 5.0 2.0 5.0
+
+[test]
+name: simple-gid-2
+kernel_name: simple_gid_2
+arg_out: 0 buffer float[8] 4.0 5.0 3.0 5.0 2.0 5.0 1.0 5.0
+
+[test]
+name: indirection-2
+kernel_name: indirection_2
+arg_in: 1 buffer uchar[4] 3 2 1 0
+arg_out: 0 buffer float[8] 1.0 5.0 2.0 5.0 3.0 5.0 4.0 5.0
 
 !*/
 
@@ -55,3 +100,52 @@ __kernel void indirection(__global float *out, __global uchar *in) {
 	int i = get_global_id(0);
 	out[i] = arr[in[i]];
 }
+
+__constant char arrc[] = { 4, 3, 2, 1, };
+
+__kernel void simple_constant_char(__global char *out) {
+	int i = get_global_id(0);
+	out[i] = arrc[1];
+}
+
+__kernel void given_constant_char(__global char *out, int c) {
+	int i = get_global_id(0);
+	out[i] = arrc[c];
+}
+
+__kernel void simple_gid_char(__global char *out) {
+	int i = get_global_id(0);
+	out[i] = arrc[i];
+}
+
+__kernel void indirection_char(__global char *out, __global uchar *in) {
+	int i = get_global_id(0);
+	out[i] = arrc[in[i]];
+}
+
+__constant float2 arr2[] = {
+{4.0f, 5.0f},
+{3.0f, 5.0f},
+{2.0f, 5.0f},
+{1.0f, 5.0f}
+};
+
+__kernel void simple_constant_2(__global float2 *out) {
+	int i = get_global_id(0);
+	out[i] = arr2[1];
+}
+
+__kernel void given_constant_2(__global float2 *out, int c) {
+	int i = get_global_id(0);
+	out[i] = arr2[c];
+}
+
+__kernel void simple_gid_2(__global float2 *out) {
+	int i = get_global_id(0);
+	out[i] = arr2[i];
+}
+
+__kernel void indirection_2(__global float2 *out, __global uchar *in) {
+	int i = get_global_id(0);
+	out[i] = arr2[in[i]];
+}
-- 
2.5.0