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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
