[CUDA] Enable variadic argument support in front-end

[LuoYuanke]

Variadice argument for NVPTX as been support in
486d00e
We can enable it in front-end.

[llvmbot]

@llvm/pr-subscribers-clang

Author: Luo, Yuanke (LuoYuanke)

Changes

Variadice argument for NVPTX as been support in
486d00e
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 --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

[jhuber6]

There are some flags like -fcuda-enable-variadic that we pass in places. These should likely be deprecated as the variadic support on GPUs has been well tested for over a year now. Just keep in mind the ABI is different between AMDGPU and NVPTX.

[LuoYuanke]

There are some flags like -fcuda-enable-variadic that we pass in places. These should likely be deprecated as the variadic support on GPUs has been well tested for over a year now. Just keep in mind the ABI is different between AMDGPU and NVPTX.

I can help to remove the code for sema check on variadic declaration and deprecate -fcuda-enable-variadic option with another patch. I don't have commit access. Could you help to land this patch?