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

The patch in D155211 added basic support for the `.alias` keyword in PTX. This means we should be able to permit use of this in clang. Reviewed By: tra Differential Revision: https://reviews.llvm.org/D156014
llvm · Aug 7, 2023 · 0ba9aec · 0ba9aec
1 parent f0e8bda
commit 0ba9aec
Show file tree

Hide file tree

Showing 3 changed files with 25 additions and 11 deletions.
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -8658,8 +8658,8 @@ def warn_kern_is_inline : Warning<
 def err_variadic_device_fn : Error<
   "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">;
+"CUDA device code does not support va_arg">;
+def err_alias_not_supported_on_nvptx : Error<"CUDA older than 10.0 does not support .alias">;
 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__ "

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -23,6 +23,7 @@
 #include "clang/AST/RecursiveASTVisitor.h"
 #include "clang/AST/Type.h"
 #include "clang/Basic/CharInfo.h"
+#include "clang/Basic/Cuda.h"
 #include "clang/Basic/DarwinSDKInfo.h"
 #include "clang/Basic/HLSLRuntime.h"
 #include "clang/Basic/LangOptions.h"
@@ -1992,8 +1993,12 @@ static void handleAliasAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
     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);
+    CudaVersion Version =
+        ToCudaVersion(S.Context.getTargetInfo().getSDKVersion());
+    if (Version != CudaVersion::UNKNOWN && Version < CudaVersion::CUDA_100)
+      S.Diag(AL.getLoc(), diag::err_alias_not_supported_on_nvptx);
   }
 
   // Aliases should be on declarations, not definitions.

diff --git a/clang/test/CodeGenCUDA/alias.cu b/clang/test/CodeGenCUDA/alias.cu
@@ -4,17 +4,26 @@
 
 // RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -emit-llvm \
 // RUN:   -o - %s | FileCheck %s
-// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn -emit-llvm \
+// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -emit-llvm -target-sdk-version=10.1 \
+// RUN:   -o - %s | FileCheck %s
+// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \
 // RUN:   -o - %s | FileCheck %s
 
 #include "Inputs/cuda.h"
 
-// Check that we don't generate an alias from "foo" to the mangled name for
-// ns::foo() -- nvptx doesn't support aliases.
-
-namespace ns {
 extern "C" {
-// CHECK-NOT: @foo = internal alias
-__device__ __attribute__((used)) static int foo() { return 0; }
-}
+__device__ int foo() { return 1; }
 }
+
+[[gnu::alias("foo")]] __device__ int alias();
+
+// CHECK: @_Z5aliasv = alias i32 (), ptr @foo
+//
+//      CHECK: define dso_local i32 @foo() #[[ATTR0:[0-9]+]] {
+// CHECK-NEXT: entry:
+//      CHECK:   ret i32 1
+// CHECK-NEXT: }
+
+// RUN: not %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -emit-llvm -target-sdk-version=9.0 \
+// RUN:   -o - %s 2>&1 | FileCheck %s --check-prefix=NO_SUPPORT
+// NO_SUPPORT: CUDA older than 10.0 does not support .alias