diff options
author | Matt Arsenault <arsenm2@gmail.com> | 2018-08-22 03:41:37 -0700 |
---|---|---|
committer | Jan Vesely <jan.vesely@rutgers.edu> | 2018-08-30 22:45:57 -0400 |
commit | a9752f23ef3531c6c44fc60cfff9811862fd594a (patch) | |
tree | c311d594ecee2100ad7870e79b37e92ebe11575b | |
parent | a1a70991889eb45a2757a7301e7c3b0be9503e82 (diff) |
cl: Add tests for calls with special inputs
Also fixes apparently missing coverage for special
input arguments not passed in registers.
Reviewed-by: Jan Vesely <jan.vesely@rutgers.edu>
-rw-r--r-- | tests/cl/program/execute/calls-workitem-id.cl | 136 |
1 files changed, 136 insertions, 0 deletions
diff --git a/tests/cl/program/execute/calls-workitem-id.cl b/tests/cl/program/execute/calls-workitem-id.cl index 7edfad7e9..b42c85959 100644 --- a/tests/cl/program/execute/calls-workitem-id.cl +++ b/tests/cl/program/execute/calls-workitem-id.cl @@ -38,6 +38,56 @@ arg_out: 2 buffer uint[64] \ 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 \ 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 +[test] +name: Callee function stack passed get_local_id +kernel_name: kernel_call_too_many_argument_regs_get_local_id_012 +dimensions: 3 +global_size: 8 4 2 +local_size: 8 4 2 + +arg_out: 0 buffer uint[64] \ + 0 1 2 3 4 5 6 7 0 1 2 3 4 5 6 7 \ + 0 1 2 3 4 5 6 7 0 1 2 3 4 5 6 7 \ + 0 1 2 3 4 5 6 7 0 1 2 3 4 5 6 7 \ + 0 1 2 3 4 5 6 7 0 1 2 3 4 5 6 7 + +arg_out: 1 buffer uint[64] \ + 0 0 0 0 0 0 0 0 1 1 1 1 1 1 1 1 \ + 2 2 2 2 2 2 2 2 3 3 3 3 3 3 3 3 \ + 0 0 0 0 0 0 0 0 1 1 1 1 1 1 1 1 \ + 2 2 2 2 2 2 2 2 3 3 3 3 3 3 3 3 + +arg_out: 2 buffer uint[64] \ + 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 \ + 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 \ + 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 \ + 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 + +[test] +name: Callee function stack passed get_local_id with byval +kernel_name: kernel_call_too_many_argument_regs_byval_get_local_id_012 +dimensions: 3 +global_size: 8 4 2 +local_size: 8 4 2 + +arg_out: 0 buffer uint[64] \ + 45 46 47 48 49 50 51 52 45 46 47 48 49 50 51 52 \ + 45 46 47 48 49 50 51 52 45 46 47 48 49 50 51 52 \ + 45 46 47 48 49 50 51 52 45 46 47 48 49 50 51 52 \ + 45 46 47 48 49 50 51 52 45 46 47 48 49 50 51 52 + +arg_out: 1 buffer uint[64] \ + 47 47 47 47 47 47 47 47 48 48 48 48 48 48 48 48 \ + 49 49 49 49 49 49 49 49 50 50 50 50 50 50 50 50 \ + 47 47 47 47 47 47 47 47 48 48 48 48 48 48 48 48 \ + 49 49 49 49 49 49 49 49 50 50 50 50 50 50 50 50 + +arg_out: 2 buffer uint[64] \ + 50 50 50 50 50 50 50 50 50 50 50 50 50 50 50 50 \ + 50 50 50 50 50 50 50 50 50 50 50 50 50 50 50 50 \ + 51 51 51 51 51 51 51 51 51 51 51 51 51 51 51 51 \ + 51 51 51 51 51 51 51 51 51 51 51 51 51 51 51 51 + !*/ #define NOINLINE __attribute__((noinline)) @@ -75,3 +125,89 @@ kernel void kernel_call_pass_get_global_id_012(global uint *out0, { func_get_global_id_012(out0, out1, out2); } + +// On amdgcn, this will require the workitem IDs be passed as values +// on the stack after the arguments. +NOINLINE +uint3 too_many_argument_regs_get_local_id_012( + int arg0, int arg1, int arg2, int arg3, + int arg4, int arg5, int arg6, int arg7, + int arg8, int arg9, int arg10, int arg11, + int arg12, int arg13, int arg14, int arg15, + int arg16, int arg17, int arg18, int arg19, + int arg20, int arg21, int arg22, int arg23, + int arg24, int arg25, int arg26, int arg27, + int arg28, int arg29, int arg30, int arg31) +{ + uint3 id; + id.x = get_local_id(0); + id.y = get_local_id(1); + id.z = get_local_id(2); + return id; +} + +kernel void kernel_call_too_many_argument_regs_get_local_id_012(global uint* out0, global uint* out1, global uint* out2) +{ + uint id0 = get_global_id(0); + uint id1 = get_global_id(1); + uint id2 = get_global_id(2); + uint flat_id = (id2 * get_global_size(1) + id1) * get_global_size(0) + id0; + + uint3 result = too_many_argument_regs_get_local_id_012( + 1234, 999, 42, 5555, 8888, 9009, 777, 4242, + 202020, 6359, 8344, 1443, 552323, 33424, 666, 98765, + 2222, 232556, 57777, 934121, 94991, 1337, 0xdead, 0xbeef, + 0x5555, 0x3333, 0x666, 0x4141, 0x1234, 0x8888, 0xaaaa, 0xbbbb); + + out0[flat_id] = result.x; + out1[flat_id] = result.y; + out2[flat_id] = result.z; +} + + +typedef struct ByValStruct { + long array[9]; +} ByValStruct; + +// Same as previous, with an additional byval passed argument. +NOINLINE +uint3 too_many_argument_regs_byval_get_local_id_012( + ByValStruct byval_arg, + int arg0, int arg1, int arg2, int arg3, + int arg4, int arg5, int arg6, int arg7, + int arg8, int arg9, int arg10, int arg11, + int arg12, int arg13, int arg14, int arg15, + int arg16, int arg17, int arg18, int arg19, + int arg20, int arg21, int arg22, int arg23, + int arg24, int arg25, int arg26, int arg27, + int arg28, int arg29, int arg30, int arg31) +{ + uint3 id; + id.x = get_local_id(0) + byval_arg.array[3]; // + 42 + 3 + id.y = get_local_id(1) + byval_arg.array[5]; // + 42 + 5 + id.z = get_local_id(2) + byval_arg.array[8]; // + 42 + 8 + return id; +} + +kernel void kernel_call_too_many_argument_regs_byval_get_local_id_012(global uint* out0, global uint* out1, global uint* out2) +{ + uint id0 = get_global_id(0); + uint id1 = get_global_id(1); + uint id2 = get_global_id(2); + uint flat_id = (id2 * get_global_size(1) + id1) * get_global_size(0) + id0; + + ByValStruct byval; + for (int i = 0; i < 9; ++i) + byval.array[i] = 42 + i; + + uint3 result = too_many_argument_regs_byval_get_local_id_012( + byval, + 1234, 999, 42, 5555, 8888, 9009, 777, 4242, + 202020, 6359, 8344, 1443, 552323, 33424, 666, 98765, + 2222, 232556, 57777, 934121, 94991, 1337, 0xdead, 0xbeef, + 0x5555, 0x3333, 0x666, 0x4141, 0x1234, 0x8888, 0xaaaa, 0xbbbb); + + out0[flat_id] = result.x; + out1[flat_id] = result.y; + out2[flat_id] = result.z; +} |