[Piglit] [PATCH] cl: Add tests for load hi16 instructions
Jan Vesely
jan.vesely at rutgers.edu
Fri Oct 6 23:22:51 UTC 2017
On Tue, 2017-09-19 at 22:02 -0700, Matt Arsenault wrote:
> v2: Fix some formatting
Reviewed-by: Jan Vesely <jan.vesely at rutgers.edu>
sorry for the delay, feel free to cc me on cl piglit patches. I don't
think anyone else is interested in piglit cl.
Are these targeting specific gcn instructions as well?
Jan
> ---
> tests/cl/program/execute/load-hi16-generic.cl | 96 +++++++++
> tests/cl/program/execute/load-hi16.cl | 280 ++++++++++++++++++++++++++
> 2 files changed, 376 insertions(+)
> create mode 100644 tests/cl/program/execute/load-hi16-generic.cl
> create mode 100644 tests/cl/program/execute/load-hi16.cl
>
> diff --git a/tests/cl/program/execute/load-hi16-generic.cl b/tests/cl/program/execute/load-hi16-generic.cl
> new file mode 100644
> index 000000000..ac1cb7a4d
> --- /dev/null
> +++ b/tests/cl/program/execute/load-hi16-generic.cl
> @@ -0,0 +1,96 @@
> +/*!
> +
> +[config]
> +name: load into high 16-bits of 32-bit register with generic addressing
> +clc_version_min: 20
> +dimensions: 1
> +
> +[test]
> +name: load hi16 generic
> +kernel_name: load_hi16_generic
> +global_size: 4 0 0
> +local_size: 4 0 0
> +
> +arg_out: 0 buffer uint[4] \
> + 0xabcd9999 0x12343333 0x11114444 0xdeadbeef
> +
> +arg_in: 1 buffer ushort[4] \
> + 0x9999 0x3333 0x4444 0xbeef
> +
> +arg_in: 2 buffer ushort[4] \
> + 0xabcd 0x1234 0x1111 0xdead
> +
> +
> +[test]
> +name: zextloadi8 hi16 generic
> +kernel_name: zextloadi8_hi16_generic
> +global_size: 4 0 0
> +local_size: 4 0 0
> +
> +arg_out: 0 buffer uint[4] \
> + 0x00ab0099 0x00120033 0x00110044 0x00de00be
> +
> +arg_in: 1 buffer uchar[4] \
> + 0x99 0x33 0x44 0xbe
> +
> +arg_in: 2 buffer uchar[4] \
> + 0xab 0x12 0x11 0xde
> +
> +
> +[test]
> +name: sextloadi8 hi16 generic
> +kernel_name: sextloadi8_hi16_generic
> +global_size: 4 0 0
> +local_size: 4 0 0
> +
> +arg_out: 0 buffer uint[4] \
> + 0xffabff99 0x00120033 0x00110044 0xffdeffbe
> +
> +arg_in: 1 buffer char[4] \
> + 0x99 0x33 0x44 0xbe
> +
> +arg_in: 2 buffer char[4] \
> + 0xab 0x12 0x11 0xde
> +
> +!*/
> +
> +kernel void load_hi16_generic(volatile global uint* out,
> + volatile global ushort* in0,
> + volatile global ushort* in1)
> +{
> + volatile generic ushort* generic_in0 = (volatile generic ushort*)in0;
> + volatile generic ushort* generic_in1 = (volatile generic ushort*)in1;
> + int gid = get_global_id(0);
> + ushort lo = generic_in0[gid];
> + ushort hi = generic_in1[gid];
> + ushort2 vec = { lo, hi };
> + out[gid] = as_uint(vec);
> +}
> +
> +kernel void zextloadi8_hi16_generic(volatile global uint* out,
> + volatile global uchar* in0,
> + volatile global uchar* in1)
> +{
> + volatile generic uchar* generic_in0 = (volatile generic uchar*)in0;
> + volatile generic uchar* generic_in1 = (volatile generic uchar*)in1;
> +
> + int gid = get_global_id(0);
> + ushort lo = (ushort)generic_in0[gid];
> + ushort hi = (ushort)generic_in1[gid];
> + ushort2 vec = { lo, hi };
> + out[gid] = as_uint(vec);
> +}
> +
> +kernel void sextloadi8_hi16_generic(volatile global uint* out,
> + volatile global char* in0,
> + volatile global char* in1)
> +{
> + volatile generic char* generic_in0 = (volatile generic char*)in0;
> + volatile generic char* generic_in1 = (volatile generic char*)in1;
> +
> + int gid = get_global_id(0);
> + short lo = (short)generic_in0[gid];
> + short hi = (short)generic_in1[gid];
> + short2 vec = { lo, hi };
> + out[gid] = as_uint(vec);
> +}
> diff --git a/tests/cl/program/execute/load-hi16.cl b/tests/cl/program/execute/load-hi16.cl
> new file mode 100644
> index 000000000..f57e0e886
> --- /dev/null
> +++ b/tests/cl/program/execute/load-hi16.cl
> @@ -0,0 +1,280 @@
> +/*!
> +
> +[config]
> + name: load into high 16-bits of 32-bit register
> + clc_version_min: 10
> + dimensions: 1
> +
> +[test]
> + name: load hi16 global
> + kernel_name: load_hi16_global
> + global_size: 4 0 0
> + local_size: 4 0 0
> +
> + arg_out: 0 buffer uint[4] \
> + 0xabcd9999 0x12343333 0x11114444 0xdeadbeef
> +
> + arg_in: 1 buffer ushort[4] \
> + 0x9999 0x3333 0x4444 0xbeef
> +
> + arg_in: 2 buffer ushort[4] \
> + 0xabcd 0x1234 0x1111 0xdead
> +
> +
> +[test]
> + name: load hi16 local
> + kernel_name: load_hi16_local
> + global_size: 4 0 0
> + local_size: 4 0 0
> +
> + arg_out: 0 buffer uint[4] \
> + 0xabcd9999 0x12343333 0x11114444 0xdeadbeef
> +
> + arg_in: 1 buffer ushort[4] \
> + 0x9999 0x3333 0x4444 0xbeef
> +
> + arg_in: 2 buffer ushort[4] \
> + 0xabcd 0x1234 0x1111 0xdead
> +
> +
> +[test]
> + name: load hi16 private
> + kernel_name: load_hi16_private
> + global_size: 4 0 0
> + local_size: 4 0 0
> +
> + arg_out: 0 buffer uint[4] \
> + 0xabcd9999 0x12343333 0x11114444 0xdeadbeef
> +
> + arg_in: 1 buffer ushort[4] \
> + 0x9999 0x3333 0x4444 0xbeef
> +
> + arg_in: 2 buffer ushort[4] \
> + 0xabcd 0x1234 0x1111 0xdead
> +
> +
> +[test]
> + name: zextloadi8 hi16 global
> + kernel_name: zextloadi8_hi16_global
> + global_size: 4 0 0
> + local_size: 4 0 0
> +
> + arg_out: 0 buffer uint[4] \
> + 0x00ab0099 0x00120033 0x00110044 0x00de00be
> +
> + arg_in: 1 buffer uchar[4] \
> + 0x99 0x33 0x44 0xbe
> +
> + arg_in: 2 buffer uchar[4] \
> + 0xab 0x12 0x11 0xde
> +
> +
> +[test]
> + name: sextloadi8 hi16 global
> + kernel_name: sextloadi8_hi16_global
> + global_size: 4 0 0
> + local_size: 4 0 0
> +
> + arg_out: 0 buffer uint[4] \
> + 0xffabff99 0x00120033 0x00110044 0xffdeffbe
> +
> + arg_in: 1 buffer char[4] \
> + 0x99 0x33 0x44 0xbe
> +
> + arg_in: 2 buffer char[4] \
> + 0xab 0x12 0x11 0xde
> +
> +
> +[test]
> + name: zextloadi8 hi16 local
> + kernel_name: zextloadi8_hi16_local
> + global_size: 4 0 0
> + local_size: 4 0 0
> +
> + arg_out: 0 buffer uint[4] \
> + 0x00ab0099 0x00120033 0x00110044 0x00de00be
> +
> + arg_in: 1 buffer uchar[4] \
> + 0x99 0x33 0x44 0xbe
> +
> + arg_in: 2 buffer uchar[4] \
> + 0xab 0x12 0x11 0xde
> +
> +
> +[test]
> + name: sextloadi8 hi16 local
> + kernel_name: sextloadi8_hi16_local
> + global_size: 4 0 0
> + local_size: 4 0 0
> +
> + arg_out: 0 buffer uint[4] \
> + 0xffabff99 0x00120033 0x00110044 0xffdeffbe
> +
> + arg_in: 1 buffer char[4] \
> + 0x99 0x33 0x44 0xbe
> +
> + arg_in: 2 buffer char[4] \
> + 0xab 0x12 0x11 0xde
> +
> +
> +[test]
> + name: zextloadi8 hi16 private
> + kernel_name: zextloadi8_hi16_private
> + global_size: 4 0 0
> + local_size: 4 0 0
> +
> + arg_out: 0 buffer uint[4] \
> + 0x00ab0099 0x00120033 0x00110044 0x00de00be
> +
> + arg_in: 1 buffer uchar[4] \
> + 0x99 0x33 0x44 0xbe
> +
> + arg_in: 2 buffer uchar[4] \
> + 0xab 0x12 0x11 0xde
> +
> +
> +[test]
> + name: sextloadi8 hi16 private
> + kernel_name: sextloadi8_hi16_private
> + global_size: 4 0 0
> + local_size: 4 0 0
> +
> + arg_out: 0 buffer uint[4] \
> + 0xffabff99 0x00120033 0x00110044 0xffdeffbe
> +
> + arg_in: 1 buffer char[4] \
> + 0x99 0x33 0x44 0xbe
> +
> + arg_in: 2 buffer char[4] \
> + 0xab 0x12 0x11 0xde
> +
> +!*/
> +
> +kernel void load_hi16_global(volatile global uint* out,
> + volatile global ushort* in0,
> + volatile global ushort* in1)
> +{
> + int gid = get_global_id(0);
> + ushort lo = in0[gid];
> + ushort hi = in1[gid];
> + ushort2 vec = { lo, hi };
> + out[gid] = as_uint(vec);
> +}
> +
> +kernel void load_hi16_local(volatile global uint* out,
> + volatile global ushort* in0,
> + volatile global ushort* in1)
> +{
> + int lid = get_local_id(0);
> + int gid = get_global_id(0);
> +
> + volatile local ushort lds0[64];
> + volatile local ushort lds1[64];
> +
> + lds0[lid] = in0[gid];
> + lds1[lid] = in1[gid];
> +
> + ushort lo = lds0[lid];
> + ushort hi = lds1[lid];
> + ushort2 vec = { lo, hi };
> + out[gid] = as_uint(vec);
> +}
> +
> +kernel void load_hi16_private(volatile global uint* out,
> + volatile global ushort* in0,
> + volatile global ushort* in1)
> +{
> + int lid = get_local_id(0);
> + int gid = get_global_id(0);
> +
> + volatile private ushort stack0 = in0[gid];
> + volatile private ushort stack1 = in1[gid];
> +
> + ushort2 vec = { stack0, stack1 };
> + out[gid] = as_uint(vec);
> +}
> +
> +kernel void zextloadi8_hi16_global(volatile global uint* out,
> + volatile global uchar* in0,
> + volatile global uchar* in1)
> +{
> + int gid = get_global_id(0);
> + ushort lo = (ushort)in0[gid];
> + ushort hi = (ushort)in1[gid];
> + ushort2 vec = { lo, hi };
> + out[gid] = as_uint(vec);
> +}
> +
> +kernel void sextloadi8_hi16_global(volatile global uint* out,
> + volatile global char* in0,
> + volatile global char* in1)
> +{
> + int gid = get_global_id(0);
> + short lo = (short)in0[gid];
> + short hi = (short)in1[gid];
> + short2 vec = { lo, hi };
> + out[gid] = as_uint(vec);
> +}
> +
> +kernel void zextloadi8_hi16_local(volatile global uint* out,
> + volatile global uchar* in0,
> + volatile global uchar* in1)
> +{
> + int lid = get_local_id(0);
> + int gid = get_global_id(0);
> +
> + volatile local uchar lds0[64];
> + volatile local uchar lds1[64];
> +
> + lds0[lid] = in0[gid];
> + lds1[lid] = in1[gid];
> +
> + ushort lo = lds0[lid];
> + ushort hi = lds1[lid];
> + ushort2 vec = { lo, hi };
> + out[gid] = as_uint(vec);
> +}
> +
> +kernel void sextloadi8_hi16_local(volatile global uint* out,
> + volatile global char* in0,
> + volatile global char* in1)
> +{
> + int lid = get_local_id(0);
> + int gid = get_global_id(0);
> +
> + volatile local char lds0[64];
> + volatile local char lds1[64];
> +
> + lds0[lid] = in0[gid];
> + lds1[lid] = in1[gid];
> +
> + short lo = lds0[lid];
> + short hi = lds1[lid];
> + ushort2 vec = { lo, hi };
> + out[gid] = as_uint(vec);
> +}
> +
> +kernel void zextloadi8_hi16_private(volatile global uint* out,
> + volatile global uchar* in0,
> + volatile global uchar* in1)
> +{
> + int gid = get_global_id(0);
> +
> + volatile private uchar stack0 = in0[gid];
> + volatile private uchar stack1 = in1[gid];
> +
> + ushort2 vec = { (ushort)stack0, (ushort)stack1 };
> + out[gid] = as_uint(vec);
> +}
> +
> +kernel void sextloadi8_hi16_private(volatile global uint* out,
> + volatile global char* in0,
> + volatile global char* in1)
> +{
> + int gid = get_global_id(0);
> + volatile private char stack0 = in0[gid];
> + volatile private char stack1 = in1[gid];
> +
> + ushort2 vec = { (short)stack0, (short)stack1 };
> + out[gid] = as_uint(vec);
> +}
-------------- next part --------------
A non-text attachment was scrubbed...
Name: signature.asc
Type: application/pgp-signature
Size: 833 bytes
Desc: This is a digitally signed message part
URL: <https://lists.freedesktop.org/archives/piglit/attachments/20171006/d2da7c4e/attachment-0001.sig>
More information about the Piglit
mailing list