https://github.com/AdityaSinha149 updated https://github.com/llvm/llvm-project/pull/218659
>From 936f77bb7de51dc71263029aaacb17f888bacfaf 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 +++-- .../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>>>(); _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
