Skip to content

[SYCL] Fix not found kernel due to empty kernel name when using set_arg(s) #2181

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.

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
Show file tree
Hide file tree
Changes from 2 commits
Commits
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
26 changes: 17 additions & 9 deletions sycl/include/CL/sycl/handler.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -726,6 +726,7 @@ class __SYCL_EXPORT handler {
MNDRDesc.set(std::move(NumWorkItems));
MCGType = detail::CG::KERNEL;
extractArgsAndReqs();
MKernelName = getKernelName();
}

#ifdef __SYCL_DEVICE_ONLY__
Expand Down Expand Up @@ -1185,6 +1186,7 @@ class __SYCL_EXPORT handler {
MKernel = detail::getSyclObjImpl(std::move(Kernel));
MCGType = detail::CG::KERNEL;
extractArgsAndReqs();
MKernelName = getKernelName();
}

void parallel_for(range<1> NumWorkItems, kernel Kernel) {
Expand Down Expand Up @@ -1218,6 +1220,7 @@ class __SYCL_EXPORT handler {
MNDRDesc.set(std::move(NumWorkItems), std::move(WorkItemOffset));
MCGType = detail::CG::KERNEL;
extractArgsAndReqs();
MKernelName = getKernelName();
}

/// Defines and invokes a SYCL kernel function for the specified range and
Expand All @@ -1238,6 +1241,7 @@ class __SYCL_EXPORT handler {
MNDRDesc.set(std::move(NDRange));
MCGType = detail::CG::KERNEL;
extractArgsAndReqs();
MKernelName = getKernelName();
}

/// Defines and invokes a SYCL kernel function.
Expand All @@ -1260,9 +1264,10 @@ class __SYCL_EXPORT handler {
MNDRDesc.set(range<1>{1});
MKernel = detail::getSyclObjImpl(std::move(Kernel));
MCGType = detail::CG::KERNEL;
if (!MIsHost && !lambdaAndKernelHaveEqualName<NameT>())
if (!MIsHost && !lambdaAndKernelHaveEqualName<NameT>()) {
extractArgsAndReqs();
else
MKernelName = getKernelName();
} else
StoreLambda<NameT, KernelType, /*Dims*/ 0, void>(std::move(KernelFunc));
#endif
}
Expand Down Expand Up @@ -1300,9 +1305,10 @@ class __SYCL_EXPORT handler {
MNDRDesc.set(std::move(NumWorkItems));
MKernel = detail::getSyclObjImpl(std::move(Kernel));
MCGType = detail::CG::KERNEL;
if (!MIsHost && !lambdaAndKernelHaveEqualName<NameT>())
if (!MIsHost && !lambdaAndKernelHaveEqualName<NameT>()) {
extractArgsAndReqs();
else
MKernelName = getKernelName();
} else
StoreLambda<NameT, KernelType, Dims, LambdaArgType>(
std::move(KernelFunc));
#endif
Expand Down Expand Up @@ -1336,9 +1342,10 @@ class __SYCL_EXPORT handler {
MNDRDesc.set(std::move(NumWorkItems), std::move(WorkItemOffset));
MKernel = detail::getSyclObjImpl(std::move(Kernel));
MCGType = detail::CG::KERNEL;
if (!MIsHost && !lambdaAndKernelHaveEqualName<NameT>())
if (!MIsHost && !lambdaAndKernelHaveEqualName<NameT>()) {
extractArgsAndReqs();
else
MKernelName = getKernelName();
} else
StoreLambda<NameT, KernelType, Dims, LambdaArgType>(
std::move(KernelFunc));
#endif
Expand Down Expand Up @@ -1373,9 +1380,10 @@ class __SYCL_EXPORT handler {
MNDRDesc.set(std::move(NDRange));
MKernel = detail::getSyclObjImpl(std::move(Kernel));
MCGType = detail::CG::KERNEL;
if (!MIsHost && !lambdaAndKernelHaveEqualName<NameT>())
if (!MIsHost && !lambdaAndKernelHaveEqualName<NameT>()) {
extractArgsAndReqs();
else
MKernelName = getKernelName();
} else
StoreLambda<NameT, KernelType, Dims, LambdaArgType>(
std::move(KernelFunc));
#endif
Expand Down Expand Up @@ -1791,7 +1799,7 @@ class __SYCL_EXPORT handler {
unique_ptr_class<detail::HostKernelBase> MHostKernel;
/// Storage for lambda/function when using HostTask
unique_ptr_class<detail::HostTask> MHostTask;
detail::OSModuleHandle MOSModuleHandle;
detail::OSModuleHandle MOSModuleHandle = detail::OSUtil::ExeModuleHandle;
// Storage for a lambda or function when using InteropTasks
std::unique_ptr<detail::InteropTask> MInteropTask;
/// The list of events that order this operation.
Expand Down
255 changes: 255 additions & 0 deletions sycl/test/basic_tests/handler/handler_set_args.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,255 @@
// RUN: %clangxx -fsycl %s -o %t.out
// RUN: %CPU_RUN_PLACEHOLDER %t.out
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: %ACC_RUN_PLACEHOLDER %t.out
//==--------------- handler_set_args.cpp -------------------==//
//
// 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
//
//===----------------------------------------------------------------------===//
#include <CL/sycl.hpp>
#include <cassert>

struct use_offset {
static const int no = 0;
static const int yes = 1;
};

using accessor_t =
cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write,
cl::sycl::access::target::global_buffer>;

struct single_task_functor {
single_task_functor(accessor_t acc) : acc(acc) {}

void operator()() { acc[0] = 10; }

accessor_t acc;
};

struct single_task_new_functor {
single_task_new_functor(accessor_t acc) : acc(acc) {}

void operator()() { acc[0] = 10; }

accessor_t acc;
};

template <int useOffset> struct parallel_for_range_id_functor {
parallel_for_range_id_functor(accessor_t acc) : acc(acc) {}

void operator()(cl::sycl::id<1> id) { acc[0] = 10; }

accessor_t acc;
};

template <int useOffset> struct parallel_for_range_item_functor {
parallel_for_range_item_functor(accessor_t acc) : acc(acc) {}

void operator()(cl::sycl::item<1> item) { acc[0] = 10; }

accessor_t acc;
};

struct parallel_for_nd_range_functor {
parallel_for_nd_range_functor(accessor_t acc) : acc(acc) {}

void operator()(cl::sycl::nd_item<1> ndItem) { acc[0] = 10; }

accessor_t acc;
};

template <class kernel_name>
cl::sycl::kernel get_prebuilt_kernel(cl::sycl::queue &queue) {
cl::sycl::program program(queue.get_context());
program.build_with_kernel_type<kernel_name>();
return program.get_kernel<kernel_name>();
}

const cl::sycl::range<1> range = 1;

template <class kernel_wrapper>
void check_api_call(cl::sycl::queue &queue, kernel_wrapper &&kernelWrapper) {
int result = 0;
{
auto buf = cl::sycl::buffer<int, 1>(&result, range);
queue.submit([&](cl::sycl::handler &cgh) {
auto acc = buf.get_access<cl::sycl::access::mode::read_write>(cgh);
kernelWrapper(cgh, acc);
});
}
assert(result == 10);
}

int main() {
cl::sycl::queue queue;
const cl::sycl::id<1> offset(0);

const cl::sycl::nd_range<1> ndRange(range, range);

check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
cgh.single_task(single_task_functor(acc));
});

check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
cgh.parallel_for(range, parallel_for_range_id_functor<use_offset::no>(acc));
});

check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
cgh.parallel_for(range, offset,
parallel_for_range_id_functor<use_offset::yes>(acc));
});

check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
cgh.parallel_for(range,
parallel_for_range_item_functor<use_offset::no>(acc));
});

check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
cgh.parallel_for(range, offset,
parallel_for_range_item_functor<use_offset::yes>(acc));
});

check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
cgh.parallel_for(ndRange, parallel_for_nd_range_functor(acc));
});

{
auto preBuiltKernel = get_prebuilt_kernel<single_task_functor>(queue);

check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
cgh.set_args(acc);
cgh.single_task(preBuiltKernel);
});
}

{
auto preBuiltKernel =
get_prebuilt_kernel<parallel_for_range_id_functor<use_offset::no>>(
queue);

check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
cgh.set_args(acc);
cgh.parallel_for(range, preBuiltKernel);
});
}

{
auto preBuiltKernel =
get_prebuilt_kernel<parallel_for_range_id_functor<use_offset::yes>>(
queue);

check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
cgh.set_args(acc);
cgh.parallel_for(range, offset, preBuiltKernel);
});
}

{
auto preBuiltKernel =
get_prebuilt_kernel<parallel_for_range_item_functor<use_offset::no>>(
queue);

check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
cgh.set_args(acc);
cgh.parallel_for(range, preBuiltKernel);
});
}

{
auto preBuiltKernel =
get_prebuilt_kernel<parallel_for_range_item_functor<use_offset::yes>>(
queue);

check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
cgh.set_args(acc);
cgh.parallel_for(range, offset, preBuiltKernel);
});
}

{
auto preBuiltKernel =
get_prebuilt_kernel<parallel_for_nd_range_functor>(queue);

check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
cgh.set_args(acc);
cgh.parallel_for(ndRange, preBuiltKernel);
});
}

{
auto preBuiltKernel = get_prebuilt_kernel<single_task_functor>(queue);

check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
cgh.set_args(acc);
cgh.single_task<class other_kernel_name1>(preBuiltKernel,
[=]() { acc[0] = 10; });
});
}

{
auto preBuiltKernel =
get_prebuilt_kernel<parallel_for_range_id_functor<use_offset::no>>(
queue);

check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
cgh.set_args(acc);
cgh.parallel_for<class other_kernel_name2>(
preBuiltKernel, range, [=](cl::sycl::id<1> id) { acc[0] = 10; });
});
}

{
auto preBuiltKernel =
get_prebuilt_kernel<parallel_for_range_id_functor<use_offset::yes>>(
queue);

check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
cgh.set_args(acc);
cgh.parallel_for<class other_kernel_name3>(
preBuiltKernel, range, offset,
[=](cl::sycl::id<1> id) { acc[0] = 10; });
});
}

{
auto preBuiltKernel =
get_prebuilt_kernel<parallel_for_range_item_functor<use_offset::no>>(
queue);

check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
cgh.set_args(acc);
cgh.parallel_for<class other_kernel_name4>(
preBuiltKernel, range, [=](cl::sycl::item<1> item) { acc[0] = 10; });
});
}

{
auto preBuiltKernel =
get_prebuilt_kernel<parallel_for_range_item_functor<use_offset::yes>>(
queue);

check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
cgh.set_args(acc);
cgh.parallel_for<class other_kernel_name5>(
preBuiltKernel, range, offset,
[=](cl::sycl::item<1> item) { acc[0] = 10; });
});
}

{
auto preBuiltKernel =
get_prebuilt_kernel<parallel_for_nd_range_functor>(queue);

check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
cgh.set_args(acc);
cgh.parallel_for<class other_kernel_name6>(
preBuiltKernel, ndRange,
[=](cl::sycl::nd_item<1> ndItem) { acc[0] = 10; });
});
}

return 0;
}