From 128e4c550f0f20ff7cb7bf284c0b3718a3e56b87 Mon Sep 17 00:00:00 2001 From: Ilya Mashkov Date: Mon, 21 Oct 2019 20:42:23 +0300 Subject: [PATCH] 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 --- lib/SPIRV/SPIRVInternal.h | 1 + test/OpFNegate.spvasm | 19 ------------------- test/transcoding/TransFNeg.cl | 26 ++++++++++++++++++++++++++ 3 files changed, 27 insertions(+), 19 deletions(-) delete mode 100644 test/OpFNegate.spvasm create mode 100644 test/transcoding/TransFNeg.cl diff --git a/lib/SPIRV/SPIRVInternal.h b/lib/SPIRV/SPIRVInternal.h index 2e9b3acea0..3d863e1aad 100644 --- a/lib/SPIRV/SPIRVInternal.h +++ b/lib/SPIRV/SPIRVInternal.h @@ -98,6 +98,7 @@ template <> inline void SPIRVMap::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 deleted file mode 100644 index 559a4d292f..0000000000 --- a/test/OpFNegate.spvasm +++ /dev/null @@ -1,19 +0,0 @@ -; REQUIRES: spirv-as -; RUN: spirv-as --target-env spv1.0 -o %t.spv %s -; RUN: llvm-spirv -r -o - %t.spv | llvm-dis | FileCheck %s - OpCapability Addresses - OpCapability Kernel - OpMemoryModel Physical32 OpenCL - OpEntryPoint Kernel %1 "testFNegate" - OpName %a "a" - %void = OpTypeVoid - %float = OpTypeFloat 32 - %5 = OpTypeFunction %void %float - %1 = OpFunction %void None %5 - %a = OpFunctionParameter %float - %6 = OpLabel - %7 = OpFNegate %float %a - OpReturn - OpFunctionEnd - -; CHECK: fsub float -0.000000e+00, %a diff --git a/test/transcoding/TransFNeg.cl b/test/transcoding/TransFNeg.cl new file mode 100644 index 0000000000..c1ea327391 --- /dev/null +++ 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> , % + +#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; +}