113 lines
4.6 KiB
Plaintext
113 lines
4.6 KiB
Plaintext
/*
|
|
* Copyright (C) 2020-2022 Intel Corporation
|
|
*
|
|
* SPDX-License-Identifier: MIT
|
|
*
|
|
*/
|
|
|
|
R"===(
|
|
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) {
|
|
dst[currentOffset + 2] = globalStart;
|
|
dst[currentOffset + 3] = globalEnd;
|
|
} else {
|
|
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) {
|
|
for(uint i = 1; i < packetUsed; i++) {
|
|
uint timestampsOffsets = 4 * 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) {
|
|
uint gid = get_global_id(0);
|
|
uint currentOffset = offsets[gid] / 8;
|
|
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) {
|
|
for(uint i = 1; i < packetUsed; i++) {
|
|
uint timestampsOffsets = 4 * 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);
|
|
}
|
|
)==="
|