[Piglit] [PATCH] cl: Add bigger versions of calls with struct tests
Jan Vesely
jan.vesely at rutgers.edu
Mon Aug 13 22:03:01 UTC 2018
On Mon, 2018-08-13 at 13:41 -0700, Matt Arsenault wrote:
> These are just bigger versions of the existing struct
> calls tests so that they stress using byval/sret. The
> existing call with struct tests are now passed directly
> in registers.
>
> v2: Rename struct member
passes on POCL.
Reviewed-by: Jan Vesely <jan.vesely at rutgers.edu>
I'd help I you cc'ed me on patches you want me to merge directly,
rather than sending a ping later. sifting through the ML requires
extra time.
Jan
> ---
> .../cl/program/execute/calls-large-struct.cl | 156 ++++++++++++++++++
> tests/cl/program/execute/calls-struct.cl | 96 +++++------
> 2 files changed, 204 insertions(+), 48 deletions(-)
> create mode 100644 tests/cl/program/execute/calls-large-struct.cl
>
> diff --git a/tests/cl/program/execute/calls-large-struct.cl b/tests/cl/program/execute/calls-large-struct.cl
> new file mode 100644
> index 000000000..c10458f37
> --- /dev/null
> +++ b/tests/cl/program/execute/calls-large-struct.cl
> @@ -0,0 +1,156 @@
> +/*!
> +
> +[config]
> +name: calls with large structs
> +clc_version_min: 10
> +
> +[test]
> +name: byval struct
> +kernel_name: call_i32_func_byval_Char_IntArray
> +dimensions: 1
> +global_size: 16 0 0
> +
> +arg_out: 0 buffer int[16] \
> + 1021 1022 1023 1024 1025 1026 1027 1028 \
> + 1029 1030 1031 1032 1033 1034 1035 1036
> +
> +arg_out: 1 buffer int[16] \
> + 14 14 14 14 \
> + 14 14 14 14 \
> + 14 14 14 14 \
> + 14 14 14 14 \
> +
> +arg_in: 2 buffer int[16] \
> + 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15
> +
> +
> +[test]
> +name: sret struct
> +kernel_name: call_sret_Char_IntArray_func
> +dimensions: 1
> +global_size: 16 0 0
> +
> +arg_out: 0 buffer int[16] \
> + 921 922 923 924 925 926 927 928 \
> + 929 930 931 932 933 934 935 936
> +
> +arg_in: 1 buffer int[16] \
> + 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15
> +
> +!*/
> +
> +#define NOINLINE __attribute__((noinline))
> +
> +typedef struct ByVal_Char_IntArray {
> + char c;
> + int i32_arr[32];
> +} ByVal_Char_IntArray;
> +
> +NOINLINE
> +int i32_func_byval_Char_IntArray(ByVal_Char_IntArray st)
> +{
> + st.i32_arr[0] += 100;
> +
> + int sum = 0;
> + for (int i = 0; i < 32; ++i)
> + {
> + sum += st.i32_arr[i];
> + }
> +
> + sum += st.c;
> + return sum;
> +}
> +
> +kernel void call_i32_func_byval_Char_IntArray(global int* out0,
> + global int* out1,
> + global int* input)
> +{
> + ByVal_Char_IntArray st;
> + st.c = 15;
> +
> + int id = get_global_id(0);
> +
> + int val = input[id];
> +
> +
> + st.i32_arr[0] = 14;
> + st.i32_arr[1] = -8;
> + st.i32_arr[2] = val;
> + st.i32_arr[3] = 900;
> +
> + for (int i = 4; i < 32; ++i)
> + {
> + st.i32_arr[i] = 0;
> + }
> +
> + volatile int stack_object[16];
> + for (int i = 0; i < 16; ++i)
> + {
> + const int test_val = 0x07080900 | i;
> + stack_object[i] = test_val;
> + }
> +
> + int result = i32_func_byval_Char_IntArray(st);
> +
> + // Check for stack corruption
> + for (int i = 0; i < 16; ++i)
> + {
> + const int test_val = 0x07080900 | i;
> + if (stack_object[i] != test_val)
> + result = -1;
> + }
> +
> + out0[id] = result;
> + out1[id] = st.i32_arr[0];
> +}
> +
> +NOINLINE
> +ByVal_Char_IntArray sret_Char_IntArray_func(global int* input, int id)
> +{
> + ByVal_Char_IntArray st;
> + st.c = 15;
> +
> + int val = input[id];
> + st.i32_arr[0] = 14;
> + st.i32_arr[1] = -8;
> + st.i32_arr[2] = val;
> + st.i32_arr[3] = 900;
> +
> + for (int i = 4; i < 32; ++i)
> + {
> + st.i32_arr[i] = 0;
> + }
> +
> + return st;
> +}
> +
> +kernel void call_sret_Char_IntArray_func(global int* output, global int* input)
> +{
> + volatile int stack_object[16];
> + for (int i = 0; i < 16; ++i)
> + {
> + const int test_val = 0x04030200 | i;
> + stack_object[i] = test_val;
> + }
> +
> + int id = get_global_id(0);
> + ByVal_Char_IntArray st = sret_Char_IntArray_func(input, id);
> +
> + int sum = 0;
> + for (int i = 0; i < 32; ++i)
> + {
> + sum += st.i32_arr[i];
> + }
> +
> + sum += st.c;
> +
> + // Check for stack corruption
> + for (int i = 0; i < 16; ++i)
> + {
> + const int test_val = 0x04030200 | i;
> + if (stack_object[i] != test_val)
> + sum = -1;
> + }
> +
> + output[id] = sum;
> +}
> diff --git a/tests/cl/program/execute/calls-struct.cl b/tests/cl/program/execute/calls-struct.cl
> index 04f769dac..3e1fa6a85 100644
> --- a/tests/cl/program/execute/calls-struct.cl
> +++ b/tests/cl/program/execute/calls-struct.cl
> @@ -1,12 +1,12 @@
> /*!
>
> [config]
> -name: calls with structs
> +name: calls with structs passed in registers on amdgcn
> clc_version_min: 10
>
> [test]
> -name: byval struct
> -kernel_name: call_i32_func_byval_Char_IntArray
> +name: regs struct
> +kernel_name: call_i32_func_small_struct_regs_Char_IntArray
> dimensions: 1
> global_size: 16 0 0
>
> @@ -25,8 +25,8 @@ arg_in: 2 buffer int[16] \
>
>
> [test]
> -name: sret struct
> -kernel_name: call_sret_Char_IntArray_func
> +name: struct_smallregs struct
> +kernel_name: call_struct_smallregs_Char_IntArray_func
> dimensions: 1
> global_size: 16 0 0
>
> @@ -39,8 +39,8 @@ arg_in: 1 buffer int[16] \
>
>
> [test]
> -name: byval struct and sret struct
> -kernel_name: call_sret_Char_IntArray_func_byval_Char_IntArray
> +name: small struct in regs
> +kernel_name: call_struct_smallregs_Char_IntArray_func_small_struct_regs_Char_IntArray
> dimensions: 1
> global_size: 16 0 0
>
> @@ -63,70 +63,70 @@ arg_in: 2 buffer int[16] \
>
> #define NOINLINE __attribute__((noinline))
>
> -typedef struct ByVal_Char_IntArray {
> +typedef struct SmallStruct_Char_IntArray {
> char c;
> - int i[4];
> -} ByVal_Char_IntArray;
> + int i32_arr[4];
> +} SmallStruct_Char_IntArray;
>
> NOINLINE
> -int i32_func_byval_Char_IntArray(ByVal_Char_IntArray st)
> +int i32_func_small_struct_regs_Char_IntArray(SmallStruct_Char_IntArray st)
> {
> - st.i[0] += 100;
> + st.i32_arr[0] += 100;
>
> int sum = 0;
> for (int i = 0; i < 4; ++i)
> {
> - sum += st.i[i];
> + sum += st.i32_arr[i];
> }
>
> sum += st.c;
> return sum;
> }
>
> -kernel void call_i32_func_byval_Char_IntArray(global int* out0,
> - global int* out1,
> - global int* input)
> +kernel void call_i32_func_small_struct_regs_Char_IntArray(global int* out0,
> + global int* out1,
> + global int* input)
> {
> - ByVal_Char_IntArray st;
> + SmallStruct_Char_IntArray st;
> st.c = 15;
>
> int id = get_global_id(0);
>
> int val = input[id];
> - st.i[0] = 14;
> - st.i[1] = -8;
> - st.i[2] = val;
> - st.i[3] = 900;
> + st.i32_arr[0] = 14;
> + st.i32_arr[1] = -8;
> + st.i32_arr[2] = val;
> + st.i32_arr[3] = 900;
>
> - int result = i32_func_byval_Char_IntArray(st);
> + int result = i32_func_small_struct_regs_Char_IntArray(st);
> out0[id] = result;
> - out1[id] = st.i[0];
> + out1[id] = st.i32_arr[0];
> }
>
> NOINLINE
> -ByVal_Char_IntArray sret_Char_IntArray_func(global int* input, int id)
> +SmallStruct_Char_IntArray struct_smallregs_Char_IntArray_func(global int* input, int id)
> {
> - ByVal_Char_IntArray st;
> + SmallStruct_Char_IntArray st;
> st.c = 15;
>
> int val = input[id];
> - st.i[0] = 14;
> - st.i[1] = -8;
> - st.i[2] = val;
> - st.i[3] = 900;
> + st.i32_arr[0] = 14;
> + st.i32_arr[1] = -8;
> + st.i32_arr[2] = val;
> + st.i32_arr[3] = 900;
>
> return st;
> }
>
> -kernel void call_sret_Char_IntArray_func(global int* output, global int* input)
> +kernel void call_struct_smallregs_Char_IntArray_func(global int* output, global int* input)
> {
> int id = get_global_id(0);
> - ByVal_Char_IntArray st = sret_Char_IntArray_func(input, id);
> + SmallStruct_Char_IntArray st = struct_smallregs_Char_IntArray_func(input, id);
>
> int sum = 0;
> for (int i = 0; i < 4; ++i)
> {
> - sum += st.i[i];
> + sum += st.i32_arr[i];
> }
>
> sum += st.c;
> @@ -134,41 +134,41 @@ kernel void call_sret_Char_IntArray_func(global int* output, global int* input)
> }
>
> NOINLINE
> -ByVal_Char_IntArray sret_Char_IntArray_func_byval_Char_IntArray(ByVal_Char_IntArray st)
> +SmallStruct_Char_IntArray struct_smallregs_Char_IntArray_func_small_struct_regs_Char_IntArray(SmallStruct_Char_IntArray st)
> {
> st.c += 15;
>
> - st.i[0] += 14;
> - st.i[1] -= 8;
> - st.i[2] += 9;
> - st.i[3] += 18;
> + st.i32_arr[0] += 14;
> + st.i32_arr[1] -= 8;
> + st.i32_arr[2] += 9;
> + st.i32_arr[3] += 18;
>
> return st;
> }
>
> -kernel void call_sret_Char_IntArray_func_byval_Char_IntArray(global int* output0,
> - global int* output1,
> - global int* input)
> +kernel void call_struct_smallregs_Char_IntArray_func_small_struct_regs_Char_IntArray(global int* output0,
> + global int* output1,
> + global int* input)
> {
> int id = get_global_id(0);
>
> - volatile ByVal_Char_IntArray st0;
> + volatile SmallStruct_Char_IntArray st0;
> st0.c = -20;
>
> int val = input[id];
> - st0.i[0] = 14;
> - st0.i[1] = -8;
> - st0.i[2] = val;
> - st0.i[3] = 100;
> + st0.i32_arr[0] = 14;
> + st0.i32_arr[1] = -8;
> + st0.i32_arr[2] = val;
> + st0.i32_arr[3] = 100;
>
> - ByVal_Char_IntArray st1 = sret_Char_IntArray_func_byval_Char_IntArray(st0);
> + SmallStruct_Char_IntArray st1 = struct_smallregs_Char_IntArray_func_small_struct_regs_Char_IntArray(st0);
>
> int sum0 = 0;
> int sum1 = 0;
> for (int i = 0; i < 4; ++i)
> {
> - sum0 += st0.i[i];
> - sum1 += st1.i[i];
> + sum0 += st0.i32_arr[i];
> + sum1 += st1.i32_arr[i];
> }
>
> sum0 += st0.c;
--
Jan Vesely <jan.vesely at rutgers.edu>
-------------- next part --------------
A non-text attachment was scrubbed...
Name: signature.asc
Type: application/pgp-signature
Size: 833 bytes
Desc: This is a digitally signed message part
URL: <https://lists.freedesktop.org/archives/piglit/attachments/20180813/30bd9dd1/attachment.sig>
More information about the Piglit
mailing list