diff options
author | Matt Arsenault <arsenm2@gmail.com> | 2018-08-22 03:41:28 -0700 |
---|---|---|
committer | Jan Vesely <jan.vesely@rutgers.edu> | 2018-08-30 22:45:57 -0400 |
commit | a1a70991889eb45a2757a7301e7c3b0be9503e82 (patch) | |
tree | 253fc0e61dad2da997918faf494344e66cd259ae | |
parent | 47af7fcad070c755327cd3a075a8cbc524f4d37d (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.cl | 36 |
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); +} |