Magic-Mirror/llvm-capstone
1
0
Fork 0
mirror of https://github.com/capstone-engine/llvm-capstone.git synced 2025-01-08 17:11:33 +00:00
Code Issues Actions 2 Packages Projects Releases Wiki Activity

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

Browse Source 
Reviewers: yaxunl, arsenm

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

Tags: #clang

Differential Revision: https://reviews.llvm.org/D81849
This commit is contained in:
Michael Liao 2020-06-02 15:39:52 -04:00
parent 14b947f306
commit e830fa260d
2 changed files with 18 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
clang/lib/Basic/Targets/AMDGPU.h
View File

@ -219,6 +219,8 @@ public:
ArrayRef<Builtin::Info> getTargetBuiltins() const override;
bool useFP16ConversionIntrinsics() const override { return false; }
void getTargetDefines(const LangOptions &Opts,
MacroBuilder &Builder) const override;

16
clang/test/CodeGenHIP/half.hip Normal file
View File

@ -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;
}