On Mon, 2017-10-02 at 10:32 -0700, Matt Arsenault wrote: > ping Reviewed-by: Jan Vesely <jan.ves...@rutgers.edu>
Jan > > > On Sep 19, 2017, at 19:51, Matt Arsenault <arse...@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@lists.freedesktop.org > https://lists.freedesktop.org/mailman/listinfo/piglit
signature.asc
Description: This is a digitally signed message part
_______________________________________________ Piglit mailing list Piglit@lists.freedesktop.org https://lists.freedesktop.org/mailman/listinfo/piglit