https://github.com/AdityaSinha149 updated https://github.com/llvm/llvm-project/pull/218659
>From bb98a18c321808d5b7501119d2c97e817b04c044 Mon Sep 17 00:00:00 2001 From: AdityaSinha149 <[email protected]> Date: Tue, 25 Aug 2026 11:29:00 +0530 Subject: [PATCH] [clang][Sema] Allow null caller for HIP kernel launch in incremental mode --- clang/lib/Sema/SemaCUDA.cpp | 8 +++++--- .../SemaCUDA/hip-incremental-toplevel-launch.hip | 15 +++++++++++++++ 2 files changed, 20 insertions(+), 3 deletions(-) create mode 100644 clang/test/SemaCUDA/hip-incremental-toplevel-launch.hip diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp index a2a088ab7c3ab..29cacd6cf9af9 100644 --- a/clang/lib/Sema/SemaCUDA.cpp +++ b/clang/lib/Sema/SemaCUDA.cpp @@ -61,9 +61,8 @@ ExprResult SemaCUDA::ActOnExecConfigExpr(Scope *S, SourceLocation LLLLoc, case CUDAFunctionTarget::HostDevice: if (getLangOpts().CUDAIsDevice) { IsDeviceKernelCall = true; - if (FunctionDecl *Caller = - SemaRef.getCurFunctionDecl(/*AllowLambda=*/true); - Caller && isImplicitHostDeviceFunction(Caller)) { + FunctionDecl *Caller = SemaRef.getCurFunctionDecl(/*AllowLambda=*/true); + if (Caller && isImplicitHostDeviceFunction(Caller)) { // Under the device compilation, config call under an HD function should // be treated as a device kernel call. But, for implicit HD ones (such // as lambdas), need to check whether RDC is enabled or not. @@ -73,6 +72,9 @@ ExprResult SemaCUDA::ActOnExecConfigExpr(Scope *S, SourceLocation LLLLoc, // the host-side kernel call. if (getLangOpts().HIP) IsDeviceKernelCall = false; + } else if (!Caller && getLangOpts().HIP && + getLangOpts().IncrementalExtensions) { + IsDeviceKernelCall = false; } } break; diff --git a/clang/test/SemaCUDA/hip-incremental-toplevel-launch.hip b/clang/test/SemaCUDA/hip-incremental-toplevel-launch.hip new file mode 100644 index 0000000000000..c21eaa30014d2 --- /dev/null +++ b/clang/test/SemaCUDA/hip-incremental-toplevel-launch.hip @@ -0,0 +1,15 @@ +// In incremental mode (clang-repl) statements may appear at the top level, so a +// HIP kernel launch can have no enclosing caller function. During device +// compilation such a launch must be accepted (treated as a host-side launch) +// rather than rejected as an unsupported device-side kernel launch. + +// RUN: %clang_cc1 -fsyntax-only -triple amdgcn-amd-amdhsa -fcud a-is-device -x hip -fincremental-extensions -verify %s + +// expected-no-diagnostics + +#include "Inputs/cuda.h" + +__global__ void kernel() {} + +// Top-level kernel launch with no surrounding function, i.e. a null caller. +kernel<<<1, 1>>>(); _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
