diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp --- a/clang/lib/AST/ASTContext.cpp +++ b/clang/lib/AST/ASTContext.cpp @@ -2171,9 +2171,14 @@ Align = Target->getLongFractAlign(); break; case BuiltinType::BFloat16: + // HIP does not currently support bf16, so in that case allow querying the + // auxiliary target. if (Target->hasBFloat16Type()) { Width = Target->getBFloat16Width(); Align = Target->getBFloat16Align(); + } else if (getLangOpts().HIP && AuxTarget->hasBFloat16Type()) { + Width = AuxTarget->getBFloat16Width(); + Align = AuxTarget->getBFloat16Align(); } break; case BuiltinType::Float16: diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp --- a/clang/lib/Sema/SemaType.cpp +++ b/clang/lib/Sema/SemaType.cpp @@ -1518,7 +1518,11 @@ break; case DeclSpec::TST_half: Result = Context.HalfTy; break; case DeclSpec::TST_BFloat16: - if (!S.Context.getTargetInfo().hasBFloat16Type()) + // HIP does not currently support bf16. Avoid diagnosing uses of bf16 + // if the auxiliary target supports it. + if (!S.Context.getTargetInfo().hasBFloat16Type() && + !(S.getLangOpts().HIP && + S.Context.getAuxTargetInfo()->hasBFloat16Type())) S.Diag(DS.getTypeSpecTypeLoc(), diag::err_type_unsupported) << "__bf16"; Result = Context.BFloat16Ty; diff --git a/clang/test/SemaCUDA/amdgpu-bf16.cu b/clang/test/SemaCUDA/amdgpu-bf16.cu new file mode 100644 --- /dev/null +++ b/clang/test/SemaCUDA/amdgpu-bf16.cu @@ -0,0 +1,16 @@ +// RUN: %clang_cc1 -x hip -fsyntax-only -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-linux-gnu -verify %s +// expected-no-diagnostics + +// If AMDGPU is the main target and X86 the aux target, ensure we +// don't complain about unsupported BF16 types in x86 code. + +#include "Inputs/cuda.h" + +__device__ void devicefn() { +} + +__bf16 hostfn(__bf16 a) { + return a; +} + +typedef __bf16 foo __attribute__((__vector_size__(16), __aligned__(16)));