[clang] e830fa2 - [clang][amdgpu] Prefer not using `fp16` conversion intrinsics.

[Michael Liao via cfe-commits]

Author: Michael Liao
Date: 2020-06-16T10:21:56-04:00
New Revision: e830fa260da9d3bbe99c4176b4ddb6aa5e6229dd

URL: https://github.com/llvm/llvm-project/commit/e830fa260da9d3bbe99c4176b4ddb6aa5e6229dd
DIFF: https://github.com/llvm/llvm-project/commit/e830fa260da9d3bbe99c4176b4ddb6aa5e6229dd.diff

LOG: [clang][amdgpu] Prefer not using `fp16` conversion intrinsics.

Reviewers: yaxunl, arsenm

Subscribers: kzhuravl, jvesely, wdng, nhaehnle, dstuttard, tpr, t-tye, kerbowa, cfe-commits

Tags: #clang

Differential Revision: https://reviews.llvm.org/D81849

Added: 
    clang/test/CodeGenHIP/half.hip

Modified: 
    clang/lib/Basic/Targets/AMDGPU.h

Removed: 
    


################################################################################
diff  --git a/clang/lib/Basic/Targets/AMDGPU.h b/clang/lib/Basic/Targets/AMDGPU.h
index e4194a881e3f..387b91abb537 100644
--- a/clang/lib/Basic/Targets/AMDGPU.h
+++ b/clang/lib/Basic/Targets/AMDGPU.h
@@ -219,6 +219,8 @@ class LLVM_LIBRARY_VISIBILITY AMDGPUTargetInfo final : public TargetInfo {
 
   ArrayRef<Builtin::Info> getTargetBuiltins() const override;
 
+  bool useFP16ConversionIntrinsics() const override { return false; }
+
   void getTargetDefines(const LangOptions &Opts,
                         MacroBuilder &Builder) const override;
 

diff  --git a/clang/test/CodeGenHIP/half.hip b/clang/test/CodeGenHIP/half.hip
new file mode 100644
index 000000000000..d5dd43d51fdf
--- /dev/null
+++ b/clang/test/CodeGenHIP/half.hip
@@ -0,0 +1,16 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -emit-llvm -fcuda-is-device -o - %s | FileCheck %s
+
+#define __device__ __attribute__((device))
+
+// CHECK-LABEL: @_Z2d0DF16_
+// CHECK: fpext
+__device__ float d0(_Float16 x) {
+  return x;
+}
+
+// CHECK-LABEL: @_Z2d1f
+// CHECK: fptrunc
+__device__ _Float16 d1(float x) {
+  return x;
+}