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

[Michael Liao via Phabricator via cfe-commits]

hliao created this revision.
hliao added reviewers: yaxunl, arsenm.
Herald added subscribers: cfe-commits, kerbowa, t-tye, tpr, dstuttard, nhaehnle, wdng, jvesely, kzhuravl.
Herald added a project: clang.

Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D81849

Files:
  clang/lib/Basic/Targets/AMDGPU.h
  clang/test/CodeGenHIP/half.hip


Index: clang/test/CodeGenHIP/half.hip
===================================================================
--- /dev/null
+++ 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;
+}
Index: clang/lib/Basic/Targets/AMDGPU.h
===================================================================
--- clang/lib/Basic/Targets/AMDGPU.h
+++ clang/lib/Basic/Targets/AMDGPU.h
@@ -219,6 +219,8 @@
 
   ArrayRef<Builtin::Info> getTargetBuiltins() const override;
 
+  bool useFP16ConversionIntrinsics() const override { return false; }
+
   void getTargetDefines(const LangOptions &Opts,
                         MacroBuilder &Builder) const override;
 


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D81849.270777.patch
Type: text/x-patch
Size: 1015 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20200615/0fdaac42/attachment.bin>