[CUDA] Remove CUDAAllowVariadicFunctions option and its sema check (#161350)

Variadic argument for NVPTX has been support in
486d00e
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]>

Author: LuoYuanke

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")

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">>;

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);

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
 }