[CUDA] Remove sema check of function declaration with variadic argument #161350
Conversation
llvmbot
added
clang
|
@llvm/pr-subscribers-clang Author: Luo, Yuanke (LuoYuanke) ChangesVariadic argument for NVPTX has been support in Full diff: https://github.com/llvm/llvm-project/pull/161350.diff 3 Files Affected:
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
}
|
Variadic argument for NVPTX has been support in llvm@486d00e We can remove the sema check in front-end.
LuoYuanke
force-pushed
the
varadic-arg-decl
branch
from
e2001a6 to
708e050
Compare
There was a problem hiding this comment.
I'd expect this to completely remove usage of CUDAAllowVariadicFunctions and remove the associated driver handling code. That way passing it will result in the generic 'argument unused' warning. Though, I don't know if there was anything special about the old handling, the 'new' handling is ABI compatible with it, so I'd assume not but just in case I'll refer to @Artem-B
|
This seems legitimate to me. The IR pass is enabled on nvptx and turns LLVM variadics into functions which match the PTX variadic calling convention as far as I know. I think the guard is still there on CUDA because Joseph and I were mostly thinking in terms of freestanding C++ when we wrote this. |
If we completely remove |
The driver and |
|
I'd expect the -f flag behavior to be independent of what the backend supports. What does nvcc do? I'd expect it to require an opt-in or language standard flag |
|
Pretty sure nvcc CUDA handles it by default https://godbolt.org/z/abqWz6ans |
|
ABI compatibility is likely not a major issue as most of CUDA compilations consist of a single TU, and the relatively rare RDC compilations are usually withing the same library compiled with the same compiler. Before the new clang driver showed up, there were virtually no users who'd be affected by the ABI break, as RDC compilation with clang required a lot of additional external build changes. I'm aware of only one actual use case (build of NCCL inside of XLA https://github.com/openxla/xla/blob/45947e2a819102aef04454fc24f311f8a50e1c6a/third_party/nccl/build_defs.bzl.tpl#L232) and that's the "one library, one compiler" scenario that's not affected by the ABI change.
It evolved over time. Initially it allowed accepting variadic functions during parsing, but did not allow to generate any code for them. This was needed as some CUDA code relied on variadic function declarations as a wildcard for the template instantiations. |
|
There is no abi break here as far as I can tell. The variadic lowering pass lays things out the same way cuda's printf works and matches the documentation on ptx. As in if you've written a variadic function in ptx, calling it from IR is expected to work. Also, it's going from refuse to compile to does compile, so there's nominally no pre-existing code to be compatible with. There may be bugs in the lowering pass. I really should revive the x86/aarch64 implementations and/or chase wasm to get it running on non-gpu targets. |
Yes, I think that's what this PR does. |
Am I missing something, we still have |
I understand now. Revised. |
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/27/builds/17063 Here is the relevant piece of the build log for the reference |
…lvm#161350) Variadic argument for NVPTX has been support in llvm@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]>
Variadic argument for NVPTX has been support in
486d00e
We can remove the sema check in front-end.