https://github.com/AdityaSinha149 updated https://github.com/llvm/llvm-project/pull/218659
>From b0ef43425df8e47ba586a57def57c9db8fd11ddc 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 | 34 +++++++++++++++++++ 2 files changed, 39 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..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/SemaCUDA/hip-incremental-toplevel-launch.hip b/clang/test/SemaCUDA/hip-incremental-toplevel-launch.hip new file mode 100644 index 0000000000000..2ac46c01c0f57 --- /dev/null +++ b/clang/test/SemaCUDA/hip-incremental-toplevel-launch.hip @@ -0,0 +1,34 @@ +// 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: cat %s | clang-repl --hip | FileCheck %s + +#include <hip/hip_runtime.h> + +extern "C" int printf(const char*, ...); + +__global__ void set_kernel(int* value, int x) { *value = x; } +__global__ void inc_kernel(int* value) { *value += 1; } + +int var; +int* devptr = nullptr; +printf("hipMalloc: %d\n", hipMalloc((void **) &devptr, sizeof(int))); +// CHECK: hipMalloc: 0 + +set_kernel<<<1,1>>>(devptr, 42); +printf("HIP Error: %d\n", hipGetLastError()); +// CHECK-NEXT: HIP Error: 0 + +inc_kernel<<<1,1>>>(devptr); +printf("HIP Error: %d\n", hipGetLastError()); +// CHECK-NEXT: HIP Error: 0 + +printf("hipMemcpy: %d\n", hipMemcpy(&var, devptr, sizeof(int), hipMemcpyDeviceToHost)); +// CHECK-NEXT: hipMemcpy: 0 + +printf("Value: %d\n", var); +// CHECK-NEXT: Value: 43 + +%quit _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
