Skip to content

Commit b6faefb

Browse files
steffenlarsensteffenlarsen
authored
[SYCL] Add host_task image accessor support (#9718)
This commit adds the support for SYCL 2020 image accessors in host_task commands as well as implementations of their member functions on host. --------- Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
1 parent a230a62 commit b6faefb

17 files changed

+996
-180
lines changed

‎sycl/include/sycl/accessor.hpp

Lines changed: 177 additions & 59 deletions
Large diffs are not rendered by default.

‎sycl/include/sycl/detail/handler_proxy.hpp

Lines changed: 8 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -19,15 +19,22 @@ class handler;
1919
namespace detail {
2020

2121
class AccessorBaseHost;
22+
class UnsampledImageAccessorBaseHost;
23+
class SampledImageAccessorBaseHost;
2224

2325
#ifdef __SYCL_DEVICE_ONLY__
24-
// In device compilation accessor isn't inherited from AccessorBaseHost, so
26+
// In device compilation accessor isn't inherited from host base classes, so
2527
// can't detect by it. Since we don't expect it to be ever called in device
2628
// execution, just use blind void *.
2729
inline void associateWithHandler(handler &, void *, access::target) {}
30+
inline void associateWithHandler(handler &, void *, image_target) {}
2831
#else
2932
__SYCL_EXPORT void associateWithHandler(handler &, AccessorBaseHost *,
3033
access::target);
34+
__SYCL_EXPORT void
35+
associateWithHandler(handler &, UnsampledImageAccessorBaseHost *, image_target);
36+
__SYCL_EXPORT void
37+
associateWithHandler(handler &, SampledImageAccessorBaseHost *, image_target);
3138
#endif
3239
} // namespace detail
3340
} // __SYCL_INLINE_VER_NAMESPACE(_V1)

‎sycl/include/sycl/ext/oneapi/owner_less.hpp

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -102,11 +102,23 @@ template <typename DataT, int Dimensions>
102102
struct owner_less<local_accessor<DataT, Dimensions>>
103103
: public detail::owner_less_base<local_accessor<DataT, Dimensions>> {};
104104

105+
template <typename DataT, int Dimensions, access_mode AccessMode,
106+
image_target AccessTarget>
107+
struct owner_less<
108+
unsampled_image_accessor<DataT, Dimensions, AccessMode, AccessTarget>>
109+
: public detail::owner_less_base<unsampled_image_accessor<
110+
DataT, Dimensions, AccessMode, AccessTarget>> {};
111+
105112
template <typename DataT, int Dimensions, access_mode AccessMode>
106113
struct owner_less<host_unsampled_image_accessor<DataT, Dimensions, AccessMode>>
107114
: public detail::owner_less_base<
108115
host_unsampled_image_accessor<DataT, Dimensions, AccessMode>> {};
109116

117+
template <typename DataT, int Dimensions, image_target AccessTarget>
118+
struct owner_less<sampled_image_accessor<DataT, Dimensions, AccessTarget>>
119+
: public detail::owner_less_base<
120+
sampled_image_accessor<DataT, Dimensions, AccessTarget>> {};
121+
110122
template <typename DataT, int Dimensions>
111123
struct owner_less<host_sampled_image_accessor<DataT, Dimensions>>
112124
: public detail::owner_less_base<

‎sycl/include/sycl/handler.hpp

Lines changed: 12 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -454,13 +454,20 @@ class __SYCL_EXPORT handler {
454454
bool is_host() { return MIsHost; }
455455

456456
#ifdef __SYCL_DEVICE_ONLY__
457-
// In device compilation accessor isn't inherited from AccessorBaseHost, so
457+
// In device compilation accessor isn't inherited from host base classes, so
458458
// can't detect by it. Since we don't expect it to be ever called in device
459459
// execution, just use blind void *.
460460
void associateWithHandler(void *AccBase, access::target AccTarget);
461+
void associateWithHandler(void *AccBase, image_target AccTarget);
461462
#else
463+
void associateWithHandlerCommon(detail::AccessorImplPtr AccImpl,
464+
int AccTarget);
462465
void associateWithHandler(detail::AccessorBaseHost *AccBase,
463466
access::target AccTarget);
467+
void associateWithHandler(detail::UnsampledImageAccessorBaseHost *AccBase,
468+
image_target AccTarget);
469+
void associateWithHandler(detail::SampledImageAccessorBaseHost *AccBase,
470+
image_target AccTarget);
464471
#endif
465472

466473
// Recursively calls itself until arguments pack is fully processed.
@@ -2935,6 +2942,10 @@ class __SYCL_EXPORT handler {
29352942
friend void detail::associateWithHandler(handler &,
29362943
detail::AccessorBaseHost *,
29372944
access::target);
2945+
friend void detail::associateWithHandler(
2946+
handler &, detail::UnsampledImageAccessorBaseHost *, image_target);
2947+
friend void detail::associateWithHandler(
2948+
handler &, detail::SampledImageAccessorBaseHost *, image_target);
29382949
#endif
29392950

29402951
friend class ::MockHandler;

‎sycl/source/detail/handler_proxy.cpp

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -19,6 +19,16 @@ void associateWithHandler(handler &CGH, AccessorBaseHost *Acc,
1919
CGH.associateWithHandler(Acc, Target);
2020
}
2121

22+
void associateWithHandler(handler &CGH, UnsampledImageAccessorBaseHost *Acc,
23+
image_target Target) {
24+
CGH.associateWithHandler(Acc, Target);
25+
}
26+
27+
void associateWithHandler(handler &CGH, SampledImageAccessorBaseHost *Acc,
28+
image_target Target) {
29+
CGH.associateWithHandler(Acc, Target);
30+
}
31+
2232
} // namespace detail
2333
} // __SYCL_INLINE_VER_NAMESPACE(_V1)
2434
} // namespace sycl

‎sycl/source/handler.cpp

Lines changed: 21 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -376,9 +376,8 @@ void handler::addReduction(const std::shared_ptr<const void> &ReduObj) {
376376
MImpl->MAuxiliaryResources.push_back(ReduObj);
377377
}
378378

379-
void handler::associateWithHandler(detail::AccessorBaseHost *AccBase,
380-
access::target AccTarget) {
381-
detail::AccessorImplPtr AccImpl = detail::getSyclObjImpl(*AccBase);
379+
void handler::associateWithHandlerCommon(detail::AccessorImplPtr AccImpl,
380+
int AccTarget) {
382381
detail::Requirement *Req = AccImpl.get();
383382
// Add accessor to the list of requirements.
384383
CGData.MRequirements.push_back(Req);
@@ -387,8 +386,25 @@ void handler::associateWithHandler(detail::AccessorBaseHost *AccBase,
387386
// Add an accessor to the handler list of associated accessors.
388387
// For associated accessors index does not means nothing.
389388
MAssociatedAccesors.emplace_back(detail::kernel_param_kind_t::kind_accessor,
390-
Req, static_cast<int>(AccTarget),
391-
/*index*/ 0);
389+
Req, AccTarget, /*index*/ 0);
390+
}
391+
392+
void handler::associateWithHandler(detail::AccessorBaseHost *AccBase,
393+
access::target AccTarget) {
394+
associateWithHandlerCommon(detail::getSyclObjImpl(*AccBase),
395+
static_cast<int>(AccTarget));
396+
}
397+
398+
void handler::associateWithHandler(
399+
detail::UnsampledImageAccessorBaseHost *AccBase, image_target AccTarget) {
400+
associateWithHandlerCommon(detail::getSyclObjImpl(*AccBase),
401+
static_cast<int>(AccTarget));
402+
}
403+
404+
void handler::associateWithHandler(
405+
detail::SampledImageAccessorBaseHost *AccBase, image_target AccTarget) {
406+
associateWithHandlerCommon(detail::getSyclObjImpl(*AccBase),
407+
static_cast<int>(AccTarget));
392408
}
393409

394410
static void addArgsForGlobalAccessor(detail::Requirement *AccImpl, size_t Index,

‎sycl/test-e2e/Basic/sycl_2020_images/common.hpp

Lines changed: 117 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -200,7 +200,7 @@ CoordT<ImgT, Dims> DelinearizeToCoord(size_t Idx, range<Dims> ImageRange,
200200
} else if constexpr (Dims == 2) {
201201
Out = CoordT<ImgT, Dims>{Idx % ImageRange[0], Idx / ImageRange[0]};
202202
} else {
203-
Out = CoordT<ImgT, Dims>{Idx % ImageRange[0] % ImageRange[1],
203+
Out = CoordT<ImgT, Dims>{Idx % ImageRange[0],
204204
Idx / ImageRange[0] % ImageRange[1],
205205
Idx / ImageRange[0] / ImageRange[1], 0};
206206
}
@@ -328,3 +328,119 @@ ApplyAddressingMode(CoordT<ImageType::Sampled, Dims> Coord,
328328
}
329329
}
330330
}
331+
332+
template <image_format Format> static constexpr size_t getMaxInt() {
333+
using rep_elem_type = typename FormatTraits<Format>::rep_elem_type;
334+
return static_cast<size_t>(std::numeric_limits<rep_elem_type>::max());
335+
}
336+
337+
template <image_format Format>
338+
typename FormatTraits<Format>::pixel_type PickNewColor(size_t I,
339+
size_t AccSize) {
340+
using PixelType = typename FormatTraits<Format>::pixel_type;
341+
size_t Idx = I * 4;
342+
343+
// Pick a new color. Make sure it isn't too big for the data type.
344+
PixelType NewColor{Idx, Idx + 1, Idx + 2, Idx + 3};
345+
NewColor = sycl::min(NewColor, PixelType{getMaxInt<Format>()});
346+
if constexpr (FormatTraits<Format>::Normalized)
347+
NewColor /= AccSize * 4;
348+
return NewColor;
349+
}
350+
351+
// Implemented as specified by the OpenCL 1.2 specification for
352+
// CLK_FILTER_NEAREST.
353+
template <image_format Format, addressing_mode AddrMode, int Dims>
354+
typename FormatTraits<Format>::pixel_type
355+
ReadNearest(typename FormatTraits<Format>::rep_elem_type *RefData,
356+
CoordT<ImageType::Sampled, Dims> Coord, range<2> ImagePitch,
357+
range<Dims> ImageRange, bool Normalized) {
358+
CoordT<ImageType::Sampled, Dims> AdjCoord = Coord;
359+
if constexpr (AddrMode == addressing_mode::repeat) {
360+
assert(Normalized);
361+
AdjCoord -= sycl::floor(AdjCoord);
362+
AdjCoord *= RangeToCoord<ImageType::Sampled, Dims>(ImageRange);
363+
AdjCoord = sycl::floor(AdjCoord);
364+
} else if constexpr (AddrMode == addressing_mode::mirrored_repeat) {
365+
assert(Normalized);
366+
AdjCoord = 2.0f * sycl::rint(0.5f * Coord);
367+
AdjCoord = sycl::fabs(Coord - AdjCoord);
368+
AdjCoord *= RangeToCoord<ImageType::Sampled, Dims>(ImageRange);
369+
AdjCoord = sycl::floor(AdjCoord);
370+
} else {
371+
if (Normalized)
372+
AdjCoord *= RangeToCoord<ImageType::Sampled, Dims>(ImageRange);
373+
AdjCoord = sycl::floor(AdjCoord);
374+
}
375+
AdjCoord = ApplyAddressingMode<AddrMode>(AdjCoord, ImageRange);
376+
return SimulateRead<Format, ImageType::Sampled>(RefData, AdjCoord, ImagePitch,
377+
ImageRange, false);
378+
}
379+
380+
// Implemented as specified by the OpenCL 1.2 specification for
381+
// CLK_FILTER_LINEAR.
382+
template <image_format Format, addressing_mode AddrMode, int Dims>
383+
float4 CalcLinearRead(typename FormatTraits<Format>::rep_elem_type *RefData,
384+
CoordT<ImageType::Sampled, Dims> Coord,
385+
range<2> ImagePitch, range<Dims> ImageRange,
386+
bool Normalized) {
387+
using UpscaledCoordT = CoordT<ImageType::Sampled, 3>;
388+
389+
auto Read = [&](UpscaledCoordT UpCoord) {
390+
auto DownCoord = DownscaleCoord<Dims>(UpCoord);
391+
return SimulateRead<Format, ImageType::Sampled>(
392+
RefData, DownCoord, ImagePitch, ImageRange, false);
393+
};
394+
395+
CoordT<ImageType::Sampled, Dims> AdjCoord = Coord;
396+
if constexpr (AddrMode == addressing_mode::repeat) {
397+
assert(Normalized);
398+
AdjCoord -= floor(AdjCoord);
399+
AdjCoord *= RangeToCoord<ImageType::Sampled, Dims>(ImageRange);
400+
} else if constexpr (AddrMode == addressing_mode::mirrored_repeat) {
401+
assert(Normalized);
402+
AdjCoord = 2.0f * sycl::rint(0.5f * Coord);
403+
AdjCoord = sycl::fabs(Coord - AdjCoord);
404+
AdjCoord *= RangeToCoord<ImageType::Sampled, Dims>(ImageRange);
405+
} else {
406+
if (Normalized)
407+
AdjCoord *= RangeToCoord<ImageType::Sampled, Dims>(ImageRange);
408+
}
409+
410+
auto Prev = sycl::floor(AdjCoord - 0.5f);
411+
auto Next = Prev + 1;
412+
auto CA = (AdjCoord - 0.5f) - Prev;
413+
414+
Prev = ApplyAddressingMode<AddrMode>(Prev, ImageRange);
415+
Next = ApplyAddressingMode<AddrMode>(Next, ImageRange);
416+
417+
auto UPrev = UpscaleCoord<Dims>(Prev);
418+
auto UNext = UpscaleCoord<Dims>(Next);
419+
auto UCA = UpscaleCoord<Dims>(CA, 1);
420+
421+
auto CA000 = Read(UpscaledCoordT{UPrev[0], UPrev[1], UPrev[2], 0})
422+
.template convert<float>() *
423+
(1 - UCA[0]) * (1 - UCA[1]) * (1 - UCA[2]);
424+
auto CA100 = Read(UpscaledCoordT{UNext[0], UPrev[1], UPrev[2], 0})
425+
.template convert<float>() *
426+
UCA[0] * (1 - UCA[1]) * (1 - UCA[2]);
427+
auto CA010 = Read(UpscaledCoordT{UPrev[0], UNext[1], UPrev[2], 0})
428+
.template convert<float>() *
429+
(1 - UCA[0]) * UCA[1] * (1 - UCA[2]);
430+
auto CA110 = Read(UpscaledCoordT{UNext[0], UNext[1], UPrev[2], 0})
431+
.template convert<float>() *
432+
UCA[0] * UCA[1] * (1 - UCA[2]);
433+
auto CA001 = Read(UpscaledCoordT{UPrev[0], UPrev[1], UNext[2], 0})
434+
.template convert<float>() *
435+
(1 - UCA[0]) * (1 - UCA[1]) * UCA[2];
436+
auto CA101 = Read(UpscaledCoordT{UNext[0], UPrev[1], UNext[2], 0})
437+
.template convert<float>() *
438+
UCA[0] * (1 - UCA[1]) * UCA[2];
439+
auto CA011 = Read(UpscaledCoordT{UPrev[0], UNext[1], UNext[2], 0})
440+
.template convert<float>() *
441+
(1 - UCA[0]) * UCA[1] * UCA[2];
442+
auto CA111 = Read(UpscaledCoordT{UNext[0], UNext[1], UNext[2], 0})
443+
.template convert<float>() *
444+
UCA[0] * UCA[1] * UCA[2];
445+
return CA000 + CA100 + CA010 + CA110 + CA001 + CA101 + CA011 + CA111;
446+
}

‎sycl/test-e2e/Basic/sycl_2020_images/host_sampled_image_read_linear.cpp

Lines changed: 0 additions & 68 deletions
Original file line numberDiff line numberDiff line change
@@ -14,74 +14,6 @@ constexpr size_t IMAGE_DEPTH = 2;
1414
constexpr size_t IMAGE_PITCH_WIDTH = 7;
1515
constexpr size_t IMAGE_PITCH_HEIGHT = 5 * IMAGE_PITCH_WIDTH;
1616

17-
// Implemented as specified by the OpenCL 1.2 specification for
18-
// CLK_FILTER_LINEAR.
19-
template <image_format Format, addressing_mode AddrMode, int Dims>
20-
float4 CalcLinearRead(typename FormatTraits<Format>::rep_elem_type *RefData,
21-
CoordT<ImageType::Sampled, Dims> Coord,
22-
range<2> ImagePitch, range<Dims> ImageRange,
23-
bool Normalized) {
24-
using UpscaledCoordT = CoordT<ImageType::Sampled, 3>;
25-
26-
auto Read = [&](UpscaledCoordT UpCoord) {
27-
auto DownCoord = DownscaleCoord<Dims>(UpCoord);
28-
return SimulateRead<Format, ImageType::Sampled>(
29-
RefData, DownCoord, ImagePitch, ImageRange, false);
30-
};
31-
32-
CoordT<ImageType::Sampled, Dims> AdjCoord = Coord;
33-
if constexpr (AddrMode == addressing_mode::repeat) {
34-
assert(Normalized);
35-
AdjCoord -= floor(AdjCoord);
36-
AdjCoord *= RangeToCoord<ImageType::Sampled, Dims>(ImageRange);
37-
} else if constexpr (AddrMode == addressing_mode::mirrored_repeat) {
38-
assert(Normalized);
39-
AdjCoord = 2.0f * sycl::rint(0.5f * Coord);
40-
AdjCoord = sycl::fabs(Coord - AdjCoord);
41-
AdjCoord *= RangeToCoord<ImageType::Sampled, Dims>(ImageRange);
42-
} else {
43-
if (Normalized)
44-
AdjCoord *= RangeToCoord<ImageType::Sampled, Dims>(ImageRange);
45-
}
46-
47-
auto Prev = sycl::floor(AdjCoord - 0.5f);
48-
auto Next = Prev + 1;
49-
auto CA = (AdjCoord - 0.5f) - Prev;
50-
51-
Prev = ApplyAddressingMode<AddrMode>(Prev, ImageRange);
52-
Next = ApplyAddressingMode<AddrMode>(Next, ImageRange);
53-
54-
auto UPrev = UpscaleCoord<Dims>(Prev);
55-
auto UNext = UpscaleCoord<Dims>(Next);
56-
auto UCA = UpscaleCoord<Dims>(CA, 1);
57-
58-
auto CA000 = Read(UpscaledCoordT{UPrev[0], UPrev[1], UPrev[2], 0})
59-
.template convert<float>() *
60-
(1 - UCA[0]) * (1 - UCA[1]) * (1 - UCA[2]);
61-
auto CA100 = Read(UpscaledCoordT{UNext[0], UPrev[1], UPrev[2], 0})
62-
.template convert<float>() *
63-
UCA[0] * (1 - UCA[1]) * (1 - UCA[2]);
64-
auto CA010 = Read(UpscaledCoordT{UPrev[0], UNext[1], UPrev[2], 0})
65-
.template convert<float>() *
66-
(1 - UCA[0]) * UCA[1] * (1 - UCA[2]);
67-
auto CA110 = Read(UpscaledCoordT{UNext[0], UNext[1], UPrev[2], 0})
68-
.template convert<float>() *
69-
UCA[0] * UCA[1] * (1 - UCA[2]);
70-
auto CA001 = Read(UpscaledCoordT{UPrev[0], UPrev[1], UNext[2], 0})
71-
.template convert<float>() *
72-
(1 - UCA[0]) * (1 - UCA[1]) * UCA[2];
73-
auto CA101 = Read(UpscaledCoordT{UNext[0], UPrev[1], UNext[2], 0})
74-
.template convert<float>() *
75-
UCA[0] * (1 - UCA[1]) * UCA[2];
76-
auto CA011 = Read(UpscaledCoordT{UPrev[0], UNext[1], UNext[2], 0})
77-
.template convert<float>() *
78-
(1 - UCA[0]) * UCA[1] * UCA[2];
79-
auto CA111 = Read(UpscaledCoordT{UNext[0], UNext[1], UNext[2], 0})
80-
.template convert<float>() *
81-
UCA[0] * UCA[1] * UCA[2];
82-
return CA000 + CA100 + CA010 + CA110 + CA001 + CA101 + CA011 + CA111;
83-
}
84-
8517
template <image_format Format, addressing_mode AddrMode,
8618
coordinate_normalization_mode CoordNormMode, int Dims>
8719
bool checkSampledImageHostReadLinear(

‎sycl/test-e2e/Basic/sycl_2020_images/host_sampled_image_read_nearest.cpp

Lines changed: 0 additions & 29 deletions
Original file line numberDiff line numberDiff line change
@@ -14,35 +14,6 @@ constexpr size_t IMAGE_DEPTH = 2;
1414
constexpr size_t IMAGE_PITCH_WIDTH = 7;
1515
constexpr size_t IMAGE_PITCH_HEIGHT = 5 * IMAGE_PITCH_WIDTH;
1616

17-
// Implemented as specified by the OpenCL 1.2 specification for
18-
// CLK_FILTER_NEAREST.
19-
template <image_format Format, addressing_mode AddrMode, int Dims>
20-
typename FormatTraits<Format>::pixel_type
21-
ReadNearest(typename FormatTraits<Format>::rep_elem_type *RefData,
22-
CoordT<ImageType::Sampled, Dims> Coord, range<2> ImagePitch,
23-
range<Dims> ImageRange, bool Normalized) {
24-
CoordT<ImageType::Sampled, Dims> AdjCoord = Coord;
25-
if constexpr (AddrMode == addressing_mode::repeat) {
26-
assert(Normalized);
27-
AdjCoord -= sycl::floor(AdjCoord);
28-
AdjCoord *= RangeToCoord<ImageType::Sampled, Dims>(ImageRange);
29-
AdjCoord = sycl::floor(AdjCoord);
30-
} else if constexpr (AddrMode == addressing_mode::mirrored_repeat) {
31-
assert(Normalized);
32-
AdjCoord = 2.0f * sycl::rint(0.5f * Coord);
33-
AdjCoord = sycl::fabs(Coord - AdjCoord);
34-
AdjCoord *= RangeToCoord<ImageType::Sampled, Dims>(ImageRange);
35-
AdjCoord = sycl::floor(AdjCoord);
36-
} else {
37-
if (Normalized)
38-
AdjCoord *= RangeToCoord<ImageType::Sampled, Dims>(ImageRange);
39-
AdjCoord = sycl::floor(AdjCoord);
40-
}
41-
AdjCoord = ApplyAddressingMode<AddrMode>(AdjCoord, ImageRange);
42-
return SimulateRead<Format, ImageType::Sampled>(RefData, AdjCoord, ImagePitch,
43-
ImageRange, false);
44-
}
45-
4617
template <image_format Format, addressing_mode AddrMode,
4718
coordinate_normalization_mode CoordNormMode, int Dims>
4819
bool checkSampledImageHostReadNearest(

0 commit comments

Comments
 (0)