Skip to content

Commit b4983f2

Browse files
narenbagriaCompute-Runtime-Automation
narenbagria
authored and
Compute-Runtime-Automation
committed
fix: update CopyBufferRectBytes*2d* builtins
Related-To: NEO-16155 Signed-off-by: Narendra Bagria <[email protected]>
1 parent 7833d62 commit b4983f2

File tree

11 files changed

+73
-71
lines changed

11 files changed

+73
-71
lines changed

‎manifests/manifest.yml‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -41,7 +41,7 @@ components:
4141
dest_dir: kernels_bin
4242
type: git
4343
branch: kernels_bin
44-
revision: 3748-5703
44+
revision: 3748-5704
4545
level_zero:
4646
asset_name: level_zero
4747
dest_dir: level_zero

‎opencl/source/built_ins/builtins_dispatch_builder.cpp‎

Lines changed: 14 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -222,6 +222,8 @@ class BuiltInOp<EBuiltInOps::copyBufferRect> : public BuiltinDispatchInfoBuilder
222222

223223
const uint32_t rootDeviceIndex = clDevice.getRootDeviceIndex();
224224
const int dimensions = is3D ? 3 : 2;
225+
const size_t originSize = is3D ? sizeof(OffsetType) * 4 : sizeof(OffsetType) * 2;
226+
const size_t pitchSize = is3D ? sizeof(OffsetType) * 2 : sizeof(OffsetType);
225227

226228
if (this->clDevice.getProductHelper().isCopyBufferRectSplitSupported()) {
227229
DispatchInfoBuilder<SplitDispatch::Dim::d1D, SplitDispatch::SplitMode::kernelSplit> kernelSplit3DBuilder(clDevice);
@@ -287,27 +289,27 @@ class BuiltInOp<EBuiltInOps::copyBufferRect> : public BuiltinDispatchInfoBuilder
287289

288290
// arg2 = srcOrigin
289291
OffsetType kSrcOrigin[4] = {static_cast<OffsetType>(operationParams.srcOffset.x + srcOffsetFromAlignedPtr), static_cast<OffsetType>(operationParams.srcOffset.y), static_cast<OffsetType>(operationParams.srcOffset.z), 0};
290-
kernelSplit3DBuilder.setArg(SplitDispatch::RegionCoordX::left, 2, sizeof(OffsetType) * 4, kSrcOrigin);
292+
kernelSplit3DBuilder.setArg(SplitDispatch::RegionCoordX::left, 2, originSize, kSrcOrigin);
291293
kSrcOrigin[0] += static_cast<uint32_t>(leftSize);
292-
kernelSplit3DBuilder.setArg(SplitDispatch::RegionCoordX::middle, 2, sizeof(OffsetType) * 4, kSrcOrigin);
294+
kernelSplit3DBuilder.setArg(SplitDispatch::RegionCoordX::middle, 2, originSize, kSrcOrigin);
293295
kSrcOrigin[0] += static_cast<uint32_t>(middleSizeBytes);
294-
kernelSplit3DBuilder.setArg(SplitDispatch::RegionCoordX::right, 2, sizeof(OffsetType) * 4, kSrcOrigin);
296+
kernelSplit3DBuilder.setArg(SplitDispatch::RegionCoordX::right, 2, originSize, kSrcOrigin);
295297

296298
// arg3 = dstOrigin
297299
OffsetType kDstOrigin[4] = {static_cast<OffsetType>(operationParams.dstOffset.x + dstOffsetFromAlignedPtr), static_cast<OffsetType>(operationParams.dstOffset.y), static_cast<OffsetType>(operationParams.dstOffset.z), 0};
298-
kernelSplit3DBuilder.setArg(SplitDispatch::RegionCoordX::left, 3, sizeof(OffsetType) * 4, kDstOrigin);
300+
kernelSplit3DBuilder.setArg(SplitDispatch::RegionCoordX::left, 3, originSize, kDstOrigin);
299301
kDstOrigin[0] += static_cast<uint32_t>(leftSize);
300-
kernelSplit3DBuilder.setArg(SplitDispatch::RegionCoordX::middle, 3, sizeof(OffsetType) * 4, kDstOrigin);
302+
kernelSplit3DBuilder.setArg(SplitDispatch::RegionCoordX::middle, 3, originSize, kDstOrigin);
301303
kDstOrigin[0] += static_cast<uint32_t>(middleSizeBytes);
302-
kernelSplit3DBuilder.setArg(SplitDispatch::RegionCoordX::right, 3, sizeof(OffsetType) * 4, kDstOrigin);
304+
kernelSplit3DBuilder.setArg(SplitDispatch::RegionCoordX::right, 3, originSize, kDstOrigin);
303305

304306
// arg4 = srcPitch
305307
OffsetType kSrcPitch[2] = {static_cast<OffsetType>(operationParams.srcRowPitch), static_cast<OffsetType>(operationParams.srcSlicePitch)};
306-
kernelSplit3DBuilder.setArg(4, sizeof(OffsetType) * 2, kSrcPitch);
308+
kernelSplit3DBuilder.setArg(4, pitchSize, kSrcPitch);
307309

308310
// arg5 = dstPitch
309311
OffsetType kDstPitch[2] = {static_cast<OffsetType>(operationParams.dstRowPitch), static_cast<OffsetType>(operationParams.dstSlicePitch)};
310-
kernelSplit3DBuilder.setArg(5, sizeof(OffsetType) * 2, kDstPitch);
312+
kernelSplit3DBuilder.setArg(5, pitchSize, kDstPitch);
311313

312314
// Set-up work sizes
313315
kernelSplit3DBuilder.setDispatchGeometry(SplitDispatch::RegionCoordX::left, Vec3<size_t>{leftSize, operationParams.size.y, operationParams.size.z}, Vec3<size_t>{0, 0, 0}, Vec3<size_t>{0, 0, 0});
@@ -350,19 +352,19 @@ class BuiltInOp<EBuiltInOps::copyBufferRect> : public BuiltinDispatchInfoBuilder
350352

351353
// arg2 = srcOrigin
352354
OffsetType kSrcOrigin[4] = {static_cast<OffsetType>(operationParams.srcOffset.x + srcOffsetFromAlignedPtr), static_cast<OffsetType>(operationParams.srcOffset.y), static_cast<OffsetType>(operationParams.srcOffset.z), 0};
353-
kernelNoSplit3DBuilder.setArg(2, sizeof(OffsetType) * 4, kSrcOrigin);
355+
kernelNoSplit3DBuilder.setArg(2, originSize, kSrcOrigin);
354356

355357
// arg3 = dstOrigin
356358
OffsetType kDstOrigin[4] = {static_cast<OffsetType>(operationParams.dstOffset.x + dstOffsetFromAlignedPtr), static_cast<OffsetType>(operationParams.dstOffset.y), static_cast<OffsetType>(operationParams.dstOffset.z), 0};
357-
kernelNoSplit3DBuilder.setArg(3, sizeof(OffsetType) * 4, kDstOrigin);
359+
kernelNoSplit3DBuilder.setArg(3, originSize, kDstOrigin);
358360

359361
// arg4 = srcPitch
360362
OffsetType kSrcPitch[2] = {static_cast<OffsetType>(operationParams.srcRowPitch), static_cast<OffsetType>(operationParams.srcSlicePitch)};
361-
kernelNoSplit3DBuilder.setArg(4, sizeof(OffsetType) * 2, kSrcPitch);
363+
kernelNoSplit3DBuilder.setArg(4, pitchSize, kSrcPitch);
362364

363365
// arg5 = dstPitch
364366
OffsetType kDstPitch[2] = {static_cast<OffsetType>(operationParams.dstRowPitch), static_cast<OffsetType>(operationParams.dstSlicePitch)};
365-
kernelNoSplit3DBuilder.setArg(5, sizeof(OffsetType) * 2, kDstPitch);
367+
kernelNoSplit3DBuilder.setArg(5, pitchSize, kDstPitch);
366368

367369
// Set-up work sizes
368370
kernelNoSplit3DBuilder.setDispatchGeometry(operationParams.size, Vec3<size_t>{0, 0, 0}, Vec3<size_t>{0, 0, 0});

‎opencl/test/unit_test/command_queue/enqueue_read_buffer_rect_tests.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -619,12 +619,12 @@ HWTEST_F(EnqueueReadWriteBufferRectDispatch, givenOffsetResultingInMisalignedPtr
619619
}
620620
}
621621

622-
if (kernelInfo.getArgDescriptorAt(3).as<ArgDescValue>().elements[0].size == 4 * sizeof(uint32_t)) { // size of uint4 DstOrigin
622+
if (kernelInfo.getArgDescriptorAt(3).as<ArgDescValue>().elements[0].size == 2 * sizeof(uint32_t)) { // size of uint2 DstOrigin
623623
auto dstOffset = (uint32_t *)(kernel->getCrossThreadData() +
624624
kernelInfo.getArgDescriptorAt(3).as<ArgDescValue>().elements[0].offset);
625625
EXPECT_EQ(hostOffset.x + ptrDiff(misalignedDstPtr, alignDown(misalignedDstPtr, 4)), *dstOffset);
626626
} else {
627-
// DstOrigin arg should be 16 bytes in size, if that changes, above if path should be modified
627+
// DstOrigin arg should be 8 bytes in size, if that changes, above if path should be modified
628628
EXPECT_TRUE(false);
629629
}
630630
}

‎opencl/test/unit_test/command_queue/enqueue_write_buffer_rect_tests.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -624,12 +624,12 @@ HWTEST_F(EnqueueReadWriteBufferRectDispatch, givenOffsetResultingInMisalignedPtr
624624
}
625625
}
626626

627-
if (kernelInfo.getArgDescriptorAt(2).as<ArgDescValue>().elements[0].size == 4 * sizeof(uint32_t)) { // size of uint4 SrcOrigin
627+
if (kernelInfo.getArgDescriptorAt(2).as<ArgDescValue>().elements[0].size == 2 * sizeof(uint32_t)) { // size of uint2 SrcOrigin
628628
auto dstOffset = (uint32_t *)(kernel->getCrossThreadData() +
629629
kernelInfo.getArgDescriptorAt(2).as<ArgDescValue>().elements[0].offset);
630630
EXPECT_EQ(hostOffset.x + ptrDiff(misalignedHostPtr, alignDown(misalignedHostPtr, 4)), *dstOffset);
631631
} else {
632-
// SrcOrigin arg should be 16 bytes in size, if that changes, above if path should be modified
632+
// SrcOrigin arg should be 8 bytes in size, if that changes, above if path should be modified
633633
EXPECT_TRUE(false);
634634
}
635635
}

‎shared/source/built_ins/kernels/copy_buffer_rect.builtin_kernel‎

Lines changed: 13 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
/*
2-
* Copyright (C) 2018-2024 Intel Corporation
2+
* Copyright (C) 2018-2025 Intel Corporation
33
*
44
* SPDX-License-Identifier: MIT
55
*
@@ -10,17 +10,17 @@ R"===(
1010
__kernel void CopyBufferRectBytes2d(
1111
__global const char* src,
1212
__global char* dst,
13-
uint4 SrcOrigin,
14-
uint4 DstOrigin,
15-
uint2 SrcPitch,
16-
uint2 DstPitch )
13+
uint2 SrcOrigin,
14+
uint2 DstOrigin,
15+
uint SrcPitch,
16+
uint DstPitch )
1717

1818
{
1919
int x = get_global_id(0);
2020
int y = get_global_id(1);
2121

22-
uint LSrcOffset = x + SrcOrigin.x + ( ( y + SrcOrigin.y ) * SrcPitch.x );
23-
uint LDstOffset = x + DstOrigin.x + ( ( y + DstOrigin.y ) * DstPitch.x );
22+
uint LSrcOffset = x + SrcOrigin.x + ( ( y + SrcOrigin.y ) * SrcPitch );
23+
uint LDstOffset = x + DstOrigin.x + ( ( y + DstOrigin.y ) * DstPitch );
2424

2525
*( dst + LDstOffset ) = *( src + LSrcOffset );
2626

@@ -29,17 +29,17 @@ __kernel void CopyBufferRectBytes2d(
2929
__kernel void CopyBufferRectBytesMiddle2d(
3030
const __global uint* src,
3131
__global uint* dst,
32-
uint4 SrcOrigin,
33-
uint4 DstOrigin,
34-
uint2 SrcPitch,
35-
uint2 DstPitch )
32+
uint2 SrcOrigin,
33+
uint2 DstOrigin,
34+
uint SrcPitch,
35+
uint DstPitch )
3636

3737
{
3838
int x = get_global_id(0);
3939
int y = get_global_id(1);
4040

41-
uint LSrcOffset = SrcOrigin.x + ( ( y + SrcOrigin.y ) * SrcPitch.x );
42-
uint LDstOffset = DstOrigin.x + ( ( y + DstOrigin.y ) * DstPitch.x );
41+
uint LSrcOffset = SrcOrigin.x + ( ( y + SrcOrigin.y ) * SrcPitch );
42+
uint LDstOffset = DstOrigin.x + ( ( y + DstOrigin.y ) * DstPitch );
4343

4444
src += LSrcOffset >> 2;
4545
dst += LDstOffset >> 2;

‎shared/source/built_ins/kernels/copy_buffer_rect_stateless.builtin_kernel‎

Lines changed: 13 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
/*
2-
* Copyright (C) 2018-2024 Intel Corporation
2+
* Copyright (C) 2018-2025 Intel Corporation
33
*
44
* SPDX-License-Identifier: MIT
55
*
@@ -10,17 +10,17 @@ R"===(
1010
__kernel void CopyBufferRectBytes2dStateless(
1111
__global const char* src,
1212
__global char* dst,
13-
ulong4 SrcOrigin,
14-
ulong4 DstOrigin,
15-
ulong2 SrcPitch,
16-
ulong2 DstPitch )
13+
ulong2 SrcOrigin,
14+
ulong2 DstOrigin,
15+
ulong SrcPitch,
16+
ulong DstPitch )
1717

1818
{
1919
size_t x = get_global_id(0);
2020
size_t y = get_global_id(1);
2121

22-
size_t LSrcOffset = x + SrcOrigin.x + ( ( y + SrcOrigin.y ) * SrcPitch.x );
23-
size_t LDstOffset = x + DstOrigin.x + ( ( y + DstOrigin.y ) * DstPitch.x );
22+
size_t LSrcOffset = x + SrcOrigin.x + ( ( y + SrcOrigin.y ) * SrcPitch );
23+
size_t LDstOffset = x + DstOrigin.x + ( ( y + DstOrigin.y ) * DstPitch );
2424

2525
*( dst + LDstOffset ) = *( src + LSrcOffset );
2626

@@ -29,17 +29,17 @@ __kernel void CopyBufferRectBytes2dStateless(
2929
__kernel void CopyBufferRectBytesMiddle2dStateless(
3030
const __global uint* src,
3131
__global uint* dst,
32-
ulong4 SrcOrigin,
33-
ulong4 DstOrigin,
34-
ulong2 SrcPitch,
35-
ulong2 DstPitch )
32+
ulong2 SrcOrigin,
33+
ulong2 DstOrigin,
34+
ulong SrcPitch,
35+
ulong DstPitch )
3636

3737
{
3838
size_t x = get_global_id(0);
3939
size_t y = get_global_id(1);
4040

41-
size_t LSrcOffset = SrcOrigin.x + ( ( y + SrcOrigin.y ) * SrcPitch.x );
42-
size_t LDstOffset = DstOrigin.x + ( ( y + DstOrigin.y ) * DstPitch.x );
41+
size_t LSrcOffset = SrcOrigin.x + ( ( y + SrcOrigin.y ) * SrcPitch );
42+
size_t LDstOffset = DstOrigin.x + ( ( y + DstOrigin.y ) * DstPitch );
4343

4444
src += LSrcOffset >> 2;
4545
dst += LDstOffset >> 2;

‎shared/test/common/helpers/kernel_binary_helper_hash_value.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -7,5 +7,5 @@
77

88
#include "shared/test/common/helpers/kernel_binary_helper.h"
99

10-
const std::string KernelBinaryHelper::BUILT_INS("15672580764041246108");
11-
const std::string KernelBinaryHelper::BUILT_INS_WITH_IMAGES("2205520382307710565_images");
10+
const std::string KernelBinaryHelper::BUILT_INS("6133084427540774618");
11+
const std::string KernelBinaryHelper::BUILT_INS_WITH_IMAGES("15342443153856668610_images");

shared/test/common/test_files/2205520382307710565_images.cl renamed to shared/test/common/test_files/15342443153856668610_images.cl

Lines changed: 12 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -224,17 +224,17 @@ __kernel void FillBufferSSHOffset(
224224
__kernel void CopyBufferRectBytes2d(
225225
__global const char* src,
226226
__global char* dst,
227-
uint4 SrcOrigin,
228-
uint4 DstOrigin,
229-
uint2 SrcPitch,
230-
uint2 DstPitch )
227+
uint2 SrcOrigin,
228+
uint2 DstOrigin,
229+
uint SrcPitch,
230+
uint DstPitch )
231231

232232
{
233233
int x = get_global_id(0);
234234
int y = get_global_id(1);
235235

236-
uint LSrcOffset = x + SrcOrigin.x + ( ( y + SrcOrigin.y ) * SrcPitch.x );
237-
uint LDstOffset = x + DstOrigin.x + ( ( y + DstOrigin.y ) * DstPitch.x );
236+
uint LSrcOffset = x + SrcOrigin.x + ( ( y + SrcOrigin.y ) * SrcPitch );
237+
uint LDstOffset = x + DstOrigin.x + ( ( y + DstOrigin.y ) * DstPitch );
238238

239239
*( dst + LDstOffset ) = *( src + LSrcOffset );
240240

@@ -243,17 +243,17 @@ __kernel void CopyBufferRectBytes2d(
243243
__kernel void CopyBufferRectBytesMiddle2d(
244244
const __global uint* src,
245245
__global uint* dst,
246-
uint4 SrcOrigin,
247-
uint4 DstOrigin,
248-
uint2 SrcPitch,
249-
uint2 DstPitch )
246+
uint2 SrcOrigin,
247+
uint2 DstOrigin,
248+
uint SrcPitch,
249+
uint DstPitch )
250250

251251
{
252252
int x = get_global_id(0);
253253
int y = get_global_id(1);
254254

255-
uint LSrcOffset = SrcOrigin.x + ( ( y + SrcOrigin.y ) * SrcPitch.x );
256-
uint LDstOffset = DstOrigin.x + ( ( y + DstOrigin.y ) * DstPitch.x );
255+
uint LSrcOffset = SrcOrigin.x + ( ( y + SrcOrigin.y ) * SrcPitch );
256+
uint LDstOffset = DstOrigin.x + ( ( y + DstOrigin.y ) * DstPitch );
257257

258258
src += LSrcOffset >> 2;
259259
dst += LDstOffset >> 2;

shared/test/common/test_files/2205520382307710565_images_options.txt renamed to shared/test/common/test_files/15342443153856668610_images_options.txt

File renamed without changes.

shared/test/common/test_files/15672580764041246108.cl renamed to shared/test/common/test_files/6133084427540774618.cl

Lines changed: 13 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
/*
2-
* Copyright (C) 2024 Intel Corporation
2+
* Copyright (C) 2024-2025 Intel Corporation
33
*
44
* SPDX-License-Identifier: MIT
55
*
@@ -224,17 +224,17 @@ __kernel void FillBufferSSHOffset(
224224
__kernel void CopyBufferRectBytes2d(
225225
__global const char* src,
226226
__global char* dst,
227-
uint4 SrcOrigin,
228-
uint4 DstOrigin,
229-
uint2 SrcPitch,
230-
uint2 DstPitch )
227+
uint2 SrcOrigin,
228+
uint2 DstOrigin,
229+
uint SrcPitch,
230+
uint DstPitch )
231231

232232
{
233233
int x = get_global_id(0);
234234
int y = get_global_id(1);
235235

236-
uint LSrcOffset = x + SrcOrigin.x + ( ( y + SrcOrigin.y ) * SrcPitch.x );
237-
uint LDstOffset = x + DstOrigin.x + ( ( y + DstOrigin.y ) * DstPitch.x );
236+
uint LSrcOffset = x + SrcOrigin.x + ( ( y + SrcOrigin.y ) * SrcPitch );
237+
uint LDstOffset = x + DstOrigin.x + ( ( y + DstOrigin.y ) * DstPitch );
238238

239239
*( dst + LDstOffset ) = *( src + LSrcOffset );
240240

@@ -243,17 +243,17 @@ __kernel void CopyBufferRectBytes2d(
243243
__kernel void CopyBufferRectBytesMiddle2d(
244244
const __global uint* src,
245245
__global uint* dst,
246-
uint4 SrcOrigin,
247-
uint4 DstOrigin,
248-
uint2 SrcPitch,
249-
uint2 DstPitch )
246+
uint2 SrcOrigin,
247+
uint2 DstOrigin,
248+
uint SrcPitch,
249+
uint DstPitch )
250250

251251
{
252252
int x = get_global_id(0);
253253
int y = get_global_id(1);
254254

255-
uint LSrcOffset = SrcOrigin.x + ( ( y + SrcOrigin.y ) * SrcPitch.x );
256-
uint LDstOffset = DstOrigin.x + ( ( y + DstOrigin.y ) * DstPitch.x );
255+
uint LSrcOffset = SrcOrigin.x + ( ( y + SrcOrigin.y ) * SrcPitch );
256+
uint LDstOffset = DstOrigin.x + ( ( y + DstOrigin.y ) * DstPitch );
257257

258258
src += LSrcOffset >> 2;
259259
dst += LDstOffset >> 2;

0 commit comments

Comments
 (0)