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

[Jan Vesely]

On Mon, 2017-10-02 at 10:32 -0700, Matt Arsenault wrote:
> ping

Reviewed-by: Jan Vesely <jan.vesely at rutgers.edu>

Jan

> 
> > 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
> > 
> 
> _______________________________________________
> Piglit mailing list
> Piglit at lists.freedesktop.org
> https://lists.freedesktop.org/mailman/listinfo/piglit
-------------- next part --------------
A non-text attachment was scrubbed...
Name: signature.asc
Type: application/pgp-signature
Size: 833 bytes
Desc: This is a digitally signed message part
URL: <https://lists.freedesktop.org/archives/piglit/attachments/20171006/4ddd1d45/attachment.sig>