Skip to content

Commit 43af08d

Browse files
iburyliburyl
authored
[SYCL] Host task accessor deduction guides (#2055)
This patch adds host task functionality of [SYCL_INTEL_accessor_simplification](#1498) extension. `handler` type is now valid argument for `host_accessor` and `buffer.get_host_access()`, which enables non-blocking behavior, while accessing memory on the host.
1 parent d3c7b20 commit 43af08d

File tree

5 files changed

+225
-20
lines changed

5 files changed

+225
-20
lines changed

‎sycl/include/CL/sycl/accessor.hpp

100644100755
Lines changed: 85 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -987,9 +987,9 @@ class accessor :
987987

988988
template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
989989
typename TagT,
990-
typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>() &&
991-
IsValidTag<TagT>() && IsPlaceH &&
992-
(IsGlobalBuf || IsConstantBuf)>>
990+
typename = detail::enable_if_t<
991+
IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && IsPlaceH &&
992+
(IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
993993
accessor(buffer<T, Dims, AllocatorT> &BufferRef, TagT,
994994
const property_list &PropertyList = {})
995995
: accessor(BufferRef, PropertyList) {}
@@ -1024,9 +1024,9 @@ class accessor :
10241024

10251025
template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
10261026
typename TagT,
1027-
typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>() &&
1028-
IsValidTag<TagT>() && !IsPlaceH &&
1029-
(IsGlobalBuf || IsConstantBuf)>>
1027+
typename = detail::enable_if_t<
1028+
IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && !IsPlaceH &&
1029+
(IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
10301030
accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
10311031
TagT, const property_list &PropertyList = {})
10321032
: accessor(BufferRef, CommandGroupHandler, PropertyList) {}
@@ -1058,9 +1058,9 @@ class accessor :
10581058
#endif
10591059

10601060
template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1061-
typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>() &&
1062-
(!IsPlaceH &&
1063-
(IsGlobalBuf || IsConstantBuf))>>
1061+
typename = detail::enable_if_t<
1062+
IsSameAsBuffer<T, Dims>() &&
1063+
(!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
10641064
accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
10651065
range<Dimensions> AccessRange,
10661066
const property_list &PropertyList = {})
@@ -1071,9 +1071,9 @@ class accessor :
10711071

10721072
template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
10731073
typename TagT,
1074-
typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>() &&
1075-
IsValidTag<TagT>() && !IsPlaceH &&
1076-
(IsGlobalBuf || IsConstantBuf)>>
1074+
typename = detail::enable_if_t<
1075+
IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && !IsPlaceH &&
1076+
(IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
10771077
accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
10781078
range<Dimensions> AccessRange, TagT,
10791079
const property_list &PropertyList = {})
@@ -1123,9 +1123,9 @@ class accessor :
11231123
#endif
11241124

11251125
template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1126-
typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>() &&
1127-
(!IsPlaceH &&
1128-
(IsGlobalBuf || IsConstantBuf))>>
1126+
typename = detail::enable_if_t<
1127+
IsSameAsBuffer<T, Dims>() &&
1128+
(!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
11291129
accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
11301130
range<Dimensions> AccessRange, id<Dimensions> AccessOffset,
11311131
const property_list &PropertyList = {})
@@ -1151,9 +1151,9 @@ class accessor :
11511151

11521152
template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
11531153
typename TagT,
1154-
typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>() &&
1155-
IsValidTag<TagT>() && !IsPlaceH &&
1156-
(IsGlobalBuf || IsConstantBuf)>>
1154+
typename = detail::enable_if_t<
1155+
IsSameAsBuffer<T, Dims>() && IsValidTag<TagT>() && !IsPlaceH &&
1156+
(IsGlobalBuf || IsConstantBuf || IsHostBuf)>>
11571157
accessor(buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
11581158
range<Dimensions> AccessRange, id<Dimensions> AccessOffset, TagT,
11591159
const property_list &PropertyList = {})
@@ -1675,8 +1675,6 @@ class host_accessor
16751675
// buffer | handler | range | id | | property_list
16761676
// buffer | handler | range | id | mode_tag | property_list
16771677
// -------+---------+-------+----+----------+--------------
1678-
// host_accessor with handler argument will be added later
1679-
// to facilitate non-blocking accessor use case
16801678

16811679
template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
16821680
typename = typename detail::enable_if_t<
@@ -1699,6 +1697,24 @@ class host_accessor
16991697
mode_tag_t<AccessMode>, const property_list &PropertyList = {})
17001698
: host_accessor(BufferRef, PropertyList) {}
17011699

1700+
#endif
1701+
1702+
template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1703+
typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
1704+
host_accessor(buffer<T, Dims, AllocatorT> &BufferRef,
1705+
handler &CommandGroupHandler,
1706+
const property_list &PropertyList = {})
1707+
: AccessorT(BufferRef, CommandGroupHandler, PropertyList) {}
1708+
1709+
#if __cplusplus > 201402L
1710+
1711+
template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1712+
typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
1713+
host_accessor(buffer<DataT, Dimensions, AllocatorT> &BufferRef,
1714+
handler &CommandGroupHandler, mode_tag_t<AccessMode>,
1715+
const property_list &PropertyList = {})
1716+
: host_accessor(BufferRef, CommandGroupHandler, PropertyList) {}
1717+
17021718
#endif
17031719

17041720
template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
@@ -1717,6 +1733,26 @@ class host_accessor
17171733
const property_list &PropertyList = {})
17181734
: host_accessor(BufferRef, AccessRange, {}, PropertyList) {}
17191735

1736+
#endif
1737+
1738+
template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1739+
typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
1740+
host_accessor(buffer<DataT, Dimensions, AllocatorT> &BufferRef,
1741+
handler &CommandGroupHandler, range<Dimensions> AccessRange,
1742+
const property_list &PropertyList = {})
1743+
: AccessorT(BufferRef, CommandGroupHandler, AccessRange, {},
1744+
PropertyList) {}
1745+
1746+
#if __cplusplus > 201402L
1747+
1748+
template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1749+
typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
1750+
host_accessor(buffer<DataT, Dimensions, AllocatorT> &BufferRef,
1751+
handler &CommandGroupHandler, range<Dimensions> AccessRange,
1752+
mode_tag_t<AccessMode>, const property_list &PropertyList = {})
1753+
: host_accessor(BufferRef, CommandGroupHandler, AccessRange, {},
1754+
PropertyList) {}
1755+
17201756
#endif
17211757

17221758
template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
@@ -1735,6 +1771,28 @@ class host_accessor
17351771
mode_tag_t<AccessMode>, const property_list &PropertyList = {})
17361772
: host_accessor(BufferRef, AccessRange, AccessOffset, PropertyList) {}
17371773

1774+
#endif
1775+
1776+
template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1777+
typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
1778+
host_accessor(buffer<DataT, Dimensions, AllocatorT> &BufferRef,
1779+
handler &CommandGroupHandler, range<Dimensions> AccessRange,
1780+
id<Dimensions> AccessOffset,
1781+
const property_list &PropertyList = {})
1782+
: AccessorT(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
1783+
PropertyList) {}
1784+
1785+
#if __cplusplus > 201402L
1786+
1787+
template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
1788+
typename = detail::enable_if_t<IsSameAsBuffer<T, Dims>()>>
1789+
host_accessor(buffer<DataT, Dimensions, AllocatorT> &BufferRef,
1790+
handler &CommandGroupHandler, range<Dimensions> AccessRange,
1791+
id<Dimensions> AccessOffset, mode_tag_t<AccessMode>,
1792+
const property_list &PropertyList = {})
1793+
: host_accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
1794+
PropertyList) {}
1795+
17381796
#endif
17391797
};
17401798

@@ -1767,6 +1825,13 @@ host_accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3, Type4)
17671825
->host_accessor<DataT, Dimensions,
17681826
detail::deduceAccessMode<Type3, Type4>()>;
17691827

1828+
template <typename DataT, int Dimensions, typename AllocatorT, typename Type1,
1829+
typename Type2, typename Type3, typename Type4, typename Type5>
1830+
host_accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3, Type4,
1831+
Type5)
1832+
->host_accessor<DataT, Dimensions,
1833+
detail::deduceAccessMode<Type4, Type5>()>;
1834+
17701835
#endif
17711836

17721837
} // namespace sycl

‎sycl/include/CL/sycl/buffer.hpp

100644100755
Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -292,6 +292,11 @@ class buffer {
292292
return host_accessor{*this, args...};
293293
}
294294

295+
template <typename... Ts>
296+
auto get_host_access(handler &commandGroupHandler, Ts... args) {
297+
return host_accessor{*this, commandGroupHandler, args...};
298+
}
299+
295300
#endif
296301

297302
template <typename Destination = std::nullptr_t>

‎sycl/test/basic_tests/accessor/Inputs/host_task_accessor.cpp

Lines changed: 125 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,125 @@
1+
//==-------- host_task_accessor.cpp - SYCL accessor basic test -------------==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
#include <CL/sycl.hpp>
9+
#include <cassert>
10+
11+
int main() {
12+
// Non-placeholder accessors.
13+
{
14+
int data[9] = {1, 2, 3, 4, 5, 6, 7, 8, 9};
15+
16+
sycl::buffer<int, 1> buf_data(data, sycl::range<1>(9),
17+
{cl::sycl::property::buffer::use_host_ptr()});
18+
19+
sycl::queue Queue;
20+
21+
Queue.submit([&](sycl::handler &cgh) {
22+
23+
#if defined(accessor_new_api_test)
24+
sycl::host_accessor acc_1(buf_data, cgh);
25+
sycl::host_accessor acc_2(buf_data, cgh, sycl::range<1>(8));
26+
sycl::host_accessor acc_3(buf_data, cgh, sycl::range<1>(8), sycl::id<1>(1));
27+
sycl::host_accessor acc_4(buf_data, cgh, sycl::read_only);
28+
sycl::host_accessor acc_5(buf_data, cgh, sycl::range<1>(8), sycl::read_only);
29+
sycl::host_accessor acc_6(buf_data, cgh, sycl::range<1>(8), sycl::id<1>(1),
30+
sycl::read_only);
31+
sycl::host_accessor acc_7(buf_data, cgh, sycl::write_only);
32+
sycl::host_accessor acc_8(buf_data, cgh, sycl::range<1>(8), sycl::write_only);
33+
sycl::host_accessor acc_9(buf_data, cgh, sycl::range<1>(8), sycl::id<1>(1),
34+
sycl::write_only);
35+
#elif defined(buffer_new_api_test)
36+
auto acc_1 = buf_data.get_host_access(cgh);
37+
auto acc_2 = buf_data.get_host_access(cgh, sycl::range<1>(8));
38+
auto acc_3 = buf_data.get_host_access(cgh, sycl::range<1>(8), sycl::id<1>(1));
39+
auto acc_4 = buf_data.get_host_access(cgh, sycl::read_only);
40+
auto acc_5 = buf_data.get_host_access(cgh, sycl::range<1>(8), sycl::read_only);
41+
auto acc_6 = buf_data.get_host_access(cgh, sycl::range<1>(8), sycl::id<1>(1),
42+
sycl::read_only);
43+
auto acc_7 = buf_data.get_host_access(cgh, sycl::write_only);
44+
auto acc_8 = buf_data.get_host_access(cgh, sycl::range<1>(8), sycl::write_only);
45+
auto acc_9 = buf_data.get_host_access(cgh, sycl::range<1>(8), sycl::id<1>(1),
46+
sycl::write_only);
47+
#endif
48+
49+
cgh.codeplay_host_task(
50+
[=]() {
51+
acc_7[6] = acc_1[0];
52+
acc_8[7] = acc_2[1];
53+
acc_9[7] = acc_3[1];
54+
acc_1[0] = acc_4[3];
55+
acc_2[1] = acc_5[4];
56+
acc_3[1] = acc_6[4];
57+
});
58+
});
59+
Queue.wait();
60+
61+
#if defined(accessor_new_api_test)
62+
sycl::host_accessor host_acc(buf_data, sycl::read_only);
63+
#elif defined(buffer_new_api_test)
64+
auto host_acc = buf_data.get_host_access(sycl::read_only);
65+
#endif
66+
assert(host_acc[0] == 4 && host_acc[1] == 5 && host_acc[2] == 6);
67+
assert(host_acc[3] == 4 && host_acc[4] == 5 && host_acc[5] == 6);
68+
assert(host_acc[6] == 1 && host_acc[7] == 2 && host_acc[8] == 3);
69+
}
70+
71+
// noinit accessors.
72+
{
73+
int data[9] = {1, 2, 3, 4, 5, 6, 7, 8, 9};
74+
75+
sycl::buffer<int, 1> buf_data(data, sycl::range<1>(9),
76+
{cl::sycl::property::buffer::use_host_ptr()});
77+
78+
sycl::queue Queue;
79+
80+
Queue.submit([&](sycl::handler &cgh) {
81+
82+
#if defined(accessor_new_api_test)
83+
sycl::host_accessor acc_1(buf_data, cgh, sycl::noinit);
84+
sycl::host_accessor acc_2(buf_data, cgh, sycl::range<1>(8), sycl::noinit);
85+
sycl::host_accessor acc_3(buf_data, cgh, sycl::range<1>(8), sycl::id<1>(1),
86+
sycl::noinit);
87+
sycl::host_accessor acc_7(buf_data, cgh, sycl::write_only, sycl::noinit);
88+
sycl::host_accessor acc_8(buf_data, cgh, sycl::range<1>(8), sycl::write_only,
89+
sycl::noinit);
90+
sycl::host_accessor acc_9(buf_data, cgh, sycl::range<1>(8), sycl::id<1>(1),
91+
sycl::write_only, sycl::noinit);
92+
#elif defined(buffer_new_api_test)
93+
auto acc_1 = buf_data.get_host_access(cgh, sycl::noinit);
94+
auto acc_2 = buf_data.get_host_access(cgh, sycl::range<1>(8), sycl::noinit);
95+
auto acc_3 = buf_data.get_host_access(cgh, sycl::range<1>(8), sycl::id<1>(1),
96+
sycl::noinit);
97+
auto acc_7 = buf_data.get_host_access(cgh, sycl::write_only, sycl::noinit);
98+
auto acc_8 = buf_data.get_host_access(cgh, sycl::range<1>(8), sycl::write_only,
99+
sycl::noinit);
100+
auto acc_9 = buf_data.get_host_access(cgh, sycl::range<1>(8), sycl::id<1>(1),
101+
sycl::write_only, sycl::noinit);
102+
#endif
103+
104+
cgh.codeplay_host_task(
105+
[=]() {
106+
acc_7[6] = acc_1[0];
107+
acc_8[7] = acc_2[1];
108+
acc_9[7] = acc_3[1];
109+
acc_1[0] = 4;
110+
acc_2[1] = 5;
111+
acc_3[1] = 6;
112+
});
113+
});
114+
Queue.wait();
115+
116+
#if defined(accessor_new_api_test)
117+
sycl::host_accessor host_acc(buf_data, sycl::read_only);
118+
#elif defined(buffer_new_api_test)
119+
auto host_acc = buf_data.get_host_access(sycl::read_only);
120+
#endif
121+
assert(host_acc[0] == 4 && host_acc[1] == 5 && host_acc[2] == 6);
122+
assert(host_acc[3] == 4 && host_acc[4] == 5 && host_acc[5] == 6);
123+
assert(host_acc[6] == 1 && host_acc[7] == 2 && host_acc[8] == 3);
124+
}
125+
}

‎sycl/test/basic_tests/accessor/get_host_task_access_deduction.cpp

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,5 @@
1+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -Dbuffer_new_api_test -std=c++17 %S/Inputs/host_task_accessor.cpp -o %t.out
2+
// RUN: env SYCL_DEVICE_TYPE=HOST %t.out
3+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
4+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
5+
// RUN: %ACC_RUN_PLACEHOLDER %t.out

‎sycl/test/basic_tests/accessor/host_task_accessor_deduction.cpp

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,5 @@
1+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -Daccessor_new_api_test -std=c++17 %S/Inputs/host_task_accessor.cpp -o %t.out
2+
// RUN: env SYCL_DEVICE_TYPE=HOST %t.out
3+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
4+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
5+
// RUN: %ACC_RUN_PLACEHOLDER %t.out

0 commit comments

Comments
 (0)