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

[Aditya Sinha via cfe-commits]

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

>From 936f77bb7de51dc71263029aaacb17f888bacfaf 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 +++--
 .../hip-incremental-toplevel-launch.hip       | 29 +++++++++++++++++++
 2 files changed, 34 insertions(+), 3 deletions(-)
 create mode 100644 clang/test/SemaHIP/hip-incremental-toplevel-launch.hip

diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index a2a088ab7c3ab..684f7a8db24cf 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 (getLangOpts().HIP && getLangOpts().IncrementalExtensions &&
+                 isa<TopLevelStmtDecl>(SemaRef.getCurLexicalContext())) {
+        IsDeviceKernelCall = false;
       }
     }
     break;
diff --git a/clang/test/SemaHIP/hip-incremental-toplevel-launch.hip b/clang/test/SemaHIP/hip-incremental-toplevel-launch.hip
new file mode 100644
index 0000000000000..53945587584ef
--- /dev/null
+++ b/clang/test/SemaHIP/hip-incremental-toplevel-launch.hip
@@ -0,0 +1,29 @@
+// 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
+
+#define __global__ __attribute__((global))
+#define __host__ __attribute__((host))
+#define __device__ __attribute__((device))
+
+typedef unsigned long size_t;
+
+struct dim3 {
+  unsigned x, y, z;
+  __host__ __device__ dim3(unsigned x, unsigned y = 1, unsigned z = 1)
+      : x(x), y(y), z(z) {}
+};
+
+typedef struct hipStream *hipStream_t;
+typedef enum hipError {} hipError_t;
+int hipConfigureCall(dim3 gridSize, dim3 blockSize, size_t sharedSize = 0,
+                     hipStream_t stream = 0);
+
+__global__ void kernel() {}
+
+// Top-level kernel launch with no surrounding function, i.e. a null caller.
+kernel<<<1, 1>>>();