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

[arsenm]

clang/test/CodeGenHIP/half.hip

7

Making sure the argument type is "half" wouldn't hurt either

clang/lib/Basic/Targets/AMDGPU.h

public:
 
   bool
   initFeatureMap(llvm::StringMap<bool> &Features, DiagnosticsEngine &Diags,
                  StringRef CPU,
                  const std::vector<std::string> &FeatureVec) const override;
 
   ArrayRef<Builtin::Info> getTargetBuiltins() const override;
 
+  bool useFP16ConversionIntrinsics() const override { return false; }
+
   void getTargetDefines(const LangOptions &Opts,
                         MacroBuilder &Builder) const override;
 
   BuiltinVaListKind getBuiltinVaListKind() const override {
     return TargetInfo::CharPtrBuiltinVaList;
   }
 
   bool isValidCPUName(StringRef Name) const override {

clang/test/CodeGenHIP/half.hip

+// 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;
+}