1/*
2 * Copyright (C) 2020-2021 Intel Corporation
3 *
4 * SPDX-License-Identifier: MIT
5 *
6 */
7
8R"===(
9void SetDstData(__global ulong* dst, uint currentOffset, ulong contextStart, ulong globalStart, ulong contextEnd, ulong globalEnd, uint useOnlyGlobalTimestamps) {
10    dst[currentOffset] = globalStart;
11    dst[currentOffset + 1] = globalEnd;
12    if (useOnlyGlobalTimestamps != 0) {
13        dst[currentOffset + 2] = globalStart;
14        dst[currentOffset + 3] = globalEnd;
15    } else {
16        dst[currentOffset + 2] = contextStart;
17        dst[currentOffset + 3] = contextEnd;
18    }
19}
20
21ulong GetTimestampValue(ulong srcPtr, ulong timestampSizeInDw, uint index) {
22    if(timestampSizeInDw == 1) {
23        __global uint *src = (__global uint *) srcPtr;
24        return src[index];
25    } else if(timestampSizeInDw == 2) {
26        __global ulong *src = (__global ulong *) srcPtr;
27        return src[index];
28    }
29
30    return 0;
31}
32
33__kernel void QueryKernelTimestamps(__global ulong* srcEvents, __global ulong* dst, uint useOnlyGlobalTimestamps) {
34    uint gid = get_global_id(0);
35    uint currentOffset = gid * 4;
36    dst[currentOffset] = 0;
37    dst[currentOffset + 1] = 0;
38    dst[currentOffset + 2] = 0;
39    dst[currentOffset + 3] = 0;
40
41    uint eventOffsetData = 3 * gid;
42
43    ulong srcPtr = srcEvents[eventOffsetData];
44    ulong packetUsed = srcEvents[eventOffsetData + 1];
45    ulong timestampSizeInDw = srcEvents[eventOffsetData + 2];
46
47    ulong contextStart = GetTimestampValue(srcPtr, timestampSizeInDw, 0);
48    ulong globalStart = GetTimestampValue(srcPtr, timestampSizeInDw, 1);
49    ulong contextEnd = GetTimestampValue(srcPtr, timestampSizeInDw, 2);
50    ulong globalEnd = GetTimestampValue(srcPtr, timestampSizeInDw, 3);
51
52    if(packetUsed > 1) {
53        uint timestampsOffsets = 4;
54        for(uint i = 1; i < packetUsed; i++) {
55            timestampsOffsets += i;
56            if(contextStart > GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets)) {
57              contextStart = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets);
58            }
59            if(globalStart > GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets + 1)) {
60              globalStart = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets + 1);
61            }
62            if(contextEnd < GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets + 2)) {
63              contextEnd = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets + 2);
64            }
65            if(globalEnd < GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets + 3)) {
66              globalEnd = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets + 3);
67        }
68      }
69    }
70
71    SetDstData(dst, currentOffset, contextStart, globalStart, contextEnd, globalEnd, useOnlyGlobalTimestamps);
72}
73
74__kernel void QueryKernelTimestampsWithOffsets(__global ulong* srcEvents, __global ulong* dst, __global ulong *offsets, uint useOnlyGlobalTimestamps) {
75    uint gid = get_global_id(0);
76    uint currentOffset = offsets[gid] / 8;
77    dst[currentOffset] = 0;
78    dst[currentOffset + 1] = 0;
79    dst[currentOffset + 2] = 0;
80    dst[currentOffset + 3] = 0;
81
82    uint eventOffsetData = 3 * gid;
83
84    ulong srcPtr = srcEvents[eventOffsetData];
85    ulong packetUsed = srcEvents[eventOffsetData + 1];
86    ulong timestampSizeInDw = srcEvents[eventOffsetData + 2];
87
88    ulong contextStart = GetTimestampValue(srcPtr, timestampSizeInDw, 0);
89    ulong globalStart = GetTimestampValue(srcPtr, timestampSizeInDw, 1);
90    ulong contextEnd = GetTimestampValue(srcPtr, timestampSizeInDw, 2);
91    ulong globalEnd = GetTimestampValue(srcPtr, timestampSizeInDw, 3);
92
93    if(packetUsed > 1) {
94        uint timestampsOffsets = 4;
95        for(uint i = 1; i < packetUsed; i++) {
96            timestampsOffsets += i;
97            if(contextStart > GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets)) {
98              contextStart = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets);
99            }
100            if(globalStart > GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets + 1)) {
101              globalStart = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets + 1);
102            }
103            if(contextEnd < GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets + 2)) {
104              contextEnd = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets + 2);
105            }
106            if(globalEnd < GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets + 3)) {
107              globalEnd = GetTimestampValue(srcPtr, timestampSizeInDw, timestampsOffsets + 3);
108        }
109      }
110    }
111
112    SetDstData(dst, currentOffset, contextStart, globalStart, contextEnd, globalEnd, useOnlyGlobalTimestamps);
113}
114)==="
115