[PATCH] D156014: [Clang][NVPTX] Permit use of the alias attribute for NVPTX targets

[Joseph Huber via Phabricator via cfe-commits]

jhuber6 created this revision.
jhuber6 added reviewers: tra, arsenm, jlebar, kushanam, aaron.ballman, yaxunl, jdoerfert.
Herald added subscribers: mattd, gchakrabarti, asavonic, jeroen.dobbelaere.
Herald added a project: All.
jhuber6 requested review of this revision.
Herald added subscribers: cfe-commits, wangpc, wdng.
Herald added a project: clang.

The patch in D155211 <https://reviews.llvm.org/D155211> added basic support for the `.alias` keyword in
PTX. This means we should be able to permit use of this in clang.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D156014

Files:
  clang/include/clang/Basic/DiagnosticSemaKinds.td
  clang/lib/Sema/SemaDeclAttr.cpp
  clang/test/SemaCUDA/alias.cu


Index: clang/test/SemaCUDA/alias.cu
===================================================================
--- clang/test/SemaCUDA/alias.cu
+++ /dev/null
@@ -1,11 +0,0 @@
-// RUN: %clang_cc1 -triple nvptx-unknown-cuda -fsyntax-only -fcuda-is-device -verify -DEXPECT_ERR %s
-// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s
-
-// The alias attribute is not allowed in CUDA device code.
-void bar();
-__attribute__((alias("bar"))) void foo();
-#ifdef EXPECT_ERR
-// expected-error at -2 {{CUDA does not support aliases}}
-#else
-// expected-no-diagnostics
-#endif
Index: clang/lib/Sema/SemaDeclAttr.cpp
===================================================================
--- clang/lib/Sema/SemaDeclAttr.cpp
+++ clang/lib/Sema/SemaDeclAttr.cpp
@@ -1992,9 +1992,6 @@
     S.Diag(AL.getLoc(), diag::err_alias_not_supported_on_darwin);
     return;
   }
-  if (S.Context.getTargetInfo().getTriple().isNVPTX()) {
-    S.Diag(AL.getLoc(), diag::err_alias_not_supported_on_nvptx);
-  }
 
   // Aliases should be on declarations, not definitions.
   if (const auto *FD = dyn_cast<FunctionDecl>(D)) {
Index: clang/include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -8646,7 +8646,6 @@
   "CUDA device code does not support variadic functions">;
 def err_va_arg_in_device : Error<
   "CUDA device code does not support va_arg">;
-def err_alias_not_supported_on_nvptx : Error<"CUDA does not support aliases">;
 def err_cuda_unattributed_constexpr_cannot_overload_device : Error<
   "constexpr function %0 without __host__ or __device__ attributes cannot "
   "overload __device__ function with same signature.  Add a __host__ "


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D156014.543138.patch
Type: text/x-patch
Size: 1805 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20230722/d15472de/attachment.bin>