e621757365
Pursuant to discussions at https://discourse.llvm.org/t/rfc-c-23-p1467r9-extended-floating-point-types-and-standard-names/70033/22, this commit enhances the handling of the __bf16 type in Clang. - Firstly, it upgrades __bf16 from a storage-only type to an arithmetic type. - Secondly, it changes the mangling of __bf16 to DF16b on all architectures except ARM. This change has been made in accordance with the finalization of the mangling for the std::bfloat16_t type, as discussed at https://github.com/itanium-cxx-abi/cxx-abi/pull/147. - Finally, this commit extends the existing excess precision support to the __bf16 type. This applies to hardware architectures that do not natively support bfloat16 arithmetic. Appropriate tests have been added to verify the effects of these changes and ensure no regressions in other areas of the compiler. Reviewed By: rjmccall, pengfei, zahiraam Differential Revision: https://reviews.llvm.org/D150913
47 lines
1.6 KiB
Plaintext
47 lines
1.6 KiB
Plaintext
// REQUIRES: nvptx-registered-target
|
|
// REQUIRES: x86-registered-target
|
|
|
|
// RUN: %clang_cc1 "-aux-triple" "x86_64-unknown-linux-gnu" "-triple" "nvptx64-nvidia-cuda" \
|
|
// RUN: -fcuda-is-device "-aux-target-cpu" "x86-64" -S -o - %s | FileCheck %s
|
|
|
|
#include "Inputs/cuda.h"
|
|
|
|
// CHECK-LABEL: .visible .func _Z8test_argPDF16bDF16b(
|
|
// CHECK: .param .b64 _Z8test_argPDF16bDF16b_param_0,
|
|
// CHECK: .param .b16 _Z8test_argPDF16bDF16b_param_1
|
|
//
|
|
__device__ void test_arg(__bf16 *out, __bf16 in) {
|
|
// CHECK: ld.param.b16 %{{h.*}}, [_Z8test_argPDF16bDF16b_param_1];
|
|
__bf16 bf16 = in;
|
|
*out = bf16;
|
|
// CHECK: st.b16
|
|
// CHECK: ret;
|
|
}
|
|
|
|
|
|
// CHECK-LABEL: .visible .func (.param .b32 func_retval0) _Z8test_retDF16b(
|
|
// CHECK: .param .b16 _Z8test_retDF16b_param_0
|
|
__device__ __bf16 test_ret( __bf16 in) {
|
|
// CHECK: ld.param.b16 %h{{.*}}, [_Z8test_retDF16b_param_0];
|
|
return in;
|
|
// CHECK: st.param.b16 [func_retval0+0], %h
|
|
// CHECK: ret;
|
|
}
|
|
|
|
// CHECK-LABEL: .visible .func (.param .b32 func_retval0) _Z9test_callDF16b(
|
|
// CHECK: .param .b16 _Z9test_callDF16b_param_0
|
|
__device__ __bf16 test_call( __bf16 in) {
|
|
// CHECK: ld.param.b16 %h{{.*}}, [_Z9test_callDF16b_param_0];
|
|
// CHECK: st.param.b16 [param0+0], %h2;
|
|
// CHECK: .param .b32 retval0;
|
|
// CHECK: call.uni (retval0),
|
|
// CHECK-NEXT: _Z8test_retDF16b,
|
|
// CHECK-NEXT: (
|
|
// CHECK-NEXT: param0
|
|
// CHECK-NEXT );
|
|
// CHECK: ld.param.b16 %h{{.*}}, [retval0+0];
|
|
return test_ret(in);
|
|
// CHECK: st.param.b16 [func_retval0+0], %h
|
|
// CHECK: ret;
|
|
}
|