llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang Author: Luo, Yuanke (LuoYuanke) <details> <summary>Changes</summary> Variadice argument for NVPTX as been support in https://github.com/llvm/llvm-project/commit/486d00eca6b6ab470e8324b52cdf9f32023c1c9a We can enable it in front-end. --- Full diff: https://github.com/llvm/llvm-project/pull/161305.diff 2 Files Affected: - (modified) clang/lib/Sema/SemaExpr.cpp (+2-3) - (modified) clang/test/SemaCUDA/vararg.cu (+1-1) ``````````diff 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<FunctionDecl>(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 <stdarg.h> #include "Inputs/cuda.h" -__device__ void foo() { +__global__ void foo() { va_list list; va_arg(list, int); #ifdef EXPECT_VA_ARG_ERR `````````` </details> https://github.com/llvm/llvm-project/pull/161305 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
