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