intel/compute-runtime
1
0
Fork 0
mirror of https://github.com/intel/compute-runtime.git synced 2026-01-03 23:03:02 +08:00
Code Issues Packages Projects Releases Wiki Activity

Add extended functionality for timestamps at Event

Browse Source 
Related-To: NEO-4584
Signed-off-by: Daria Hinz <daria.hinz@intel.com>
This commit is contained in:
Daria Hinz
2020-09-15 09:33:12 +02:00
committed by Compute-Runtime-Automation
parent 55fb319517
commit 1ef9a1c35f
21 changed files with 406 additions and 156 deletions
Download Patch File Download Diff File Expand all files Collapse all files

94
shared/source/built_ins/kernels/copy_kernel_timestamps.builtin_kernel
View File

@@ -8,45 +8,99 @@
R"===(
__kernel void QueryKernelTimestamps(__global ulong* srcEvents, __global ulong* dst, uint useOnlyGlobalTimestamps) {
uint gid = get_global_id(0);
const ulong tsMask = (1ull << 32) - 1;
uint currentOffset = gid * 4;
dst[currentOffset] = 0;
dst[currentOffset + 1] = 0;
dst[currentOffset + 2] = 0;
dst[currentOffset + 3] = 0;
ulong srcPtr = srcEvents[gid];
uint eventOffsetData = 2*gid;
ulong srcPtr = srcEvents[eventOffsetData];
__global uint *src = (__global uint *) srcPtr;
dst[currentOffset] = src[1] & tsMask;
dst[currentOffset + 1] = src[3] & tsMask;
if (useOnlyGlobalTimestamps != 0) {
dst[currentOffset + 2] = src[1] & tsMask;
dst[currentOffset + 3] = src[3] & tsMask;
} else {
dst[currentOffset + 2] = src[0] & tsMask;
dst[currentOffset + 3] = src[2] & tsMask;
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);
const ulong tsMask = (1ull << 32) - 1;
uint currentOffset = offsets[gid] / 8;
dst[currentOffset] = 0;
dst[currentOffset + 1] = 0;
dst[currentOffset + 2] = 0;
dst[currentOffset + 3] = 0;
ulong srcPtr = srcEvents[gid];
uint eventOffsetData = 2*gid;
ulong srcPtr = srcEvents[eventOffsetData];
__global uint *src = (__global uint *) srcPtr;
dst[currentOffset] = src[1] & tsMask;
dst[currentOffset + 1] = src[3] & tsMask;
if (useOnlyGlobalTimestamps != 0) {
dst[currentOffset + 2] = src[1] & tsMask;
dst[currentOffset + 3] = src[3] & tsMask;
} else {
dst[currentOffset + 2] = src[0] & tsMask;
dst[currentOffset + 3] = src[2] & tsMask;
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;
}
}
)==="

4
shared/source/command_stream/command_stream_receiver.h
View File

@@ -44,7 +44,9 @@ class OSInterface;
class ScratchSpaceController;
struct HwPerfCounter;
struct HwTimeStamps;
struct TimestampPacketStorage;
template <typename TSize>
struct TimestampPackets;
template <typename T1>
class TagAllocator;

6
shared/source/helpers/blit_commands_helper.h
View File

@@ -22,16 +22,20 @@ class CommandStreamReceiver;
class Device;
class GraphicsAllocation;
class LinearStream;
struct TimestampPacketStorage;
struct RootDeviceEnvironment;
template <typename TagType>
struct TagNode;
template <typename TSize>
struct TimestampPackets;
struct BlitProperties;
struct HardwareInfo;
struct TimestampPacketDependencies;
using BlitPropertiesContainer = StackVec<BlitProperties, 16>;
using TimestampPacketStorage = TimestampPackets<uint32_t>;
struct BlitProperties {
static BlitProperties constructPropertiesForReadWriteBuffer(BlitterConstants::BlitDirection blitDirection,

14
shared/source/helpers/timestamp_packet.h
View File

@@ -29,12 +29,13 @@ constexpr uint32_t preferredPacketCount = 16u;
}
#pragma pack(1)
struct TimestampPacketStorage {
template <typename TSize>
struct TimestampPackets {
struct Packet {
uint32_t contextStart = 1u;
uint32_t globalStart = 1u;
uint32_t contextEnd = 1u;
uint32_t globalEnd = 1u;
TSize contextStart = 1u;
TSize globalStart = 1u;
TSize contextEnd = 1u;
TSize globalEnd = 1u;
};
static GraphicsAllocation::AllocationType getAllocationType() {
@@ -74,12 +75,15 @@ struct TimestampPacketStorage {
};
#pragma pack()
using TimestampPacketStorage = TimestampPackets<uint32_t>;
static_assert(((4 * TimestampPacketSizeControl::preferredPacketCount + 2) * sizeof(uint32_t)) == sizeof(TimestampPacketStorage),
"This structure is consumed by GPU and has to follow specific restrictions for padding and size");
class TimestampPacketContainer : public NonCopyableClass {
public:
using Node = TagNode<TimestampPacketStorage>;
TimestampPacketContainer() = default;
TimestampPacketContainer(TimestampPacketContainer &&) = default;
TimestampPacketContainer &operator=(TimestampPacketContainer &&) = default;

94
shared/test/unit_test/test_files/17475704436818206941.cl → shared/test/unit_test/test_files/17034883254829439342.cl
View File

@@ -200,44 +200,98 @@ __kernel void CopyBufferRectBytes3d(
__kernel void QueryKernelTimestamps(__global ulong* srcEvents, __global ulong* dst, uint useOnlyGlobalTimestamps) {
uint gid = get_global_id(0);
const ulong tsMask = (1ull << 32) - 1;
uint currentOffset = gid * 4;
dst[currentOffset] = 0;
dst[currentOffset + 1] = 0;
dst[currentOffset + 2] = 0;
dst[currentOffset + 3] = 0;
ulong srcPtr = srcEvents[gid];
uint eventOffsetData = 2*gid;
ulong srcPtr = srcEvents[eventOffsetData];
__global uint *src = (__global uint *) srcPtr;
dst[currentOffset] = src[1] & tsMask;
dst[currentOffset + 1] = src[3] & tsMask;
if (useOnlyGlobalTimestamps != 0) {
dst[currentOffset + 2] = src[1] & tsMask;
dst[currentOffset + 3] = src[3] & tsMask;
} else {
dst[currentOffset + 2] = src[0] & tsMask;
dst[currentOffset + 3] = src[2] & tsMask;
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);
const ulong tsMask = (1ull << 32) - 1;
uint currentOffset = offsets[gid] / 8;
dst[currentOffset] = 0;
dst[currentOffset + 1] = 0;
dst[currentOffset + 2] = 0;
dst[currentOffset + 3] = 0;
ulong srcPtr = srcEvents[gid];
uint eventOffsetData = 2*gid;
ulong srcPtr = srcEvents[eventOffsetData];
__global uint *src = (__global uint *) srcPtr;
dst[currentOffset] = src[1] & tsMask;
dst[currentOffset + 1] = src[3] & tsMask;
if (useOnlyGlobalTimestamps != 0) {
dst[currentOffset + 2] = src[1] & tsMask;
dst[currentOffset + 3] = src[3] & tsMask;
} else {
dst[currentOffset + 2] = src[0] & tsMask;
dst[currentOffset + 3] = src[2] & tsMask;
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;
}
}

0
shared/test/unit_test/test_files/13936718464051189346_images_options.txt → shared/test/unit_test/test_files/17034883254829439342_options.txt
View File

94
shared/test/unit_test/test_files/13936718464051189346_images.cl → shared/test/unit_test/test_files/5842966610900731025_images.cl
View File

@@ -200,46 +200,100 @@ __kernel void CopyBufferRectBytes3d(
__kernel void QueryKernelTimestamps(__global ulong* srcEvents, __global ulong* dst, uint useOnlyGlobalTimestamps) {
uint gid = get_global_id(0);
const ulong tsMask = (1ull << 32) - 1;
uint currentOffset = gid * 4;
dst[currentOffset] = 0;
dst[currentOffset + 1] = 0;
dst[currentOffset + 2] = 0;
dst[currentOffset + 3] = 0;
ulong srcPtr = srcEvents[gid];
uint eventOffsetData = 2*gid;
ulong srcPtr = srcEvents[eventOffsetData];
__global uint *src = (__global uint *) srcPtr;
dst[currentOffset] = src[1] & tsMask;
dst[currentOffset + 1] = src[3] & tsMask;
if (useOnlyGlobalTimestamps != 0) {
dst[currentOffset + 2] = src[1] & tsMask;
dst[currentOffset + 3] = src[3] & tsMask;
} else {
dst[currentOffset + 2] = src[0] & tsMask;
dst[currentOffset + 3] = src[2] & tsMask;
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);
const ulong tsMask = (1ull << 32) - 1;
uint currentOffset = offsets[gid] / 8;
dst[currentOffset] = 0;
dst[currentOffset + 1] = 0;
dst[currentOffset + 2] = 0;
dst[currentOffset + 3] = 0;
ulong srcPtr = srcEvents[gid];
uint eventOffsetData = 2*gid;
ulong srcPtr = srcEvents[eventOffsetData];
__global uint *src = (__global uint *) srcPtr;
dst[currentOffset] = src[1] & tsMask;
dst[currentOffset + 1] = src[3] & tsMask;
if (useOnlyGlobalTimestamps != 0) {
dst[currentOffset + 2] = src[1] & tsMask;
dst[currentOffset + 3] = src[3] & tsMask;
} else {
dst[currentOffset + 2] = src[0] & tsMask;
dst[currentOffset + 3] = src[2] & tsMask;
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 FillImage1d(

0
shared/test/unit_test/test_files/17475704436818206941_options.txt → shared/test/unit_test/test_files/5842966610900731025_images_options.txt
View File