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@-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__ "
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@-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__ "
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits