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>

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;