summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorMatt Arsenault <arsenm2@gmail.com>2018-08-22 03:41:28 -0700
committerJan Vesely <jan.vesely@rutgers.edu>2018-08-30 22:45:57 -0400
commita1a70991889eb45a2757a7301e7c3b0be9503e82 (patch)
tree253fc0e61dad2da997918faf494344e66cd259ae
parent47af7fcad070c755327cd3a075a8cbc524f4d37d (diff)
cl: Add test for respecting byval alignment in call setup
Reviewed-by: Jan Vesely <jan.vesely@rutgers.edu>
-rw-r--r--tests/cl/program/execute/calls-large-struct.cl36
1 files changed, 36 insertions, 0 deletions
diff --git a/tests/cl/program/execute/calls-large-struct.cl b/tests/cl/program/execute/calls-large-struct.cl
index c10458f37..0eac4d470 100644
--- a/tests/cl/program/execute/calls-large-struct.cl
+++ b/tests/cl/program/execute/calls-large-struct.cl
@@ -37,6 +37,15 @@ arg_out: 0 buffer int[16] \
arg_in: 1 buffer int[16] \
0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15
+
+[test]
+name: byval struct align 8
+kernel_name: kernel_call_byval_struct_align8
+dimensions: 1
+global_size: 1 0 0
+
+arg_out: 0 buffer int[1] 1
+
!*/
#define NOINLINE __attribute__((noinline))
@@ -154,3 +163,30 @@ kernel void call_sret_Char_IntArray_func(global int* output, global int* input)
output[id] = sum;
}
+
+typedef struct ByVal_Struct_Align8 {
+ long xs[9];
+} ByVal_Struct_Align8;
+
+__attribute__((noinline))
+int func(ByVal_Struct_Align8 val)
+{
+ for (int i = 0; i < 9; ++i)
+ {
+ long ld = val.xs[i];
+ if (ld != i)
+ return 0;
+ }
+ return 1;
+}
+
+__kernel void kernel_call_byval_struct_align8(__global uint* result)
+{
+ struct ByVal_Struct_Align8 val = { { 0x1337 } };
+ for (int i = 0; i < 9; ++i)
+ {
+ val.xs[i] = i;
+ }
+
+ *result = func(val);
+}