DPC++ Runtime
Runtime libraries for oneAPI DPC++
memcpy2d.hpp
Go to the documentation of this file.
1 //==----- memcpy2d.hpp --- SYCL 2D memcpy 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 #pragma once
9 
10 #include <sycl/handler.hpp>
11 #include <sycl/queue.hpp>
12 #include <sycl/usm/usm_enums.hpp>
14 
15 namespace sycl {
16 inline namespace _V1 {
17 template <typename T, typename>
18 void handler::ext_oneapi_memcpy2d(void *Dest, size_t DestPitch, const void *Src,
19  size_t SrcPitch, size_t Width,
20  size_t Height) {
21  throwIfGraphAssociated<
22  ext::oneapi::experimental::detail::UnsupportedGraphFeatures::
23  sycl_ext_oneapi_memcpy2d>();
24  throwIfActionIsCreated();
25  if (Width > DestPitch)
27  "Destination pitch must be greater than or equal "
28  "to the width specified in 'ext_oneapi_memcpy2d'");
29  if (Width > SrcPitch)
31  "Source pitch must be greater than or equal "
32  "to the width specified in 'ext_oneapi_memcpy2d'");
33 
34  // Get the type of the pointers.
35  context Ctx = detail::createSyclObjFromImpl<context>(getContextImplPtr());
36  usm::alloc SrcAllocType = get_pointer_type(Src, Ctx);
37  usm::alloc DestAllocType = get_pointer_type(Dest, Ctx);
38  bool SrcIsHost =
39  SrcAllocType == usm::alloc::unknown || SrcAllocType == usm::alloc::host;
40  bool DestIsHost =
41  DestAllocType == usm::alloc::unknown || DestAllocType == usm::alloc::host;
42 
43  // Do the following:
44  // 1. If both are host, use host_task to copy.
45  // 2. If either pointer is host or the backend supports native memcpy2d, use
46  // special command.
47  // 3. Otherwise, launch a kernel for copying.
48  if (SrcIsHost && DestIsHost) {
49  commonUSMCopy2DFallbackHostTask<T>(Src, SrcPitch, Dest, DestPitch, Width,
50  Height);
51  } else if (SrcIsHost || DestIsHost || supportsUSMMemcpy2D()) {
52  ext_oneapi_memcpy2d_impl(Dest, DestPitch, Src, SrcPitch, Width, Height);
53  } else {
54  commonUSMCopy2DFallbackKernel<T>(Src, SrcPitch, Dest, DestPitch, Width,
55  Height);
56  }
57 }
58 
59 template <typename T>
60 void handler::ext_oneapi_copy2d(const T *Src, size_t SrcPitch, T *Dest,
61  size_t DestPitch, size_t Width, size_t Height) {
62  if (Width > DestPitch)
64  "Destination pitch must be greater than or equal "
65  "to the width specified in 'ext_oneapi_copy2d'");
66  if (Width > SrcPitch)
68  "Source pitch must be greater than or equal "
69  "to the width specified in 'ext_oneapi_copy2d'");
70 
71  // Get the type of the pointers.
72  context Ctx = detail::createSyclObjFromImpl<context>(getContextImplPtr());
73  usm::alloc SrcAllocType = get_pointer_type(Src, Ctx);
74  usm::alloc DestAllocType = get_pointer_type(Dest, Ctx);
75  bool SrcIsHost =
76  SrcAllocType == usm::alloc::unknown || SrcAllocType == usm::alloc::host;
77  bool DestIsHost =
78  DestAllocType == usm::alloc::unknown || DestAllocType == usm::alloc::host;
79 
80  // Do the following:
81  // 1. If both are host, use host_task to copy.
82  // 2. If either pointer is host or of the backend supports native memcpy2d,
83  // use special command.
84  // 3. Otherwise, launch a kernel for copying.
85  if (SrcIsHost && DestIsHost) {
86  commonUSMCopy2DFallbackHostTask<T>(Src, SrcPitch, Dest, DestPitch, Width,
87  Height);
88  } else if (SrcIsHost || DestIsHost || supportsUSMMemcpy2D()) {
89  ext_oneapi_memcpy2d_impl(Dest, DestPitch * sizeof(T), Src,
90  SrcPitch * sizeof(T), Width * sizeof(T), Height);
91  } else {
92  commonUSMCopy2DFallbackKernel<T>(Src, SrcPitch, Dest, DestPitch, Width,
93  Height);
94  }
95 }
96 
97 template <typename T, typename>
98 void handler::ext_oneapi_memset2d(void *Dest, size_t DestPitch, int Value,
99  size_t Width, size_t Height) {
100  throwIfActionIsCreated();
101  if (Width > DestPitch)
103  "Destination pitch must be greater than or equal "
104  "to the width specified in 'ext_oneapi_memset2d'");
105  T CharVal = static_cast<T>(Value);
106 
107  context Ctx = detail::createSyclObjFromImpl<context>(getContextImplPtr());
108  usm::alloc DestAllocType = get_pointer_type(Dest, Ctx);
109 
110  // If the backends supports 2D fill we use that. Otherwise we use a fallback
111  // kernel. If the target is on host we will always do the operation on host.
112  if (DestAllocType == usm::alloc::unknown || DestAllocType == usm::alloc::host)
113  commonUSMFill2DFallbackHostTask(Dest, DestPitch, CharVal, Width, Height);
114  else if (supportsUSMMemset2D())
115  ext_oneapi_memset2d_impl(Dest, DestPitch, Value, Width, Height);
116  else
117  commonUSMFill2DFallbackKernel(Dest, DestPitch, CharVal, Width, Height);
118 }
119 
120 template <typename T>
121 void handler::ext_oneapi_fill2d(void *Dest, size_t DestPitch, const T &Pattern,
122  size_t Width, size_t Height) {
123  throwIfActionIsCreated();
124  static_assert(is_device_copyable<T>::value,
125  "Pattern must be device copyable");
126  if (Width > DestPitch)
128  "Destination pitch must be greater than or equal "
129  "to the width specified in 'ext_oneapi_fill2d'");
130 
131  context Ctx = detail::createSyclObjFromImpl<context>(getContextImplPtr());
132  usm::alloc DestAllocType = get_pointer_type(Dest, Ctx);
133 
134  // If the backends supports 2D fill we use that. Otherwise we use a fallback
135  // kernel. If the target is on host we will always do the operation on host.
136  if (DestAllocType == usm::alloc::unknown || DestAllocType == usm::alloc::host)
137  commonUSMFill2DFallbackHostTask(Dest, DestPitch, Pattern, Width, Height);
138  else if (supportsUSMFill2D())
139  ext_oneapi_fill2d_impl(Dest, DestPitch, &Pattern, sizeof(T), Width, Height);
140  else
141  commonUSMFill2DFallbackKernel(Dest, DestPitch, Pattern, Width, Height);
142 }
143 
144 template <typename T, typename>
145 event queue::ext_oneapi_memcpy2d(void *Dest, size_t DestPitch, const void *Src,
146  size_t SrcPitch, size_t Width, size_t Height,
147  event DepEvent,
148  const detail::code_location &CodeLoc) {
149  return submit(
150  [=](handler &CGH) {
151  CGH.depends_on(DepEvent);
152  CGH.ext_oneapi_memcpy2d<T>(Dest, DestPitch, Src, SrcPitch, Width,
153  Height);
154  },
155  CodeLoc);
156 }
157 
158 template <typename T, typename>
159 event queue::ext_oneapi_memcpy2d(void *Dest, size_t DestPitch, const void *Src,
160  size_t SrcPitch, size_t Width, size_t Height,
161  const std::vector<event> &DepEvents,
162  const detail::code_location &CodeLoc) {
163  return submit(
164  [=](handler &CGH) {
165  CGH.depends_on(DepEvents);
166  CGH.ext_oneapi_memcpy2d<T>(Dest, DestPitch, Src, SrcPitch, Width,
167  Height);
168  },
169  CodeLoc);
170 }
171 
172 template <typename T>
173 event queue::ext_oneapi_copy2d(const T *Src, size_t SrcPitch, T *Dest,
174  size_t DestPitch, size_t Width, size_t Height,
175  const detail::code_location &CodeLoc) {
176  return submit(
177  [=](handler &CGH) {
178  CGH.ext_oneapi_copy2d<T>(Src, SrcPitch, Dest, DestPitch, Width, Height);
179  },
180  CodeLoc);
181 }
182 
183 template <typename T>
184 event queue::ext_oneapi_copy2d(const T *Src, size_t SrcPitch, T *Dest,
185  size_t DestPitch, size_t Width, size_t Height,
186  event DepEvent,
187  const detail::code_location &CodeLoc) {
188  return submit(
189  [=](handler &CGH) {
190  CGH.depends_on(DepEvent);
191  CGH.ext_oneapi_copy2d<T>(Src, SrcPitch, Dest, DestPitch, Width, Height);
192  },
193  CodeLoc);
194 }
195 
196 template <typename T>
197 event queue::ext_oneapi_copy2d(const T *Src, size_t SrcPitch, T *Dest,
198  size_t DestPitch, size_t Width, size_t Height,
199  const std::vector<event> &DepEvents,
200  const detail::code_location &CodeLoc) {
201  return submit(
202  [=](handler &CGH) {
203  CGH.depends_on(DepEvents);
204  CGH.ext_oneapi_copy2d<T>(Src, SrcPitch, Dest, DestPitch, Width, Height);
205  },
206  CodeLoc);
207 }
208 
209 template <typename T, typename>
210 event queue::ext_oneapi_memset2d(void *Dest, size_t DestPitch, int Value,
211  size_t Width, size_t Height,
212  const detail::code_location &CodeLoc) {
213  return submit(
214  [=](handler &CGH) {
215  CGH.ext_oneapi_memset2d<T>(Dest, DestPitch, Value, Width, Height);
216  },
217  CodeLoc);
218 }
219 
220 template <typename T, typename>
221 event queue::ext_oneapi_memset2d(void *Dest, size_t DestPitch, int Value,
222  size_t Width, size_t Height, event DepEvent,
223  const detail::code_location &CodeLoc) {
224  return submit(
225  [=](handler &CGH) {
226  CGH.depends_on(DepEvent);
227  CGH.ext_oneapi_memset2d<T>(Dest, DestPitch, Value, Width, Height);
228  },
229  CodeLoc);
230 }
231 
232 template <typename T, typename>
233 event queue::ext_oneapi_memset2d(void *Dest, size_t DestPitch, int Value,
234  size_t Width, size_t Height,
235  const std::vector<event> &DepEvents,
236  const detail::code_location &CodeLoc) {
237  return submit(
238  [=](handler &CGH) {
239  CGH.depends_on(DepEvents);
240  CGH.ext_oneapi_memset2d<T>(Dest, DestPitch, Value, Width, Height);
241  },
242  CodeLoc);
243 }
244 
245 template <typename T>
246 event queue::ext_oneapi_fill2d(void *Dest, size_t DestPitch, const T &Pattern,
247  size_t Width, size_t Height,
248  const detail::code_location &CodeLoc) {
249  return submit(
250  [=](handler &CGH) {
251  CGH.ext_oneapi_fill2d<T>(Dest, DestPitch, Pattern, Width, Height);
252  },
253  CodeLoc);
254 }
255 
256 template <typename T>
257 event queue::ext_oneapi_fill2d(void *Dest, size_t DestPitch, const T &Pattern,
258  size_t Width, size_t Height, event DepEvent,
259  const detail::code_location &CodeLoc) {
260  return submit(
261  [=](handler &CGH) {
262  CGH.depends_on(DepEvent);
263  CGH.ext_oneapi_fill2d<T>(Dest, DestPitch, Pattern, Width, Height);
264  },
265  CodeLoc);
266 }
267 
268 template <typename T>
269 event queue::ext_oneapi_fill2d(void *Dest, size_t DestPitch, const T &Pattern,
270  size_t Width, size_t Height,
271  const std::vector<event> &DepEvents,
272  const detail::code_location &CodeLoc) {
273  return submit(
274  [=](handler &CGH) {
275  CGH.depends_on(DepEvents);
276  CGH.ext_oneapi_fill2d<T>(Dest, DestPitch, Pattern, Width, Height);
277  },
278  CodeLoc);
279 }
280 } // namespace _V1
281 } // namespace sycl
The context class represents a SYCL context on which kernel functions may be executed.
Definition: context.hpp:50
An event object can be used to synchronize memory transfers, enqueues of kernels and signaling barrie...
Definition: event.hpp:44
Command group handler class.
Definition: handler.hpp:468
void depends_on(event Event)
Registers event dependencies on this command group.
Definition: handler.cpp:1491
void ext_oneapi_memset2d(void *Dest, size_t DestPitch, int Value, size_t Width, size_t Height)
Fills the memory pointed by a USM pointer with the value specified.
Definition: memcpy2d.hpp:98
void ext_oneapi_fill2d(void *Dest, size_t DestPitch, const T &Pattern, size_t Width, size_t Height)
Fills the memory pointed by a USM pointer with the value specified.
Definition: memcpy2d.hpp:121
void ext_oneapi_copy2d(const T *Src, size_t SrcPitch, T *Dest, size_t DestPitch, size_t Width, size_t Height)
Copies data from one 2D memory region to another, both pointed by USM pointers.
Definition: memcpy2d.hpp:60
void ext_oneapi_memcpy2d(void *Dest, size_t DestPitch, const void *Src, size_t SrcPitch, size_t Width, size_t Height)
Copies data from one 2D memory region to another, both pointed by USM pointers.
Definition: memcpy2d.hpp:18
event ext_oneapi_copy2d(const T *Src, size_t SrcPitch, T *Dest, size_t DestPitch, size_t Width, size_t Height, const detail::code_location &CodeLoc=detail::code_location::current())
Copies data from one 2D memory region to another, both pointed by USM pointers.
Definition: memcpy2d.hpp:173
event ext_oneapi_memset2d(void *Dest, size_t DestPitch, int Value, size_t Width, size_t Height, const detail::code_location &CodeLoc=detail::code_location::current())
Fills the memory pointed by a USM pointer with the value specified.
Definition: memcpy2d.hpp:210
event ext_oneapi_memcpy2d(void *Dest, size_t DestPitch, const void *Src, size_t SrcPitch, size_t Width, size_t Height, const detail::code_location &CodeLoc=detail::code_location::current())
Copies data from one 2D memory region to another, both pointed by USM pointers.
Definition: queue.hpp:801
event ext_oneapi_fill2d(void *Dest, size_t DestPitch, const T &Pattern, size_t Width, size_t Height, const detail::code_location &CodeLoc=detail::code_location::current())
Fills the memory pointed by a USM pointer with the value specified.
Definition: memcpy2d.hpp:246
std::enable_if_t< std::is_invocable_r_v< void, T, handler & >, event > submit(T CGF, const detail::code_location &CodeLoc=detail::code_location::current())
Submits a command group function object to the queue, in order to be scheduled for execution on the d...
Definition: queue.hpp:340
usm::alloc get_pointer_type(const void *ptr, const context &ctxt)
Query the allocation type from a USM pointer.
Definition: usm_impl.cpp:536
std::error_code make_error_code(sycl::errc E) noexcept
Constructs an error code using e and sycl_category()
Definition: exception.cpp:64
Definition: access.hpp:18
is_device_copyable is a user specializable class template to indicate that a type T is device copyabl...