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

[Matt Arsenault]

---
 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