Skip to content

Commit 4d953fd

Browse files
KamilKoprykIntelCompute-Runtime-Automation
KamilKoprykIntel
authored and
Compute-Runtime-Automation
committed
Correct timestamps offset calculation
Related-To: NEO-6653 Signed-off-by: Kamil Kopryk <[email protected]>
1 parent 8f85d4b commit 4d953fd

9 files changed

+23
-35
lines changed

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

Lines changed: 3 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
/*
2-
* Copyright (C) 2020-2021 Intel Corporation
2+
* Copyright (C) 2020-2022 Intel Corporation
33
*
44
* SPDX-License-Identifier: MIT
55
*
@@ -50,9 +50,8 @@ __kernel void QueryKernelTimestamps(__global ulong* srcEvents, __global ulong* d
5050
ulong globalEnd = GetTimestampValue(srcPtr, timestampSizeInDw, 3);
5151

5252
if(packetUsed > 1) {
53-
uint timestampsOffsets = 4;
5453
for(uint i = 1; i < packetUsed; i++) {
55-
timestampsOffsets += i;
54+
uint timestampsOffsets = 4 * i;
5655
if(contextStart > GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets)) {
5756
contextStart = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets);
5857
}
@@ -91,9 +90,8 @@ __kernel void QueryKernelTimestampsWithOffsets(__global ulong* srcEvents, __glob
9190
ulong globalEnd = GetTimestampValue(srcPtr, timestampSizeInDw, 3);
9291

9392
if(packetUsed > 1) {
94-
uint timestampsOffsets = 4;
9593
for(uint i = 1; i < packetUsed; i++) {
96-
timestampsOffsets += i;
94+
uint timestampsOffsets = 4 * i;
9795
if(contextStart > GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets)) {
9896
contextStart = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets);
9997
}

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

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,11 +1,11 @@
11
/*
2-
* Copyright (C) 2020-2021 Intel Corporation
2+
* Copyright (C) 2020-2022 Intel Corporation
33
*
44
* SPDX-License-Identifier: MIT
55
*
66
*/
77

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

10-
const std::string KernelBinaryHelper::BUILT_INS("7836643072362489210");
11-
const std::string KernelBinaryHelper::BUILT_INS_WITH_IMAGES("10613492505254921609_images");
10+
const std::string KernelBinaryHelper::BUILT_INS("7998916142903730155");
11+
const std::string KernelBinaryHelper::BUILT_INS_WITH_IMAGES("16526264370178379440_images");

shared/test/common/test_files/10613492505254921609_images.cl renamed to shared/test/common/test_files/16526264370178379440_images.cl

Lines changed: 3 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
/*
2-
* Copyright (C) 2020-2021 Intel Corporation
2+
* Copyright (C) 2022 Intel Corporation
33
*
44
* SPDX-License-Identifier: MIT
55
*
@@ -263,9 +263,8 @@ __kernel void QueryKernelTimestamps(__global ulong* srcEvents, __global ulong* d
263263
ulong globalEnd = GetTimestampValue(srcPtr, timestampSizeInDw, 3);
264264

265265
if(packetUsed > 1) {
266-
uint timestampsOffsets = 4;
267266
for(uint i = 1; i < packetUsed; i++) {
268-
timestampsOffsets += i;
267+
uint timestampsOffsets = 4 * i;
269268
if(contextStart > GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets)) {
270269
contextStart = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets);
271270
}
@@ -304,9 +303,8 @@ __kernel void QueryKernelTimestampsWithOffsets(__global ulong* srcEvents, __glob
304303
ulong globalEnd = GetTimestampValue(srcPtr, timestampSizeInDw, 3);
305304

306305
if(packetUsed > 1) {
307-
uint timestampsOffsets = 4;
308306
for(uint i = 1; i < packetUsed; i++) {
309-
timestampsOffsets += i;
307+
uint timestampsOffsets = 4 * i;
310308
if(contextStart > GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets)) {
311309
contextStart = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets);
312310
}

shared/test/common/test_files/10613492505254921609_images_options.txt renamed to shared/test/common/test_files/16526264370178379440_images_options.txt

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
/*
2-
* Copyright (C) 2021 Intel Corporation
2+
* Copyright (C) 2021-2022 Intel Corporation
33
*
44
* SPDX-License-Identifier: MIT
55
*

shared/test/common/test_files/7836643072362489210.cl renamed to shared/test/common/test_files/7998916142903730155.cl

Lines changed: 3 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
/*
2-
* Copyright (C) 2020-2021 Intel Corporation
2+
* Copyright (C) 2020-2022 Intel Corporation
33
*
44
* SPDX-License-Identifier: MIT
55
*
@@ -263,9 +263,8 @@ __kernel void QueryKernelTimestamps(__global ulong* srcEvents, __global ulong* d
263263
ulong globalEnd = GetTimestampValue(srcPtr, timestampSizeInDw, 3);
264264

265265
if(packetUsed > 1) {
266-
uint timestampsOffsets = 4;
267266
for(uint i = 1; i < packetUsed; i++) {
268-
timestampsOffsets += i;
267+
uint timestampsOffsets = 4 * i;
269268
if(contextStart > GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets)) {
270269
contextStart = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets);
271270
}
@@ -304,9 +303,8 @@ __kernel void QueryKernelTimestampsWithOffsets(__global ulong* srcEvents, __glob
304303
ulong globalEnd = GetTimestampValue(srcPtr, timestampSizeInDw, 3);
305304

306305
if(packetUsed > 1) {
307-
uint timestampsOffsets = 4;
308306
for(uint i = 1; i < packetUsed; i++) {
309-
timestampsOffsets += i;
307+
uint timestampsOffsets = 4 * i;
310308
if(contextStart > GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets)) {
311309
contextStart = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets);
312310
}

shared/test/common/test_files/7836643072362489210_options.txt renamed to shared/test/common/test_files/7998916142903730155_options.txt

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
/*
2-
* Copyright (C) 2021 Intel Corporation
2+
* Copyright (C) 2021-2022 Intel Corporation
33
*
44
* SPDX-License-Identifier: MIT
55
*

‎shared/test/common/test_files/builtin_copyfill.cl‎

Lines changed: 3 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
/*
2-
* Copyright (C) 2020-2021 Intel Corporation
2+
* Copyright (C) 2020-2022 Intel Corporation
33
*
44
* SPDX-License-Identifier: MIT
55
*
@@ -263,9 +263,8 @@ __kernel void QueryKernelTimestamps(__global ulong* srcEvents, __global ulong* d
263263
ulong globalEnd = GetTimestampValue(srcPtr, timestampSizeInDw, 3);
264264

265265
if(packetUsed > 1) {
266-
uint timestampsOffsets = 4;
267266
for(uint i = 1; i < packetUsed; i++) {
268-
timestampsOffsets *= i;
267+
uint timestampsOffsets = 4 * i;
269268
if(contextStart > GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets)) {
270269
contextStart = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets);
271270
}
@@ -304,9 +303,8 @@ __kernel void QueryKernelTimestampsWithOffsets(__global ulong* srcEvents, __glob
304303
ulong globalEnd = GetTimestampValue(srcPtr, timestampSizeInDw, 3);
305304

306305
if(packetUsed > 1) {
307-
uint timestampsOffsets = 4;
308306
for(uint i = 1; i < packetUsed; i++) {
309-
timestampsOffsets *= i;
307+
uint timestampsOffsets = 4 * i;
310308
if(contextStart > GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets)) {
311309
contextStart = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets);
312310
}

‎shared/test/common/test_files/builtin_copyfill_stateless.cl‎

Lines changed: 3 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
/*
2-
* Copyright (C) 2020-2021 Intel Corporation
2+
* Copyright (C) 2020-2022 Intel Corporation
33
*
44
* SPDX-License-Identifier: MIT
55
*
@@ -238,9 +238,8 @@ __kernel void QueryKernelTimestamps(__global ulong* srcEvents, __global ulong* d
238238
uint globalEnd = src[3];
239239

240240
if(packetUsed > 1) {
241-
uint timestampsOffsets = 4;
242241
for(uint i = 1; i < packetUsed; i++) {
243-
timestampsOffsets *= i;
242+
uint timestampsOffsets = 4 * i;
244243
if(contextStart > src[timestampsOffsets]) {
245244
contextStart = src[timestampsOffsets];
246245
}
@@ -287,9 +286,8 @@ __kernel void QueryKernelTimestampsWithOffsets(__global ulong* srcEvents, __glob
287286
uint globalEnd = src[3];
288287

289288
if(packetUsed > 1) {
290-
uint timestampsOffsets = 4;
291289
for(uint i = 1; i < packetUsed; i++) {
292-
timestampsOffsets *= i;
290+
uint timestampsOffsets = 4 * i;
293291
if(contextStart > src[timestampsOffsets]) {
294292
contextStart = src[timestampsOffsets];
295293
}

‎shared/test/common/test_files/builtin_images.cl‎

Lines changed: 3 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
/*
2-
* Copyright (C) 2020-2021 Intel Corporation
2+
* Copyright (C) 2020-2022 Intel Corporation
33
*
44
* SPDX-License-Identifier: MIT
55
*
@@ -263,9 +263,8 @@ __kernel void QueryKernelTimestamps(__global ulong* srcEvents, __global ulong* d
263263
ulong globalEnd = GetTimestampValue(srcPtr, timestampSizeInDw, 3);
264264

265265
if(packetUsed > 1) {
266-
uint timestampsOffsets = 4;
267266
for(uint i = 1; i < packetUsed; i++) {
268-
timestampsOffsets *= i;
267+
uint timestampsOffsets = 4 * i;
269268
if(contextStart > GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets)) {
270269
contextStart = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets);
271270
}
@@ -304,9 +303,8 @@ __kernel void QueryKernelTimestampsWithOffsets(__global ulong* srcEvents, __glob
304303
ulong globalEnd = GetTimestampValue(srcPtr, timestampSizeInDw, 3);
305304

306305
if(packetUsed > 1) {
307-
uint timestampsOffsets = 4;
308306
for(uint i = 1; i < packetUsed; i++) {
309-
timestampsOffsets *= i;
307+
uint timestampsOffsets = 4 * i;
310308
if(contextStart > GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets)) {
311309
contextStart = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets);
312310
}

0 commit comments

Comments
 (0)