Add support for 'fneg' when translate from LLVM to SPIR-V

It adds support for the fneg operation (added in LLVM 8.0.0) when translating from LLVM IR to SPIR-V, using the SPIR-V FNegate operation. Signed-off-by: Ilya Mashkov ilya.mashkov@intel.com Fix #359
KhronosGroup · Oct 28, 2019 · 128e4c5 · 128e4c5
1 parent 5de3935
commit 128e4c5
Show file tree

Hide file tree

Showing 3 changed files with 27 additions and 19 deletions.
diff --git a/lib/SPIRV/SPIRVInternal.h b/lib/SPIRV/SPIRVInternal.h
@@ -98,6 +98,7 @@ template <> inline void SPIRVMap<unsigned, Op>::init() {
  _SPIRV_OP(BitCast, Bitcast)
  _SPIRV_OP(AddrSpaceCast, GenericCastToPtr)
  _SPIRV_OP(GetElementPtr, AccessChain)
+ _SPIRV_OP(FNeg, FNegate)
  /*Binary*/
  _SPIRV_OP(And, BitwiseAnd)
  _SPIRV_OP(Or, BitwiseOr)

diff --git a/test/OpFNegate.spvasm b/test/OpFNegate.spvasm
diff --git a/test/transcoding/TransFNeg.cl b/test/transcoding/TransFNeg.cl
@@ -0,0 +1,26 @@
+// RUN: %clang_cc1 -triple spir-unknown-unknown -O0 -cl-std=CL2.0 -finclude-default-header -emit-llvm-bc %s -o %t.bc
+// RUN: llvm-spirv %t.bc -spirv-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+// RUN: llvm-spirv %t.bc -o %t.spv
+// RUN: spirv-val %t.spv
+// RUN: llvm-spirv -r %t.spv -o %t.rev.bc
+// RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
+
+// CHECK-SPIRV: FNegate
+// CHECK-SPIRV: FNegate
+// CHECK-SPIRV: FNegate
+// CHECK-SPIRV: FNegate
+
+// CHECK-LLVM: fsub half 0xH8000, %
+// CHECK-LLVM: fsub float -0.000000e+00, %
+// CHECK-LLVM: fsub double -0.000000e+00, % 
+// CHECK-LLVM: fsub <8 x double> <double -0.000000e+00, double -0.000000e+00, double -0.000000e+00, double -0.000000e+00, double -0.000000e+00, double -0.000000e+00, double -0.000000e+00, double -0.000000e+00>, %
+
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+__kernel void foo(double a1, __global half *h, __global float *b0, __global double *b1, __global double8 *d) {
+ *h = -*h;
+ *b0 = -*b0;
+ *b1 = -a1;
+ *d = -*d;
+}