Skip to content

[SYCL] Implement a builtin to mark a sycl kernel #3894

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 9 commits into from
Jun 9, 2021
Merged

[SYCL] Implement a builtin to mark a sycl kernel #3894

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
9 commits
Select commit Hold shift + click to select a range
c2b0991
[SYCL] implement a builtin to mark a sycl kernel
Jun 7, 2021
ddb3ae7
Fix the clang-format I'm willing to do
Jun 7, 2021
0962bf6
Document kernel_desc.hpp use of this builtin
Jun 7, 2021
e6c5f14
Fix a lang-extensions comment, update the diagnostics, add tests
Jun 7, 2021
6669b37
Fixes suggested by aaron, clarify language-extensions document
Jun 7, 2021
e9aa506
Change tests as Elizabeth requested
Jun 8, 2021
234bd2e
Add regression test requested by AlexeySachkov
Jun 8, 2021
7f61ef1
Update tests as @elizabethandrews requested in review
Jun 8, 2021
4e94a62
Accept @bader 's suggestion toremove sycldevice
Jun 8, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 3 additions & 0 deletions clang/test/CodeGenSYCL/Inputs/sycl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -187,6 +187,7 @@ class accessor {
template <int dimensions, access::mode accessmode, access::target accesstarget>
struct opencl_image_type;

#ifdef __SYCL_DEVICE_ONLY__
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Note the CFE only provides these types in device mode, so this change lets us uses this in header properly in host mode.

#define IMAGETY_DEFINE(dim, accessmode, amsuffix, Target, ifarray_) \
template <> \
struct opencl_image_type<dim, access::mode::accessmode, \
Expand Down Expand Up @@ -218,6 +219,8 @@ IMAGETY_WRITE_3_DIM_IMAGE
IMAGETY_READ_2_DIM_IARRAY
IMAGETY_WRITE_2_DIM_IARRAY

#endif

template <int dim, access::mode accessmode, access::target accesstarget>
struct _ImageImplT {
#ifdef __SYCL_DEVICE_ONLY__
Expand Down
19 changes: 12 additions & 7 deletions clang/test/CodeGenSYCL/mark-kernel-name.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,14 +1,19 @@
// RUN: %clang_cc1 -triple x86_64-linux-pc -fsycl-is-host -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
// RUN: %clang_cc1 -triple x86_64-linux-pc -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
// RUN: %clang_cc1 -triple x86_64-linux-pc -fsycl-is-host -disable-llvm-passes -fdeclare-spirv-builtins -emit-llvm %s -o - | FileCheck %s
// RUN: %clang_cc1 -triple spir64-sycldevice -aux-triple x86_64-linux-pc -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s

template<typename KN, typename Func>
__attribute__((sycl_kernel)) void kernel(const Func &F) {
F();
}
#include "Inputs/sycl.hpp"

// This test validates that the use of __builtin_sycl_mark_kernel_name alters
// the code-gen'ed value of __builtin_unique_stable_name. In this case, lambda1
// emits the unmodified version like we do typically, while lambda2 is 'marked',
// so it should follow kernel naming (that is, using the E10000 naming). Note
// that the top level kernel lambda (the E10000 in common) is automatically part
// of a kernel name, since it is passed to the kernel function (which is
// necessary sot hat the 'device' build actually emits the builtins.

int main() {

kernel<class K>([]() {
cl::sycl::kernel_single_task<class K>([]() {
auto lambda1 = []() {};
auto lambda2 = []() {};

Expand Down
2 changes: 1 addition & 1 deletion clang/test/SemaSYCL/Inputs/sycl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -228,7 +228,7 @@ template <typename Type> struct get_kernel_wrapper_name_t {

#define ATTR_SYCL_KERNEL __attribute__((sycl_kernel))
template <typename KernelName = auto_name, typename KernelType>
ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc) {
ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc) { // #KernelSingleTaskFunc
kernelFunc(); // #KernelSingleTaskKernelFuncCall
}
template <typename KernelName = auto_name, typename KernelType>
Expand Down
16 changes: 10 additions & 6 deletions clang/test/SemaSYCL/mark-kernel-name.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,9 +1,13 @@
// RUN: %clang_cc1 %s -std=c++17 -triple x86_64-linux-gnu -Wno-sycl-2020-compat -fsycl-is-device -verify -fsyntax-only -Wno-unused

template <typename KernelName, typename KernelType>
[[clang::sycl_kernel]] void kernel_single_task(const KernelType &kernelFunc) { // #kernelSingleTask
kernelFunc();
}
#include "Inputs/sycl.hpp"

// Test to validate that __builtin_sycl_mark_kernel_name properly updates the
// constexpr checking for __builtin_sycl_unique_stable_name. We need to make
// sure that the KernelInfo change in the library both still stays broken, and
// is then 'fixed', so the definitions below help ensure that is the case.
// We also validate that this works in the event that we have a wrapper that
// first calls for the KernelInfo type, then instantiates a kernel.

template <typename KN>
struct KernelInfo {
Expand All @@ -24,7 +28,7 @@ template <template <typename> class KI,
typename KernelType>
void wrapper(KernelType KernelFunc) {
(void)KI<KernelName>::c;
kernel_single_task<KernelName>(KernelFunc); // #SingleTaskInst
cl::sycl::kernel_single_task<KernelName>(KernelFunc); // #SingleTaskInst
}

int main() {
Expand All @@ -37,7 +41,7 @@ int main() {
}();

[]() {
// expected-error@#kernelSingleTask {{kernel instantiation changes the result of an evaluated '__builtin_sycl_unique_stable_name'}}
// expected-error@#KernelSingleTaskFunc {{kernel instantiation changes the result of an evaluated '__builtin_sycl_unique_stable_name'}}
// expected-note@#SingleTaskInst {{in instantiation of function template}}
// expected-note@+2 {{in instantiation of function template}}
// expected-note@#KI_USN {{'__builtin_sycl_unique_stable_name' evaluated here}}
Expand Down