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

[Tom Stellard]

On Fri, Jan 15, 2016 at 05:41:08PM -0500, Jan Vesely wrote:
> 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
> 

Reviewed-by: Tom Stellard <thomas.stellard at amd.com>
>  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
> 
> _______________________________________________
> Piglit mailing list
> Piglit at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/piglit