Skip to content

Commit e7139b0

Browse files
KornevNikitaKornevNikita
authored
[SYCL] Implement interface of sycl_ext_oneapi_prefetch (#11458)
Spec: https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_prefetch.asciidoc Properties are not yet fully functional and being ignored, they require other changes in the SW stack to be properly passed through SPIR-V layer. Will be done in follow-up patches.
1 parent 8c481bd commit e7139b0

File tree

5 files changed

+338
-5
lines changed

5 files changed

+338
-5
lines changed

‎llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -628,13 +628,13 @@ bool CompileTimePropertiesPass::transformSYCLPropertiesAnnotation(
628628
// Read the annotation values and create the new annotation string.
629629
std::string NewAnnotString = "";
630630
auto Properties = parseSYCLPropertiesString(M, IntrInst);
631-
for (auto &Property : Properties) {
631+
for (const auto &[PropName, PropVal] : Properties) {
632632
// sycl-alignment is converted to align on
633633
// previous parseAlignmentAndApply(), dropping here
634-
if (*Property.first == "sycl-alignment")
634+
if (PropName == "sycl-alignment")
635635
continue;
636636

637-
auto DecorIt = SpirvDecorMap.find(*Property.first);
637+
auto DecorIt = SpirvDecorMap.find(*PropName);
638638
if (DecorIt == SpirvDecorMap.end())
639639
continue;
640640
uint32_t DecorCode = DecorIt->second.Code;
@@ -644,8 +644,8 @@ bool CompileTimePropertiesPass::transformSYCLPropertiesAnnotation(
644644
// string values are handled correctly. Note that " around values are
645645
// always valid, even if the decoration parameters are not strings.
646646
NewAnnotString += "{" + std::to_string(DecorCode);
647-
if (Property.second)
648-
NewAnnotString += ":\"" + Property.second->str() + "\"";
647+
if (PropVal)
648+
NewAnnotString += ":\"" + PropVal->str() + "\"";
649649
NewAnnotString += "}";
650650
}
651651

‎sycl/include/sycl/ext/oneapi/experimental/prefetch.hpp

Lines changed: 269 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,269 @@
1+
//==--------------- prefetch.hpp --- SYCL prefetch extension ---------------==//
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+
9+
#pragma once
10+
11+
#include <CL/__spirv/spirv_ops.hpp>
12+
#include <sycl/ext/oneapi/properties/properties.hpp>
13+
14+
namespace sycl {
15+
inline namespace _V1 {
16+
namespace ext::oneapi::experimental {
17+
18+
enum class cache_level { L1 = 0, L2 = 1, L3 = 2, L4 = 3 };
19+
20+
struct nontemporal;
21+
22+
struct prefetch_hint_key {
23+
template <cache_level Level, typename Hint>
24+
using value_t =
25+
property_value<prefetch_hint_key,
26+
std::integral_constant<cache_level, Level>, Hint>;
27+
};
28+
29+
template <cache_level Level, typename Hint>
30+
inline constexpr prefetch_hint_key::value_t<Level, Hint> prefetch_hint;
31+
32+
inline constexpr prefetch_hint_key::value_t<cache_level::L1, void>
33+
prefetch_hint_L1;
34+
inline constexpr prefetch_hint_key::value_t<cache_level::L2, void>
35+
prefetch_hint_L2;
36+
inline constexpr prefetch_hint_key::value_t<cache_level::L3, void>
37+
prefetch_hint_L3;
38+
inline constexpr prefetch_hint_key::value_t<cache_level::L4, void>
39+
prefetch_hint_L4;
40+
41+
inline constexpr prefetch_hint_key::value_t<cache_level::L1, nontemporal>
42+
prefetch_hint_L1_nt;
43+
inline constexpr prefetch_hint_key::value_t<cache_level::L2, nontemporal>
44+
prefetch_hint_L2_nt;
45+
inline constexpr prefetch_hint_key::value_t<cache_level::L3, nontemporal>
46+
prefetch_hint_L3_nt;
47+
inline constexpr prefetch_hint_key::value_t<cache_level::L4, nontemporal>
48+
prefetch_hint_L4_nt;
49+
50+
namespace detail {
51+
template <> struct IsCompileTimeProperty<prefetch_hint_key> : std::true_type {};
52+
53+
template <cache_level Level, typename Hint>
54+
struct PropertyMetaInfo<prefetch_hint_key::value_t<Level, Hint>> {
55+
static constexpr const char *name = std::is_same_v<Hint, nontemporal>
56+
? "sycl-prefetch-hint-nt"
57+
: "sycl-prefetch-hint";
58+
static constexpr int value = static_cast<int>(Level);
59+
};
60+
61+
template <access::address_space AS>
62+
inline constexpr bool check_prefetch_AS =
63+
AS == access::address_space::global_space ||
64+
AS == access::address_space::generic_space;
65+
66+
template <access_mode mode>
67+
inline constexpr bool check_prefetch_acc_mode =
68+
mode == access_mode::read || mode == access_mode::read_write;
69+
70+
template <typename T, typename Properties>
71+
void prefetch_impl(T *ptr, size_t bytes, Properties properties) {
72+
#ifdef __SYCL_DEVICE_ONLY__
73+
auto *ptrGlobalAS = __SYCL_GenericCastToPtrExplicit_ToGlobal<const char>(ptr);
74+
const __attribute__((opencl_global)) char *ptrAnnotated = nullptr;
75+
if constexpr (!properties.template has_property<prefetch_hint_key>()) {
76+
ptrAnnotated = __builtin_intel_sycl_ptr_annotation(
77+
ptrGlobalAS, "sycl-prefetch-hint", static_cast<int>(cache_level::L1));
78+
} else {
79+
auto prop = properties.template get_property<prefetch_hint_key>();
80+
ptrAnnotated = __builtin_intel_sycl_ptr_annotation(
81+
ptrGlobalAS, PropertyMetaInfo<decltype(prop)>::name,
82+
PropertyMetaInfo<decltype(prop)>::value);
83+
}
84+
__spirv_ocl_prefetch(ptrAnnotated, bytes);
85+
#else
86+
std::ignore = ptr;
87+
std::ignore = bytes;
88+
std::ignore = properties;
89+
#endif
90+
}
91+
92+
template <typename Group, typename T, typename Properties>
93+
void joint_prefetch_impl(Group g, T *ptr, size_t bytes, Properties properties) {
94+
// Although calling joint_prefetch is functionally equivalent to calling
95+
// prefetch from every work-item in a group, native suppurt may be added to to
96+
// issue cooperative prefetches more efficiently on some hardware.
97+
std::ignore = g;
98+
prefetch_impl(ptr, bytes, properties);
99+
}
100+
} // namespace detail
101+
102+
template <typename Properties = empty_properties_t>
103+
void prefetch(void *ptr, Properties properties = {}) {
104+
detail::prefetch_impl(ptr, 1, properties);
105+
}
106+
107+
template <typename Properties = empty_properties_t>
108+
void prefetch(void *ptr, size_t bytes, Properties properties = {}) {
109+
detail::prefetch_impl(ptr, bytes, properties);
110+
}
111+
112+
template <typename T, typename Properties = empty_properties_t>
113+
void prefetch(T *ptr, Properties properties = {}) {
114+
detail::prefetch_impl(ptr, sizeof(T), properties);
115+
}
116+
117+
template <typename T, typename Properties = empty_properties_t>
118+
void prefetch(T *ptr, size_t count, Properties properties = {}) {
119+
detail::prefetch_impl(ptr, count * sizeof(T), properties);
120+
}
121+
122+
template <access::address_space AddressSpace, access::decorated IsDecorated,
123+
typename Properties = empty_properties_t>
124+
std::enable_if_t<detail::check_prefetch_AS<AddressSpace>>
125+
prefetch(multi_ptr<void, AddressSpace, IsDecorated> ptr,
126+
Properties properties = {}) {
127+
detail::prefetch_impl(ptr.get(), 1, properties);
128+
}
129+
130+
template <access::address_space AddressSpace, access::decorated IsDecorated,
131+
typename Properties = empty_properties_t>
132+
std::enable_if_t<detail::check_prefetch_AS<AddressSpace>>
133+
prefetch(multi_ptr<void, AddressSpace, IsDecorated> ptr, size_t bytes,
134+
Properties properties = {}) {
135+
detail::prefetch_impl(ptr.get(), bytes, properties);
136+
}
137+
138+
template <typename T, access::address_space AddressSpace,
139+
access::decorated IsDecorated,
140+
typename Properties = empty_properties_t>
141+
std::enable_if_t<detail::check_prefetch_AS<AddressSpace>>
142+
prefetch(multi_ptr<T, AddressSpace, IsDecorated> ptr,
143+
Properties properties = {}) {
144+
detail::prefetch_impl(ptr.get(), sizeof(T), properties);
145+
}
146+
147+
template <typename T, access::address_space AddressSpace,
148+
access::decorated IsDecorated,
149+
typename Properties = empty_properties_t>
150+
std::enable_if_t<detail::check_prefetch_AS<AddressSpace>>
151+
prefetch(multi_ptr<T, AddressSpace, IsDecorated> ptr, size_t count,
152+
Properties properties = {}) {
153+
detail::prefetch_impl(ptr.get(), count * sizeof(T), properties);
154+
}
155+
156+
template <typename DataT, int Dimensions, access_mode AccessMode,
157+
access::placeholder IsPlaceholder,
158+
typename Properties = empty_properties_t>
159+
std::enable_if_t<detail::check_prefetch_acc_mode<AccessMode> &&
160+
(Dimensions > 0)>
161+
prefetch(
162+
accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
163+
id<Dimensions> offset, Properties properties = {}) {
164+
detail::prefetch_impl(&acc[offset], sizeof(DataT), properties);
165+
}
166+
167+
template <typename DataT, int Dimensions, access_mode AccessMode,
168+
access::placeholder IsPlaceholder,
169+
typename Properties = empty_properties_t>
170+
std::enable_if_t<detail::check_prefetch_acc_mode<AccessMode> &&
171+
(Dimensions > 0)>
172+
prefetch(
173+
accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
174+
size_t offset, size_t count, Properties properties = {}) {
175+
detail::prefetch_impl(&acc[offset], count * sizeof(DataT), properties);
176+
}
177+
178+
template <typename Group, typename Properties = empty_properties_t>
179+
std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>>
180+
joint_prefetch(Group g, void *ptr, Properties properties = {}) {
181+
detail::joint_prefetch_impl(g, ptr, 1, properties);
182+
}
183+
184+
template <typename Group, typename Properties = empty_properties_t>
185+
std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>>
186+
joint_prefetch(Group g, void *ptr, size_t bytes, Properties properties = {}) {
187+
detail::joint_prefetch_impl(g, ptr, bytes, properties);
188+
}
189+
190+
template <typename Group, typename T, typename Properties = empty_properties_t>
191+
std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>>
192+
joint_prefetch(Group g, T *ptr, Properties properties = {}) {
193+
detail::joint_prefetch_impl(g, ptr, sizeof(T), properties);
194+
}
195+
196+
template <typename Group, typename T, typename Properties = empty_properties_t>
197+
std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>>
198+
joint_prefetch(Group g, T *ptr, size_t count, Properties properties = {}) {
199+
detail::joint_prefetch_impl(g, ptr, count * sizeof(T), properties);
200+
}
201+
202+
template <typename Group, access::address_space AddressSpace,
203+
access::decorated IsDecorated,
204+
typename Properties = empty_properties_t>
205+
std::enable_if_t<detail::check_prefetch_AS<AddressSpace> &&
206+
sycl::is_group_v<std::decay_t<Group>>>
207+
joint_prefetch(Group g, multi_ptr<void, AddressSpace, IsDecorated> ptr,
208+
Properties properties = {}) {
209+
detail::joint_prefetch_impl(g, ptr.get(), 1, properties);
210+
}
211+
212+
template <typename Group, access::address_space AddressSpace,
213+
access::decorated IsDecorated,
214+
typename Properties = empty_properties_t>
215+
std::enable_if_t<detail::check_prefetch_AS<AddressSpace> &&
216+
sycl::is_group_v<std::decay_t<Group>>>
217+
joint_prefetch(Group g, multi_ptr<void, AddressSpace, IsDecorated> ptr,
218+
size_t bytes, Properties properties = {}) {
219+
detail::joint_prefetch_impl(g, ptr.get(), bytes, properties);
220+
}
221+
222+
template <typename Group, typename T, access::address_space AddressSpace,
223+
access::decorated IsDecorated,
224+
typename Properties = empty_properties_t>
225+
std::enable_if_t<detail::check_prefetch_AS<AddressSpace> &&
226+
sycl::is_group_v<std::decay_t<Group>>>
227+
joint_prefetch(Group g, multi_ptr<T, AddressSpace, IsDecorated> ptr,
228+
Properties properties = {}) {
229+
detail::joint_prefetch_impl(g, ptr.get(), sizeof(T), properties);
230+
}
231+
232+
template <typename Group, typename T, access::address_space AddressSpace,
233+
access::decorated IsDecorated,
234+
typename Properties = empty_properties_t>
235+
std::enable_if_t<detail::check_prefetch_AS<AddressSpace> &&
236+
sycl::is_group_v<std::decay_t<Group>>>
237+
joint_prefetch(Group g, multi_ptr<T, AddressSpace, IsDecorated> ptr,
238+
size_t count, Properties properties = {}) {
239+
detail::joint_prefetch_impl(g, ptr.get(), count * sizeof(T), properties);
240+
}
241+
242+
template <typename Group, typename DataT, int Dimensions,
243+
access_mode AccessMode, access::placeholder IsPlaceholder,
244+
typename Properties = empty_properties_t>
245+
std::enable_if_t<detail::check_prefetch_acc_mode<AccessMode> &&
246+
(Dimensions > 0) && sycl::is_group_v<std::decay_t<Group>>>
247+
joint_prefetch(
248+
Group g,
249+
accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
250+
size_t offset, Properties properties = {}) {
251+
detail::joint_prefetch_impl(g, &acc[offset], sizeof(DataT), properties);
252+
}
253+
254+
template <typename Group, typename DataT, int Dimensions,
255+
access_mode AccessMode, access::placeholder IsPlaceholder,
256+
typename Properties = empty_properties_t>
257+
std::enable_if_t<detail::check_prefetch_acc_mode<AccessMode> &&
258+
(Dimensions > 0) && sycl::is_group_v<std::decay_t<Group>>>
259+
joint_prefetch(
260+
Group g,
261+
accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
262+
size_t offset, size_t count, Properties properties = {}) {
263+
detail::joint_prefetch_impl(g, &acc[offset], count * sizeof(DataT),
264+
properties);
265+
}
266+
267+
} // namespace ext::oneapi::experimental
268+
} // namespace _V1
269+
} // namespace sycl

‎sycl/include/sycl/sycl.hpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -82,6 +82,7 @@
8282
#include <sycl/ext/oneapi/experimental/cuda/barrier.hpp>
8383
#include <sycl/ext/oneapi/experimental/fixed_size_group.hpp>
8484
#include <sycl/ext/oneapi/experimental/opportunistic_group.hpp>
85+
#include <sycl/ext/oneapi/experimental/prefetch.hpp>
8586
#include <sycl/ext/oneapi/experimental/tangle_group.hpp>
8687
#include <sycl/ext/oneapi/filter_selector.hpp>
8788
#include <sycl/ext/oneapi/functional.hpp>

‎sycl/source/feature_test.hpp.in

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -90,6 +90,7 @@ inline namespace _V1 {
9090
#define SYCL_EXT_CODEPLAY_MAX_REGISTERS_PER_WORK_GROUP_QUERY 1
9191
#define SYCL_EXT_ONEAPI_DEVICE_GLOBAL 1
9292
#define SYCL_EXT_INTEL_QUEUE_IMMEDIATE_COMMAND_LIST 1
93+
#define SYCL_EXT_ONEAPI_PREFETCH 1
9394

9495
#ifndef __has_include
9596
#define __has_include(x) 0

‎sycl/test/extensions/prefetch.cpp

Lines changed: 62 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,62 @@
1+
// RUN: %clangxx -fsycl -fsyntax-only %s
2+
3+
#include <sycl/sycl.hpp>
4+
5+
int data[] = {0, 1, 2, 3};
6+
7+
int main() {
8+
namespace syclex = sycl::ext::oneapi::experimental;
9+
void *dataPtrVoid = data;
10+
int *dataPtrInt = data;
11+
auto prop = syclex::properties{syclex::prefetch_hint_L1};
12+
13+
{
14+
sycl::buffer<int, 1> buf(data, 4);
15+
sycl::queue q;
16+
q.submit([&](sycl::handler &h) {
17+
auto acc = buf.get_access<sycl::access_mode::read>(h);
18+
h.parallel_for<class Kernel>(
19+
sycl::nd_range<1>(1, 1), ([=](sycl::nd_item<1> index) {
20+
syclex::prefetch(dataPtrVoid, prop);
21+
syclex::prefetch(dataPtrVoid, 16, prop);
22+
23+
syclex::prefetch(dataPtrInt, prop);
24+
syclex::prefetch(dataPtrInt, 4, prop);
25+
26+
auto mPtrVoid = sycl::address_space_cast<
27+
sycl::access::address_space::global_space,
28+
sycl::access::decorated::yes>(dataPtrVoid);
29+
syclex::prefetch(mPtrVoid, prop);
30+
syclex::prefetch(mPtrVoid, 16, prop);
31+
32+
auto mPtrInt = sycl::address_space_cast<
33+
sycl::access::address_space::global_space,
34+
sycl::access::decorated::yes>(dataPtrInt);
35+
syclex::prefetch(mPtrInt, prop);
36+
syclex::prefetch(mPtrInt, 8, prop);
37+
38+
syclex::prefetch(acc, sycl::id(0), prop);
39+
syclex::prefetch(acc, sycl::id(0), 4, prop);
40+
41+
auto g = index.get_group();
42+
syclex::joint_prefetch(g, dataPtrVoid, prop);
43+
syclex::joint_prefetch(g, dataPtrVoid, 16, prop);
44+
45+
syclex::joint_prefetch(g, dataPtrInt, prop);
46+
syclex::joint_prefetch(g, dataPtrInt, 4, prop);
47+
48+
syclex::joint_prefetch(g, mPtrVoid, prop);
49+
syclex::joint_prefetch(g, mPtrVoid, 16, prop);
50+
51+
syclex::joint_prefetch(g, mPtrInt, prop);
52+
syclex::joint_prefetch(g, mPtrInt, 8, prop);
53+
54+
syclex::joint_prefetch(g, acc, sycl::id(0), prop);
55+
syclex::joint_prefetch(g, acc, sycl::id(0), 4, prop);
56+
}));
57+
});
58+
q.wait();
59+
}
60+
61+
return 0;
62+
}

0 commit comments

Comments
 (0)