Skip to content

Commit 9808525

Browse files
maximdimakovmaximdimakov
authored
[SYCL] Pass buffer_location property to buffer (#5604)
There is the accessor property "buffer_location" that allows to allocate buffer in definite location (spec: https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/supported/sycl_ext_intel_buffer_location.asciidoc) Current implementation doesn't implement allocating a buffer at the passed location and the buffer will be re-sided when a kernel is enqueued. It leads to problems of various kinds. The proposed solution implies adding new buffer property in order to store it in a buffer and use it when the buffer is allocated. If the property is not supported by device it will be ignored.
1 parent a2583ec commit 9808525

15 files changed

+384
-16
lines changed

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

Lines changed: 62 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -1195,7 +1195,9 @@ class __SYCL_SPECIAL_CLASS accessor :
11951195
buffer<T, Dims, AllocatorT> &BufferRef, TagT,
11961196
const property_list &PropertyList = {},
11971197
const detail::code_location CodeLoc = detail::code_location::current())
1198-
: accessor(BufferRef, PropertyList, CodeLoc) {}
1198+
: accessor(BufferRef, PropertyList, CodeLoc) {
1199+
adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
1200+
}
11991201

12001202
template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
12011203
typename TagT, typename... PropTypes,
@@ -1208,7 +1210,9 @@ class __SYCL_SPECIAL_CLASS accessor :
12081210
const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
12091211
{},
12101212
const detail::code_location CodeLoc = detail::code_location::current())
1211-
: accessor(BufferRef, PropertyList, CodeLoc) {}
1213+
: accessor(BufferRef, PropertyList, CodeLoc) {
1214+
adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
1215+
}
12121216
#endif
12131217

12141218
template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
@@ -1285,7 +1289,9 @@ class __SYCL_SPECIAL_CLASS accessor :
12851289
buffer<T, Dims, AllocatorT> &BufferRef, handler &CommandGroupHandler,
12861290
TagT, const property_list &PropertyList = {},
12871291
const detail::code_location CodeLoc = detail::code_location::current())
1288-
: accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {}
1292+
: accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {
1293+
adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
1294+
}
12891295

12901296
template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
12911297
typename TagT, typename... PropTypes,
@@ -1299,7 +1305,9 @@ class __SYCL_SPECIAL_CLASS accessor :
12991305
const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
13001306
{},
13011307
const detail::code_location CodeLoc = detail::code_location::current())
1302-
: accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {}
1308+
: accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {
1309+
adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
1310+
}
13031311

13041312
#endif
13051313

@@ -1341,7 +1349,9 @@ class __SYCL_SPECIAL_CLASS accessor :
13411349
buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
13421350
TagT, const property_list &PropertyList = {},
13431351
const detail::code_location CodeLoc = detail::code_location::current())
1344-
: accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
1352+
: accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {
1353+
adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
1354+
}
13451355

13461356
template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
13471357
typename TagT, typename... PropTypes,
@@ -1355,7 +1365,9 @@ class __SYCL_SPECIAL_CLASS accessor :
13551365
const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
13561366
{},
13571367
const detail::code_location CodeLoc = detail::code_location::current())
1358-
: accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {}
1368+
: accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {
1369+
adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
1370+
}
13591371
#endif
13601372

13611373
template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
@@ -1399,7 +1411,9 @@ class __SYCL_SPECIAL_CLASS accessor :
13991411
const property_list &PropertyList = {},
14001412
const detail::code_location CodeLoc = detail::code_location::current())
14011413
: accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
1402-
CodeLoc) {}
1414+
CodeLoc) {
1415+
adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
1416+
}
14031417

14041418
template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
14051419
typename TagT, typename... PropTypes,
@@ -1414,7 +1428,9 @@ class __SYCL_SPECIAL_CLASS accessor :
14141428
{},
14151429
const detail::code_location CodeLoc = detail::code_location::current())
14161430
: accessor(BufferRef, CommandGroupHandler, AccessRange, {}, PropertyList,
1417-
CodeLoc) {}
1431+
CodeLoc) {
1432+
adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
1433+
}
14181434
#endif
14191435

14201436
template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
@@ -1508,7 +1524,9 @@ class __SYCL_SPECIAL_CLASS accessor :
15081524
buffer<T, Dims, AllocatorT> &BufferRef, range<Dimensions> AccessRange,
15091525
id<Dimensions> AccessOffset, TagT, const property_list &PropertyList = {},
15101526
const detail::code_location CodeLoc = detail::code_location::current())
1511-
: accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {}
1527+
: accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {
1528+
adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
1529+
}
15121530

15131531
template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
15141532
typename TagT, typename... PropTypes,
@@ -1522,7 +1540,9 @@ class __SYCL_SPECIAL_CLASS accessor :
15221540
const ext::oneapi::accessor_property_list<PropTypes...> &PropertyList =
15231541
{},
15241542
const detail::code_location CodeLoc = detail::code_location::current())
1525-
: accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {}
1543+
: accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {
1544+
adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
1545+
}
15261546
#endif
15271547

15281548
template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
@@ -1617,7 +1637,9 @@ class __SYCL_SPECIAL_CLASS accessor :
16171637
const property_list &PropertyList = {},
16181638
const detail::code_location CodeLoc = detail::code_location::current())
16191639
: accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
1620-
PropertyList, CodeLoc) {}
1640+
PropertyList, CodeLoc) {
1641+
adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
1642+
}
16211643

16221644
template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
16231645
typename TagT, typename... PropTypes,
@@ -1632,7 +1654,9 @@ class __SYCL_SPECIAL_CLASS accessor :
16321654
{},
16331655
const detail::code_location CodeLoc = detail::code_location::current())
16341656
: accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
1635-
PropertyList, CodeLoc) {}
1657+
PropertyList, CodeLoc) {
1658+
adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get());
1659+
}
16361660
#endif
16371661

16381662
template <typename... NewPropsT>
@@ -1802,6 +1826,32 @@ class __SYCL_SPECIAL_CLASS accessor :
18021826
PI_INVALID_VALUE);
18031827
}
18041828
}
1829+
1830+
#if __cplusplus >= 201703L
1831+
template <typename... PropTypes>
1832+
void adjustAccPropsInBuf(detail::SYCLMemObjI *SYCLMemObject) {
1833+
if constexpr (PropertyListT::template has_property<
1834+
sycl::ext::intel::property::buffer_location>()) {
1835+
auto location = (PropertyListT::template get_property<
1836+
sycl::ext::intel::property::buffer_location>())
1837+
.get_location();
1838+
property_list PropList{
1839+
sycl::property::buffer::detail::buffer_location(location)};
1840+
detail::SYCLMemObjT *SYCLMemObjectT =
1841+
dynamic_cast<detail::SYCLMemObjT *>(SYCLMemObject);
1842+
SYCLMemObjectT->addOrReplaceAccessorProperties(PropList);
1843+
} else {
1844+
deleteAccPropsFromBuf(SYCLMemObject);
1845+
}
1846+
}
1847+
1848+
void deleteAccPropsFromBuf(detail::SYCLMemObjI *SYCLMemObject) {
1849+
detail::SYCLMemObjT *SYCLMemObjectT =
1850+
dynamic_cast<detail::SYCLMemObjT *>(SYCLMemObject);
1851+
SYCLMemObjectT->deleteAccessorProperty(
1852+
sycl::detail::PropWithDataKind::AccPropBufferLocation);
1853+
}
1854+
#endif
18051855
};
18061856

18071857
#if __cplusplus >= 201703L

‎sycl/include/CL/sycl/detail/pi.h

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -591,6 +591,8 @@ constexpr pi_map_flags PI_MAP_WRITE_INVALIDATE_REGION =
591591
// make the translation to OpenCL transparent.
592592
using pi_mem_properties = pi_bitfield;
593593
constexpr pi_mem_properties PI_MEM_PROPERTIES_CHANNEL = CL_MEM_CHANNEL_INTEL;
594+
constexpr pi_mem_properties PI_MEM_PROPERTIES_ALLOC_BUFFER_LOCATION =
595+
CL_MEM_ALLOC_BUFFER_LOCATION_INTEL;
594596

595597
// NOTE: this is made 64-bit to match the size of cl_mem_properties_intel to
596598
// make the translation to OpenCL transparent.

‎sycl/include/CL/sycl/detail/property_helper.hpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -48,7 +48,7 @@ enum PropWithDataKind {
4848
ImageContextBound = 3,
4949
BufferMemChannel = 4,
5050
AccPropBufferLocation = 5,
51-
PropWithDataKindSize = 6
51+
PropWithDataKindSize = 6,
5252
};
5353

5454
// Base class for dataless properties, needed to check that the type of an

‎sycl/include/CL/sycl/detail/property_list_base.hpp

Lines changed: 24 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -103,6 +103,30 @@ class PropertyListBase {
103103
PI_INVALID_VALUE);
104104
}
105105

106+
void add_or_replace_accessor_properties_helper(
107+
const std::vector<std::shared_ptr<PropertyWithDataBase>> &PropsWithData) {
108+
for (auto &Prop : PropsWithData) {
109+
if (Prop->isSame(sycl::detail::PropWithDataKind::AccPropBufferLocation)) {
110+
delete_accessor_property_helper(
111+
sycl::detail::PropWithDataKind::AccPropBufferLocation);
112+
MPropsWithData.push_back(Prop);
113+
break;
114+
}
115+
}
116+
}
117+
118+
void delete_accessor_property_helper(const PropWithDataKind &Kind) {
119+
auto It = MPropsWithData.begin();
120+
for (; It != MPropsWithData.end(); ++It) {
121+
if ((*It)->isSame(Kind))
122+
break;
123+
}
124+
if (It != MPropsWithData.end()) {
125+
std::iter_swap(It, MPropsWithData.end() - 1);
126+
MPropsWithData.pop_back();
127+
}
128+
}
129+
106130
// Stores enabled/disabled for simple properties
107131
std::bitset<DataLessPropKind::DataLessPropKindSize> MDataLessProps;
108132
// Stores shared_ptrs to complex properties

‎sycl/include/CL/sycl/detail/sycl_mem_obj_t.hpp

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -111,6 +111,15 @@ class __SYCL_EXPORT SYCLMemObjT : public SYCLMemObjI {
111111
return MProps.get_property<propertyT>();
112112
}
113113

114+
__SYCL_DLL_LOCAL void
115+
addOrReplaceAccessorProperties(const property_list &PropertyList) {
116+
MProps.add_or_replace_accessor_properties(PropertyList);
117+
}
118+
119+
__SYCL_DLL_LOCAL void deleteAccessorProperty(const PropWithDataKind &Kind) {
120+
MProps.delete_accessor_property(Kind);
121+
}
122+
114123
template <typename AllocatorT>
115124
__SYCL_DLL_LOCAL AllocatorT get_allocator() const {
116125
return MAllocator->getAllocator<AllocatorT>();

‎sycl/include/CL/sycl/properties/accessor_properties.hpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -58,6 +58,7 @@ struct buffer_location {
5858
constexpr bool operator!=(const buffer_location::instance<B> &) const {
5959
return A != B;
6060
}
61+
int get_location() { return A; }
6162
};
6263
};
6364
} // namespace property

‎sycl/include/CL/sycl/properties/buffer_properties.hpp

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -51,6 +51,18 @@ class mem_channel : public detail::PropertyWithData<
5151
uint32_t MChannel;
5252
};
5353

54+
namespace detail {
55+
class buffer_location
56+
: public sycl::detail::PropertyWithData<
57+
sycl::detail::PropWithDataKind::AccPropBufferLocation> {
58+
public:
59+
buffer_location(uint64_t Location) : MLocation(Location) {}
60+
uint64_t get_buffer_location() const { return MLocation; }
61+
62+
private:
63+
uint64_t MLocation;
64+
};
65+
} // namespace detail
5466
} // namespace buffer
5567
} // namespace property
5668

@@ -75,6 +87,9 @@ template <>
7587
struct is_property<property::buffer::use_host_ptr> : std::true_type {};
7688
template <> struct is_property<property::buffer::use_mutex> : std::true_type {};
7789
template <>
90+
struct is_property<property::buffer::detail::buffer_location> : std::true_type {
91+
};
92+
template <>
7893
struct is_property<property::buffer::context_bound> : std::true_type {};
7994
template <>
8095
struct is_property<property::buffer::mem_channel> : std::true_type {};
@@ -91,6 +106,10 @@ struct is_property_of<property::buffer::use_mutex,
91106
buffer<T, Dimensions, AllocatorT, void>>
92107
: std::true_type {};
93108
template <typename T, int Dimensions, typename AllocatorT>
109+
struct is_property_of<property::buffer::detail::buffer_location,
110+
buffer<T, Dimensions, AllocatorT, void>>
111+
: std::true_type {};
112+
template <typename T, int Dimensions, typename AllocatorT>
94113
struct is_property_of<property::buffer::context_bound,
95114
buffer<T, Dimensions, AllocatorT, void>>
96115
: std::true_type {};

‎sycl/include/CL/sycl/property_list.hpp

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -51,6 +51,13 @@ class property_list : protected detail::PropertyListBase {
5151
return has_property_helper<PropT>();
5252
}
5353

54+
void add_or_replace_accessor_properties(const property_list &PropertyList) {
55+
add_or_replace_accessor_properties_helper(PropertyList.MPropsWithData);
56+
}
57+
void delete_accessor_property(const sycl::detail::PropWithDataKind &Kind) {
58+
delete_accessor_property_helper(Kind);
59+
}
60+
5461
template <typename... T> operator ext::oneapi::accessor_property_list<T...>();
5562

5663
private:

‎sycl/source/detail/context_impl.cpp

Lines changed: 19 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -28,15 +28,17 @@ namespace detail {
2828
context_impl::context_impl(const device &Device, async_handler AsyncHandler,
2929
const property_list &PropList)
3030
: MAsyncHandler(AsyncHandler), MDevices(1, Device), MContext(nullptr),
31-
MPlatform(), MPropList(PropList), MHostContext(Device.is_host()) {
31+
MPlatform(), MPropList(PropList), MHostContext(Device.is_host()),
32+
MSupportBufferLocationByDevices(NotChecked) {
3233
MKernelProgramCache.setContextPtr(this);
3334
}
3435

3536
context_impl::context_impl(const std::vector<cl::sycl::device> Devices,
3637
async_handler AsyncHandler,
3738
const property_list &PropList)
3839
: MAsyncHandler(AsyncHandler), MDevices(Devices), MContext(nullptr),
39-
MPlatform(), MPropList(PropList), MHostContext(false) {
40+
MPlatform(), MPropList(PropList), MHostContext(false),
41+
MSupportBufferLocationByDevices(NotChecked) {
4042
MPlatform = detail::getSyclObjImpl(MDevices[0].get_platform());
4143
std::vector<RT::PiDevice> DeviceIds;
4244
for (const auto &D : MDevices) {
@@ -66,7 +68,7 @@ context_impl::context_impl(const std::vector<cl::sycl::device> Devices,
6668
context_impl::context_impl(RT::PiContext PiContext, async_handler AsyncHandler,
6769
const plugin &Plugin)
6870
: MAsyncHandler(AsyncHandler), MDevices(), MContext(PiContext), MPlatform(),
69-
MHostContext(false) {
71+
MHostContext(false), MSupportBufferLocationByDevices(NotChecked) {
7072

7173
std::vector<RT::PiDevice> DeviceIds;
7274
size_t DevicesNum = 0;
@@ -206,6 +208,20 @@ pi_native_handle context_impl::getNative() const {
206208
return Handle;
207209
}
208210

211+
bool context_impl::isBufferLocationSupported() const {
212+
if (MSupportBufferLocationByDevices != NotChecked)
213+
return MSupportBufferLocationByDevices == Supported ? true : false;
214+
// Check that devices within context have support of buffer location
215+
MSupportBufferLocationByDevices = Supported;
216+
for (auto &Device : MDevices) {
217+
if (!Device.has_extension("cl_intel_mem_alloc_buffer_location")) {
218+
MSupportBufferLocationByDevices = NotSupported;
219+
break;
220+
}
221+
}
222+
return MSupportBufferLocationByDevices == Supported ? true : false;
223+
}
224+
209225
} // namespace detail
210226
} // namespace sycl
211227
} // __SYCL_INLINE_NAMESPACE(cl)

‎sycl/source/detail/context_impl.hpp

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -167,6 +167,11 @@ class context_impl {
167167
/// \return a native handle.
168168
pi_native_handle getNative() const;
169169

170+
// Returns true if buffer_location property is supported by devices
171+
bool isBufferLocationSupported() const;
172+
173+
enum PropertySupport { NotSupported = 0, Supported = 1, NotChecked = 2 };
174+
170175
private:
171176
async_handler MAsyncHandler;
172177
std::vector<device> MDevices;
@@ -177,6 +182,7 @@ class context_impl {
177182
std::map<std::pair<DeviceLibExt, RT::PiDevice>, RT::PiProgram>
178183
MCachedLibPrograms;
179184
mutable KernelProgramCache MKernelProgramCache;
185+
mutable PropertySupport MSupportBufferLocationByDevices;
180186
};
181187

182188
} // namespace detail

0 commit comments

Comments
 (0)