https://github.com/LuoYuanke updated https://github.com/llvm/llvm-project/pull/161350
>From 708e050a5e4097e505f0215be67f7bbb42749763 Mon Sep 17 00:00:00 2001 From: Yuanke Luo <[email protected]> Date: Tue, 30 Sep 2025 18:39:15 +0800 Subject: [PATCH 1/2] [CUDA] Remove sema check of function declaration with variadic argument Variadic argument for NVPTX has been support in https://github.com/llvm/llvm-project/commit/486d00eca6b6ab470e8324b52cdf9f32023c1c9a We can remove the sema check in front-end. --- clang/include/clang/Driver/Options.td | 2 +- clang/lib/Sema/SemaDecl.cpp | 11 ----------- clang/test/SemaCUDA/vararg.cu | 25 +++++-------------------- 3 files changed, 6 insertions(+), 32 deletions(-) diff --git a/clang/include/clang/Driver/Options.td b/clang/include/clang/Driver/Options.td index 6245cf33a0719..323ffea5afa59 100644 --- a/clang/include/clang/Driver/Options.td +++ b/clang/include/clang/Driver/Options.td @@ -8729,7 +8729,7 @@ def fcuda_include_gpubinary : Separate<["-"], "fcuda-include-gpubinary">, HelpText<"Incorporate CUDA device-side binary into host object file.">, MarshallingInfoString<CodeGenOpts<"CudaGpuBinaryFileName">>; def fcuda_allow_variadic_functions : Flag<["-"], "fcuda-allow-variadic-functions">, - HelpText<"Allow variadic functions in CUDA device code.">, + HelpText<"Deprecated; Allow variadic functions in CUDA device code.">, MarshallingInfoFlag<LangOpts<"CUDAAllowVariadicFunctions">>; def fno_cuda_host_device_constexpr : Flag<["-"], "fno-cuda-host-device-constexpr">, HelpText<"Don't treat unattributed constexpr functions as __host__ __device__.">, diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 9ef7a2698913d..357af2a50e75b 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -11041,17 +11041,6 @@ Sema::ActOnFunctionDeclarator(Scope *S, Declarator &D, DeclContext *DC, << CUDA().getConfigureFuncName(); Context.setcudaConfigureCallDecl(NewFD); } - - // Variadic functions, other than a *declaration* of printf, are not allowed - // in device-side CUDA code, unless someone passed - // -fcuda-allow-variadic-functions. - if (!getLangOpts().CUDAAllowVariadicFunctions && NewFD->isVariadic() && - (NewFD->hasAttr<CUDADeviceAttr>() || - NewFD->hasAttr<CUDAGlobalAttr>()) && - !(II && II->isStr("printf") && NewFD->isExternC() && - !D.isFunctionDefinition())) { - Diag(NewFD->getLocation(), diag::err_variadic_device_fn); - } } MarkUnusedFileScopedDecl(NewFD); diff --git a/clang/test/SemaCUDA/vararg.cu b/clang/test/SemaCUDA/vararg.cu index 0238f42dc40a9..62693e1d4a0af 100644 --- a/clang/test/SemaCUDA/vararg.cu +++ b/clang/test/SemaCUDA/vararg.cu @@ -1,11 +1,9 @@ // REQUIRES: x86-registered-target // REQUIRES: nvptx-registered-target // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -fsyntax-only \ -// RUN: -verify -DEXPECT_VA_ARG_ERR -DEXPECT_VARARG_ERR %s +// RUN: -verify -DEXPECT_VA_ARG_ERR %s // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -fsyntax-only \ // RUN: -fcuda-allow-variadic-functions -verify -DEXPECT_VA_ARG_ERR %s -// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify \ -// RUN: -DEXPECT_VARARG_ERR %s #include <stdarg.h> #include "Inputs/cuda.h" @@ -30,28 +28,15 @@ __device__ void baz() { #endif } -__device__ void vararg(const char* x, ...) {} -#ifdef EXPECT_VARARG_ERR -// expected-error@-2 {{CUDA device code does not support variadic functions}} -#endif +__device__ void vararg(const char* x, ...) {} // OK template <typename T> -__device__ void vararg(T t, ...) {} -#ifdef EXPECT_VARARG_ERR -// expected-error@-2 {{CUDA device code does not support variadic functions}} -#endif +__device__ void vararg(T t, ...) {} // OK extern "C" __device__ int printf(const char* fmt, ...); // OK, special case. -// Definition of printf not allowed. -extern "C" __device__ int printf(const char* fmt, ...) { return 0; } -#ifdef EXPECT_VARARG_ERR -// expected-error@-2 {{CUDA device code does not support variadic functions}} -#endif +extern "C" __device__ int printf(const char* fmt, ...) { return 0; } // OK namespace ns { -__device__ int printf(const char* fmt, ...); -#ifdef EXPECT_VARARG_ERR -// expected-error@-2 {{CUDA device code does not support variadic functions}} -#endif +__device__ int printf(const char* fmt, ...); // OK } >From 8a69808c5bd0b49bae29f3ab8025cc5182c38f6d Mon Sep 17 00:00:00 2001 From: Yuanke Luo <[email protected]> Date: Sun, 5 Oct 2025 10:37:59 +0800 Subject: [PATCH 2/2] [Clang] Add hidden flag --- clang/include/clang/Driver/Options.td | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/include/clang/Driver/Options.td b/clang/include/clang/Driver/Options.td index 323ffea5afa59..3cfb36c33ad55 100644 --- a/clang/include/clang/Driver/Options.td +++ b/clang/include/clang/Driver/Options.td @@ -8729,7 +8729,7 @@ def fcuda_include_gpubinary : Separate<["-"], "fcuda-include-gpubinary">, HelpText<"Incorporate CUDA device-side binary into host object file.">, MarshallingInfoString<CodeGenOpts<"CudaGpuBinaryFileName">>; def fcuda_allow_variadic_functions : Flag<["-"], "fcuda-allow-variadic-functions">, - HelpText<"Deprecated; Allow variadic functions in CUDA device code.">, + HelpText<"Deprecated; Allow variadic functions in CUDA device code.">, Flags<[HelpHidden]>, MarshallingInfoFlag<LangOpts<"CUDAAllowVariadicFunctions">>; def fno_cuda_host_device_constexpr : Flag<["-"], "fno-cuda-host-device-constexpr">, HelpText<"Don't treat unattributed constexpr functions as __host__ __device__.">, _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
