[Piglit] [PATCH] CL/math: Add frexp builtin tests

Aaron Watry awatry at gmail.com
Mon Feb 1 20:41:47 PST 2016


Add some non-generated float/double tests for the frexp builtin function.

float frexp(float x, int* exp);

The tests are in separate files so that devices that don't support doubles
can still have the float variant tested.

Because the function signature is so unique, I elected to not auto-generate
tests for now.

The double tests don't handle inf/nan, as I couldn't get piglit's CL tester
to interpret those values correctly... Advice welcome.

Signed-off-by: Aaron Watry <awatry at gmail.com>
---
 tests/cl/program/execute/builtin/math/frexp.cl     | 98 ++++++++++++++++++++++
 .../program/execute/builtin/math/frexp_double.cl   | 95 +++++++++++++++++++++
 2 files changed, 193 insertions(+)
 create mode 100644 tests/cl/program/execute/builtin/math/frexp.cl
 create mode 100644 tests/cl/program/execute/builtin/math/frexp_double.cl

diff --git a/tests/cl/program/execute/builtin/math/frexp.cl b/tests/cl/program/execute/builtin/math/frexp.cl
new file mode 100644
index 0000000..66b0bdf
--- /dev/null
+++ b/tests/cl/program/execute/builtin/math/frexp.cl
@@ -0,0 +1,98 @@
+/*!
+[config]
+name: frexp
+clc_version_min: 10
+dimensions: 1
+
+#import struct
+#def float_to_hex(f):
+#    return hex(struct.unpack('<I', struct.pack('<f', f))[0])
+#
+
+[test]
+name: float scalar simple
+kernel_name: frexp_float_scalar
+global_size: 1 0 0
+arg_out: 0 buffer float[1] 0.602783203125
+arg_in: 1 buffer float[1] 1234.5
+arg_out: 2 buffer int[1] 11
+
+[test]
+name: float scalar
+kernel_name: frexp_float_scalar
+global_size: 7 0 0
+arg_out: 0 buffer float[7] 0.602783203125 0.5 nan nan  inf -inf 0.0
+arg_in: 1 buffer float[7]  1234.5         1.0 nan -nan inf -inf 0.0
+arg_out: 2 buffer int[7]   11             1   0   0    0   0    0
+
+[test]
+name: float vector 2 simple
+kernel_name: frexp_float_vec2
+global_size: 1 0 0
+arg_out: 0 buffer float[2] repeat 0.602783203125
+arg_in: 1 buffer float[2] repeat 1234.5
+arg_out: 2 buffer int[2] repeat 11
+
+[test]
+name: float vector 3 complex
+kernel_name: frexp_float_vec3
+global_size: 1 0 0
+arg_out: 0 buffer float[3] 0.602783203125 0.0 0.5
+arg_in: 1 buffer float[3] 1234.5 0.0 1.0
+arg_out: 2 buffer int[3] 11 0 1
+
+[test]
+name: float vector 4 complex
+kernel_name: frexp_float_vec4
+global_size: 1 0 0
+arg_out: 0 buffer float[4] 0.602783203125 0.0 0.5 nan
+arg_in: 1 buffer float[4] 1234.5 0.0 1.0 nan
+arg_out: 2 buffer int[4] 11 0 1 0
+
+[test]
+name: float vector 8 complex
+kernel_name: frexp_float_vec8
+global_size: 1 0 0
+arg_out: 0 buffer float[8] 0.602783203125 0.0 0.5 nan 0.5 nan 0.602783203125 0.0
+arg_in: 1 buffer float[8] 1234.5 0.0 1.0 nan 1.0 nan 1234.5 0.0
+arg_out: 2 buffer int[8] 11 0 1 0 1 0 11 0
+
+[test]
+name: float vector 16 simple
+kernel_name: frexp_float_vec16
+global_size: 1 0 0
+arg_out: 0 buffer float[16] repeat 0.602783203125
+arg_in: 1 buffer float[16] repeat 1234.5
+arg_out: 2 buffer int[16] repeat 11
+
+!*/
+
+kernel void frexp_float_scalar(global float* out, global float* in1, global int* out2) {
+    size_t id = get_global_id(0);
+    out[id] = frexp(in1[id], &out2[id]);
+}
+
+kernel void frexp_float_vec2(global float2* out, global float2* in1, global int2* out2) {
+    size_t id = get_global_id(0);
+    out[id] = frexp(in1[id], &out2[id]);
+}
+
+kernel void frexp_float_vec3(global float3* out, global float3* in1, global int3* out2) {
+    size_t id = get_global_id(0);
+    out[id] = frexp(in1[id], &out2[id]);
+}
+
+kernel void frexp_float_vec4(global float4* out, global float4* in1, global int4* out2) {
+    size_t id = get_global_id(0);
+    out[id] = frexp(in1[id], &out2[id]);
+}
+
+kernel void frexp_float_vec8(global float8* out, global float8* in1, global int8* out2) {
+    size_t id = get_global_id(0);
+    out[id] = frexp(in1[id], &out2[id]);
+}
+
+kernel void frexp_float_vec16(global float16* out, global float16* in1, global int16* out2) {
+    size_t id = get_global_id(0);
+    out[id] = frexp(in1[id], &out2[id]);
+}
diff --git a/tests/cl/program/execute/builtin/math/frexp_double.cl b/tests/cl/program/execute/builtin/math/frexp_double.cl
new file mode 100644
index 0000000..2a3f458
--- /dev/null
+++ b/tests/cl/program/execute/builtin/math/frexp_double.cl
@@ -0,0 +1,95 @@
+#pragma OPENCL EXTENSION cl_khr_fp64: enable
+/*!
+[config]
+name: frexp
+clc_version_min: 10
+dimensions: 1
+require_device_extensions: cl_khr_fp64
+
+[test]
+name: double scalar simple
+kernel_name: frexp_double_scalar
+global_size: 1 0 0
+arg_out: 0 buffer double[1] 0.602783203125
+arg_in: 1 buffer double[1] 1234.5
+arg_out: 2 buffer int[1] 11
+
+[test]
+name: double scalar
+kernel_name: frexp_double_scalar
+global_size: 7 0 0
+arg_out: 0 buffer double[7] 0.602783203125 0.5 nan nan  inf -inf 0.0
+arg_in: 1 buffer double[7]  1234.5         1.0 nan -nan inf -inf 0.0
+arg_out: 2 buffer int[7]    11             1   0   0    0   0    0
+
+[test]
+name: double vector 2 simple
+kernel_name: frexp_double_vec2
+global_size: 1 0 0
+arg_out: 0 buffer double[2] repeat 0.602783203125
+arg_in: 1 buffer double[2] repeat 1234.5
+arg_out: 2 buffer int[2] repeat 11
+
+[test]
+name: double vector 3 complex
+kernel_name: frexp_double_vec3
+global_size: 1 0 0
+arg_out: 0 buffer double[3] 0.602783203125 0.0 0.5
+arg_in: 1 buffer double[3] 1234.5 0.0 1.0
+arg_out: 2 buffer int[3] 11 0 1
+
+[test]
+name: double vector 4 complex
+kernel_name: frexp_double_vec4
+global_size: 1 0 0
+arg_out: 0 buffer double[4] 0.602783203125 0.0 0.5 0.512500
+arg_in: 1 buffer double[4] 1234.5 0.0 1.0 16.4
+arg_out: 2 buffer int[4] 11 0 1 5
+
+[test]
+name: double vector 8 complex
+kernel_name: frexp_double_vec8
+global_size: 1 0 0
+arg_out: 0 buffer double[8] 0.602783203125 0.0 0.5 0.512500 0.5 0.512500 0.602783203125 0.0
+arg_in: 1 buffer double[8] 1234.5 0.0 1.0 16.4 1.0 16.4 1234.5 0.0
+arg_out: 2 buffer int[8] 11 0 1 5 1 5 11 0
+
+[test]
+name: double vector 16 simple
+kernel_name: frexp_double_vec16
+global_size: 1 0 0
+arg_out: 0 buffer double[16] repeat 0.602783203125
+arg_in: 1 buffer double[16] repeat 1234.5
+arg_out: 2 buffer int[16] repeat 11
+
+!*/
+
+kernel void frexp_double_scalar(global double* out, global double* in1, global int* out2) {
+    size_t id = get_global_id(0);
+    out[id] = frexp(in1[id], &out2[id]);
+}
+
+kernel void frexp_double_vec2(global double2* out, global double2* in1, global int2* out2) {
+    size_t id = get_global_id(0);
+    out[id] = frexp(in1[id], &out2[id]);
+}
+
+kernel void frexp_double_vec3(global double3* out, global double3* in1, global int3* out2) {
+    size_t id = get_global_id(0);
+    out[id] = frexp(in1[id], &out2[id]);
+}
+
+kernel void frexp_double_vec4(global double4* out, global double4* in1, global int4* out2) {
+    size_t id = get_global_id(0);
+    out[id] = frexp(in1[id], &out2[id]);
+}
+
+kernel void frexp_double_vec8(global double8* out, global double8* in1, global int8* out2) {
+    size_t id = get_global_id(0);
+    out[id] = frexp(in1[id], &out2[id]);
+}
+
+kernel void frexp_double_vec16(global double16* out, global double16* in1, global int16* out2) {
+    size_t id = get_global_id(0);
+    out[id] = frexp(in1[id], &out2[id]);
+}
-- 
2.5.0



More information about the Piglit mailing list