Author: Aditya Sinha Date: 2026-09-02T13:30:12+05:30 New Revision: 058c7282003a68d6e2cb3bd923e7510c9433393a
URL: https://github.com/llvm/llvm-project/commit/058c7282003a68d6e2cb3bd923e7510c9433393a DIFF: https://github.com/llvm/llvm-project/commit/058c7282003a68d6e2cb3bd923e7510c9433393a.diff LOG: [clang][Sema] Allow null caller for HIP kernel launch in incremental mode (#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. Assisted by Claude Opus 4.8 Added: clang/test/SemaHIP/hip-incremental-toplevel-launch.hip Modified: clang/lib/Sema/SemaCUDA.cpp Removed: ################################################################################ diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp index 78cfa2c1a8662..553df2db01869 100644 --- a/clang/lib/Sema/SemaCUDA.cpp +++ b/clang/lib/Sema/SemaCUDA.cpp @@ -62,9 +62,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. @@ -74,6 +73,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
