[Piglit] [PATCH] cl: Add test for bswap
arsenm2 at gmail.com
arsenm2 at gmail.com
Thu Feb 13 17:01:27 UTC 2020
From: Matt Arsenault <matt at MattBook.local>
Surprisingly, OpenCL seems to not define a function for this, but I do
want to test codegen for it. The vector shuffle gets idiom recognized
into a bswap.
---
tests/cl/program/execute/bswap.cl | 119 ++++++++++++++++++++++++++++++
1 file changed, 119 insertions(+)
create mode 100644 tests/cl/program/execute/bswap.cl
diff --git a/tests/cl/program/execute/bswap.cl b/tests/cl/program/execute/bswap.cl
new file mode 100644
index 000000000..bc1825cfd
--- /dev/null
+++ b/tests/cl/program/execute/bswap.cl
@@ -0,0 +1,119 @@
+/*!
+
+[config]
+name: bswap
+clc_version_min: 10
+dimensions: 1
+
+[test]
+name: v_bswap_u32
+kernel_name: v_bswap_u32
+global_size: 24 0 0
+
+arg_out: 0 buffer uint[24] \
+ 0x00000000 0xffffffff 0xddccbbaa 0xaabbccdd \
+ 0x01000000 0x00010000 0x00000100 0x00000001 \
+ 0x10000000 0x00100000 0x00001000 0x00000010 \
+ 0x000000ff 0x0000ff00 0x00ff0000 0xff000000 \
+ 0x01020304 0x04030201 0x78563412 0x21436587 \
+ 0x0000ffff 0xffff0000 0x00ffff00 0x0000ffff
+
+arg_in: 1 buffer uint[24] \
+ 0x00000000 0xffffffff 0xaabbccdd 0xddccbbaa \
+ 0x00000001 0x00000100 0x00010000 0x01000000 \
+ 0x00000010 0x00001000 0x00100000 0x10000000 \
+ 0xff000000 0x00ff0000 0x0000ff00 0x000000ff \
+ 0x04030201 0x01020304 0x12345678 0x87654321 \
+ 0xffff0000 0x0000ffff 0x00ffff00 0xffff0000
+
+
+[test]
+name: s_bswap_u32
+kernel_name: s_bswap_u32
+global_size: 1 0 0
+
+arg_out: 0 buffer uint[6] \
+ 0x00000000 0xffffffff 0x01020304 0x04030201 \
+ 0xaabbccdd 0x78563412
+
+arg_in: 1 uint 0x00000000
+arg_in: 2 uint 0xffffffff
+arg_in: 3 uint 0x04030201
+arg_in: 4 uint 0x01020304
+arg_in: 5 uint 0xddccbbaa
+arg_in: 6 uint 0x12345678
+
+
+[test]
+name: v_bswap_u64
+kernel_name: v_bswap_u64
+global_size: 4 0 0
+
+arg_out: 0 buffer ulong[4] \
+ 0x0000000000000000 0xffffffffffffffff 0x8877665544332211 \
+ 0x2131415161718191
+
+arg_in: 1 buffer ulong[4] \
+ 0x0000000000000000 0xffffffffffffffff 0x1122334455667788 \
+ 0x9181716151413121
+
+
+[test]
+name: s_bswap_u64
+kernel_name: s_bswap_u64
+global_size: 1 0 0
+
+arg_out: 0 buffer ulong[4] \
+ 0x0000000000000000 0xffffffffffffffff 0x8877665544332211 \
+ 0x2131415161718191
+
+arg_in: 1 ulong 0x0000000000000000
+arg_in: 2 ulong 0xffffffffffffffff
+arg_in: 3 ulong 0x1122334455667788
+arg_in: 4 ulong 0x9181716151413121
+
+!*/
+
+uint bswap32(uint src) {
+ uchar4 vec = as_uchar4(src);
+ return as_uint(vec.wzyx);
+}
+
+ulong bswap64(ulong src) {
+ uchar8 vec = as_uchar8(src);
+ return as_ulong(vec.s76543210);
+}
+
+kernel void v_bswap_u32(global uint* out, global uint* in)
+{
+ int id = get_global_id(0);
+ out[id] = bswap32(in[id]);
+}
+
+kernel void s_bswap_u32(global uint* out,
+ uint arg0, uint arg1, uint arg2,
+ uint arg3, uint arg4, uint arg5) {
+ uint array[] = { arg0, arg1, arg2, arg3, arg4, arg5 };
+ const int n = sizeof(array) / sizeof(array[0]);
+
+ #pragma unroll
+ for (int i = 0; i < n; ++i)
+ out[i] = bswap32(array[i]);
+}
+
+kernel void v_bswap_u64(global ulong* out, global ulong* in)
+{
+ int id = get_global_id(0);
+ out[id] = bswap64(in[id]);
+}
+
+kernel void s_bswap_u64(global ulong* out,
+ ulong arg0, ulong arg1, ulong arg2,
+ ulong arg3) {
+ ulong array[] = { arg0, arg1, arg2, arg3 };
+ const int n = sizeof(array) / sizeof(array[0]);
+
+ #pragma unroll
+ for (int i = 0; i < n; ++i)
+ out[i] = bswap64(array[i]);
+}
--
2.20.1
More information about the Piglit
mailing list