[PATCH] D138651: [CUDA][HIP] Don't diagnose use for __bf16

[Pierre van Houtryve via Phabricator via cfe-commits]

Pierre-vh updated this revision to Diff 478823.
Pierre-vh marked 5 inline comments as done.
Pierre-vh added a comment.

- Recentering the patch around HIP only.
  - I was using too much from D57369 <https://reviews.llvm.org/D57369> and was involving OpenMP when there's no reason to. Just checking if HIP is being used should be enough.


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D138651/new/

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,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)));
Index: clang/lib/Sema/SemaType.cpp
===================================================================
--- clang/lib/Sema/SemaType.cpp
+++ 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;
Index: clang/lib/AST/ASTContext.cpp
===================================================================
--- clang/lib/AST/ASTContext.cpp
+++ 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:


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D138651.478823.patch
Type: text/x-patch
Size: 2131 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20221130/d5359f26/attachment-0001.bin>