[CUDA][HIP] Don't diagnose use for __bf16

[Pierre-vh]

Summary

e0fb01e97b6b7d2fe66b17b36eeb98aa78c6e3bb 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 but for bf16 which prevents Clang from diagnosing uses of bf16 when compiling heterogenous applications.

[Pierre-vh]

Add newline at end of file

[Pierre-vh]

Need to fix a test crash

[Pierre-vh]

Not all targets have bf16 and AuxTarget may not be available all the time so I changed the condition slightly

[Pierre-vh]

Fixing condition, adding new test case

[tra]

Builds were failing because "__bf16" wasn't allowed on the target.

For CUDA/NVPTX we've solved the issue by implementing storage-only support for NVPTX: https://reviews.llvm.org/D136311
It's a bit more work, but I think it would be a better fix long-term for AMDGPU, too.

clang/lib/AST/ASTContext.cpp

2177

Is there a particular reason we're singling out OpenMP here and not HIP/CUDA?

2177

Nit: I'd use if (Target->hasBFloat16Type() && !(getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice))

clang/lib/Sema/SemaType.cpp

1521–1525

CUDA/NVPTX does have __bf16 support since https://reviews.llvm.org/rG0e8a414ab3d330ebb2996ec95d8141618ee0278b

On the other hand, OpenMP is special-cased, but is not mentioned.

clang/test/SemaCUDA/amdgpu-bf16.cu

2

The changes do affect OpenMP, but there are no OpenMP tests. I think we could use some.

[yaxunl]

In D138651#3954179, @tra wrote:

Builds were failing because "__bf16" wasn't allowed on the target.

For CUDA/NVPTX we've solved the issue by implementing storage-only support for NVPTX: https://reviews.llvm.org/D136311
It's a bit more work, but I think it would be a better fix long-term for AMDGPU, too.

Agree supportting bf16 in backend is a better solution. Limited by resource, overall this patch LGTM.

clang/test/SemaCUDA/amdgpu-bf16.cu

17

missing new line

[Pierre-vh]

• Recentering the patch around HIP only.
  • I was using too much from D57369 and was involving OpenMP when there's no reason to. Just checking if HIP is being used should be enough.

[tra]

clang/test/SemaCUDA/amdgpu-bf16.cu

9

We should probably also have a case verifying that actual attempt to use __bf16 in device code is still diagnosed.

[Pierre-vh]

clang/test/SemaCUDA/amdgpu-bf16.cu

9

Good catch, it's currently no longer diagnosed.
What can I use in ConvertDeclSpecToType to make it diagnose only if the current function is a device function? If I understand correctly, LangOpts are for the whole TU so I can't use that (e.g. CUDAIsDevice), right?

[tra]

clang/test/SemaCUDA/amdgpu-bf16.cu

9

I don't have a good answer. That was one of the reasons I ended up implementing __bf16 support in NVPTX. :-)
You may need to use deferred diagnostics, instead of disabling them) and trigger them if we need to codegen anything that uses __bf16.

[Pierre-vh]

I'll take a look at handling bf16 storage-only for AMDGPU. Looks like our Backend already handles it and converts it to i16 so maybe it'll be really easy.

[Pierre-vh]

Added bf16 storage support instead: D139398

clang/lib/AST/ASTContext.cpp

case Type::Builtin:
     case BuiltinType::LongFract:
     case BuiltinType::ULongFract:
     case BuiltinType::SatLongFract:
     case BuiltinType::SatULongFract:
       Width = Target->getLongFractWidth();
       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:
     case BuiltinType::Half:
       if (Target->hasFloat16Type() || !getLangOpts().OpenMP ||
           !getLangOpts().OpenMPIsDevice) {
         Width = Target->getHalfWidth();
         Align = Target->getHalfAlign();

clang/lib/Sema/SemaType.cpp

case DeclSpec::TST_float16:
     if (!S.Context.getTargetInfo().hasFloat16Type() && !S.getLangOpts().CUDA &&
         !(S.getLangOpts().OpenMP && S.getLangOpts().OpenMPIsDevice))
       S.Diag(DS.getTypeSpecTypeLoc(), diag::err_type_unsupported)
         << "_Float16";
     Result = Context.Float16Ty;
     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;
     break;
   case DeclSpec::TST_float:   Result = Context.FloatTy; break;
   case DeclSpec::TST_double:
     if (DS.getTypeSpecWidth() == TypeSpecifierWidth::Long)
       Result = Context.LongDoubleTy;

clang/test/SemaCUDA/amdgpu-bf16.cu

+// 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)));