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

[via cfe-commits]

llvmorg-github-actions[bot] wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-clang

Author: Aditya Sinha (AdityaSinha149)

<details>
<summary>Changes</summary>

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.

---
Full diff: https://github.com/llvm/llvm-project/pull/218659.diff


2 Files Affected:

- (modified) clang/lib/Sema/SemaCUDA.cpp (+5-3) 
- (added) clang/test/SemaCUDA/hip-incremental-toplevel-launch.hip (+15) 


``````````diff
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>>>();

``````````

</details>


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