Author: Luo, Yuanke Date: 2025-10-06T14:07:08+08:00 New Revision: a406eb460c05e1171971ed1dace2546e3901eb61
URL: https://github.com/llvm/llvm-project/commit/a406eb460c05e1171971ed1dace2546e3901eb61 DIFF: https://github.com/llvm/llvm-project/commit/a406eb460c05e1171971ed1dace2546e3901eb61.diff LOG: [CUDA] Remove CUDAAllowVariadicFunctions option and its sema check (#161350) Variadic argument for NVPTX has been support in https://github.com/llvm/llvm-project/commit/486d00eca6b6ab470e8324b52cdf9f32023c1c9a We can remove `CUDAAllowVariadicFunctions` option and its sema check. The CC1 option `fcuda_allow_variadic_functions` is retained to not break the existing code building. --------- Co-authored-by: Yuanke Luo <[email protected]> Added: Modified: clang/include/clang/Basic/LangOptions.def clang/include/clang/Driver/Options.td clang/lib/Sema/SemaDecl.cpp clang/test/SemaCUDA/vararg.cu Removed: ################################################################################ diff --git a/clang/include/clang/Basic/LangOptions.def b/clang/include/clang/Basic/LangOptions.def index 84f5ab3443a59..9e850089ad87f 100644 --- a/clang/include/clang/Basic/LangOptions.def +++ b/clang/include/clang/Basic/LangOptions.def @@ -245,7 +245,6 @@ LANGOPT(HLSLStrictAvailability, 1, 0, NotCompatible, LANGOPT(HLSLSpvUseUnknownImageFormat, 1, 0, NotCompatible, "For storage images and texel buffers, sets the default format to 'Unknown' when not specified via the `vk::image_format` attribute. If this option is not used, the format is inferred from the resource's data type.") LANGOPT(CUDAIsDevice , 1, 0, NotCompatible, "compiling for CUDA device") -LANGOPT(CUDAAllowVariadicFunctions, 1, 0, NotCompatible, "allowing variadic functions in CUDA device code") LANGOPT(CUDAHostDeviceConstexpr, 1, 1, NotCompatible, "treating unattributed constexpr functions as __host__ __device__") LANGOPT(GPUDeviceApproxTranscendentals, 1, 0, NotCompatible, "using approximate transcendental functions") LANGOPT(GPURelocatableDeviceCode, 1, 0, NotCompatible, "generate relocatable device code") diff --git a/clang/include/clang/Driver/Options.td b/clang/include/clang/Driver/Options.td index 5a48f0bcf65e5..9bfa1dd52effe 100644 --- a/clang/include/clang/Driver/Options.td +++ b/clang/include/clang/Driver/Options.td @@ -8733,8 +8733,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.">, - MarshallingInfoFlag<LangOpts<"CUDAAllowVariadicFunctions">>; + HelpText<"Deprecated; Allow variadic functions in CUDA device code.">; def fno_cuda_host_device_constexpr : Flag<["-"], "fno-cuda-host-device-constexpr">, HelpText<"Don't treat unattributed constexpr functions as __host__ __device__.">, MarshallingInfoNegativeFlag<LangOpts<"CUDAHostDeviceConstexpr">>; diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 0069b08f1991a..6eaf7b9435491 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 } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
