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

Add appendQueryKernelTimestamps API

Browse Source 
Change-Id: I44b1d34a822b74a5c3404da62962ec324079744b
This commit is contained in:
Kamil Diedrich
2020-08-05 13:17:08 +02:00
committed by sys_ocldev
parent 8c64d65867
commit 0105960e13
22 changed files with 919 additions and 489 deletions
Download Patch File Download Diff File Expand all files Collapse all files

1
shared/source/built_ins/CMakeLists.txt
View File

@@ -27,6 +27,7 @@ set(NEO_CORE_SRCS_BUILT_IN_KERNELS
${CMAKE_CURRENT_SOURCE_DIR}/kernels/copy_buffer_to_buffer_stateless.builtin_kernel
${CMAKE_CURRENT_SOURCE_DIR}/kernels/copy_buffer_to_image3d.builtin_kernel
${CMAKE_CURRENT_SOURCE_DIR}/kernels/copy_buffer_to_image3d_stateless.builtin_kernel
${CMAKE_CURRENT_SOURCE_DIR}/kernels/copy_kernel_timestamps.builtin_kernel
${CMAKE_CURRENT_SOURCE_DIR}/kernels/copy_image3d_to_buffer.builtin_kernel
${CMAKE_CURRENT_SOURCE_DIR}/kernels/copy_image3d_to_buffer_stateless.builtin_kernel
${CMAKE_CURRENT_SOURCE_DIR}/kernels/copy_image_to_image1d.builtin_kernel

3
shared/source/built_ins/built_in_ops_base.h
View File

@@ -30,8 +30,9 @@ constexpr Type CopyImageToImage3d{13};
constexpr Type FillImage1d{14};
constexpr Type FillImage2d{15};
constexpr Type FillImage3d{16};
constexpr Type QueryKernelTimestamps{17};
constexpr Type MaxBaseValue{16};
constexpr Type MaxBaseValue{17};
constexpr Type COUNT{64};
} // namespace EBuiltInOps
} // namespace NEO

2
shared/source/built_ins/built_ins_storage.cpp
View File

@@ -59,6 +59,8 @@ const char *getBuiltinAsString(EBuiltInOps::Type builtin) {
return "fill_image2d.builtin_kernel";
case EBuiltInOps::FillImage3d:
return "fill_image3d.builtin_kernel";
case EBuiltInOps::QueryKernelTimestamps:
return "copy_kernel_timestamps.builtin_kernel";
};
}

1
shared/source/built_ins/builtins_binary.cmake
View File

@@ -14,6 +14,7 @@ set(GENERATED_BUILTINS
"copy_buffer_rect"
"copy_buffer_to_buffer"
"copy_buffer_to_image3d"
"copy_kernel_timestamps"
"fill_buffer"
"fill_image1d"
"fill_image2d"

42
shared/source/built_ins/kernels/copy_kernel_timestamps.builtin_kernel Normal file
View File

@@ -0,0 +1,42 @@
/*
* Copyright (C) 2020 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
*/
R"===(
__kernel void QueryKernelTimestamps(__global ulong* srcEvents, __global ulong* dst) {
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];
__global uint *src = (__global uint *) srcPtr;
dst[currentOffset] = src[1] & tsMask;
dst[currentOffset + 1] = src[3] & tsMask;
dst[currentOffset + 2] = src[0] & tsMask;
dst[currentOffset + 3] = src[2] & tsMask;
}
__kernel void QueryKernelTimestampsWithOffsets(__global ulong* srcEvents, __global ulong* dst, __global ulong *offsets) {
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];
__global uint *src = (__global uint *) srcPtr;
dst[currentOffset] = src[1] & tsMask;
dst[currentOffset + 1] = src[3] & tsMask;
dst[currentOffset + 2] = src[0] & tsMask;
dst[currentOffset + 3] = src[2] & tsMask;
}
)==="

9
shared/source/built_ins/registry/register_copy_kernels_source.cpp
View File

@@ -164,4 +164,13 @@ static RegisterEmbeddedResource registerAuxTranslationSrc(
#include "shared/source/built_ins/kernels/aux_translation.builtin_kernel"
));
static RegisterEmbeddedResource registerCopyKernelTimestampsSrc(
createBuiltinResourceName(
EBuiltInOps::FillImage3d,
BuiltinCode::getExtension(BuiltinCode::ECodeType::Source))
.c_str(),
std::string(
#include "shared/source/built_ins/kernels/copy_kernel_timestamps.builtin_kernel"
));
} // namespace NEO