23e4fe040b
The `sycl_kernel_entry_point` attribute facilitates the generation of an
offload kernel entry point function based on the parameters and body
of the attributed function. This change extends the behavior of that
attribute to support integration with a SYCL runtime library through
an interface that communicates symbol names and kernel arguments
for the generated offload kernel entry point functions.
Consider the following function declared with the
`sycl_kernel_entry_point` attribute with a call to this function
occurring in the implementation of a SYCL kernel invocation function
such as `sycl::handler::single_task()`.
```c++
template<typename KernelName, typename KernelType>
[[clang::sycl_kernel_entry_point(KernelName)]]
void kernel_entry_point(KernelType kernel) {
kernel();
}
```
The body of the above function specifies the parameters and body of the
generated offload kernel entry point. Clearly, a call to the above
function by a SYCL kernel invocation function is not intended to execute
the body as written. Previously, code generation emitted an empty
function body so that calls to the function had no effect other than to
trigger the generation of the offload kernel entry point. The function
body is therefore available to hook for SYCL library support and is now
substituted with a call to a (SYCL library provided) function template
or variable template named `sycl_kernel_launch()` with the kernel
name type passed as the first template argument, the symbol name
of the offload kernel entry point passed as a string literal for the first
function argument, and the function parameters passed as the
remaining explicit function arguments. Given a call like this:
```c++
kernel_entry_point<struct KN>([]{})
```
the body of the instantiated `kernel_entry_point()` specialization would
be substituted as follows with "kernel-symbol-name" substituted for the
generated symbol name and `kernel` forwarded.
```c++
sycl_kernel_launch<KN>("kernel-symbol-name", kernel)
```
Name lookup and overload resolution for the `sycl_kernel_launch()`
function is performed at the point of definition of the
`sycl_kernel_entry_point` attributed function (or the point of
instantiation for an instantiated function template specialization). If
overload resolution fails, the program is ill-formed.
Implementation of the `sycl_kernel_launch()` function might require
additional information provided by the SYCL library. This is facilitated
by removing the previous prohibition against use of the
`sycl_kernel_entry_point` attribute with a non-static member function.
If the `sycl_kernel_entry_point` attributed function is a non-static
member function, then overload resolution for the `sycl_kernel_launch()`
function template may select a non-static member function in which case,
`this` will be implicitly passed as the implicit object argument.
If a `sycl_kernel_entry_point` attributed function is a non-static
member function, use of `this` in a potentially evaluated expression is
prohibited in the definition since `this` is not a kernel argument and
will not be available within the generated offload kernel entry point
function. The attribute cannot be applied to a function with an
explicit object parameter.
---------
Co-authored-by: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>
86 lines
3.6 KiB
C++
86 lines
3.6 KiB
C++
//===--------- CodeGenSYCL.cpp - Code for SYCL kernel generation ----------===//
|
|
//
|
|
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
|
|
// See https://llvm.org/LICENSE.txt for license information.
|
|
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
|
|
//
|
|
//===----------------------------------------------------------------------===//
|
|
//
|
|
// This contains code required for generation of SYCL kernel caller offload
|
|
// entry point functions.
|
|
//
|
|
//===----------------------------------------------------------------------===//
|
|
|
|
#include "CodeGenFunction.h"
|
|
#include "CodeGenModule.h"
|
|
#include <cassert>
|
|
|
|
using namespace clang;
|
|
using namespace CodeGen;
|
|
|
|
void CodeGenFunction::EmitSYCLKernelCallStmt(const SYCLKernelCallStmt &S) {
|
|
// SYCLKernelCallStmt instances are only injected in the definitions of
|
|
// functions declared with the sycl_kernel_entry_point attribute. ODR-use of
|
|
// such a function in code emitted during device compilation should be
|
|
// diagnosed. Thus, any attempt to emit a SYCLKernelCallStmt during device
|
|
// compilation indicates a missing diagnostic.
|
|
assert(!getLangOpts().SYCLIsDevice &&
|
|
"Attempt to emit a SYCL kernel call statement during device"
|
|
" compilation");
|
|
EmitStmt(S.getKernelLaunchStmt());
|
|
}
|
|
|
|
static void SetSYCLKernelAttributes(llvm::Function *Fn, CodeGenFunction &CGF) {
|
|
// SYCL 2020 device language restrictions require forward progress and
|
|
// disallow recursion.
|
|
Fn->setDoesNotRecurse();
|
|
if (CGF.checkIfFunctionMustProgress())
|
|
Fn->addFnAttr(llvm::Attribute::MustProgress);
|
|
}
|
|
|
|
void CodeGenModule::EmitSYCLKernelCaller(const FunctionDecl *KernelEntryPointFn,
|
|
ASTContext &Ctx) {
|
|
assert(Ctx.getLangOpts().SYCLIsDevice &&
|
|
"SYCL kernel caller offload entry point functions can only be emitted"
|
|
" during device compilation");
|
|
|
|
const auto *KernelEntryPointAttr =
|
|
KernelEntryPointFn->getAttr<SYCLKernelEntryPointAttr>();
|
|
assert(KernelEntryPointAttr && "Missing sycl_kernel_entry_point attribute");
|
|
assert(!KernelEntryPointAttr->isInvalidAttr() &&
|
|
"sycl_kernel_entry_point attribute is invalid");
|
|
|
|
// Find the SYCLKernelCallStmt.
|
|
SYCLKernelCallStmt *KernelCallStmt =
|
|
cast<SYCLKernelCallStmt>(KernelEntryPointFn->getBody());
|
|
|
|
// Retrieve the SYCL kernel caller parameters from the OutlinedFunctionDecl.
|
|
FunctionArgList Args;
|
|
const OutlinedFunctionDecl *OutlinedFnDecl =
|
|
KernelCallStmt->getOutlinedFunctionDecl();
|
|
Args.append(OutlinedFnDecl->param_begin(), OutlinedFnDecl->param_end());
|
|
|
|
// Compute the function info and LLVM function type.
|
|
const CGFunctionInfo &FnInfo =
|
|
getTypes().arrangeDeviceKernelCallerDeclaration(Ctx.VoidTy, Args);
|
|
llvm::FunctionType *FnTy = getTypes().GetFunctionType(FnInfo);
|
|
|
|
// Retrieve the generated name for the SYCL kernel caller function.
|
|
CanQualType KernelNameType =
|
|
Ctx.getCanonicalType(KernelEntryPointAttr->getKernelName());
|
|
const SYCLKernelInfo &KernelInfo = Ctx.getSYCLKernelInfo(KernelNameType);
|
|
auto *Fn = llvm::Function::Create(FnTy, llvm::Function::ExternalLinkage,
|
|
KernelInfo.GetKernelName(), &getModule());
|
|
|
|
// Emit the SYCL kernel caller function.
|
|
CodeGenFunction CGF(*this);
|
|
SetLLVMFunctionAttributes(GlobalDecl(), FnInfo, Fn, false);
|
|
SetSYCLKernelAttributes(Fn, CGF);
|
|
CGF.StartFunction(GlobalDecl(), Ctx.VoidTy, Fn, FnInfo, Args,
|
|
SourceLocation(), SourceLocation());
|
|
CGF.EmitFunctionBody(OutlinedFnDecl->getBody());
|
|
setDSOLocal(Fn);
|
|
SetLLVMFunctionAttributesForDefinition(cast<Decl>(OutlinedFnDecl), Fn);
|
|
CGF.FinishFunction();
|
|
}
|