From db658c7032473a291afb5201b70a7f463719bff5 Mon Sep 17 00:00:00 2001 From: Yuanke Luo Date: Tue, 30 Sep 2025 10:04:46 +0800 Subject: [PATCH] [CUDA] Enable variadic argument support in front-end Variadice argument for NVPTX as been support in https://github.com/llvm/llvm-project/commit/486d00eca6b6ab470e8324b52cdf9f32023c1c9a We can enable it in front-end. --- clang/lib/Sema/SemaExpr.cpp | 5 ++--- clang/test/SemaCUDA/vararg.cu | 2 +- 2 files changed, 3 insertions(+), 4 deletions(-) diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index 3b267c1b1693d..2ab14cf0906a7 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -16791,12 +16791,11 @@ ExprResult Sema::BuildVAArgExpr(SourceLocation BuiltinLoc, Expr *OrigExpr = E; bool IsMS = false; - // CUDA device code does not support varargs. + // CUDA device global function does not support varargs. if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice) { if (const FunctionDecl *F = dyn_cast(CurContext)) { CUDAFunctionTarget T = CUDA().IdentifyTarget(F); - if (T == CUDAFunctionTarget::Global || T == CUDAFunctionTarget::Device || - T == CUDAFunctionTarget::HostDevice) + if (T == CUDAFunctionTarget::Global) return ExprError(Diag(E->getBeginLoc(), diag::err_va_arg_in_device)); } } diff --git a/clang/test/SemaCUDA/vararg.cu b/clang/test/SemaCUDA/vararg.cu index 34ef367d89820..0238f42dc40a9 100644 --- a/clang/test/SemaCUDA/vararg.cu +++ b/clang/test/SemaCUDA/vararg.cu @@ -10,7 +10,7 @@ #include #include "Inputs/cuda.h" -__device__ void foo() { +__global__ void foo() { va_list list; va_arg(list, int); #ifdef EXPECT_VA_ARG_ERR