diff options
author | Matt Arsenault <arsenm2@gmail.com> | 2018-08-13 13:41:19 -0700 |
---|---|---|
committer | Jan Vesely <jan.vesely@rutgers.edu> | 2018-08-13 18:15:57 -0400 |
commit | 57859e15dc8ba4034348b04d0b72f213b74d6347 (patch) | |
tree | 7c4ef8acf6e93c8fa9ba1e3cd424762e761ae778 | |
parent | 6dd406fcee57babfbb05ec9b5fa8ca8d6d8d6eb4 (diff) |
cl: Add bigger versions of calls with struct tests
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
Reviewed-by: Jan Vesely <jan.vesely@rutgers.edu>
-rw-r--r-- | tests/cl/program/execute/calls-large-struct.cl | 156 | ||||
-rw-r--r-- | tests/cl/program/execute/calls-struct.cl | 96 |
2 files changed, 204 insertions, 48 deletions
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; |