[Piglit] [PATCH] cl: Add bigger versions of calls with struct tests
Matt Arsenault
arsenm2 at gmail.com
Mon Aug 13 20:39:27 UTC 2018
> On Mar 23, 2018, at 23:07, Jan Vesely <jan.vesely at rutgers.edu> wrote:
>
> On Thu, 2018-03-15 at 11:41 -0400, Matt Arsenault wrote:
>> ping
>>
>>> On Oct 12, 2017, at 16:19, Matt Arsenault <arsenm2 at gmail.com> 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.
>>> ---
>>> tests/cl/program/execute/calls-large-struct.cl | 156 +++++++++++++++++++++++++
>>> tests/cl/program/execute/calls-struct.cl | 50 ++++----
>>> 2 files changed, 181 insertions(+), 25 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..46d84760d
>>> --- /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 i[32];
>>> +} ByVal_Char_IntArray;
>>> +
>>> +NOINLINE
>>> +int i32_func_byval_Char_IntArray(ByVal_Char_IntArray st)
>>> +{
>>> + st.i[0] += 100;
>>> +
>>> + int sum = 0;
>>> + for (int i = 0; i < 32; ++i)
>>> + {
>>> + sum += st.i[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.i[0] = 14;
>>> + st.i[1] = -8;
>>> + st.i[2] = val;
>>> + st.i[3] = 900;
>
> are these just some arbitrary numbers or do they have a specific
> meaning?
They’re arbitrary
>
>>> +
>>> + for (int i = 4; i < 32; ++i)
>>> + {
>>> + st.i[i] = 0;
>>> + }
>>> +
>>> + volatile int stack_object[16];
>>> + for (int i = 0; i < 16; ++i)
>>> + {
>>> + const int test_val = 0x07080900 | i;
> same here
Just arbitrary values to test against
>>> + 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.i[0];
>>> +}
>>> +
>>> +NOINLINE
>>> +ByVal_Char_IntArray sret_Char_IntArray_func(global int* input, int id)
>
> why is it called sret? is it "stack return"? why not spell it out, the
> test is using ridiculously long names anyway.
Because this will use the llvm sret attributed pointer argument to return the struct, rather than the direct value. sret is more specific than just returning a structure
>
>>> +{
>>> + ByVal_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;
>>> +
>>> + for (int i = 4; i < 32; ++i)
>>> + {
>>> + st.i[i] = 0;
>
> Can you use other iteration variable than "i", using the same name as
> struct members is confusing.
>
>>> + }
>>> +
>>> + 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.i[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..5d52e5587 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
>
> passed in registers where? is it gcn only? cudacl, beignet, xillinx?
> It'd be nice to mark tests that require more than OpenCL specs
> knowledge.
>
> Jan
>
>>> 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,13 +63,13 @@ 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;
>>> +} 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;
>>>
>>> @@ -83,11 +83,11 @@ int i32_func_byval_Char_IntArray(ByVal_Char_IntArray st)
>>> 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);
>>> @@ -98,15 +98,15 @@ kernel void call_i32_func_byval_Char_IntArray(global int* out0,
>>> st.i[2] = val;
>>> st.i[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];
>>> }
>>>
>>> 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];
>>> @@ -118,10 +118,10 @@ ByVal_Char_IntArray sret_Char_IntArray_func(global int* input, int id)
>>> 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)
>>> @@ -134,7 +134,7 @@ 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;
>>>
>>> @@ -146,13 +146,13 @@ ByVal_Char_IntArray sret_Char_IntArray_func_byval_Char_IntArray(ByVal_Char_IntAr
>>> 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];
>>> @@ -161,7 +161,7 @@ kernel void call_sret_Char_IntArray_func_byval_Char_IntArray(global int* output0
>>> st0.i[2] = val;
>>> st0.i[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;
>>> --
>>> 2.11.0
>>>
>>
>>
>
> --
> Jan Vesely <jan.vesely at rutgers.edu>
More information about the Piglit
mailing list