[clang] [clang][Sema] Allow null caller for HIP kernel launch in incremental mode (PR #218659)

[Aditya Sinha via cfe-commits]

https://github.com/AdityaSinha149 created https://github.com/llvm/llvm-project/pull/218659

This PR allows a HIP kernel to be launched without a caller in incremental mode. This is necessary because in incremental mode(clang-repl), there isn't a `main` or any function that calls the kernel.

Also adds a Sema test that compiles a top-level HIP kernel launch in incremental device mode and verifies it is accepted.

>From c528473a068b062d1e7f42d3086c7d0568b5c362 Mon Sep 17 00:00:00 2001
From: AdityaSinha149 <adsinha at amd.com>
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..2f09baba2405a
--- /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 -fcuda-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>>>();