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

[jhuber6]

Summary

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.

[tra]

clang/lib/Sema/SemaDeclAttr.cpp

1995

Allowing or not noreturn depends on the CUDA version we're building with (or rather on the PTX version we need for .noreturn instruction).

We would still need to issue the diagnostics if we're using CUDA older than 10.1.

[tra]

clang/lib/Sema/SemaDeclAttr.cpp

1995

Make it .alias and CUDA older than 10.0.

[jhuber6]

clang/lib/Sema/SemaDeclAttr.cpp

1995

Do we do any similar diagnostics checks on the CUDA version? I thought that was more of a clang driver thing and we'd just let the backend handle the failure, since we can emit LLVM-IR that can be compiled irrespective of the CUDA version used to make it.

[jhuber6]

clang/lib/Sema/SemaDeclAttr.cpp

1995

I checked and I don't think we pass in any CUDA version information to the -cc1 compiler. In this case if the user didn't have sufficient utilities it would simply fail in the backend or in PTX. We have semi-helpful messages there and it would be a good indicator to update CUDA. Is this fine given that?

[jhuber6]

ping

[tra]

clang/lib/Sema/SemaDeclAttr.cpp

1995

We do pass it via -target-sdk-version=... https://github.com/llvm/llvm-project/blob/1b74459df8a6d960f7387f0c8379047e42811f58/clang/lib/Driver/ToolChains/Clang.cpp#L4707

And then check with getSDKVersion. E.g. https://github.com/llvm/llvm-project/blob/1b74459df8a6d960f7387f0c8379047e42811f58/clang/lib/CodeGen/CGCUDANV.cpp#L317

[jhuber6]

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.

[steven_wu]

This breaks macOS bot: https://green.lab.llvm.org/green/job/clang-stage1-cmake-RA-incremental/36900/testReport/junit/Clang/SemaCUDA/alias_cu/

[jhuber6]

In D156014#4567363, @steven_wu wrote:

This breaks macOS bot: https://green.lab.llvm.org/green/job/clang-stage1-cmake-RA-incremental/36900/testReport/junit/Clang/SemaCUDA/alias_cu/

I should've fixed that already. Is it still broken?

[steven_wu]

In D156014#4567446, @jhuber6 wrote:

In D156014#4567363, @steven_wu wrote:

This breaks macOS bot: https://green.lab.llvm.org/green/job/clang-stage1-cmake-RA-incremental/36900/testReport/junit/Clang/SemaCUDA/alias_cu/

I should've fixed that already. Is it still broken?

9e99a4f0db0e21b68e9aab9ad6f32f34d42eb460 ? That is not yet built but looks promising. Thanks.

clang/include/clang/Basic/DiagnosticSemaKinds.td

def warn_kern_is_method : Extension<
   "kernel function %0 is a member function; this may not be accepted by nvcc">,
   InGroup<CudaCompat>;
 def warn_kern_is_inline : Warning<
   "ignored 'inline' attribute on kernel function %0">,
   InGroup<CudaCompat>;
 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">;
+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__ "
   "attribute, or build with -fno-cuda-host-device-constexpr.">;
 def note_cuda_conflicting_device_function_declared_here : Note<
   "conflicting __device__ function declared here">;
 def err_cuda_device_exceptions : Error<
   "cannot use '%0' in "

clang/lib/Sema/SemaDeclAttr.cpp

 #include "clang/AST/DeclObjC.h"
 #include "clang/AST/DeclTemplate.h"
 #include "clang/AST/Expr.h"
 #include "clang/AST/ExprCXX.h"
 #include "clang/AST/Mangle.h"
 #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"
 #include "clang/Basic/SourceLocation.h"
 #include "clang/Basic/SourceManager.h"
 #include "clang/Basic/TargetBuiltins.h"
 #include "clang/Basic/TargetInfo.h"
 #include "clang/Lex/Preprocessor.h"
static void handleAliasAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
   StringRef Str;
   if (!S.checkStringLiteralArgumentAttr(AL, 0, Str))
     return;
 
   if (S.Context.getTargetInfo().getTriple().isOSDarwin()) {
     S.Diag(AL.getLoc(), diag::err_alias_not_supported_on_darwin);
     return;
   }
+
   if (S.Context.getTargetInfo().getTriple().isNVPTX()) {
+    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.
   if (const auto *FD = dyn_cast<FunctionDecl>(D)) {
     if (FD->isThisDeclarationADefinition()) {
       S.Diag(AL.getLoc(), diag::err_alias_is_definition) << FD << 0;
       return;
     }

clang/test/CodeGenCUDA/alias.cu

 // REQUIRES: x86-registered-target
 // REQUIRES: nvptx-registered-target
 // REQUIRES: amdgpu-registered-target
 
 // 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