diff options
author | Matt Arsenault <arsenm2@gmail.com> | 2017-09-19 19:51:56 -0700 |
---|---|---|
committer | Jan Vesely <jan.vesely@rutgers.edu> | 2017-10-12 16:49:26 -0400 |
commit | 8f2ec26445daf4c8da2e994afb4f46f46b44c117 (patch) | |
tree | 9053e4ef5fcd52087f6a73fb1018c8aaf13d0237 | |
parent | 28a77cb84dcec49b55c82663615f052efee2bc67 (diff) |
cl: Add tests for store hi16 instructions
Reviewed-by: Jan Vesely <jan.vesely@rutgers.edu>
-rw-r--r-- | tests/cl/program/execute/store-hi16-generic.cl | 51 | ||||
-rw-r--r-- | tests/cl/program/execute/store-hi16.cl | 134 |
2 files changed, 185 insertions, 0 deletions
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; +} |