1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
|
/*
* 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);
}
)==="
|
