[clang] [clang] fix half && bfloat16 convert node expr codegen (PR #89051)
via cfe-commits
cfe-commits at lists.llvm.org
Fri Apr 26 04:08:27 PDT 2024
https://github.com/JinjinLi868 updated https://github.com/llvm/llvm-project/pull/89051
>From 0afac9d8a6acedff53089f55eacb92a2880f58aa Mon Sep 17 00:00:00 2001
From: Jinjin Li <lijinjin.868 at bytedance.com>
Date: Wed, 17 Apr 2024 16:44:50 +0800
Subject: [PATCH] [clang] Fix half && bfloat16 convert node expr codegen
Data type conversion between fp16 and bf16 will generate fptrunc
and fpextend nodes, but they are actually bitcast nodes.
---
clang/lib/CodeGen/CGExprScalar.cpp | 15 +++-
.../test/CodeGen/X86/bfloat16-convert-half.c | 25 +++++++
.../test/CodeGenHIP/bfloat16-half-convert.hip | 71 +++++++++++++++++++
3 files changed, 109 insertions(+), 2 deletions(-)
create mode 100644 clang/test/CodeGen/X86/bfloat16-convert-half.c
create mode 100644 clang/test/CodeGenHIP/bfloat16-half-convert.hip
diff --git a/clang/lib/CodeGen/CGExprScalar.cpp b/clang/lib/CodeGen/CGExprScalar.cpp
index 1f18e0d5ba409a..8e35c801bc9599 100644
--- a/clang/lib/CodeGen/CGExprScalar.cpp
+++ b/clang/lib/CodeGen/CGExprScalar.cpp
@@ -1431,7 +1431,10 @@ Value *ScalarExprEmitter::EmitScalarCast(Value *Src, QualType SrcType,
return Builder.CreateFPToUI(Src, DstTy, "conv");
}
- if (DstElementTy->getTypeID() < SrcElementTy->getTypeID())
+ if ((DstElementTy->is16bitFPTy() && SrcElementTy->is16bitFPTy())) {
+ Value *FloatVal = Builder.CreateFPExt(Src, Builder.getFloatTy(), "fpext");
+ return Builder.CreateFPTrunc(FloatVal, DstTy, "fptrunc");
+ } else if (DstElementTy->getTypeID() < SrcElementTy->getTypeID())
return Builder.CreateFPTrunc(Src, DstTy, "conv");
return Builder.CreateFPExt(Src, DstTy, "conv");
}
@@ -1906,7 +1909,15 @@ Value *ScalarExprEmitter::VisitConvertVectorExpr(ConvertVectorExpr *E) {
} else {
assert(SrcEltTy->isFloatingPointTy() && DstEltTy->isFloatingPointTy() &&
"Unknown real conversion");
- if (DstEltTy->getTypeID() < SrcEltTy->getTypeID())
+ if ((DstEltTy->is16bitFPTy() && SrcEltTy->is16bitFPTy())) {
+ auto *ScrVecTy = cast<llvm::VectorType>(SrcTy);
+ Value *FloatVal = Builder.CreateFPExt(
+ Src,
+ llvm::VectorType::get(Builder.getFloatTy(),
+ ScrVecTy->getElementCount()),
+ "fpext");
+ Res = Builder.CreateFPTrunc(FloatVal, DstTy, "fptrunc");
+ } else if (DstEltTy->getTypeID() < SrcEltTy->getTypeID())
Res = Builder.CreateFPTrunc(Src, DstTy, "conv");
else
Res = Builder.CreateFPExt(Src, DstTy, "conv");
diff --git a/clang/test/CodeGen/X86/bfloat16-convert-half.c b/clang/test/CodeGen/X86/bfloat16-convert-half.c
new file mode 100644
index 00000000000000..55451dc6f092cd
--- /dev/null
+++ b/clang/test/CodeGen/X86/bfloat16-convert-half.c
@@ -0,0 +1,25 @@
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -disable-O0-optnone -emit-llvm \
+// RUN: %s -o - | opt -S -passes=mem2reg | FileCheck %s
+
+// CHECK-LABEL: define dso_local half @test_convert_from_bf16_to_fp16(
+// CHECK-SAME: bfloat noundef [[A:%.*]]) #[[ATTR0:[0-9]+]] {
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[FPEXT:%.*]] = fpext bfloat [[A]] to float
+// CHECK-NEXT: [[FPTRUNC:%.*]] = fptrunc float [[FPEXT]] to half
+// CHECK-NEXT: ret half [[FPTRUNC]]
+//
+_Float16 test_convert_from_bf16_to_fp16(__bf16 a) {
+ return (_Float16)a;
+}
+
+// CHECK-LABEL: define dso_local bfloat @test_convert_from_fp16_to_bf16(
+// CHECK-SAME: half noundef [[A:%.*]]) #[[ATTR0]] {
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[FPEXT:%.*]] = fpext half [[A]] to float
+// CHECK-NEXT: [[FPTRUNC:%.*]] = fptrunc float [[FPEXT]] to bfloat
+// CHECK-NEXT: ret bfloat [[FPTRUNC]]
+//
+__bf16 test_convert_from_fp16_to_bf16(_Float16 a) {
+ return (__bf16)a;
+}
+
diff --git a/clang/test/CodeGenHIP/bfloat16-half-convert.hip b/clang/test/CodeGenHIP/bfloat16-half-convert.hip
new file mode 100644
index 00000000000000..0ffebb44c969b4
--- /dev/null
+++ b/clang/test/CodeGenHIP/bfloat16-half-convert.hip
@@ -0,0 +1,71 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -disable-O0-optnone -emit-llvm -fcuda-is-device \
+// RUN: %s -o - | opt -S -passes=mem2reg | FileCheck %s
+
+#define __device__ __attribute__((device))
+
+typedef _Float16 half2 __attribute__((ext_vector_type(2)));
+typedef _Float16 half4 __attribute__((ext_vector_type(4)));
+
+typedef __bf16 bfloat2 __attribute__((ext_vector_type(2)));
+typedef __bf16 bfloat4 __attribute__((ext_vector_type(4)));
+
+// CHECK-LABEL: define dso_local noundef <2 x bfloat> @_Z40test_convertvector_from_half2_to_bfloat2Dv2_DF16_
+// CHECK-SAME: (<2 x half> noundef [[IN:%.*]]) #[[ATTR0:[0-9]+]] {
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[IN_ADDR:%.*]] = alloca <2 x half>, align 4, addrspace(5)
+// CHECK-NEXT: [[IN_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[IN_ADDR]] to ptr
+// CHECK-NEXT: store <2 x half> [[IN]], ptr [[IN_ADDR_ASCAST]], align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load <2 x half>, ptr [[IN_ADDR_ASCAST]], align 4
+// CHECK-NEXT: [[FPEXT:%.*]] = fpext <2 x half> [[TMP0]] to <2 x float>
+// CHECK-NEXT: [[FPTRUNC:%.*]] = fptrunc <2 x float> [[FPEXT]] to <2 x bfloat>
+// CHECK-NEXT: ret <2 x bfloat> [[FPTRUNC]]
+//
+__device__ bfloat2 test_convertvector_from_half2_to_bfloat2(half2 in) {
+ return __builtin_convertvector(in, bfloat2);
+}
+
+// CHECK-LABEL: define dso_local noundef <2 x half> @_Z40test_convertvector_from_bfloat2_to_half2Dv2_DF16b
+// CHECK-SAME: (<2 x bfloat> noundef [[IN:%.*]]) #[[ATTR0]] {
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[IN_ADDR:%.*]] = alloca <2 x bfloat>, align 4, addrspace(5)
+// CHECK-NEXT: [[IN_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[IN_ADDR]] to ptr
+// CHECK-NEXT: store <2 x bfloat> [[IN]], ptr [[IN_ADDR_ASCAST]], align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load <2 x bfloat>, ptr [[IN_ADDR_ASCAST]], align 4
+// CHECK-NEXT: [[FPEXT:%.*]] = fpext <2 x bfloat> [[TMP0]] to <2 x float>
+// CHECK-NEXT: [[FPTRUNC:%.*]] = fptrunc <2 x float> [[FPEXT]] to <2 x half>
+// CHECK-NEXT: ret <2 x half> [[FPTRUNC]]
+//
+__device__ half2 test_convertvector_from_bfloat2_to_half2(bfloat2 in) {
+ return __builtin_convertvector(in, half2);
+}
+
+// CHECK-LABEL: define dso_local noundef <4 x bfloat> @_Z40test_convertvector_from_half4_to_bfloat4Dv4_DF16_
+// CHECK-SAME: (<4 x half> noundef [[IN:%.*]]) #[[ATTR0]] {
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[IN_ADDR:%.*]] = alloca <4 x half>, align 8, addrspace(5)
+// CHECK-NEXT: [[IN_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[IN_ADDR]] to ptr
+// CHECK-NEXT: store <4 x half> [[IN]], ptr [[IN_ADDR_ASCAST]], align 8
+// CHECK-NEXT: [[TMP0:%.*]] = load <4 x half>, ptr [[IN_ADDR_ASCAST]], align 8
+// CHECK-NEXT: [[FPEXT:%.*]] = fpext <4 x half> [[TMP0]] to <4 x float>
+// CHECK-NEXT: [[FPTRUNC:%.*]] = fptrunc <4 x float> [[FPEXT]] to <4 x bfloat>
+// CHECK-NEXT: ret <4 x bfloat> [[FPTRUNC]]
+//
+__device__ bfloat4 test_convertvector_from_half4_to_bfloat4(half4 in) {
+ return __builtin_convertvector(in, bfloat4);
+}
+
+// CHECK-LABEL: define dso_local noundef <4 x half> @_Z40test_convertvector_from_bfloat4_to_half4Dv4_DF16b
+// CHECK-SAME: (<4 x bfloat> noundef [[IN:%.*]]) #[[ATTR0]] {
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[IN_ADDR:%.*]] = alloca <4 x bfloat>, align 8, addrspace(5)
+// CHECK-NEXT: [[IN_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[IN_ADDR]] to ptr
+// CHECK-NEXT: store <4 x bfloat> [[IN]], ptr [[IN_ADDR_ASCAST]], align 8
+// CHECK-NEXT: [[TMP0:%.*]] = load <4 x bfloat>, ptr [[IN_ADDR_ASCAST]], align 8
+// CHECK-NEXT: [[FPEXT:%.*]] = fpext <4 x bfloat> [[TMP0]] to <4 x float>
+// CHECK-NEXT: [[FPTRUNC:%.*]] = fptrunc <4 x float> [[FPEXT]] to <4 x half>
+// CHECK-NEXT: ret <4 x half> [[FPTRUNC]]
+//
+__device__ half4 test_convertvector_from_bfloat4_to_half4(bfloat4 in) {
+ return __builtin_convertvector(in, half4);
+}
More information about the cfe-commits
mailing list