Skip to content

Commit b936b0d

Browse files
garimagugarimagu
committed
[SYCL] Image_accessor Host Implementation - Constructors and few basic
APIs. - Enabled get_access methods in image class. - Base class for accessor to image/image_array/host_image access target for host compiler introduced with basic implementation. - Implemented a few APIs for Image accessors for host compiler - get_count and get_size. - Test code to be added after scheduler support is added for image class. Signed-off-by: Garima Gupta <[email protected]>
1 parent 2b585ff commit b936b0d

File tree

3 files changed

+300
-71
lines changed

3 files changed

+300
-71
lines changed

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

Lines changed: 229 additions & 62 deletions
Original file line numberDiff line numberDiff line change
@@ -13,10 +13,13 @@
1313
#include <CL/sycl/buffer.hpp>
1414
#include <CL/sycl/detail/accessor_impl.hpp>
1515
#include <CL/sycl/detail/common.hpp>
16+
#include <CL/sycl/detail/generic_type_traits.hpp>
1617
#include <CL/sycl/detail/image_ocl_types.hpp>
1718
#include <CL/sycl/handler.hpp>
1819
#include <CL/sycl/id.hpp>
20+
#include <CL/sycl/image.hpp>
1921
#include <CL/sycl/pointers.hpp>
22+
#include <CL/sycl/sampler.hpp>
2023

2124
// The file contains implementations of accessor class. Objects of accessor
2225
// class define a requirement to access some SYCL memory object or local memory
@@ -234,6 +237,194 @@ class accessor_common {
234237
};
235238
};
236239

240+
// Image accessor
241+
template <typename DataT, int Dimensions, access::mode AccessMode,
242+
access::target AccessTarget, access::placeholder IsPlaceholder>
243+
class image_accessor
244+
#ifndef __SYCL_DEVICE_ONLY__
245+
: public detail::AccessorBaseHost {
246+
size_t MImageSize;
247+
size_t MImageCount;
248+
#else
249+
{
250+
/*
251+
// TODO: Define the datatype here based on Dimensions, AccessMode and
252+
AccessTarget.
253+
__ocl_image[Dim]d[array/non_array]_[AM]_t MImage;
254+
__init(__ocl_imagexx_t Image) { MImage = Image; }
255+
*/
256+
#endif
257+
constexpr static bool IsHostImageAcc =
258+
(AccessTarget == access::target::host_image);
259+
260+
constexpr static bool IsImageAcc = (AccessTarget == access::target::image);
261+
262+
constexpr static bool IsImageArrayAcc =
263+
(AccessTarget == access::target::image_array);
264+
265+
constexpr static bool IsImageAccessAnyWrite =
266+
(AccessMode == access::mode::write ||
267+
AccessMode == access::mode::discard_write);
268+
269+
constexpr static bool IsImageAccessRead = (AccessMode == access::mode::read);
270+
271+
static_assert(IsImageAcc || IsHostImageAcc || IsImageArrayAcc,
272+
"Expected image type");
273+
274+
static_assert(IsPlaceholder == access::placeholder::false_t,
275+
"Expected false as Placeholder value for image accessor.");
276+
277+
static_assert(
278+
AccessMode == access::mode::read || AccessMode == access::mode::write ||
279+
AccessMode == access::mode::discard_write,
280+
"Access modes can be only read/write/discard_write for image accessor.");
281+
282+
static_assert(Dimensions > 0 && Dimensions <= 3,
283+
"Dimensions can be 1/2/3 for image accessor.");
284+
285+
public:
286+
using value_type = DataT;
287+
using reference = DataT &;
288+
using const_reference = const DataT &;
289+
290+
// image_accessor Constructors.
291+
292+
// Available only when: accessTarget == access::target::host_image
293+
// template <typename AllocatorT>
294+
// accessor(image<dimensions, AllocatorT> &imageRef);
295+
template <
296+
typename AllocatorT, int Dims = Dimensions,
297+
typename = detail::enable_if_t<(Dims > 0 && Dims <= 3) && IsHostImageAcc>>
298+
image_accessor(image<Dims, AllocatorT> &ImageRef, int ImageElementSize)
299+
#ifdef __SYCL_DEVICE_ONLY__
300+
{
301+
// TODO: Implement this function.
302+
}
303+
#else
304+
: AccessorBaseHost(id<3>(0, 0, 0) /* Offset,*/,
305+
detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
306+
detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
307+
AccessMode, detail::getSyclObjImpl(ImageRef).get(),
308+
Dimensions, ImageElementSize),
309+
MImageSize(ImageRef.get_size()), MImageCount(ImageRef.get_count()) {
310+
detail::EventImplPtr Event =
311+
detail::Scheduler::getInstance().addHostAccessor(
312+
AccessorBaseHost::impl.get());
313+
Event->wait(Event);
314+
}
315+
#endif
316+
317+
// Available only when: accessTarget == access::target::image
318+
// template <typename AllocatorT>
319+
// accessor(image<dimensions, AllocatorT> &imageRef,
320+
// handler &commandGroupHandlerRef);
321+
template <
322+
typename AllocatorT, int Dims = Dimensions,
323+
typename = detail::enable_if_t<(Dims > 0 && Dims <= 3) && IsImageAcc>>
324+
image_accessor(image<Dims, AllocatorT> &ImageRef,
325+
handler &CommandGroupHandlerRef, int ImageElementSize)
326+
#ifdef __SYCL_DEVICE_ONLY__
327+
{
328+
// TODO: Implement this function.
329+
}
330+
#else
331+
: AccessorBaseHost(id<3>(0, 0, 0) /* Offset,*/,
332+
detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
333+
detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
334+
AccessMode, detail::getSyclObjImpl(ImageRef).get(),
335+
Dimensions, ImageElementSize),
336+
MImageSize(ImageRef.get_size()), MImageCount(ImageRef.get_count()) {
337+
}
338+
#endif
339+
340+
template <typename AllocatorT, int Dims = Dimensions,
341+
typename = detail::enable_if_t<(Dims > 0) && (Dims < 3) &&
342+
IsImageArrayAcc>>
343+
image_accessor(image<Dims + 1, AllocatorT> &ImageRef,
344+
handler &CommandGroupHandlerRef, int ImageElementSize)
345+
#ifdef __SYCL_DEVICE_ONLY__
346+
{
347+
// TODO: Implement this function.
348+
}
349+
#else
350+
: AccessorBaseHost(id<3>(0, 0, 0) /* Offset,*/,
351+
detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
352+
detail::convertToArrayOfN<3, 1>(ImageRef.get_range()),
353+
AccessMode, detail::getSyclObjImpl(ImageRef).get(),
354+
Dimensions, ImageElementSize),
355+
MImageSize(ImageRef.get_size()), MImageCount(ImageRef.get_count()) {
356+
// TODO: Implement this function.
357+
}
358+
#endif
359+
360+
/* TODO -- common interface members -- */
361+
362+
#ifdef __SYCL_DEVICE_ONLY__
363+
// TODO: Define the get_size(), get_count() methods.
364+
#else
365+
size_t get_size() const { return MImageSize; };
366+
size_t get_count() const { return MImageCount; };
367+
#endif
368+
369+
template <int Dim, typename T> struct IsValidCoordDataT;
370+
template <typename T> struct IsValidCoordDataT<1, T> {
371+
constexpr static bool value =
372+
detail::is_contained<T,
373+
detail::type_list<cl_int, cl_float>>::type::value;
374+
};
375+
template <typename T> struct IsValidCoordDataT<2, T> {
376+
constexpr static bool value = detail::is_contained<
377+
T, detail::type_list<cl_int2, cl_float2>>::type::value;
378+
};
379+
template <typename T> struct IsValidCoordDataT<3, T> {
380+
constexpr static bool value = detail::is_contained<
381+
T, detail::type_list<cl_int4, cl_float4>>::type::value;
382+
};
383+
384+
// Available only when: (accessTarget == access::target::image ||
385+
// accessTarget == access::target::host_image) && accessMode ==
386+
// access::mode::read
387+
template <typename CoordT, int Dims = Dimensions,
388+
typename = detail::enable_if_t<
389+
(Dims > 0) && (IsValidCoordDataT<Dims, CoordT>::value) &&
390+
(IsImageAcc || IsHostImageAcc) && IsImageAccessRead>>
391+
DataT read(const CoordT &Coords) const {
392+
// TODO: To be implemented.
393+
throw cl::sycl::feature_not_supported("Read API is not implemented.");
394+
return;
395+
};
396+
397+
// Available only when: (accessTarget == access::target::image || accessTarget
398+
// == access::target::host_image) && accessMode == access::mode::read
399+
template <typename CoordT, int Dims = Dimensions,
400+
typename = detail::enable_if_t<
401+
(Dims > 0) && (IsValidCoordDataT<Dims, CoordT>::value) &&
402+
(IsImageAcc || IsHostImageAcc) && IsImageAccessRead>>
403+
DataT read(const CoordT &Coords, const sampler &Smpl) const {
404+
// TODO: To be implemented.
405+
throw cl::sycl::feature_not_supported("Read API is not implemented.");
406+
return;
407+
};
408+
409+
// Available only when: (accessTarget == access::target::image || accessTarget
410+
// == access::target::host_image) && accessMode == access::mode::write ||
411+
// accessMode == access::mode::discard_write
412+
template <typename CoordT, int Dims = Dimensions,
413+
typename = detail::enable_if_t<
414+
(Dims > 0) && (detail::is_intn<CoordT>::value) &&
415+
(IsValidCoordDataT<Dims, CoordT>::value) &&
416+
(IsImageAcc || IsHostImageAcc) && IsImageAccessAnyWrite>>
417+
void write(const CoordT &Coords, const DataT &Color) const {
418+
// TODO: To be implemented.
419+
throw cl::sycl::feature_not_supported("Write API is not implemented.");
420+
return;
421+
};
422+
423+
// Available only when: accessTarget == access::target::image_array &&
424+
// dimensions < 3
425+
//__image_array_slice__ operator[](size_t index) const;
426+
};
427+
237428
} // namespace detail
238429

239430
template <typename DataT, int Dimensions, access::mode AccessMode,
@@ -315,8 +506,8 @@ class accessor :
315506
#else
316507

317508
using AccessorBaseHost::getAccessRange;
318-
using AccessorBaseHost::getOffset;
319509
using AccessorBaseHost::getMemoryRange;
510+
using AccessorBaseHost::getOffset;
320511

321512
char padding[sizeof(detail::AccessorImplDevice<AdjustedDim>) +
322513
sizeof(PtrType) - sizeof(detail::AccessorBaseHost)];
@@ -720,86 +911,62 @@ class accessor<DataT, Dimensions, AccessMode, access::target::local,
720911
bool operator!=(const accessor &Rhs) const { return !(*this == Rhs); }
721912
};
722913

723-
// Image accessor
724-
template <typename DataT, int Dimensions, access::mode AccessMode,
725-
access::target AccessTarget, access::placeholder IsPlaceholder>
726-
class image_accessor {
727-
static_assert(AccessTarget == access::target::image ||
728-
AccessTarget == access::target::host_image ||
729-
AccessTarget == access::target::image_array,
730-
"Expected image type");
731-
// TODO: Check if placeholder is applicable here.
732-
public:
733-
using value_type = DataT;
734-
using reference = DataT &;
735-
using const_reference = const DataT &;
736-
737-
/* Available only when: accessTarget == access::target::host_image */
738-
// template <typename AllocatorT>
739-
// accessor(image<dimensions, AllocatorT> &imageRef);
740-
/* Available only when: accessTarget == access::target::image */
741-
// template <typename AllocatorT>
742-
// accessor(image<dimensions, AllocatorT> &imageRef,
743-
// handler &commandGroupHandlerRef);
744-
745-
/* Available only when: accessTarget == access::target::image_array &&
746-
dimensions < 3 */
747-
// template <typename AllocatorT>
748-
// accessor(image<dimensions + 1, AllocatorT> &imageRef,
749-
// handler &commandGroupHandlerRef);
750-
751-
/* TODO -- common interface members -- */
752-
// size_t get_size() const;
753-
754-
// size_t get_count() const;
755-
756-
/* Available only when: (accessTarget == access::target::image || accessTarget
757-
== access::target::host_image) && accessMode == access::mode::read */
758-
// template <typename coordT> dataT read(const coordT &coords) const;
759-
760-
/* Available only when: (accessTarget == access::target::image || accessTarget
761-
== access::target::host_image) && accessMode == access::mode::read */
762-
// template <typename coordT>
763-
// dataT read(const coordT &coords, const sampler &smpl) const;
764-
765-
/* Available only when: (accessTarget == access::target::image || accessTarget
766-
== access::target::host_image) && accessMode == access::mode::write ||
767-
accessMode == access::mode::discard_write */
768-
// template <typename coordT>
769-
// void write(const coordT &coords, const dataT &color) const;
770-
771-
/* Available only when: accessTarget == access::target::image_array &&
772-
dimensions < 3 */
773-
//__image_array_slice__ operator[](size_t index) const;
774-
};
775-
776914
// Image accessors
915+
// Available only when: accessTarget == access::target::host_image
916+
// template <typename AllocatorT>
917+
// accessor(image<dimensions, AllocatorT> &imageRef);
777918
template <typename DataT, int Dimensions, access::mode AccessMode,
778919
access::placeholder IsPlaceholder>
779920
class accessor<DataT, Dimensions, AccessMode, access::target::image,
780921
IsPlaceholder>
781-
: public image_accessor<DataT, Dimensions, AccessMode,
782-
access::target::image, IsPlaceholder> {};
922+
: public detail::image_accessor<DataT, Dimensions, AccessMode,
923+
access::target::image, IsPlaceholder> {
924+
public:
925+
template <typename AllocatorT>
926+
accessor(cl::sycl::image<Dimensions, AllocatorT> &Image,
927+
handler &CommandGroupHandler)
928+
: detail::image_accessor<DataT, Dimensions, AccessMode, access::target::image,
929+
IsPlaceholder>(
930+
Image, CommandGroupHandler,
931+
(detail::getSyclObjImpl(Image))->getElementSize()) {
932+
CommandGroupHandler.associateWithHandler(*this);
933+
}
934+
};
783935

936+
// Available only when: accessTarget == access::target::image
937+
// template <typename AllocatorT>
938+
// accessor(image<dimensions, AllocatorT> &imageRef,
939+
// handler &commandGroupHandlerRef);
784940
template <typename DataT, int Dimensions, access::mode AccessMode,
785941
access::placeholder IsPlaceholder>
786942
class accessor<DataT, Dimensions, AccessMode, access::target::host_image,
787943
IsPlaceholder>
788-
: public image_accessor<DataT, Dimensions, AccessMode,
789-
access::target::host_image, IsPlaceholder> {};
944+
: public detail::image_accessor<DataT, Dimensions, AccessMode,
945+
access::target::host_image, IsPlaceholder> {
946+
public:
947+
template <typename AllocatorT>
948+
accessor(cl::sycl::image<Dimensions, AllocatorT> &Image)
949+
: detail::image_accessor<DataT, Dimensions, AccessMode,
950+
access::target::host_image, IsPlaceholder>(
951+
Image, (detail::getSyclObjImpl(Image))->getElementSize()) {}
952+
};
790953

954+
// Available only when: accessTarget == access::target::image_array &&
955+
// dimensions < 3
956+
// template <typename AllocatorT> accessor(image<dimensions + 1,
957+
// AllocatorT> &imageRef, handler &commandGroupHandlerRef);
791958
template <typename DataT, int Dimensions, access::mode AccessMode,
792959
access::placeholder IsPlaceholder>
793960
class accessor<DataT, Dimensions, AccessMode, access::target::image_array,
794961
IsPlaceholder>
795-
: public image_accessor<DataT, Dimensions, AccessMode,
796-
access::target::image_array, IsPlaceholder> {};
962+
: public detail::image_accessor<DataT, Dimensions, AccessMode,
963+
access::target::image_array, IsPlaceholder> {
964+
// TODO: To be Implemented.
965+
};
797966

798967
} // namespace sycl
799968
} // namespace cl
800969

801-
802-
803970
namespace std {
804971
template <typename DataT, int Dimensions, cl::sycl::access::mode AccessMode,
805972
cl::sycl::access::target AccessTarget,

0 commit comments

Comments
 (0)