Skip to content

Commit e7012e2

Browse files
slawekptakslawekptak
committed
[WIP] No handler submit
1 parent e5a13e9 commit e7012e2

File tree

14 files changed

+866
-33
lines changed

14 files changed

+866
-33
lines changed

‎sycl/source/detail/kernel_name_based_cache_t.hpp‎ renamed to ‎sycl/include/sycl/detail/kernel_name_based_cache_t.hpp‎

Lines changed: 10 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -7,8 +7,8 @@
77
//===----------------------------------------------------------------------===//
88
#pragma once
99

10-
#include <detail/hashers.hpp>
11-
#include <detail/kernel_arg_mask.hpp>
10+
//#include <detail/hashers.hpp>
11+
//#include <sycl/detail/kernel_arg_mask.hpp>
1212
#include <emhash/hash_table8.hpp>
1313
#include <sycl/detail/spinlock.hpp>
1414
#include <sycl/detail/ur.hpp>
@@ -20,6 +20,7 @@ namespace sycl {
2020
inline namespace _V1 {
2121
namespace detail {
2222
using FastKernelCacheKeyT = std::pair<ur_device_handle_t, ur_context_handle_t>;
23+
using KernelArgMask = std::vector<bool>;
2324

2425
struct FastKernelCacheVal {
2526
ur_kernel_handle_t MKernelHandle; /* UR kernel handle pointer. */
@@ -29,25 +30,27 @@ struct FastKernelCacheVal {
2930
const KernelArgMask *MKernelArgMask; /* Eliminated kernel argument mask. */
3031
ur_program_handle_t MProgramHandle; /* UR program handle corresponding to
3132
this kernel. */
32-
const adapter_impl &MAdapterPtr; /* We can keep reference to the adapter
33+
/*const adapter_impl &MAdapterPtr;*/ /* We can keep reference to the adapter
3334
because during 2-stage shutdown the kernel
3435
cache is destroyed deliberately before the
3536
adapter. */
3637

3738
FastKernelCacheVal(ur_kernel_handle_t KernelHandle, std::mutex *Mutex,
3839
const KernelArgMask *KernelArgMask,
39-
ur_program_handle_t ProgramHandle,
40-
const adapter_impl &AdapterPtr)
40+
ur_program_handle_t ProgramHandle)
41+
//const adapter_impl &AdapterPtr)
4142
: MKernelHandle(KernelHandle), MMutex(Mutex),
42-
MKernelArgMask(KernelArgMask), MProgramHandle(ProgramHandle),
43-
MAdapterPtr(AdapterPtr) {}
43+
MKernelArgMask(KernelArgMask), MProgramHandle(ProgramHandle)
44+
/*MAdapterPtr(AdapterPtr)*/ {}
4445

4546
~FastKernelCacheVal() {
47+
/*
4648
if (MKernelHandle)
4749
MAdapterPtr.call<sycl::detail::UrApiKind::urKernelRelease>(MKernelHandle);
4850
if (MProgramHandle)
4951
MAdapterPtr.call<sycl::detail::UrApiKind::urProgramRelease>(
5052
MProgramHandle);
53+
*/
5154
MKernelHandle = nullptr;
5255
MMutex = nullptr;
5356
MKernelArgMask = nullptr;

‎sycl/include/sycl/khr/free_function_commands.hpp‎

Lines changed: 30 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -1,11 +1,13 @@
11
#pragma once
22

33
#include <sycl/ext/oneapi/experimental/enqueue_functions.hpp>
4+
#include <sycl/khr/requirements.hpp>
45

56
namespace sycl {
67
inline namespace _V1 {
78

8-
#ifdef __DPCPP_ENABLE_UNFINISHED_KHR_EXTENSIONS
9+
10+
//#ifdef __DPCPP_ENABLE_UNFINISHED_KHR_EXTENSIONS
911
namespace khr {
1012

1113
template <typename CommandGroupFunc>
@@ -148,32 +150,33 @@ void launch_grouped(handler &h, range<3> r, range<3> size,
148150
h.parallel_for(nd_range<3>(r, size), k);
149151
}
150152

151-
template <typename KernelType>
152-
void launch_grouped(const queue &q, range<1> r, range<1> size,
153+
template <typename KernelType, typename... Requirements>
154+
void launch_grouped(queue &q, range<1> r, range<1> size,
153155
const KernelType &k,
154156
const sycl::detail::code_location &codeLoc =
155-
sycl::detail::code_location::current()) {
156-
submit(
157-
q, [&](handler &h) { launch_grouped<KernelType>(h, r, size, k); },
158-
codeLoc);
157+
sycl::detail::code_location::current(),
158+
const requirements<Requirements...> req = {}) {
159+
(void)codeLoc;
160+
q.parallel_for_no_handler(nd_range<1>(r, size), k, req);
159161
}
160-
template <typename KernelType>
161-
void launch_grouped(const queue &q, range<2> r, range<2> size,
162+
163+
template <typename KernelType, typename... Requirements>
164+
void launch_grouped(queue &q, range<2> r, range<2> size,
162165
const KernelType &k,
163166
const sycl::detail::code_location &codeLoc =
164-
sycl::detail::code_location::current()) {
165-
submit(
166-
q, [&](handler &h) { launch_grouped<KernelType>(h, r, size, k); },
167-
codeLoc);
167+
sycl::detail::code_location::current(),
168+
const requirements<Requirements...> req = {}) {
169+
(void)codeLoc;
170+
q.parallel_for_no_handler(nd_range<2>(r, size), k, req);
168171
}
169-
template <typename KernelType>
170-
void launch_grouped(const queue &q, range<3> r, range<3> size,
172+
template <typename KernelType, typename... Requirements>
173+
void launch_grouped(queue &q, range<3> r, range<3> size,
171174
const KernelType &k,
172175
const sycl::detail::code_location &codeLoc =
173-
sycl::detail::code_location::current()) {
174-
submit(
175-
q, [&](handler &h) { launch_grouped<KernelType>(h, r, size, k); },
176-
codeLoc);
176+
sycl::detail::code_location::current(),
177+
const requirements<Requirements...> req = {}) {
178+
(void)codeLoc;
179+
q.parallel_for_no_handler(nd_range<3>(r, size), k, req);
177180
}
178181

179182
template <typename... Args>
@@ -283,7 +286,8 @@ template <typename KernelType>
283286
void launch_task(const sycl::queue &q, const KernelType &k,
284287
const sycl::detail::code_location &codeLoc =
285288
sycl::detail::code_location::current()) {
286-
submit(q, [&](handler &h) { launch_task<KernelType>(h, k); }, codeLoc);
289+
(void)codeLoc;
290+
q.single_task_no_handler(k);
287291
}
288292

289293
template <typename... Args>
@@ -298,6 +302,11 @@ void launch_task(const queue &q, const kernel &k, Args &&...args) {
298302
[&](handler &h) { launch_task(h, k, std::forward<Args>(args)...); });
299303
}
300304

305+
template <typename FuncT>
306+
void launch_host_task(queue &q, FuncT &&Func) {
307+
q.host_task_no_handler(std::move(Func));
308+
}
309+
301310
inline void memcpy(handler &h, void *dest, const void *src, size_t numBytes) {
302311
h.memcpy(dest, src, numBytes);
303312
}
@@ -520,6 +529,6 @@ inline void event_barrier(const queue &q, const std::vector<event> &events,
520529
}
521530

522531
} // namespace khr
523-
#endif
532+
//#endif
524533
} // namespace _V1
525534
} // namespace sycl

‎sycl/include/sycl/khr/requirements.hpp‎

Lines changed: 66 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,66 @@
1+
#pragma once
2+
3+
#include <vector>
4+
#include <tuple>
5+
#include <sycl/event.hpp>
6+
7+
namespace sycl {
8+
inline namespace _V1 {
9+
10+
//#ifdef __DPCPP_ENABLE_UNFINISHED_KHR_EXTENSIONS
11+
namespace khr {
12+
13+
template <typename... Requirements>
14+
class requirements;
15+
16+
template <typename... Requirements>
17+
void add_events(const requirements<Requirements...> &Reqs, std::vector<event> &Events);
18+
19+
template <typename... Requirements>
20+
class requirements {
21+
public:
22+
requirements(Requirements... r) : MRequirements(r...) {}
23+
24+
private:
25+
std::tuple<Requirements...> MRequirements;
26+
27+
template <typename... R>
28+
friend void add_events(const requirements<R...> &Reqs, std::vector<event> &Events);
29+
};
30+
31+
template <typename... Requirements>
32+
requirements(Requirements... r) -> requirements<Requirements...>;
33+
34+
template <typename Requirement, typename T>
35+
void add_requirement(std::vector<T> &ReqCont, const Requirement &Req) {
36+
if constexpr (std::is_same_v<Requirement, T>)
37+
ReqCont.push_back(Req);
38+
}
39+
40+
template <typename Requirement, typename T, typename... Requirements>
41+
void add_requirement(std::vector<T> &ReqCont, const Requirement &Req, const Requirements&... Rest) {
42+
if constexpr (std::is_same_v<Requirement, T>)
43+
ReqCont.push_back(Req);
44+
add_requirement(ReqCont, Rest...);
45+
}
46+
47+
template <typename T, typename... Requirements, size_t... Is>
48+
void add_requirements(const std::tuple<Requirements...> &ReqsTuple, std::vector<T> &ReqCont,
49+
std::index_sequence<Is...>) {
50+
add_requirement(ReqCont, std::get<Is>(ReqsTuple)...);
51+
}
52+
53+
template <typename... Requirements>
54+
void add_events(const sycl::khr::requirements<Requirements...> &Reqs, std::vector<event> &Events) {
55+
add_requirements(Reqs.MRequirements, Events, std::make_index_sequence<sizeof...(Requirements)>());
56+
}
57+
58+
template <typename... Requirements>
59+
constexpr bool has_events() {
60+
return (std::is_same_v<event, Requirements> || ...);
61+
}
62+
63+
} // namespace khr
64+
//#endif
65+
} // namespace _V1
66+
} // namespace sycl

0 commit comments

Comments
 (0)