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.ves...@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.ves...@rutgers.edu>
signature.asc
Description: This is a digitally signed message part
_______________________________________________ Piglit mailing list Piglit@lists.freedesktop.org https://lists.freedesktop.org/mailman/listinfo/piglit