/*
 * Copyright (C) 2020 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];
        }
      }
    }

    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;
    }

}

__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 = 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];
        }
      }
    }

    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;
    }

}
)==="