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

[Joseph Huber via Phabricator via cfe-commits]

jhuber6 updated this revision to Diff 547871.
jhuber6 added a comment.

Update to check the SDK version. Permit this if there is not passed in SDK version so that freestanding targets can still target CUDA and assume it's supported.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D156014

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


Index: clang/test/CodeGenCUDA/alias.cu
===================================================================
--- clang/test/CodeGenCUDA/alias.cu
+++ 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
Index: clang/lib/Sema/SemaDeclAttr.cpp
===================================================================
--- clang/lib/Sema/SemaDeclAttr.cpp
+++ 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 @@
     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.
Index: clang/include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -8648,8 +8648,8 @@
 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__ "


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D156014.547871.patch
Type: text/x-patch
Size: 3278 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20230807/4c8e8895/attachment-0001.bin>