[Piglit] [PATCH 5/5] cl: Run stack array subtest in parallel

[Jan Vesely]

Avoid duplicate subtest names

Signed-off-by: Jan Vesely <jan.vesely at rutgers.edu>
---
 tests/cl/program/execute/i32-stack-array.cl | 125 +++++++++++-----------------
 1 file changed, 49 insertions(+), 76 deletions(-)

diff --git a/tests/cl/program/execute/i32-stack-array.cl b/tests/cl/program/execute/i32-stack-array.cl
index 534ee13..c040f79 100644
--- a/tests/cl/program/execute/i32-stack-array.cl
+++ b/tests/cl/program/execute/i32-stack-array.cl
@@ -1,6 +1,6 @@
 /*!
 [config]
-name: i32 stack array read
+name: i32 stack array
 clc_version_min: 10
 
 dimensions: 1
@@ -16,26 +16,9 @@ global_size: 1 0 0
 [test]
 kernel_name: stack_array_read
 name: i32 stack array read
-arg_out: 0 buffer int[5] 4 5 6 7 8
-arg_in:  1 buffer int[5] 0 1 2 3 4
-
-[test]
-kernel_name: stack_array_read
-name: i32 stack array read
-arg_out: 0 buffer int[5] 5 5 5 5 5
-arg_in:  1 buffer int[5] 1 1 1 1 1
-
-[test]
-kernel_name: stack_array_read
-name: i32 stack array read
-arg_out: 0 buffer int[5] 8 7 5 6 4
-arg_in:  1 buffer int[5] 4 3 1 2 0
-
-[test]
-kernel_name: stack_array_read
-name: i32 stack array read
-arg_out: 0 buffer int[5] 7 5 8 4 6
-arg_in:  1 buffer int[5] 3 1 4 0 2
+arg_out: 0 buffer int[20] 4 5 6 7 8 5 5 5 5 5 8 7 5 6 4 7 5 8 4 6
+arg_in:  1 buffer int[20] 0 1 2 3 4 1 1 1 1 1 4 3 1 2 0 3 1 4 0 2
+global_size: 4 0 0
 
 ##===----------------------------------------------------------------------===##
 # INDIRECT WRITE / DIRECT READ
@@ -73,23 +56,10 @@ arg_in:  1 buffer int[1] 1
 [test]
 kernel_name: stack_array_write_if_else_indirect_read
 name: i32 stack array direct write (IF and ELSE) indirect read
-arg_out: 0 buffer int[5] 4 5 6 7 8
-arg_in:  1 buffer int[6] 1 \
-                         0 1 2 3 4
-
-[test]
-kernel_name: stack_array_write_if_else_indirect_read
-name: i32 stack array direct write (IF and ELSE) indirect read
-arg_out: 0 buffer int[5] 0 0 0 0 0
-arg_in:  1 buffer int[6] 0 \
-                         0 1 2 3 4
-
-[test]
-kernel_name: stack_array_write_if_else_indirect_read
-name: i32 stack array direct write (IF and ELSE) indirect read
-arg_out: 0 buffer int[5] 8 7 6 5 4
-arg_in:  1 buffer int[6] 1 \
-                         4 3 2 1 0
+arg_out: 0 buffer int[15] 4 5 6 7 8 0 0 0 0 0 8 7 6 5 4
+arg_in:  1 buffer int[3] 1 0 1
+arg_in:  2 buffer int[15] 0 1 2 3 4 0 1 2 3 4 4 3 2 1 0
+global_size: 3 0 0
 
 ##===----------------------------------------------------------------------===##
 # INDIRECT WRITE (IF and ELSE) / DIRECT READ
@@ -107,32 +77,28 @@ arg_in:  1 buffer int[6] 1 \
 [test]
 kernel_name: stack_array_indirect_write_if_else_indirect_read
 name: i32 stack array indirect write (IF and ELSE) indirect read
-arg_out: 0 buffer int[5] 4 5 6 7 8
-arg_in:  1 buffer int[11] 1 \
-                          0 1 2 3 4 \
-                          0 1 2 3 4
-[test]
-kernel_name: stack_array_indirect_write_if_else_indirect_read
-name: i32 stack array indirect write (IF and ELSE) indirect read
-arg_out: 0 buffer int[5] 8 7 6 5 4
-arg_in:  1 buffer int[11] 1 \
+arg_out: 0 buffer int[10] 4 5 6 7 8 8 7 6 5 4
+arg_in:  1 buffer int[2] 1 1
+arg_in:  2 buffer int[15] 0 1 2 3 4 \
                           0 1 2 3 4 \
                           4 3 2 1 0
+global_size: 2 0 0
 !*/
 
 kernel void stack_array_read(global int* out, global int *in) {
 
+	int i = get_global_id(0) * 5;
 	int stack[5];
 	stack[0] = 4;
 	stack[1] = 5;
 	stack[2] = 6;
 	stack[3] = 7;
 	stack[4] = 8;
-	out[0] = stack[in[0]];
-	out[1] = stack[in[1]];
-	out[2] = stack[in[2]];
-	out[3] = stack[in[3]];
-	out[4] = stack[in[4]];
+	out[i + 0] = stack[in[i + 0]];
+	out[i + 1] = stack[in[i + 1]];
+	out[i + 2] = stack[in[i + 2]];
+	out[i + 3] = stack[in[i + 3]];
+	out[i + 4] = stack[in[i + 4]];
 }
 
 kernel void stack_array_write(global int* out, global int *in) {
@@ -188,10 +154,11 @@ kernel void stack_array_write_if_else_read(global int *out, global int *in) {
 	out[4] = stack[4];
 }
 
-kernel void stack_array_write_if_else_indirect_read(global int *out, global int *in) {
+kernel void stack_array_write_if_else_indirect_read(global int *out, global int *sel, global int *in) {
 
+	int i = get_global_id(0);
 	int stack[5];
-	if (in[0] == 1) {
+	if (sel[i] == 1) {
 		stack[0] = 4;
 		stack[1] = 5;
 		stack[2] = 6;
@@ -201,11 +168,13 @@ kernel void stack_array_write_if_else_indirect_read(global int *out, global int
 		stack[0] = stack[1] = stack[2] = stack[3] = stack[4] = 0;
 	}
 
-	out[0] = stack[in[1]];
-	out[1] = stack[in[2]];
-	out[2] = stack[in[3]];
-	out[3] = stack[in[4]];
-	out[4] = stack[in[5]];
+	i *= 5;
+
+	out[i + 0] = stack[in[i + 0]];
+	out[i + 1] = stack[in[i + 1]];
+	out[i + 2] = stack[in[i + 2]];
+	out[i + 3] = stack[in[i + 3]];
+	out[i + 4] = stack[in[i + 4]];
 }
 
 kernel void stack_array_indirect_write_if_else_read(global int *out, global int *in) {
@@ -232,25 +201,29 @@ kernel void stack_array_indirect_write_if_else_read(global int *out, global int
 	out[4] = stack[4];
 }
 
-kernel void stack_array_indirect_write_if_else_indirect_read(global int *out, global int *in) {
+kernel void stack_array_indirect_write_if_else_indirect_read(global int *out, global int *sel, global int *in) {
+
+	int i = get_global_id(0);
 	int stack[5];
-	if (in[0] == 1) {
-		stack[in[1]] = 4;
-		stack[in[2]] = 5;
-		stack[in[3]] = 6;
-		stack[in[4]] = 7;
-		stack[in[5]] = 8;
+	if (sel[i] == 1) {
+		stack[in[0]] = 4;
+		stack[in[1]] = 5;
+		stack[in[2]] = 6;
+		stack[in[3]] = 7;
+		stack[in[4]] = 8;
 	} else {
-		stack[in[1]] = 9;
-		stack[in[2]] = 10;
-		stack[in[3]] = 11;
-		stack[in[4]] = 12;
-		stack[in[5]] = 13;
+		stack[in[0]] = 9;
+		stack[in[1]] = 10;
+		stack[in[2]] = 11;
+		stack[in[3]] = 12;
+		stack[in[4]] = 13;
 	}
 
-	out[0] = stack[in[6]];
-	out[1] = stack[in[7]];
-	out[2] = stack[in[8]];
-	out[3] = stack[in[9]];
-	out[4] = stack[in[10]];
+	i *= 5;
+
+	out[i + 0] = stack[in[i + 5]];
+	out[i + 1] = stack[in[i + 6]];
+	out[i + 2] = stack[in[i + 7]];
+	out[i + 3] = stack[in[i + 8]];
+	out[i + 4] = stack[in[i + 9]];
 }
-- 
1.9.3