[clang][CodeGen] Allow fp16 arg pass by register

[yaxunl]

Summary

HIP supports _Float16 and __fp16 types. In x86_64 host they are for storage
only. Since they have the same size and alignment as int16, they are supposed
to be passed by value in the same way as int16. Currently clang pass them
by stack when included in a struct, which is not efficient. This also causes
interoperability difficulty with gcc. On gcc since there is no _Float16 type,
int16 is used as replacement for _Float16 for passing arguments, which is
passed by register.

This patch changes x86_64 target codegen info so that _Float16 and __fp16
in structs can be passed by register.

[tra]

clang/lib/CodeGen/TargetInfo.cpp

2825

Do we need to set Hi, too? We do set it for int128.

[yaxunl]

clang/lib/CodeGen/TargetInfo.cpp

2825

Seems not. By default Hi is NoClass. I tried setting Hi to Integer and it resulted in an extra int64 argument of no use.

[pengfei]

clang/lib/CodeGen/TargetInfo.cpp

2821

It's true that AMD64 does not support _Float16, but __fp16 is supported on every target.
See https://clang.llvm.org/docs/LanguageExtensions.html#half-precision-floating-point

2822

__fp16 is storage format while _Float16 isn't.

2824

The GCC should take _Floatn as floating types, see https://gcc.gnu.org/onlinedocs/gcc-8.2.0/gcc/Floating-Types.html#Floating-Types
Can you provide the information how gcc uses as 16 bit integer for _Float16?

[yaxunl]

clang/lib/CodeGen/TargetInfo.cpp

2821

will fix comments

2822

will fix comments

2824

GCC does not allow _Float16 or __fp16 to be used for x86_64.

HIP language allows _Float16 or fp16 type to be used with x86_64. To be able to passing _Float16 or fp16 values between C++ using gcc and HIP using clang, users often write code like this:

struct half_t {
#ifdef __GNUC__
uint16_t x;
#else
_Float16 x;
#endif
};

void fun1(half_t x);

And usually fun1 is defined in HIP and compiled by clang to object file, then they call it in C++ which is compiled by gcc. Then the caller will pass half_t by register. However clang passes half_t by stack. This causes issue.

[yaxunl]

revised comments and fixed test

[pengfei]

I don't think the patch is doing right.

clang/lib/CodeGen/TargetInfo.cpp

2821

This is still not correct I think. As Clang dos says, _Float16 is not support (including load and store) unless ABI defines it.
We cannot add it before there's clear definition in the ABI.
See https://github.com/hjl-tools/x86-psABI/wiki/x86-64-psABI-1.0.pdf

2824

It looks to me this is a user scenario issue. You are using different types between GCC and Clang. There's nothing surprised if they aren't interoperable. You should use uint16_t in fun1 too instead of changing the behavior of the floating type.
Changing type behavior is dangerous, you will result in backward compatiblity problems. That's also the reason why we need a well-defined ABI :)

[yaxunl]

clang/lib/CodeGen/TargetInfo.cpp

2824

That makes sense. Thanks.

[hliao]

clang/lib/CodeGen/TargetInfo.cpp

2821

That's quite outdated spec. Please check the latest one from https://gitlab.com/x86-psABIs/x86-64-ABI, which classify _Float16 as SSE.

[hliao]

clang/lib/CodeGen/TargetInfo.cpp

2821

That's quite outdated spec. Please check the latest one from https://gitlab.com/x86-psABIs/x86-64-ABI, which classify _Float16 as SSE.

From the log

commit 71d1183e7bb95e9f8ad732e0f2b5a4f127796e2a (origin/usr/hjl/_Float16)
Author: H.J. Lu <[e-mail]>
Date: Wed Feb 20 05:45:39 2019 -0800

Add optional _Float16 support

Pass and return _Float16 values in XMM registers.

diff --git a/x86-64-ABI/low-level-sys-info.tex b/x86-64-ABI/low-level-sys-info.tex
index ca84fff..0b06c56 100644

• a/x86-64-ABI/low-level-sys-info.tex

+++ b/x86-64-ABI/low-level-sys-info.tex
@@ -25,7 +25,8 @@ object, and the term \emph{\textindex{\sixteenbyte{}}} refers to a
\subsubsection{Fundamental Types}

Figure~\ref{basic-types} shows the correspondence between ISO C's
-scalar types and the processor's. \code{int128}, \code{float80},
+scalar types and the processor's. \code{int128}, \code{_Float16},
+\code{float80},
\code{float128}, \code{m64}, \code{m128}, \code{m256} and
\code{__m512} types are optional.

[yaxunl]

So should we revive this patch by classifying _Float16 as SSE, then it will conform to the ABI. And we should be compatible with gcc assuming it also follows the ABI. @pengfei

[pengfei]

In D97318#2587866, @yaxunl wrote:

So should we revive this patch by classifying _Float16 as SSE, then it will conform to the ABI. And we should be compatible with gcc assuming it also follows the ABI. @pengfei

Thanks @hliao for the information. Yes, we should classify _Float16 as SSE. Sorry for the misleading.

clang/lib/CodeGen/TargetInfo.cpp

if (k == BuiltinType::Void) {
         Hi = SSEUp;
       } else if (LDF == &llvm::APFloat::x87DoubleExtended()) {
         Lo = X87;
         Hi = X87Up;
       } else if (LDF == &llvm::APFloat::IEEEdouble()) {
         Current = SSE;
       } else
         llvm_unreachable("unexpected long double representation!");
+    } else if (k == BuiltinType::Float16 || k == BuiltinType::Half) {
+      // AMD64 does not support operations on _Float16 or __fp16 other than
+      // load and store. For load/store operations, _Float16 and __fp16 is
+      // equivalent to 16 bit integer since they have the same size and
+      // alignment. We need this to interop with gcc where 16 bit integer
+      // is used in place of _Float16 or __fp16.
+      Lo = Integer;
     }
     // FIXME: _Decimal32 and _Decimal64 are SSE.
     // FIXME: _float128 and _Decimal128 are (SSE, SSEUp).
     return;
   }
 
   if (const EnumType *ET = Ty->getAs<EnumType>()) {
     // Classify the underlying integer type.

clang/test/CodeGenCUDA/float16.cu

+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -o - -x hip %s \
+// RUN:   -fhip-new-launch-api | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm -o - -x hip %s \
+// RUN:   -fcuda-is-device | FileCheck -check-prefix=DEV %s
+
+#include "Inputs/cuda.h"
+
+// CHECK: %struct.A = type { i16 }
+struct A {
+  short x;
+};
+
+// CHECK: %struct.B = type { half }
+struct B {
+  _Float16 x;
+};
+
+// CHECK: %struct.C = type { half }
+struct C {
+  __fp16 x;
+};
+
+// Check struct containing _Float16 is coerced and passed correctly to kernel
+// in a similar way as int16.
+
+// CHECK: define dso_local void @_Z20__device_stub__kern11A(i16 %x.coerce)
+// CHECK:  %x = alloca %struct.A, align 2
+// CHECK:  %coerce.dive = getelementptr inbounds %struct.A, %struct.A* %x, i32 0, i32 0
+// CHECK:  store i16 %x.coerce, i16* %coerce.dive, align 2
+// CHECK:  %kernel_args = alloca i8*, i64 1, align 16
+// CHECK:  %[[PTR:.*]] = bitcast %struct.A* %x to i8*
+// CHECK:  %[[ARGS:.*]] = getelementptr i8*, i8** %kernel_args, i32 0
+// CHECK:  store i8* %[[PTR]], i8** %[[ARGS]], align 8
+// DEV: define dso_local amdgpu_kernel void @_Z5kern11A(i16 %x.coerce)
+// DEV: %coerce.dive = getelementptr inbounds %struct.A, %struct.A* %x1, i32 0, i32 0
+// DEV: store i16 %x.coerce, i16* %coerce.dive, align 2
+__global__ void kern1(A x) {
+  x.x += 1;
+}
+
+// CHECK: define dso_local void @_Z20__device_stub__kern21B(i16 %x.coerce)
+// CHECK:  %x = alloca %struct.B, align 2
+// CHECK:  %coerce.dive = getelementptr inbounds %struct.B, %struct.B* %x, i32 0, i32 0
+// CHECK:  %[[PTR:.*]] = bitcast half* %coerce.dive to i16*
+// CHECK:  store i16 %x.coerce, i16* %[[PTR]], align 2
+// CHECK:  %kernel_args = alloca i8*, i64 1, align 16
+// CHECK:  %[[PTR:.*]] = bitcast %struct.B* %x to i8*
+// CHECK:  %[[ARGS:.*]] = getelementptr i8*, i8** %kernel_args, i32 0
+// CHECK:  store i8* %[[PTR]], i8** %[[ARGS]], align 8
+// DEV: define dso_local amdgpu_kernel void @_Z5kern21B(half %x.coerce)
+// DEV:  %coerce.dive = getelementptr inbounds %struct.B, %struct.B* %x1, i32 0, i32 0
+// DEV:  store half %x.coerce, half* %coerce.dive, align 2
+// DEV:  %[[HALF:.*]] = load half, half* %x2, align 2
+// DEV:  %add = fadd contract half %[[HALF]], 0xH3C00
+// DEV:  store half %add, half* %x2, align 2
+__global__ void kern2(B x) {
+  x.x += 1;
+}
+
+// CHECK: define dso_local void @_Z20__device_stub__kern31C(i16 %x.coerce)
+// CHECK:  %x = alloca %struct.C, align 2
+// CHECK:  %coerce.dive = getelementptr inbounds %struct.C, %struct.C* %x, i32 0, i32 0
+// CHECK:  %[[PTR:.*]] = bitcast half* %coerce.dive to i16*
+// CHECK:  store i16 %x.coerce, i16* %[[PTR]], align 2
+// CHECK:  %kernel_args = alloca i8*, i64 1, align 16
+// CHECK:  %[[PTR:.*]] = bitcast %struct.C* %x to i8*
+// CHECK:  %[[ARGS:.*]] = getelementptr i8*, i8** %kernel_args, i32 0
+// CHECK:  store i8* %[[PTR]], i8** %[[ARGS]], align 8
+// DEV: define dso_local amdgpu_kernel void @_Z5kern31C(half %x.coerce)
+// DEV:  %coerce.dive = getelementptr inbounds %struct.C, %struct.C* %x1, i32 0, i32 0
+// DEV:  store half %x.coerce, half* %coerce.dive, align 2
+// DEV:  %[[HALF:.*]] = load half, half* %x2, align 2
+// DEV:  %conv = fpext half %[[HALF]] to float
+// DEV:  %add = fadd contract float %conv, 1.000000e+00
+// DEV:  %[[HALF:.*]] = fptrunc float %add to half
+// DEV:  store half %[[HALF]], half* %x2, align 2
+__global__ void kern3(C x) {
+  x.x += 1;
+}
+
+// CHECK: define dso_local void @_Z4fun11A(i16 %x.coerce)
+void fun1(A x) {
+  kern1<<<1, 1>>>(x);
+}
+
+// CHECK: define dso_local void @_Z4fun21B(i16 %x.coerce)
+void fun2(B x) {
+  kern2<<<1, 1>>>(x);
+}
+
+// CHECK: define dso_local void @_Z5func31C(i16 %x.coerce)
+void func3(C x) {
+  kern3<<<1, 1>>>(x);
+}

clang/test/CodeGenOpenCL/builtins-f16.cl

 // RUN: %clang_cc1 -emit-llvm -o - -triple x86_64-darwin-apple %s | FileCheck %s
 
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
 
 // CHECK-LABEL: define{{.*}} void @test_half_builtins
 void test_half_builtins(half h0, half h1, half h2) {
   volatile half res;
 
-  // CHECK: call half @llvm.copysign.f16(half %h0, half %h1)
+  // CHECK: [[h0:%.*]] = bitcast i16 %h0.coerce to half
+  // CHECK: [[h1:%.*]] = bitcast i16 %h1.coerce to half
+  // CHECK: [[h2:%.*]] = bitcast i16 %h2.coerce to half
+
+  // CHECK: call half @llvm.copysign.f16(half [[h0]], half [[h1]])
   res = __builtin_copysignf16(h0, h1);
 
-  // CHECK: call half @llvm.fabs.f16(half %h0)
+  // CHECK: call half @llvm.fabs.f16(half [[h0]])
   res = __builtin_fabsf16(h0);
 
-  // CHECK: call half @llvm.ceil.f16(half %h0)
+  // CHECK: call half @llvm.ceil.f16(half [[h0]])
   res = __builtin_ceilf16(h0);
 
-  // CHECK: call half @llvm.cos.f16(half %h0)
+  // CHECK: call half @llvm.cos.f16(half [[h0]])
   res = __builtin_cosf16(h0);
 
-  // CHECK: call half @llvm.exp.f16(half %h0)
+  // CHECK: call half @llvm.exp.f16(half [[h0]])
   res = __builtin_expf16(h0);
 
-  // CHECK: call half @llvm.exp2.f16(half %h0)
+  // CHECK: call half @llvm.exp2.f16(half [[h0]])
   res = __builtin_exp2f16(h0);
 
-  // CHECK: call half @llvm.floor.f16(half %h0)
+  // CHECK: call half @llvm.floor.f16(half [[h0]])
   res = __builtin_floorf16(h0);
 
-  // CHECK: call half @llvm.fma.f16(half %h0, half %h1, half %h2)
+  // CHECK: call half @llvm.fma.f16(half [[h0]], half [[h1]], half [[h2]])
   res = __builtin_fmaf16(h0, h1 ,h2);
 
-  // CHECK: call half @llvm.maxnum.f16(half %h0, half %h1)
+  // CHECK: call half @llvm.maxnum.f16(half [[h0]], half [[h1]])
   res = __builtin_fmaxf16(h0, h1);
 
-  // CHECK: call half @llvm.minnum.f16(half %h0, half %h1)
+  // CHECK: call half @llvm.minnum.f16(half [[h0]], half [[h1]])
   res = __builtin_fminf16(h0, h1);
 
-  // CHECK: frem half %h0, %h1
+  // CHECK: frem half [[h0]], [[h1]]
   res = __builtin_fmodf16(h0, h1);
 
-  // CHECK: call half @llvm.pow.f16(half %h0, half %h1)
+  // CHECK: call half @llvm.pow.f16(half [[h0]], half [[h1]])
   res = __builtin_powf16(h0, h1);
 
-  // CHECK: call half @llvm.log10.f16(half %h0)
+  // CHECK: call half @llvm.log10.f16(half [[h0]])
   res = __builtin_log10f16(h0);
 
-  // CHECK: call half @llvm.log2.f16(half %h0)
+  // CHECK: call half @llvm.log2.f16(half [[h0]])
   res = __builtin_log2f16(h0);
 
-  // CHECK: call half @llvm.log.f16(half %h0)
+  // CHECK: call half @llvm.log.f16(half [[h0]])
   res = __builtin_logf16(h0);
 
-  // CHECK: call half @llvm.rint.f16(half %h0)
+  // CHECK: call half @llvm.rint.f16(half [[h0]])
   res = __builtin_rintf16(h0);
 
-  // CHECK: call half @llvm.round.f16(half %h0)
+  // CHECK: call half @llvm.round.f16(half [[h0]])
   res = __builtin_roundf16(h0);
 
-  // CHECK: call half @llvm.sin.f16(half %h0)
+  // CHECK: call half @llvm.sin.f16(half [[h0]])
   res = __builtin_sinf16(h0);
 
-  // CHECK: call half @llvm.sqrt.f16(half %h0)
+  // CHECK: call half @llvm.sqrt.f16(half [[h0]])
   res = __builtin_sqrtf16(h0);
 
-  // CHECK: call half @llvm.trunc.f16(half %h0)
+  // CHECK: call half @llvm.trunc.f16(half [[h0]])
   res = __builtin_truncf16(h0);
 
-  // CHECK: call half @llvm.canonicalize.f16(half %h0)
+  // CHECK: call half @llvm.canonicalize.f16(half [[h0]])
   res = __builtin_canonicalizef16(h0);
 }