Pierre-vh created this revision.
Pierre-vh added reviewers: arsenm, rjmccall, tra.
Herald added subscribers: kosarev, mattd, kerbowa, pengfei, tpr, yaxunl,
jvesely.
Herald added a project: All.
Pierre-vh requested review of this revision.
Herald added subscribers: cfe-commits, sstefan1, wdng.
Herald added a reviewer: jdoerfert.
Herald added a project: clang.
e0fb01e97b6b7d2fe66b17b36eeb98aa78c6e3bb
<https://reviews.llvm.org/rGe0fb01e97b6b7d2fe66b17b36eeb98aa78c6e3bb> caused
issues in some of our HIP projects. Builds were failing because "__bf16" wasn't
allowed on the target. This is because in those cases, the main target is
AMDGPU (which doesn't have bf16), and the aux target is X86 (which has bf16).
This implements a fix similar to D57369 <https://reviews.llvm.org/D57369> but
for bf16 which prevents Clang from diagnosing uses of bf16 when compiling
heterogenous applications.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D138651
Files:
clang/lib/AST/ASTContext.cpp
clang/lib/Sema/SemaType.cpp
clang/test/SemaCUDA/amdgpu-bf16.cu
Index: clang/test/SemaCUDA/amdgpu-bf16.cu
===================================================================
--- /dev/null
+++ clang/test/SemaCUDA/amdgpu-bf16.cu
@@ -0,0 +1,14 @@
+// RUN: %clang_cc1 -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;
+}
\ No newline at end of file
Index: clang/lib/Sema/SemaType.cpp
===================================================================
--- clang/lib/Sema/SemaType.cpp
+++ clang/lib/Sema/SemaType.cpp
@@ -1518,7 +1518,9 @@
break;
case DeclSpec::TST_half: Result = Context.HalfTy; break;
case DeclSpec::TST_BFloat16:
- if (!S.Context.getTargetInfo().hasBFloat16Type())
+ // Likewise, CUDA host and device may have different __bf16 support.
+ if (!S.Context.getTargetInfo().hasBFloat16Type() && !S.getLangOpts().CUDA
&&
+ !(S.getLangOpts().OpenMP && S.getLangOpts().OpenMPIsDevice))
S.Diag(DS.getTypeSpecTypeLoc(), diag::err_type_unsupported)
<< "__bf16";
Result = Context.BFloat16Ty;
Index: clang/lib/AST/ASTContext.cpp
===================================================================
--- clang/lib/AST/ASTContext.cpp
+++ clang/lib/AST/ASTContext.cpp
@@ -2171,9 +2171,15 @@
Align = Target->getLongFractAlign();
break;
case BuiltinType::BFloat16:
- if (Target->hasBFloat16Type()) {
+ if (Target->hasBFloat16Type() || !getLangOpts().OpenMP ||
+ !getLangOpts().OpenMPIsDevice) {
Width = Target->getBFloat16Width();
Align = Target->getBFloat16Align();
+ } else {
+ assert(getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice &&
+ "Expected OpenMP device compilation.");
+ Width = AuxTarget->getBFloat16Width();
+ Align = AuxTarget->getBFloat16Align();
}
break;
case BuiltinType::Float16:
Index: clang/test/SemaCUDA/amdgpu-bf16.cu
===================================================================
--- /dev/null
+++ clang/test/SemaCUDA/amdgpu-bf16.cu
@@ -0,0 +1,14 @@
+// RUN: %clang_cc1 -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;
+}
\ No newline at end of file
Index: clang/lib/Sema/SemaType.cpp
===================================================================
--- clang/lib/Sema/SemaType.cpp
+++ clang/lib/Sema/SemaType.cpp
@@ -1518,7 +1518,9 @@
break;
case DeclSpec::TST_half: Result = Context.HalfTy; break;
case DeclSpec::TST_BFloat16:
- if (!S.Context.getTargetInfo().hasBFloat16Type())
+ // Likewise, CUDA host and device may have different __bf16 support.
+ if (!S.Context.getTargetInfo().hasBFloat16Type() && !S.getLangOpts().CUDA &&
+ !(S.getLangOpts().OpenMP && S.getLangOpts().OpenMPIsDevice))
S.Diag(DS.getTypeSpecTypeLoc(), diag::err_type_unsupported)
<< "__bf16";
Result = Context.BFloat16Ty;
Index: clang/lib/AST/ASTContext.cpp
===================================================================
--- clang/lib/AST/ASTContext.cpp
+++ clang/lib/AST/ASTContext.cpp
@@ -2171,9 +2171,15 @@
Align = Target->getLongFractAlign();
break;
case BuiltinType::BFloat16:
- if (Target->hasBFloat16Type()) {
+ if (Target->hasBFloat16Type() || !getLangOpts().OpenMP ||
+ !getLangOpts().OpenMPIsDevice) {
Width = Target->getBFloat16Width();
Align = Target->getBFloat16Align();
+ } else {
+ assert(getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice &&
+ "Expected OpenMP device compilation.");
+ Width = AuxTarget->getBFloat16Width();
+ Align = AuxTarget->getBFloat16Align();
}
break;
case BuiltinType::Float16:
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
