intel/compute-runtime
1
0
Fork 0
mirror of https://github.com/intel/compute-runtime.git synced 2026-01-11 08:07:19 +08:00
Code Issues Packages Projects Releases Wiki Activity

L0::Event to support dynamic size - part 3

Browse Source 
Signed-off-by: Bartosz Dunajski <bartosz.dunajski@intel.com>
This commit is contained in:
Bartosz Dunajski
2021-05-19 15:36:43 +00:00
committed by Compute-Runtime-Automation
parent 1c9bd5f114
commit 1bca3b2ab5
11 changed files with 339 additions and 195 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -1,48 +1,12 @@
/*
* Copyright (C) 2020 Intel Corporation
* Copyright (C) 2020-2021 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
*/
R"===(
__kernel void QueryKernelTimestamps(__global ulong* srcEvents, __global ulong* dst, uint useOnlyGlobalTimestamps) {
uint gid = get_global_id(0);
uint currentOffset = gid * 4;
dst[currentOffset] = 0;
dst[currentOffset + 1] = 0;
dst[currentOffset + 2] = 0;
dst[currentOffset + 3] = 0;
uint eventOffsetData = 2*gid;
ulong srcPtr = srcEvents[eventOffsetData];
__global uint *src = (__global uint *) srcPtr;
ulong packetUsed = srcEvents[eventOffsetData + 1];
uint contextStart = src[0];
uint globalStart = src[1];
uint contextEnd = src[2];
uint globalEnd = src[3];
if(packetUsed > 1) {
uint timestampsOffsets = 4;
for(uint i = 1; i < packetUsed; i++) {
timestampsOffsets *= i;
if(contextStart > src[timestampsOffsets]) {
contextStart = src[timestampsOffsets];
}
if(globalStart > src[timestampsOffsets + 1]) {
globalStart = src[timestampsOffsets + 1];
}
if(contextEnd < src[timestampsOffsets + 2]) {
contextEnd = src[timestampsOffsets + 2];
}
if(globalEnd < src[timestampsOffsets + 3]) {
globalEnd = src[timestampsOffsets + 3];
}
}
}
void SetDstData(__global ulong* dst, uint currentOffset, ulong contextStart, ulong globalStart, ulong contextEnd, ulong globalEnd, uint useOnlyGlobalTimestamps) {
dst[currentOffset] = globalStart;
dst[currentOffset + 1] = globalEnd;
if (useOnlyGlobalTimestamps != 0) {
@@ -52,7 +16,59 @@ __kernel void QueryKernelTimestamps(__global ulong* srcEvents, __global ulong* d
dst[currentOffset + 2] = contextStart;
dst[currentOffset + 3] = contextEnd;
}
}
ulong GetTimestampValue(ulong srcPtr, ulong timestampSizeInDw, uint index) {
if(timestampSizeInDw == 1) {
__global uint *src = (__global uint *) srcPtr;
return src[index];
} else if(timestampSizeInDw == 2) {
__global ulong *src = (__global ulong *) srcPtr;
return src[index];
}
return 0;
}
__kernel void QueryKernelTimestamps(__global ulong* srcEvents, __global ulong* dst, uint useOnlyGlobalTimestamps) {
uint gid = get_global_id(0);
uint currentOffset = gid * 4;
dst[currentOffset] = 0;
dst[currentOffset + 1] = 0;
dst[currentOffset + 2] = 0;
dst[currentOffset + 3] = 0;
uint eventOffsetData = 3 * gid;
ulong srcPtr = srcEvents[eventOffsetData];
ulong packetUsed = srcEvents[eventOffsetData + 1];
ulong timestampSizeInDw = srcEvents[eventOffsetData + 2];
ulong contextStart = GetTimestampValue(srcPtr, timestampSizeInDw, 0);
ulong globalStart = GetTimestampValue(srcPtr, timestampSizeInDw, 1);
ulong contextEnd = GetTimestampValue(srcPtr, timestampSizeInDw, 2);
ulong globalEnd = GetTimestampValue(srcPtr, timestampSizeInDw, 3);
if(packetUsed > 1) {
uint timestampsOffsets = 4;
for(uint i = 1; i < packetUsed; i++) {
timestampsOffsets *= i;
if(contextStart > GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets)) {
contextStart = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets);
}
if(globalStart > GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets + 1)) {
globalStart = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets + 1);
}
if(contextEnd < GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets + 2)) {
contextEnd = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets + 2);
}
if(globalEnd < GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets + 3)) {
globalEnd = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets + 3);
}
}
}
SetDstData(dst, currentOffset, contextStart, globalStart, contextEnd, globalEnd, useOnlyGlobalTimestamps);
}
__kernel void QueryKernelTimestampsWithOffsets(__global ulong* srcEvents, __global ulong* dst, __global ulong *offsets, uint useOnlyGlobalTimestamps) {
@@ -63,44 +79,36 @@ __kernel void QueryKernelTimestampsWithOffsets(__global ulong* srcEvents, __glob
dst[currentOffset + 2] = 0;
dst[currentOffset + 3] = 0;
uint eventOffsetData = 2*gid;
ulong srcPtr = srcEvents[eventOffsetData];
__global uint *src = (__global uint *) srcPtr;
ulong packetUsed = srcEvents[eventOffsetData + 1];
uint eventOffsetData = 3 * gid;
uint contextStart = src[0];
uint globalStart = src[1];
uint contextEnd = src[2];
uint globalEnd = src[3];
ulong srcPtr = srcEvents[eventOffsetData];
ulong packetUsed = srcEvents[eventOffsetData + 1];
ulong timestampSizeInDw = srcEvents[eventOffsetData + 2];
ulong contextStart = GetTimestampValue(srcPtr, timestampSizeInDw, 0);
ulong globalStart = GetTimestampValue(srcPtr, timestampSizeInDw, 1);
ulong contextEnd = GetTimestampValue(srcPtr, timestampSizeInDw, 2);
ulong globalEnd = GetTimestampValue(srcPtr, timestampSizeInDw, 3);
if(packetUsed > 1) {
uint timestampsOffsets = 4;
uint timestampsOffsets = 4;
for(uint i = 1; i < packetUsed; i++) {
timestampsOffsets *= i;
if(contextStart > src[timestampsOffsets]) {
contextStart = src[timestampsOffsets];
if(contextStart > GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets)) {
contextStart = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets);
}
if(globalStart > src[timestampsOffsets + 1]) {
globalStart = src[timestampsOffsets + 1];
if(globalStart > GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets + 1)) {
globalStart = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets + 1);
}
if(contextEnd < src[timestampsOffsets + 2]) {
contextEnd = src[timestampsOffsets + 2];
if(contextEnd < GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets + 2)) {
contextEnd = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets + 2);
}
if(globalEnd < src[timestampsOffsets + 3]) {
globalEnd = src[timestampsOffsets + 3];
if(globalEnd < GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets + 3)) {
globalEnd = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets + 3);
}
}
}
dst[currentOffset] = globalStart;
dst[currentOffset + 1] = globalEnd;
if (useOnlyGlobalTimestamps != 0) {
dst[currentOffset + 2] = globalStart;
dst[currentOffset + 3] = globalEnd;
} else {
dst[currentOffset + 2] = contextStart;
dst[currentOffset + 3] = contextEnd;
}
SetDstData(dst, currentOffset, contextStart, globalStart, contextEnd, globalEnd, useOnlyGlobalTimestamps);
}
)==="