Media/compute-runtime
1
0
Fork 0
mirror of https://github.com/intel/compute-runtime.git synced 2025-09-15 13:01:45 +08:00
Code Issues Packages Projects Releases Wiki Activity

Correct timestamps offset calculation

Browse Source 
Related-To: NEO-6653
Signed-off-by: Kamil Kopryk <kamil.kopryk@intel.com>
This commit is contained in:
Kamil Kopryk
2022-02-02 16:20:04 +00:00
committed by Compute-Runtime-Automation
parent 8f85d4b8f8
commit 4d953fd125
9 changed files with 23 additions and 35 deletions
Download Patch File Download Diff File Expand all files Collapse all files

8
shared/source/built_ins/kernels/copy_kernel_timestamps.builtin_kernel
View File

@ -1,5 +1,5 @@
/*
* Copyright (C) 2020-2021 Intel Corporation
* Copyright (C) 2020-2022 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
@ -50,9 +50,8 @@ __kernel void QueryKernelTimestamps(__global ulong* srcEvents, __global ulong* d
ulong globalEnd = GetTimestampValue(srcPtr, timestampSizeInDw, 3);
if(packetUsed > 1) {
uint timestampsOffsets = 4;
for(uint i = 1; i < packetUsed; i++) {
timestampsOffsets += i;
uint timestampsOffsets = 4 * i;
if(contextStart > GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets)) {
contextStart = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets);
}
@ -91,9 +90,8 @@ __kernel void QueryKernelTimestampsWithOffsets(__global ulong* srcEvents, __glob
ulong globalEnd = GetTimestampValue(srcPtr, timestampSizeInDw, 3);
if(packetUsed > 1) {
uint timestampsOffsets = 4;
for(uint i = 1; i < packetUsed; i++) {
timestampsOffsets += i;
uint timestampsOffsets = 4 * i;
if(contextStart > GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets)) {
contextStart = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets);
}

6
shared/test/common/helpers/kernel_binary_helper_hash_value.cpp
View File

@ -1,5 +1,5 @@
/*
* Copyright (C) 2020-2021 Intel Corporation
* Copyright (C) 2020-2022 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
@ -7,5 +7,5 @@
#include "shared/test/common/helpers/kernel_binary_helper.h"
const std::string KernelBinaryHelper::BUILT_INS("7836643072362489210");
const std::string KernelBinaryHelper::BUILT_INS_WITH_IMAGES("10613492505254921609_images");
const std::string KernelBinaryHelper::BUILT_INS("7998916142903730155");
const std::string KernelBinaryHelper::BUILT_INS_WITH_IMAGES("16526264370178379440_images");

8
shared/test/common/test_files/10613492505254921609_images.cl → shared/test/common/test_files/16526264370178379440_images.cl
View File

@ -1,5 +1,5 @@
/*
* Copyright (C) 2020-2021 Intel Corporation
* Copyright (C) 2022 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
@ -263,9 +263,8 @@ __kernel void QueryKernelTimestamps(__global ulong* srcEvents, __global ulong* d
ulong globalEnd = GetTimestampValue(srcPtr, timestampSizeInDw, 3);
if(packetUsed > 1) {
uint timestampsOffsets = 4;
for(uint i = 1; i < packetUsed; i++) {
timestampsOffsets += i;
uint timestampsOffsets = 4 * i;
if(contextStart > GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets)) {
contextStart = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets);
}
@ -304,9 +303,8 @@ __kernel void QueryKernelTimestampsWithOffsets(__global ulong* srcEvents, __glob
ulong globalEnd = GetTimestampValue(srcPtr, timestampSizeInDw, 3);
if(packetUsed > 1) {
uint timestampsOffsets = 4;
for(uint i = 1; i < packetUsed; i++) {
timestampsOffsets += i;
uint timestampsOffsets = 4 * i;
if(contextStart > GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets)) {
contextStart = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets);
}

2
shared/test/common/test_files/10613492505254921609_images_options.txt → shared/test/common/test_files/16526264370178379440_images_options.txt
View File

@ -1,5 +1,5 @@
/*
* Copyright (C) 2021 Intel Corporation
* Copyright (C) 2021-2022 Intel Corporation
*
* SPDX-License-Identifier: MIT
*

8
shared/test/common/test_files/7836643072362489210.cl → shared/test/common/test_files/7998916142903730155.cl
View File

@ -1,5 +1,5 @@
/*
* Copyright (C) 2020-2021 Intel Corporation
* Copyright (C) 2020-2022 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
@ -263,9 +263,8 @@ __kernel void QueryKernelTimestamps(__global ulong* srcEvents, __global ulong* d
ulong globalEnd = GetTimestampValue(srcPtr, timestampSizeInDw, 3);
if(packetUsed > 1) {
uint timestampsOffsets = 4;
for(uint i = 1; i < packetUsed; i++) {
timestampsOffsets += i;
uint timestampsOffsets = 4 * i;
if(contextStart > GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets)) {
contextStart = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets);
}
@ -304,9 +303,8 @@ __kernel void QueryKernelTimestampsWithOffsets(__global ulong* srcEvents, __glob
ulong globalEnd = GetTimestampValue(srcPtr, timestampSizeInDw, 3);
if(packetUsed > 1) {
uint timestampsOffsets = 4;
for(uint i = 1; i < packetUsed; i++) {
timestampsOffsets += i;
uint timestampsOffsets = 4 * i;
if(contextStart > GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets)) {
contextStart = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets);
}

2
shared/test/common/test_files/7836643072362489210_options.txt → shared/test/common/test_files/7998916142903730155_options.txt
View File

@ -1,5 +1,5 @@
/*
* Copyright (C) 2021 Intel Corporation
* Copyright (C) 2021-2022 Intel Corporation
*
* SPDX-License-Identifier: MIT
*

8
shared/test/common/test_files/builtin_copyfill.cl
View File

@ -1,5 +1,5 @@
/*
* Copyright (C) 2020-2021 Intel Corporation
* Copyright (C) 2020-2022 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
@ -263,9 +263,8 @@ __kernel void QueryKernelTimestamps(__global ulong* srcEvents, __global ulong* d
ulong globalEnd = GetTimestampValue(srcPtr, timestampSizeInDw, 3);
if(packetUsed > 1) {
uint timestampsOffsets = 4;
for(uint i = 1; i < packetUsed; i++) {
timestampsOffsets *= i;
uint timestampsOffsets = 4 * i;
if(contextStart > GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets)) {
contextStart = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets);
}
@ -304,9 +303,8 @@ __kernel void QueryKernelTimestampsWithOffsets(__global ulong* srcEvents, __glob
ulong globalEnd = GetTimestampValue(srcPtr, timestampSizeInDw, 3);
if(packetUsed > 1) {
uint timestampsOffsets = 4;
for(uint i = 1; i < packetUsed; i++) {
timestampsOffsets *= i;
uint timestampsOffsets = 4 * i;
if(contextStart > GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets)) {
contextStart = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets);
}

8
shared/test/common/test_files/builtin_copyfill_stateless.cl
View File

@ -1,5 +1,5 @@
/*
* Copyright (C) 2020-2021 Intel Corporation
* Copyright (C) 2020-2022 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
@ -238,9 +238,8 @@ __kernel void QueryKernelTimestamps(__global ulong* srcEvents, __global ulong* d
uint globalEnd = src[3];
if(packetUsed > 1) {
uint timestampsOffsets = 4;
for(uint i = 1; i < packetUsed; i++) {
timestampsOffsets *= i;
uint timestampsOffsets = 4 * i;
if(contextStart > src[timestampsOffsets]) {
contextStart = src[timestampsOffsets];
}
@ -287,9 +286,8 @@ __kernel void QueryKernelTimestampsWithOffsets(__global ulong* srcEvents, __glob
uint globalEnd = src[3];
if(packetUsed > 1) {
uint timestampsOffsets = 4;
for(uint i = 1; i < packetUsed; i++) {
timestampsOffsets *= i;
uint timestampsOffsets = 4 * i;
if(contextStart > src[timestampsOffsets]) {
contextStart = src[timestampsOffsets];
}

8
shared/test/common/test_files/builtin_images.cl
View File

@ -1,5 +1,5 @@
/*
* Copyright (C) 2020-2021 Intel Corporation
* Copyright (C) 2020-2022 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
@ -263,9 +263,8 @@ __kernel void QueryKernelTimestamps(__global ulong* srcEvents, __global ulong* d
ulong globalEnd = GetTimestampValue(srcPtr, timestampSizeInDw, 3);
if(packetUsed > 1) {
uint timestampsOffsets = 4;
for(uint i = 1; i < packetUsed; i++) {
timestampsOffsets *= i;
uint timestampsOffsets = 4 * i;
if(contextStart > GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets)) {
contextStart = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets);
}
@ -304,9 +303,8 @@ __kernel void QueryKernelTimestampsWithOffsets(__global ulong* srcEvents, __glob
ulong globalEnd = GetTimestampValue(srcPtr, timestampSizeInDw, 3);
if(packetUsed > 1) {
uint timestampsOffsets = 4;
for(uint i = 1; i < packetUsed; i++) {
timestampsOffsets *= i;
uint timestampsOffsets = 4 * i;
if(contextStart > GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets)) {
contextStart = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets);
}