Also fixes apparently missing coverage for special
input arguments not passed in registers.
---
tests/cl/program/execute/calls-workitem-id.cl | 136 ++++++++++++++++++
1 file changed, 136 insertions(+)
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;
+}
--
2.17.1
_______________________________________________
Piglit mailing list
Piglit@lists.freedesktop.org
https://lists.freedesktop.org/mailman/listinfo/piglit
