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>
96 lines
3.7 KiB
C++
96 lines
3.7 KiB
C++
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fcxx-exceptions -fexceptions -fsycl-is-host -emit-llvm -o - %s | FileCheck %s
|
|
|
|
// Validate generation of exception handling code for functions declared
|
|
// with the sycl_kernel_entry_point attribute that implicitly call a
|
|
// sycl_kernel_launch function that may throw an exception. Exception
|
|
// handling is not relevant for the generated offload kernel entry point
|
|
// function, so device compilation is intentionally not exercised.
|
|
|
|
// A unique kernel name type is required for each declared kernel entry point.
|
|
template<int> struct KN;
|
|
|
|
// A generic kernel object type.
|
|
template<int, int = 0>
|
|
struct KT {
|
|
void operator()() const;
|
|
};
|
|
|
|
|
|
// Validate that exception handling instructions are omitted when a
|
|
// potentially throwing sycl_kernel_entry_point attributed function
|
|
// calls a potentially throwing sycl_kernel_launch function (a thrown
|
|
// exception will propagate with no explicit handling required).
|
|
namespace ns1 {
|
|
template<typename KN, typename... Ts>
|
|
void sycl_kernel_launch(const char *, Ts...);
|
|
[[clang::sycl_kernel_entry_point(KN<1>)]]
|
|
void skep(KT<1> k) {
|
|
k();
|
|
}
|
|
}
|
|
// CHECK: ; Function Attrs: mustprogress noinline optnone
|
|
// CHECK: define dso_local void @_ZN3ns14skepE2KTILi1ELi0EE() #{{[0-9]+}} {
|
|
// CHECK: call void @_ZN3ns118sycl_kernel_launchI2KNILi1EEJ2KTILi1ELi0EEEEEvPKcDpT0_(ptr noundef @.str)
|
|
// CHECK: ret void
|
|
// CHECK: }
|
|
|
|
|
|
// Validate that exception handling instructions are emitted when a
|
|
// non-throwing sycl_kernel_entry_point attributed function calls
|
|
// a potentially throwing sycl_kernel_launch function.
|
|
namespace ns2 {
|
|
template<typename KN, typename... Ts>
|
|
void sycl_kernel_launch(const char *, Ts...);
|
|
[[clang::sycl_kernel_entry_point(KN<2>)]]
|
|
void skep(KT<2> k) noexcept {
|
|
k();
|
|
}
|
|
}
|
|
// CHECK: ; Function Attrs: mustprogress noinline nounwind optnone
|
|
// CHECK: define dso_local void @_ZN3ns24skepE2KTILi2ELi0EE() #{{[0-9]+}} personality ptr @__gxx_personality_v0 {
|
|
// CHECK: invoke void @_ZN3ns218sycl_kernel_launchI2KNILi2EEJ2KTILi2ELi0EEEEEvPKcDpT0_(ptr noundef @.str.1)
|
|
// CHECK: to label %invoke.cont unwind label %terminate.lpad
|
|
// CHECK: invoke.cont:
|
|
// CHECK: ret void
|
|
// CHECK: terminate.lpad:
|
|
// CHECK: call void @__clang_call_terminate(ptr %1) #{{[0-9]+}}
|
|
// CHECK: unreachable
|
|
// CHECK: }
|
|
|
|
|
|
// Validate that exception handling instructions are omitted when a
|
|
// potentially throwing sycl_kernel_entry_point attributed function
|
|
// calls a non-throwing sycl_kernel_launch function (a thrown
|
|
// exception will terminate within sycl_kernel_launch).
|
|
namespace ns3 {
|
|
template<typename KN, typename... Ts>
|
|
void sycl_kernel_launch(const char *, Ts...) noexcept;
|
|
[[clang::sycl_kernel_entry_point(KN<3>)]]
|
|
void skep(KT<3> k) {
|
|
k();
|
|
}
|
|
}
|
|
// CHECK: ; Function Attrs: mustprogress noinline nounwind optnone
|
|
// CHECK: define dso_local void @_ZN3ns34skepE2KTILi3ELi0EE() #{{[0-9]+}} {
|
|
// CHECK: call void @_ZN3ns318sycl_kernel_launchI2KNILi3EEJ2KTILi3ELi0EEEEEvPKcDpT0_(ptr noundef @.str.2)
|
|
// CHECK: ret void
|
|
// CHECK: }
|
|
|
|
|
|
// Validate that exception handling instructions are omitted when a
|
|
// non-throwing sycl_kernel_entry_point attributed function calls a
|
|
// non-throwing sycl_kernel_launch function.
|
|
namespace ns4 {
|
|
template<typename KN, typename... Ts>
|
|
void sycl_kernel_launch(const char *, Ts...) noexcept;
|
|
[[clang::sycl_kernel_entry_point(KN<4>)]]
|
|
void skep(KT<4> k) noexcept {
|
|
k();
|
|
}
|
|
}
|
|
// CHECK: ; Function Attrs: mustprogress noinline nounwind optnone
|
|
// CHECK: define dso_local void @_ZN3ns44skepE2KTILi4ELi0EE() #{{[0-9]+}} {
|
|
// CHECK: call void @_ZN3ns418sycl_kernel_launchI2KNILi4EEJ2KTILi4ELi0EEEEEvPKcDpT0_(ptr noundef @.str.3)
|
|
// CHECK: ret void
|
|
// CHECK: }
|