[Piglit] [PATCH] cl: Add tests for store hi16 instructions

[Matt Arsenault]

ping

> On Sep 19, 2017, at 19:51, Matt Arsenault <arsenm2 at gmail.com> wrote:
> 
> ---
> tests/cl/program/execute/store-hi16-generic.cl |  51 ++++++++++
> tests/cl/program/execute/store-hi16.cl         | 134 +++++++++++++++++++++++++
> 2 files changed, 185 insertions(+)
> create mode 100644 tests/cl/program/execute/store-hi16-generic.cl
> create mode 100644 tests/cl/program/execute/store-hi16.cl
> 
> diff --git a/tests/cl/program/execute/store-hi16-generic.cl b/tests/cl/program/execute/store-hi16-generic.cl
> new file mode 100644
> index 000000000..807818971
> --- /dev/null
> +++ b/tests/cl/program/execute/store-hi16-generic.cl
> @@ -0,0 +1,51 @@
> +/*!
> +
> +[config]
> +name: store high 16-bits of 32-bit value with generic addressing
> +clc_version_min: 20
> +dimensions: 1
> +
> +[test]
> +name: store hi16 generic
> +kernel_name: store_hi16_generic
> +global_size: 4 0 0
> +local_size: 4 0 0
> +
> +arg_out: 0 buffer ushort[4] \
> +    0xabcd    0x1111   0x2222 0x3333
> +
> +arg_in: 1 buffer uint[4] \
> +   0xabcd1234    0x11112222    0x22221111    0x33334444
> +
> +
> +[test]
> +name: store hi16 generic trunc i8
> +kernel_name: truncstorei8_hi16_generic
> +global_size: 4 0 0
> +local_size: 4 0 0
> +
> +arg_out: 0 buffer uchar[4] \
> +  0xcd 0x22 0xad 0x80
> +
> +arg_in: 1 buffer uint[4] \
> +   0xabcd1234    0x11223344    0xdeadbeef  0x70809024
> +
> +!*/
> +
> +kernel void store_hi16_generic(volatile global ushort* out, volatile global uint* in)
> +{
> +    int gid = get_global_id(0);
> +    uint value = in[gid];
> +
> +    volatile generic ushort* generic_out = (volatile generic ushort*)out;
> +    generic_out[gid] = value >> 16;
> +}
> +
> +kernel void truncstorei8_hi16_generic(volatile global uchar* out, volatile global uint* in)
> +{
> +    int gid = get_global_id(0);
> +    uint value = in[gid];
> +
> +    volatile generic uchar* generic_out = (volatile generic ushort*)out;
> +    generic_out[gid] = (uchar)(value >> 16);
> +}
> diff --git a/tests/cl/program/execute/store-hi16.cl b/tests/cl/program/execute/store-hi16.cl
> new file mode 100644
> index 000000000..b734b3766
> --- /dev/null
> +++ b/tests/cl/program/execute/store-hi16.cl
> @@ -0,0 +1,134 @@
> +/*!
> +
> +[config]
> +name: store high 16-bits of 32-bit value
> +clc_version_min: 10
> +
> +dimensions: 1
> +
> +[test]
> +name: store hi16 global
> +kernel_name: store_hi16_global
> +global_size: 4 0 0
> +local_size: 4 0 0
> +
> +arg_out: 0 buffer ushort[4] \
> +    0xabcd    0x1111   0x2222 0x3333
> +
> +arg_in: 1 buffer uint[4] \
> +   0xabcd1234    0x11112222    0x22221111    0x33334444
> +
> +[test]
> +name: store hi16 local
> +kernel_name: store_hi16_local
> +global_size: 4 0 0
> +local_size: 4 0 0
> +
> +arg_out: 0 buffer ushort[4] \
> +  0xabcd    0x1111   0x2222 0x3333
> +
> +arg_in: 1 buffer uint[4] \
> +   0xabcd1234    0x11112222    0x22221111    0x33334444
> +
> +[test]
> +name: store hi16 private
> +kernel_name: store_hi16_private
> +global_size: 4 0 0
> +local_size: 4 0 0
> +
> +arg_out: 0 buffer ushort[4] \
> +  0xabcd    0x1111   0x2222 0x3333
> +
> +arg_in: 1 buffer uint[4] \
> +   0xabcd1234    0x11112222    0x22221111    0x33334444
> +
> +
> +[test]
> +name: store hi16 global trunc i8
> +kernel_name: truncstorei8_hi16_global
> +global_size: 4 0 0
> +local_size: 4 0 0
> +
> +arg_out: 0 buffer uchar[4] \
> +  0xcd 0x22 0xad 0x80
> +
> +arg_in: 1 buffer uint[4] \
> +   0xabcd1234    0x11223344    0xdeadbeef  0x70809024
> +
> +
> +[test]
> +name: store hi16 local trunc i8
> +kernel_name: truncstorei8_hi16_local
> +global_size: 4 0 0
> +local_size: 4 0 0
> +
> +arg_out: 0 buffer uchar[4] \
> +  0xcd 0x22 0xad 0x80
> +
> +arg_in: 1 buffer uint[4] \
> +   0xabcd1234    0x11223344    0xdeadbeef  0x70809024
> +
> +
> +[test]
> +name: store hi16 private trunc i8
> +kernel_name: truncstorei8_hi16_private
> +global_size: 4 0 0
> +local_size: 4 0 0
> +
> +arg_out: 0 buffer uchar[4] \
> +  0xcd 0x22 0xad 0x80
> +
> +arg_in: 1 buffer uint[4] \
> +   0xabcd1234    0x11223344    0xdeadbeef  0x70809024
> +
> +!*/
> +
> +kernel void store_hi16_global(volatile global ushort* out, volatile global uint* in)
> +{
> +    int gid = get_global_id(0);
> +    uint value = in[gid];
> +    out[gid] = value >> 16;
> +}
> +
> +kernel void store_hi16_local(volatile global ushort* out, volatile global uint* in)
> +{
> +    volatile local short lds[64];
> +    int lid = get_local_id(0);
> +    int gid = get_global_id(0);
> +
> +    uint value = in[gid];
> +    lds[lid] = value >> 16;
> +    out[gid] = lds[lid];
> +}
> +
> +kernel void store_hi16_private(volatile global ushort* out, volatile global uint* in)
> +{
> +    int gid = get_global_id(0);
> +    volatile private short stack = in[gid] >> 16;
> +    out[gid] = stack;
> +}
> +
> +kernel void truncstorei8_hi16_global(volatile global uchar* out, volatile global uint* in)
> +{
> +    int gid = get_global_id(0);
> +    uint value = in[gid];
> +    out[gid] = (uchar)(value >> 16);
> +}
> +
> +kernel void truncstorei8_hi16_local(volatile global uchar* out, volatile global uint* in)
> +{
> +    volatile local short lds[64];
> +    int lid = get_local_id(0);
> +    int gid = get_global_id(0);
> +
> +    uint value = in[gid];
> +    lds[lid] = value >> 16;
> +    out[gid] = (uchar)lds[lid];
> +}
> +
> +kernel void truncstorei8_hi16_private(volatile global uchar* out, volatile global uint* in)
> +{
> +    int gid = get_global_id(0);
> +    volatile private short stack = in[gid] >> 16;
> +    out[gid] = (uchar)stack;
> +}
> -- 
> 2.11.0
>