[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<
   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_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

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()) {
-    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;
     }
   } else {

clang/test/SemaCUDA/alias.cu

-// 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@-2 {{CUDA does not support aliases}}
-#else
-// expected-no-diagnostics
-#endif