[Piglit] [PATCH] cl: Add test for clz optimizations
Jan Vesely
jan.vesely at rutgers.edu
Tue Dec 6 23:48:37 UTC 2016
On Tue, 2016-12-06 at 11:15 -0800, arsenm2 at gmail.com wrote:
> From: Matt Arsenault <arsenm2 at gmail.com>
>
> v2: Rename test
Reviewed-by: Jan Vesely <jan.vesely at rutgers.edu>
and pushed (with slight rename followup)
Jan
> ---
> tests/cl/program/execute/clz-optimizations.cl | 389 ++++++++++++++++++++++++++
> 1 file changed, 389 insertions(+)
> create mode 100644 tests/cl/program/execute/clz-optimizations.cl
>
> diff --git a/tests/cl/program/execute/clz-optimizations.cl b/tests/cl/program/execute/clz-optimizations.cl
> new file mode 100644
> index 0000000..a09e79b
> --- /dev/null
> +++ b/tests/cl/program/execute/clz-optimizations.cl
> @@ -0,0 +1,389 @@
> +/*!
> +
> +[config]
> +name: clz
> +clc_version_min: 10
> +
> +dimensions: 1
> +global_size: 34 0 0
> +
> +## clz ##
> +
> +[test]
> +name: v_clz_u32
> +kernel_name: v_clz_u32
> +
> +arg_out: 0 buffer int[34] \
> + 32 31 30 29 28 \
> + 27 26 25 24 23 \
> + 22 21 20 19 18 \
> + 17 16 15 14 13 \
> + 12 11 10 9 8 \
> + 7 6 5 4 3 \
> + 2 1 0 0
> +
> +arg_in: 1 buffer uint[34] \
> + 0 1 2 4 8 \
> + 16 32 64 128 256 \
> + 512 1024 2048 4096 8192 \
> + 16384 32768 65536 131072 262144 \
> + 524288 1048576 2097152 4194304 8388608 \
> + 16777216 33554432 67108864 134217728 268435456 \
> + 536870912 1073741824 2147483648 4294967295
> +
> +
> +[test]
> +name: s_clz_u32
> +kernel_name: s_clz_u32
> +global_size: 6 0 0
> +
> +arg_out: 0 buffer uint[6] \
> + 42 41 40 11 10 10
> +
> +arg_in: 1 uint 0
> +arg_in: 2 uint 1
> +arg_in: 3 uint 2
> +arg_in: 4 uint 1073741824
> +arg_in: 5 uint 2147483648
> +arg_in: 6 uint 4294967295
> +
> +[test]
> +name: s_firstbit_u32
> +kernel_name: s_firstbit_u32
> +global_size: 6 0 0
> +
> +arg_out: 0 buffer uint[6] \
> + 9 41 40 11 10 10
> +
> +arg_in: 1 uint 0
> +arg_in: 2 uint 1
> +arg_in: 3 uint 2
> +arg_in: 4 uint 1073741824
> +arg_in: 5 uint 2147483648
> +arg_in: 6 uint 4294967295
> +
> +
> +[test]
> +name: s_firstbit_cmp_result_u32
> +kernel_name: s_firstbit_cmp_result_u32
> +global_size: 6 0 0
> +
> +arg_out: 0 buffer uint[6] \
> + 9 41 40 11 10 10
> +
> +arg_in: 1 uint 0
> +arg_in: 2 uint 1
> +arg_in: 3 uint 2
> +arg_in: 4 uint 1073741824
> +arg_in: 5 uint 2147483648
> +arg_in: 6 uint 4294967295
> +
> +
> +[test]
> +name: v_firstbit_u32
> +kernel_name: v_firstbit_u32
> +
> +arg_out: 0 buffer int[34] \
> + 0xffffffff 31 30 29 28 \
> + 27 26 25 24 23 \
> + 22 21 20 19 18 \
> + 17 16 15 14 13 \
> + 12 11 10 9 8 \
> + 7 6 5 4 3 \
> + 2 1 0 0
> +
> +arg_in: 1 buffer uint[34] \
> + 0 1 2 4 8 \
> + 16 32 64 128 256 \
> + 512 1024 2048 4096 8192 \
> + 16384 32768 65536 131072 262144 \
> + 524288 1048576 2097152 4194304 8388608 \
> + 16777216 33554432 67108864 134217728 268435456 \
> + 536870912 1073741824 2147483648 4294967295
> +
> +
> +[test]
> +name: v_firstbit_u32_inv
> +kernel_name: v_firstbit_u32_inv
> +
> +arg_out: 0 buffer int[34] \
> + 0xffffffff 31 30 29 28 \
> + 27 26 25 24 23 \
> + 22 21 20 19 18 \
> + 17 16 15 14 13 \
> + 12 11 10 9 8 \
> + 7 6 5 4 3 \
> + 2 1 0 0
> +
> +arg_in: 1 buffer uint[34] \
> + 0 1 2 4 8 \
> + 16 32 64 128 256 \
> + 512 1024 2048 4096 8192 \
> + 16384 32768 65536 131072 262144 \
> + 524288 1048576 2097152 4194304 8388608 \
> + 16777216 33554432 67108864 134217728 268435456 \
> + 536870912 1073741824 2147483648 4294967295
> +
> +[test]
> +name: v_firstbit_u32_cmp_result
> +kernel_name: v_firstbit_u32_cmp_result
> +
> +arg_out: 0 buffer int[34] \
> + 0xffffffff 31 30 29 28 \
> + 27 26 25 24 23 \
> + 22 21 20 19 18 \
> + 17 16 15 14 13 \
> + 12 11 10 9 8 \
> + 7 6 5 4 3 \
> + 2 1 0 0
> +
> +arg_in: 1 buffer uint[34] \
> + 0 1 2 4 8 \
> + 16 32 64 128 256 \
> + 512 1024 2048 4096 8192 \
> + 16384 32768 65536 131072 262144 \
> + 524288 1048576 2097152 4194304 8388608 \
> + 16777216 33554432 67108864 134217728 268435456 \
> + 536870912 1073741824 2147483648 4294967295
> +
> +[test]
> +name: v_firstbit_u32_cmp_result_inv
> +kernel_name: v_firstbit_u32_cmp_result_inv
> +
> +arg_out: 0 buffer int[34] \
> + 0xffffffff 31 30 29 28 \
> + 27 26 25 24 23 \
> + 22 21 20 19 18 \
> + 17 16 15 14 13 \
> + 12 11 10 9 8 \
> + 7 6 5 4 3 \
> + 2 1 0 0
> +
> +arg_in: 1 buffer uint[34] \
> + 0 1 2 4 8 \
> + 16 32 64 128 256 \
> + 512 1024 2048 4096 8192 \
> + 16384 32768 65536 131072 262144 \
> + 524288 1048576 2097152 4194304 8388608 \
> + 16777216 33554432 67108864 134217728 268435456 \
> + 536870912 1073741824 2147483648 4294967295
> +
> +
> +[test]
> +name: v_clz_u64
> +kernel_name: v_clz_u64
> +global_size: 66 0 0
> +
> +arg_out: 0 buffer int[66] \
> + 64 63 62 61 \
> + 60 59 58 57 \
> + 56 55 54 53 \
> + 52 51 50 49 \
> + 48 47 46 45 \
> + 44 43 42 41 \
> + 40 39 38 37 \
> + 36 35 34 33 \
> + 32 31 30 29 \
> + 28 27 26 25 \
> + 24 23 22 21 \
> + 20 19 18 17 \
> + 16 15 14 13 \
> + 12 11 10 9 \
> + 8 7 6 5 \
> + 4 3 2 1 \
> + 0 0
> +
> +arg_in: 1 buffer ulong[66] \
> + 0x0 0x1 0x2 0x4 \
> + 0x8 0x10 0x20 0x40 \
> + 0x80 0x100 0x200 0x400 \
> + 0x800 0x1000 0x2000 0x4000 \
> + 0x8000 0x10000 0x20000 0x40000 \
> + 0x80000 0x100000 0x200000 0x400000 \
> + 0x800000 0x1000000 0x2000000 0x4000000 \
> + 0x8000000 0x10000000 0x20000000 0x40000000 \
> + 0x80000000 0x100000000 0x200000000 0x400000000 \
> + 0x800000000 0x1000000000 0x2000000000 0x4000000000 \
> + 0x8000000000 0x10000000000 0x20000000000 0x40000000000 \
> + 0x80000000000 0x100000000000 0x200000000000 0x400000000000 \
> + 0x800000000000 0x1000000000000 0x2000000000000 0x4000000000000 \
> + 0x8000000000000 0x10000000000000 0x20000000000000 0x40000000000000 \
> + 0x80000000000000 0x100000000000000 0x200000000000000 0x400000000000000 \
> + 0x800000000000000 0x1000000000000000 0x2000000000000000 0x4000000000000000 \
> + 0x8000000000000000 \
> + 0xffffffffffffffff
> +
> +[test]
> +name: v_firstbit_u64
> +kernel_name: v_firstbit_u64
> +global_size: 66 0 0
> +
> +arg_out: 0 buffer int[66] \
> + -1 63 62 61 \
> + 60 59 58 57 \
> + 56 55 54 53 \
> + 52 51 50 49 \
> + 48 47 46 45 \
> + 44 43 42 41 \
> + 40 39 38 37 \
> + 36 35 34 33 \
> + 32 31 30 29 \
> + 28 27 26 25 \
> + 24 23 22 21 \
> + 20 19 18 17 \
> + 16 15 14 13 \
> + 12 11 10 9 \
> + 8 7 6 5 \
> + 4 3 2 1 \
> + 0 0
> +
> +arg_in: 1 buffer ulong[66] \
> + 0x0 0x1 0x2 0x4 \
> + 0x8 0x10 0x20 0x40 \
> + 0x80 0x100 0x200 0x400 \
> + 0x800 0x1000 0x2000 0x4000 \
> + 0x8000 0x10000 0x20000 0x40000 \
> + 0x80000 0x100000 0x200000 0x400000 \
> + 0x800000 0x1000000 0x2000000 0x4000000 \
> + 0x8000000 0x10000000 0x20000000 0x40000000 \
> + 0x80000000 0x100000000 0x200000000 0x400000000 \
> + 0x800000000 0x1000000000 0x2000000000 0x4000000000 \
> + 0x8000000000 0x10000000000 0x20000000000 0x40000000000 \
> + 0x80000000000 0x100000000000 0x200000000000 0x400000000000 \
> + 0x800000000000 0x1000000000000 0x2000000000000 0x4000000000000 \
> + 0x8000000000000 0x10000000000000 0x20000000000000 0x40000000000000 \
> + 0x80000000000000 0x100000000000000 0x200000000000000 0x400000000000000 \
> + 0x800000000000000 0x1000000000000000 0x2000000000000000 0x4000000000000000 \
> + 0x8000000000000000 \
> + 0xffffffffffffffff
> +
> +
> +[test]
> +name: v_clz_u16
> +kernel_name: v_clz_u16
> +
> +arg_out: 0 buffer int[18] \
> + 16 15 14 13 12 \
> + 11 10 9 8 7 \
> + 6 5 4 3 2 \
> + 1 0 0
> +
> +arg_in: 1 buffer ushort[18] \
> + 0 1 2 4 8 \
> + 16 32 64 128 256 \
> + 512 1024 2048 4096 8192 \
> + 16384 32768 0xffff
> +
> +[test]
> +name: v_firstbit_u16
> +kernel_name: v_firstbit_u16
> +
> +arg_out: 0 buffer int[18] \
> + -1 15 14 13 12 \
> + 11 10 9 8 7 \
> + 6 5 4 3 2 \
> + 1 0 0
> +
> +arg_in: 1 buffer ushort[18] \
> + 0 1 2 4 8 \
> + 16 32 64 128 256 \
> + 512 1024 2048 4096 8192 \
> + 16384 32768 0xffff
> +
> +!*/
> +
> +kernel void v_clz_u32(global int* out, global uint* in)
> +{
> + int id = get_global_id(0);
> + out[id] = clz(in[id]);
> +}
> +
> +kernel void v_clz_u16(global int* out, global ushort* in)
> +{
> + int id = get_global_id(0);
> + out[id] = clz(in[id]);
> +}
> +
> +kernel void s_clz_u32(global uint* out,
> + uint arg0, uint arg1, uint arg2,
> + uint arg3, uint arg4, uint arg5) {
> + uint array[] = { arg0, arg1, arg2, arg3, arg4, arg5 };
> +
> +#pragma unroll
> + for (int i = 0; i < sizeof(array) / sizeof(array[0]); ++i)
> + out[i] = clz(array[i]) + 10;
> +}
> +
> +kernel void v_firstbit_u32(global uint* out, global uint* in)
> +{
> + uint id = get_global_id(0);
> + uint val = in[id];
> + out[id] = (val == 0) ? -1 : clz(val);
> +}
> +
> +kernel void v_firstbit_u16(global uint* out, global ushort* in)
> +{
> + uint id = get_global_id(0);
> + ushort val = in[id];
> + out[id] = (val == 0) ? -1 : clz(val);
> +}
> +
> +kernel void v_firstbit_u32_inv(global uint* out, global uint* in)
> +{
> + uint id = get_global_id(0);
> + uint val = in[id];
> + out[id] = (val != 0) ? clz(val) : -1;
> +}
> +
> +kernel void v_firstbit_u32_cmp_result(global uint* out,
> + global uint* in)
> +{
> + uint id = get_global_id(0);
> + uint res = clz(in[id]);
> + out[id] = (res == 32) ? -1 : res;
> +}
> +
> +kernel void v_firstbit_u32_cmp_result_inv(global uint* out,
> + global uint* in)
> +{
> + uint id = get_global_id(0);
> + uint res = clz(in[id]);
> + out[id] = (res != 32) ? res : -1;
> +}
> +
> +kernel void s_firstbit_u32(global uint* out,
> + uint arg0, uint arg1, uint arg2,
> + uint arg3, uint arg4, uint arg5) {
> + uint array[] = { arg0, arg1, arg2, arg3, arg4, arg5 };
> +
> +#pragma unroll
> + for (int i = 0; i < sizeof(array) / sizeof(array[0]); ++i) {
> + int val = array[i];
> + out[i] = ((val == 0) ? -1 : clz(val)) + 10;
> + }
> +}
> +
> +kernel void s_firstbit_cmp_result_u32(global uint* out,
> + uint arg0, uint arg1, uint arg2,
> + uint arg3, uint arg4, uint arg5) {
> + uint array[] = { arg0, arg1, arg2, arg3, arg4, arg5 };
> +
> +#pragma unroll
> + for (int i = 0; i < sizeof(array) / sizeof(array[0]); ++i) {
> + int val = clz(array[i]);
> + out[i] = ((val == 32) ? -1 : val) + 10;
> + }
> +}
> +
> +kernel void v_clz_u64(global int* out, global ulong* in)
> +{
> + int id = get_global_id(0);
> + out[id] = clz(in[id]);
> +}
> +
> +kernel void v_firstbit_u64(global uint* out, global ulong* in)
> +{
> + uint id = get_global_id(0);
> + ulong val = in[id];
> + out[id] = (val == 0) ? -1 : clz(val);
> +}
-------------- next part --------------
A non-text attachment was scrubbed...
Name: signature.asc
Type: application/pgp-signature
Size: 819 bytes
Desc: This is a digitally signed message part
URL: <https://lists.freedesktop.org/archives/piglit/attachments/20161206/a2262626/attachment-0001.sig>
More information about the Piglit
mailing list