summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorMatt Arsenault <arsenm2@gmail.com>2017-01-26 15:09:52 -0800
committerJan Vesely <jan.vesely@rutgers.edu>2017-01-27 13:03:18 -0500
commit7d422a3bf6d658330c663bf45478afb325f901a3 (patch)
tree2aedd43b604afb59f4cf735426d05cbf4be31317
parent01eb7321b46dbd11d71a2ca07217ef387486db3a (diff)
cl: Add tests for stored fneg
Make sure fneg that can't be folded into a use operand is executed correctly. v2: Merge with other test v3: Remove other test Reviewed-by: Jan Vesely <jan.vesely@rutgers.edu>
-rw-r--r--tests/cl/program/execute/scalar-arithmetic-double.cl36
-rw-r--r--tests/cl/program/execute/scalar-arithmetic-float.cl42
2 files changed, 55 insertions, 23 deletions
diff --git a/tests/cl/program/execute/scalar-arithmetic-double.cl b/tests/cl/program/execute/scalar-arithmetic-double.cl
new file mode 100644
index 000000000..3c2cb6271
--- /dev/null
+++ b/tests/cl/program/execute/scalar-arithmetic-double.cl
@@ -0,0 +1,36 @@
+/*!
+[config]
+name: fneg f64
+clc_version_min: 10
+require_device_extensions: cl_khr_fp64
+dimensions: 1
+
+## Unary minus ##
+
+[test]
+name: minus
+kernel_name: minus
+global_size: 20 0 0
+
+arg_out: 0 buffer double[20] \
+ -0.0 0.0 -0.5 0.5 \
+ -1.0 1.0 -2.0 2.0 \
+ -4.0 4.0 -10.0 10.0 \
+ -inf inf nan -345.25 \
+ 345.25 -455.125 455.125 0.12345
+
+arg_in: 1 buffer double[20] \
+ 0.0 -0.0 0.5 -0.5 \
+ 1.0 -1.0 2.0 -2.0 \
+ 4.0 -4.0 10.0 -10.0 \
+ inf -inf nan 345.25 \
+ -345.25 455.125 -455.125 -0.12345
+
+!*/
+
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+
+kernel void minus(global double* out, global double* in) {
+ int id = get_global_id(0);
+ out[id] = -in[id];
+}
diff --git a/tests/cl/program/execute/scalar-arithmetic-float.cl b/tests/cl/program/execute/scalar-arithmetic-float.cl
index 84c397030..3412e1b7b 100644
--- a/tests/cl/program/execute/scalar-arithmetic-float.cl
+++ b/tests/cl/program/execute/scalar-arithmetic-float.cl
@@ -293,28 +293,23 @@ arg_out: 0 buffer float[1] -inf
## Unary minus ##
[test]
-name: -pos
+name: minus
kernel_name: minus
-arg_in: 1 float 345.25
-arg_out: 0 buffer float[1] -345.25
-
-[test]
-name: -neg
-kernel_name: minus
-arg_in: 1 float -455.125
-arg_out: 0 buffer float[1] 455.125
-
-[test]
-name: -inf
-kernel_name: minus
-arg_in: 1 float inf
-arg_out: 0 buffer float[1] -inf
-
-[test]
-name: --inf
-kernel_name: minus
-arg_in: 1 float -inf
-arg_out: 0 buffer float[1] inf
+global_size: 20 0 0
+
+arg_out: 0 buffer float[20] \
+ -0.0 0.0 -0.5 0.5 \
+ -1.0 1.0 -2.0 2.0 \
+ -4.0 4.0 -10.0 10.0 \
+ -inf inf nan -345.25 \
+ 345.25 -455.125 455.125 0.12345
+
+arg_in: 1 buffer float[20] \
+ 0.0 -0.0 0.5 -0.5 \
+ 1.0 -1.0 2.0 -2.0 \
+ 4.0 -4.0 10.0 -10.0 \
+ inf -inf nan 345.25 \
+ -345.25 455.125 -455.125 -0.12345
!*/
@@ -338,6 +333,7 @@ kernel void plus(global float* out, float in) {
out[0] = +in;
}
-kernel void minus(global float* out, float in) {
- out[0] = -in;
+kernel void minus(global float* out, global float* in) {
+ int id = get_global_id(0);
+ out[id] = -in[id];
}