Skip to content

Commit eb7e1b5

Browse files
AlexeySachkovAlexeySachkov
committed
[SYCL] Fix and re-enable bindless image copy tests on L0
1 parent 3d6a85d commit eb7e1b5

38 files changed

+379
-196
lines changed

‎sycl/source/detail/cg.hpp‎

Lines changed: 7 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -600,6 +600,7 @@ class CGCopyImage : public CG {
600600
ur_image_format_t MSrcImageFormat;
601601
ur_image_format_t MDstImageFormat;
602602
ur_exp_image_copy_flags_t MImageCopyFlags;
603+
ur_exp_image_copy_input_types_t MImageInputTypes;
603604
ur_rect_offset_t MSrcOffset;
604605
ur_rect_offset_t MDstOffset;
605606
ur_rect_region_t MCopyExtent;
@@ -609,14 +610,15 @@ class CGCopyImage : public CG {
609610
ur_image_desc_t DstImageDesc, ur_image_format_t SrcImageFormat,
610611
ur_image_format_t DstImageFormat,
611612
ur_exp_image_copy_flags_t ImageCopyFlags,
613+
ur_exp_image_copy_input_types_t ImageInputTypes,
612614
ur_rect_offset_t SrcOffset, ur_rect_offset_t DstOffset,
613615
ur_rect_region_t CopyExtent, CG::StorageInitHelper CGData,
614616
detail::code_location loc = {})
615617
: CG(CGType::CopyImage, std::move(CGData), std::move(loc)), MSrc(Src),
616618
MDst(Dst), MSrcImageDesc(SrcImageDesc), MDstImageDesc(DstImageDesc),
617619
MSrcImageFormat(SrcImageFormat), MDstImageFormat(DstImageFormat),
618-
MImageCopyFlags(ImageCopyFlags), MSrcOffset(SrcOffset),
619-
MDstOffset(DstOffset), MCopyExtent(CopyExtent) {}
620+
MImageCopyFlags(ImageCopyFlags), MImageInputTypes(ImageInputTypes),
621+
MSrcOffset(SrcOffset), MDstOffset(DstOffset), MCopyExtent(CopyExtent) {}
620622

621623
void *getSrc() const { return MSrc; }
622624
void *getDst() const { return MDst; }
@@ -625,6 +627,9 @@ class CGCopyImage : public CG {
625627
ur_image_format_t getSrcFormat() const { return MSrcImageFormat; }
626628
ur_image_format_t getDstFormat() const { return MDstImageFormat; }
627629
ur_exp_image_copy_flags_t getCopyFlags() const { return MImageCopyFlags; }
630+
ur_exp_image_copy_input_types_t getCopyInputTypes() const {
631+
return MImageInputTypes;
632+
}
628633
ur_rect_offset_t getSrcOffset() const { return MSrcOffset; }
629634
ur_rect_offset_t getDstOffset() const { return MDstOffset; }
630635
ur_rect_region_t getCopyExtent() const { return MCopyExtent; }

‎sycl/source/detail/handler_impl.hpp‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -120,6 +120,7 @@ class handler_impl {
120120
ur_image_format_t MSrcImageFormat = {};
121121
ur_image_format_t MDstImageFormat = {};
122122
ur_exp_image_copy_flags_t MImageCopyFlags = {};
123+
ur_exp_image_copy_input_types_t MImageCopyInputTypes = {};
123124

124125
ur_rect_offset_t MSrcOffset = {};
125126
ur_rect_offset_t MDestOffset = {};

‎sycl/source/detail/memory_manager.cpp‎

Lines changed: 6 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1570,8 +1570,10 @@ void MemoryManager::copy_image_bindless(
15701570
queue_impl &Queue, const void *Src, void *Dst,
15711571
const ur_image_desc_t &SrcDesc, const ur_image_desc_t &DstDesc,
15721572
const ur_image_format_t &SrcFormat, const ur_image_format_t &DstFormat,
1573-
const ur_exp_image_copy_flags_t Flags, ur_rect_offset_t SrcOffset,
1574-
ur_rect_offset_t DstOffset, ur_rect_region_t CopyExtent,
1573+
const ur_exp_image_copy_flags_t Flags,
1574+
const ur_exp_image_copy_input_types_t InputTypes,
1575+
ur_rect_offset_t SrcOffset, ur_rect_offset_t DstOffset,
1576+
ur_rect_region_t CopyExtent,
15751577
const std::vector<ur_event_handle_t> &DepEvents,
15761578
ur_event_handle_t *OutEvent) {
15771579
assert((Flags == UR_EXP_IMAGE_COPY_FLAG_HOST_TO_DEVICE ||
@@ -1594,8 +1596,8 @@ void MemoryManager::copy_image_bindless(
15941596

15951597
Adapter.call<UrApiKind::urBindlessImagesImageCopyExp>(
15961598
Queue.getHandleRef(), Src, Dst, &SrcDesc, &DstDesc, &SrcFormat,
1597-
&DstFormat, &CopyRegion, Flags, DepEvents.size(), DepEvents.data(),
1598-
OutEvent);
1599+
&DstFormat, &CopyRegion, Flags, InputTypes, DepEvents.size(),
1600+
DepEvents.data(), OutEvent);
15991601
}
16001602

16011603
} // namespace detail

‎sycl/source/detail/memory_manager.hpp‎

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -264,8 +264,10 @@ class MemoryManager {
264264
queue_impl &Queue, const void *Src, void *Dst,
265265
const ur_image_desc_t &SrcDesc, const ur_image_desc_t &DstDesc,
266266
const ur_image_format_t &SrcFormat, const ur_image_format_t &DstFormat,
267-
const ur_exp_image_copy_flags_t Flags, ur_rect_offset_t SrcOffset,
268-
ur_rect_offset_t DstOffset, ur_rect_region_t CopyExtent,
267+
const ur_exp_image_copy_flags_t Flags,
268+
const ur_exp_image_copy_input_types_t InputTypes,
269+
ur_rect_offset_t SrcOffset, ur_rect_offset_t DstOffset,
270+
ur_rect_region_t CopyExtent,
269271
const std::vector<ur_event_handle_t> &DepEvents,
270272
ur_event_handle_t *OutEvent);
271273
};

‎sycl/source/detail/scheduler/commands.cpp‎

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -3665,8 +3665,9 @@ ur_result_t ExecCGCommand::enqueueImpQueue() {
36653665
MemoryManager::copy_image_bindless, *MQueue, Copy->getSrc(),
36663666
Copy->getDst(), Copy->getSrcDesc(), Copy->getDstDesc(),
36673667
Copy->getSrcFormat(), Copy->getDstFormat(), Copy->getCopyFlags(),
3668-
Copy->getSrcOffset(), Copy->getDstOffset(), Copy->getCopyExtent(),
3669-
std::move(RawEvents), Event);
3668+
Copy->getCopyInputTypes(), Copy->getSrcOffset(),
3669+
Copy->getDstOffset(), Copy->getCopyExtent(), std::move(RawEvents),
3670+
Event);
36703671
Result != UR_RESULT_SUCCESS)
36713672
return Result;
36723673

‎sycl/source/handler.cpp‎

Lines changed: 75 additions & 52 deletions
Original file line numberDiff line numberDiff line change
@@ -239,16 +239,16 @@ fill_image_desc(const ext::oneapi::experimental::image_descriptor &ImgDesc) {
239239
return UrDesc;
240240
}
241241

242-
static void
243-
fill_copy_args(detail::handler_impl *impl,
244-
const ext::oneapi::experimental::image_descriptor &SrcImgDesc,
245-
const ext::oneapi::experimental::image_descriptor &DestImgDesc,
246-
ur_exp_image_copy_flags_t ImageCopyFlags, size_t SrcPitch,
247-
size_t DestPitch, sycl::range<3> SrcOffset = {0, 0, 0},
248-
sycl::range<3> SrcExtent = {0, 0, 0},
249-
sycl::range<3> DestOffset = {0, 0, 0},
250-
sycl::range<3> DestExtent = {0, 0, 0},
251-
sycl::range<3> CopyExtent = {0, 0, 0}) {
242+
static void fill_copy_args(
243+
detail::handler_impl *impl,
244+
const ext::oneapi::experimental::image_descriptor &SrcImgDesc,
245+
const ext::oneapi::experimental::image_descriptor &DestImgDesc,
246+
ur_exp_image_copy_flags_t ImageCopyFlags,
247+
ur_exp_image_copy_input_types_t ImageCopyInputTypes, size_t SrcPitch,
248+
size_t DestPitch, sycl::range<3> SrcOffset = {0, 0, 0},
249+
sycl::range<3> SrcExtent = {0, 0, 0}, sycl::range<3> DestOffset = {0, 0, 0},
250+
sycl::range<3> DestExtent = {0, 0, 0},
251+
sycl::range<3> CopyExtent = {0, 0, 0}) {
252252
SrcImgDesc.verify();
253253
DestImgDesc.verify();
254254

@@ -267,12 +267,13 @@ fill_copy_args(detail::handler_impl *impl,
267267
auto ZCopyExtentComponent = detail::fill_image_type(SrcImgDesc, UrSrcDesc);
268268
detail::fill_image_type(DestImgDesc, UrDestDesc);
269269

270-
// Copy args computed here are directly passed to UR. Various offsets and
271-
// extents end up passed as ur_rect_offset_t and ur_rect_region_t. Both those
272-
// structs expect their first component to be in bytes, not in pixels
273-
size_t SrcPixelSize = SrcImgDesc.num_channels * get_channel_size(SrcImgDesc);
274-
size_t DestPixelSize =
275-
DestImgDesc.num_channels * get_channel_size(DestImgDesc);
270+
// ur_rect_offset_t and ur_rect_offset_t which represent image offsets and
271+
// copy extents expect that X-axis offset and region width are specified in
272+
// bytes rather then in elements.
273+
auto SrcPixelSize =
274+
SrcImgDesc.num_channels * detail::get_channel_size(SrcImgDesc);
275+
auto DestPixelSize =
276+
DestImgDesc.num_channels * detail::get_channel_size(DestImgDesc);
276277

277278
impl->MSrcOffset = {SrcOffset[0] * SrcPixelSize, SrcOffset[1], SrcOffset[2]};
278279
impl->MDestOffset = {DestOffset[0] * DestPixelSize, DestOffset[1],
@@ -282,6 +283,7 @@ fill_copy_args(detail::handler_impl *impl,
282283
impl->MSrcImageFormat = UrSrcFormat;
283284
impl->MDstImageFormat = UrDestFormat;
284285
impl->MImageCopyFlags = ImageCopyFlags;
286+
impl->MImageCopyInputTypes = ImageCopyInputTypes;
285287

286288
if (CopyExtent.size() != 0) {
287289
impl->MCopyExtent = {CopyExtent[0] * SrcPixelSize, CopyExtent[1],
@@ -311,6 +313,7 @@ static void
311313
fill_copy_args(detail::handler_impl *impl,
312314
const ext::oneapi::experimental::image_descriptor &Desc,
313315
ur_exp_image_copy_flags_t ImageCopyFlags,
316+
ur_exp_image_copy_input_types_t ImageCopyInputTypes,
314317
sycl::range<3> SrcOffset = {0, 0, 0},
315318
sycl::range<3> SrcExtent = {0, 0, 0},
316319
sycl::range<3> DestOffset = {0, 0, 0},
@@ -320,29 +323,32 @@ fill_copy_args(detail::handler_impl *impl,
320323
size_t SrcPitch = SrcExtent[0] * Desc.num_channels * get_channel_size(Desc);
321324
size_t DestPitch = DestExtent[0] * Desc.num_channels * get_channel_size(Desc);
322325

323-
fill_copy_args(impl, Desc, Desc, ImageCopyFlags, SrcPitch, DestPitch,
324-
SrcOffset, SrcExtent, DestOffset, DestExtent, CopyExtent);
326+
fill_copy_args(impl, Desc, Desc, ImageCopyFlags, ImageCopyInputTypes,
327+
SrcPitch, DestPitch, SrcOffset, SrcExtent, DestOffset,
328+
DestExtent, CopyExtent);
325329
}
326330

327-
static void
328-
fill_copy_args(detail::handler_impl *impl,
329-
const ext::oneapi::experimental::image_descriptor &Desc,
330-
ur_exp_image_copy_flags_t ImageCopyFlags, size_t SrcPitch,
331-
size_t DestPitch, sycl::range<3> SrcOffset = {0, 0, 0},
332-
sycl::range<3> SrcExtent = {0, 0, 0},
333-
sycl::range<3> DestOffset = {0, 0, 0},
334-
sycl::range<3> DestExtent = {0, 0, 0},
335-
sycl::range<3> CopyExtent = {0, 0, 0}) {
331+
static void fill_copy_args(
332+
detail::handler_impl *impl,
333+
const ext::oneapi::experimental::image_descriptor &Desc,
334+
ur_exp_image_copy_flags_t ImageCopyFlags,
335+
ur_exp_image_copy_input_types_t ImageCopyInputTypes, size_t SrcPitch,
336+
size_t DestPitch, sycl::range<3> SrcOffset = {0, 0, 0},
337+
sycl::range<3> SrcExtent = {0, 0, 0}, sycl::range<3> DestOffset = {0, 0, 0},
338+
sycl::range<3> DestExtent = {0, 0, 0},
339+
sycl::range<3> CopyExtent = {0, 0, 0}) {
336340

337-
fill_copy_args(impl, Desc, Desc, ImageCopyFlags, SrcPitch, DestPitch,
338-
SrcOffset, SrcExtent, DestOffset, DestExtent, CopyExtent);
341+
fill_copy_args(impl, Desc, Desc, ImageCopyFlags, ImageCopyInputTypes,
342+
SrcPitch, DestPitch, SrcOffset, SrcExtent, DestOffset,
343+
DestExtent, CopyExtent);
339344
}
340345

341346
static void
342347
fill_copy_args(detail::handler_impl *impl,
343348
const ext::oneapi::experimental::image_descriptor &SrcImgDesc,
344349
const ext::oneapi::experimental::image_descriptor &DestImgDesc,
345350
ur_exp_image_copy_flags_t ImageCopyFlags,
351+
ur_exp_image_copy_input_types_t ImageCopyInputTypes,
346352
sycl::range<3> SrcOffset = {0, 0, 0},
347353
sycl::range<3> SrcExtent = {0, 0, 0},
348354
sycl::range<3> DestOffset = {0, 0, 0},
@@ -354,9 +360,9 @@ fill_copy_args(detail::handler_impl *impl,
354360
size_t DestPitch =
355361
DestExtent[0] * DestImgDesc.num_channels * get_channel_size(DestImgDesc);
356362

357-
fill_copy_args(impl, SrcImgDesc, DestImgDesc, ImageCopyFlags, SrcPitch,
358-
DestPitch, SrcOffset, SrcExtent, DestOffset, DestExtent,
359-
CopyExtent);
363+
fill_copy_args(impl, SrcImgDesc, DestImgDesc, ImageCopyFlags,
364+
ImageCopyInputTypes, SrcPitch, DestPitch, SrcOffset, SrcExtent,
365+
DestOffset, DestExtent, CopyExtent);
360366
}
361367

362368
} // namespace detail
@@ -905,8 +911,8 @@ event handler::finalize() {
905911
CommandGroup.reset(new detail::CGCopyImage(
906912
MSrcPtr, MDstPtr, impl->MSrcImageDesc, impl->MDstImageDesc,
907913
impl->MSrcImageFormat, impl->MDstImageFormat, impl->MImageCopyFlags,
908-
impl->MSrcOffset, impl->MDestOffset, impl->MCopyExtent,
909-
std::move(impl->CGData), MCodeLoc));
914+
impl->MImageCopyInputTypes, impl->MSrcOffset, impl->MDestOffset,
915+
impl->MCopyExtent, std::move(impl->CGData), MCodeLoc));
910916
break;
911917
case detail::CGType::SemaphoreWait:
912918
CommandGroup.reset(
@@ -1622,7 +1628,8 @@ void handler::ext_oneapi_copy(
16221628
MDstPtr = reinterpret_cast<void *>(Dest.raw_handle);
16231629

16241630
detail::fill_copy_args(get_impl(), DestImgDesc,
1625-
UR_EXP_IMAGE_COPY_FLAG_HOST_TO_DEVICE);
1631+
UR_EXP_IMAGE_COPY_FLAG_HOST_TO_DEVICE,
1632+
UR_EXP_IMAGE_COPY_INPUT_TYPES_MEM_TO_IMAGE);
16261633

16271634
setType(detail::CGType::CopyImage);
16281635
}
@@ -1640,7 +1647,8 @@ void handler::ext_oneapi_copy(
16401647
MDstPtr = reinterpret_cast<void *>(Dest.raw_handle);
16411648

16421649
detail::fill_copy_args(get_impl(), DestImgDesc,
1643-
UR_EXP_IMAGE_COPY_FLAG_HOST_TO_DEVICE, SrcOffset,
1650+
UR_EXP_IMAGE_COPY_FLAG_HOST_TO_DEVICE,
1651+
UR_EXP_IMAGE_COPY_INPUT_TYPES_MEM_TO_IMAGE, SrcOffset,
16441652
SrcExtent, DestOffset, {0, 0, 0}, CopyExtent);
16451653

16461654
setType(detail::CGType::CopyImage);
@@ -1657,7 +1665,8 @@ void handler::ext_oneapi_copy(
16571665
MDstPtr = Dest;
16581666

16591667
detail::fill_copy_args(get_impl(), SrcImgDesc,
1660-
UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_HOST);
1668+
UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_HOST,
1669+
UR_EXP_IMAGE_COPY_INPUT_TYPES_IMAGE_TO_MEM);
16611670

16621671
setType(detail::CGType::CopyImage);
16631672
}
@@ -1676,7 +1685,8 @@ void handler::ext_oneapi_copy(
16761685
MDstPtr = Dest;
16771686

16781687
detail::fill_copy_args(get_impl(), SrcImgDesc,
1679-
UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_HOST, SrcOffset,
1688+
UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_HOST,
1689+
UR_EXP_IMAGE_COPY_INPUT_TYPES_IMAGE_TO_MEM, SrcOffset,
16801690
{0, 0, 0}, DestOffset, DestExtent, CopyExtent);
16811691

16821692
setType(detail::CGType::CopyImage);
@@ -1705,11 +1715,13 @@ void handler::ext_oneapi_copy(
17051715
Desc.width * Desc.num_channels * detail::get_channel_size(Desc);
17061716

17071717
if (ImageCopyFlags == UR_EXP_IMAGE_COPY_FLAG_HOST_TO_DEVICE) {
1708-
detail::fill_copy_args(get_impl(), Desc, ImageCopyFlags, HostRowPitch,
1709-
DeviceRowPitch);
1718+
detail::fill_copy_args(get_impl(), Desc, ImageCopyFlags,
1719+
UR_EXP_IMAGE_COPY_INPUT_TYPES_MEM_TO_MEM,
1720+
HostRowPitch, DeviceRowPitch);
17101721
} else if (ImageCopyFlags == UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_HOST) {
1711-
detail::fill_copy_args(get_impl(), Desc, ImageCopyFlags, DeviceRowPitch,
1712-
HostRowPitch);
1722+
detail::fill_copy_args(get_impl(), Desc, ImageCopyFlags,
1723+
UR_EXP_IMAGE_COPY_INPUT_TYPES_MEM_TO_MEM,
1724+
DeviceRowPitch, HostRowPitch);
17131725
} else {
17141726
throw sycl::exception(make_error_code(errc::invalid),
17151727
"Copy Error: This copy function only performs host "
@@ -1746,10 +1758,12 @@ void handler::ext_oneapi_copy(
17461758
// Fill the host extent based on the type of copy.
17471759
if (ImageCopyFlags == UR_EXP_IMAGE_COPY_FLAG_HOST_TO_DEVICE) {
17481760
detail::fill_copy_args(get_impl(), DeviceImgDesc, ImageCopyFlags,
1761+
UR_EXP_IMAGE_COPY_INPUT_TYPES_MEM_TO_MEM,
17491762
HostRowPitch, DeviceRowPitch, SrcOffset, HostExtent,
17501763
DestOffset, {0, 0, 0}, CopyExtent);
17511764
} else if (ImageCopyFlags == UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_HOST) {
17521765
detail::fill_copy_args(get_impl(), DeviceImgDesc, ImageCopyFlags,
1766+
UR_EXP_IMAGE_COPY_INPUT_TYPES_MEM_TO_MEM,
17531767
DeviceRowPitch, HostRowPitch, SrcOffset, {0, 0, 0},
17541768
DestOffset, HostExtent, CopyExtent);
17551769
} else {
@@ -1774,7 +1788,8 @@ void handler::ext_oneapi_copy(
17741788
MDstPtr = reinterpret_cast<void *>(Dest.raw_handle);
17751789

17761790
detail::fill_copy_args(get_impl(), SrcImgDesc, DestImgDesc,
1777-
UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE);
1791+
UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE,
1792+
UR_EXP_IMAGE_COPY_INPUT_TYPES_IMAGE_TO_IMAGE);
17781793

17791794
setType(detail::CGType::CopyImage);
17801795
}
@@ -1794,8 +1809,10 @@ void handler::ext_oneapi_copy(
17941809
MDstPtr = reinterpret_cast<void *>(Dest.raw_handle);
17951810

17961811
detail::fill_copy_args(get_impl(), SrcImgDesc, DestImgDesc,
1797-
UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE, SrcOffset,
1798-
{0, 0, 0}, DestOffset, {0, 0, 0}, CopyExtent);
1812+
UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE,
1813+
UR_EXP_IMAGE_COPY_INPUT_TYPES_IMAGE_TO_IMAGE,
1814+
SrcOffset, {0, 0, 0}, DestOffset, {0, 0, 0},
1815+
CopyExtent);
17991816

18001817
setType(detail::CGType::CopyImage);
18011818
}
@@ -1813,7 +1830,8 @@ void handler::ext_oneapi_copy(
18131830
MDstPtr = Dest;
18141831

18151832
detail::fill_copy_args(get_impl(), SrcImgDesc, DestImgDesc,
1816-
UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE, 0,
1833+
UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE,
1834+
UR_EXP_IMAGE_COPY_INPUT_TYPES_IMAGE_TO_MEM, 0,
18171835
DestRowPitch);
18181836

18191837
setType(detail::CGType::CopyImage);
@@ -1834,7 +1852,8 @@ void handler::ext_oneapi_copy(
18341852
MDstPtr = Dest;
18351853

18361854
detail::fill_copy_args(get_impl(), SrcImgDesc, DestImgDesc,
1837-
UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE, 0,
1855+
UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE,
1856+
UR_EXP_IMAGE_COPY_INPUT_TYPES_IMAGE_TO_MEM, 0,
18381857
DestRowPitch, SrcOffset, {0, 0, 0}, DestOffset,
18391858
{0, 0, 0}, CopyExtent);
18401859

@@ -1854,8 +1873,9 @@ void handler::ext_oneapi_copy(
18541873
MDstPtr = reinterpret_cast<void *>(Dest.raw_handle);
18551874

18561875
detail::fill_copy_args(get_impl(), SrcImgDesc, DestImgDesc,
1857-
UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE, SrcRowPitch,
1858-
0);
1876+
UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE,
1877+
UR_EXP_IMAGE_COPY_INPUT_TYPES_MEM_TO_IMAGE,
1878+
SrcRowPitch, 0);
18591879

18601880
setType(detail::CGType::CopyImage);
18611881
}
@@ -1875,9 +1895,10 @@ void handler::ext_oneapi_copy(
18751895
MDstPtr = reinterpret_cast<void *>(Dest.raw_handle);
18761896

18771897
detail::fill_copy_args(get_impl(), SrcImgDesc, DestImgDesc,
1878-
UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE, SrcRowPitch,
1879-
0, SrcOffset, {0, 0, 0}, DestOffset, {0, 0, 0},
1880-
CopyExtent);
1898+
UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE,
1899+
UR_EXP_IMAGE_COPY_INPUT_TYPES_MEM_TO_IMAGE,
1900+
SrcRowPitch, 0, SrcOffset, {0, 0, 0}, DestOffset,
1901+
{0, 0, 0}, CopyExtent);
18811902

18821903
setType(detail::CGType::CopyImage);
18831904
}
@@ -1904,6 +1925,7 @@ void handler::ext_oneapi_copy(
19041925
if (ImageCopyFlags == UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE ||
19051926
ImageCopyFlags == UR_EXP_IMAGE_COPY_FLAG_HOST_TO_HOST) {
19061927
detail::fill_copy_args(get_impl(), SrcImgDesc, DestImgDesc, ImageCopyFlags,
1928+
UR_EXP_IMAGE_COPY_INPUT_TYPES_MEM_TO_MEM,
19071929
SrcRowPitch, DestRowPitch);
19081930
} else {
19091931
throw sycl::exception(make_error_code(errc::invalid),
@@ -1933,6 +1955,7 @@ void handler::ext_oneapi_copy(
19331955
if (ImageCopyFlags == UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE ||
19341956
ImageCopyFlags == UR_EXP_IMAGE_COPY_FLAG_HOST_TO_HOST) {
19351957
detail::fill_copy_args(get_impl(), SrcImgDesc, DestImgDesc, ImageCopyFlags,
1958+
UR_EXP_IMAGE_COPY_INPUT_TYPES_MEM_TO_MEM,
19361959
SrcRowPitch, DestRowPitch, SrcOffset, {0, 0, 0},
19371960
DestOffset, {0, 0, 0}, CopyExtent);
19381961
} else {

‎sycl/test-e2e/bindless_images/copies/copy_subregion_1D.cpp‎

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,4 @@
11
// REQUIRES: aspect-ext_oneapi_bindless_images
2-
// REQUIRES: cuda
32
// XFAIL: hip
43
// XFAIL-TRACKER: https://github.com/intel/llvm/issues/19957
54

‎sycl/test-e2e/bindless_images/copies/copy_subregion_2D.cpp‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,6 @@
11
// REQUIRES: aspect-ext_oneapi_bindless_images
2-
// REQUIRES: cuda
2+
// XFAIL: linux && arch-intel_gpu_acm_g10
3+
// XFAIL-TRACKER: https://github.com/intel/llvm/issues/20004
34
// XFAIL: hip
45
// XFAIL-TRACKER: https://github.com/intel/llvm/issues/19957
56

0 commit comments

Comments
 (0)