ping > On Sep 19, 2017, at 19:51, Matt Arsenault <[email protected]> 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 [email protected] https://lists.freedesktop.org/mailman/listinfo/piglit
