Aux translation [2/n]: Add new builtins

Each Kernel arg for aux translation needs to have own builtin kernel.
This is required to build MultiDispatchInfo before copying into SSH

Change-Id: I4b2f42518cf06bb31c31fd5f83f7da927bde99c3

runtime/built_ins/CMakeLists.txt

@@ -20,10 +20,12 @@
 
 set(RUNTIME_SRCS_BUILT_INS
   ${CMAKE_CURRENT_SOURCE_DIR}/CMakeLists.txt
+  ${CMAKE_CURRENT_SOURCE_DIR}/aux_translation_builtin.h
   ${CMAKE_CURRENT_SOURCE_DIR}/builtins_dispatch_builder.h
   ${CMAKE_CURRENT_SOURCE_DIR}/built_ins_storage.cpp
   ${CMAKE_CURRENT_SOURCE_DIR}/built_ins.cpp
   ${CMAKE_CURRENT_SOURCE_DIR}/built_ins.h
+  ${CMAKE_CURRENT_SOURCE_DIR}/built_ins.inl
   ${CMAKE_CURRENT_SOURCE_DIR}/sip.cpp
   ${CMAKE_CURRENT_SOURCE_DIR}/sip.h
   ${CMAKE_CURRENT_SOURCE_DIR}/vme_dispatch_builder.h
@@ -32,6 +34,7 @@ target_sources(${NEO_STATIC_LIB_NAME} PRIVATE ${RUNTIME_SRCS_BUILT_INS})
 set_property(GLOBAL PROPERTY RUNTIME_SRCS_BUILT_INS ${RUNTIME_SRCS_BUILT_INS})
 
 set(RUNTIME_SRCS_BUILT_IN_KERNELS
+  ${CMAKE_CURRENT_SOURCE_DIR}/kernels/aux_translation.igdrcl_built_in
   ${CMAKE_CURRENT_SOURCE_DIR}/kernels/copy_buffer_rect.igdrcl_built_in
   ${CMAKE_CURRENT_SOURCE_DIR}/kernels/copy_buffer_to_buffer.igdrcl_built_in
   ${CMAKE_CURRENT_SOURCE_DIR}/kernels/copy_buffer_to_image3d.igdrcl_built_in

runtime/built_ins/aux_translation_builtin.h

@@ -0,0 +1,44 @@
+/*
+ * Copyright (c) 2018, Intel Corporation
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and associated documentation files (the "Software"),
+ * to deal in the Software without restriction, including without limitation
+ * the rights to use, copy, modify, merge, publish, distribute, sublicense,
+ * and/or sell copies of the Software, and to permit persons to whom the
+ * Software is furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS
+ * OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
+ * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR
+ * OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE,
+ * ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
+ * OTHER DEALINGS IN THE SOFTWARE.
+ */
+
+#pragma once
+#include "runtime/built_ins/built_ins.h"
+#include "runtime/built_ins/builtins_dispatch_builder.h"
+#include "runtime/helpers/dispatch_info_builder.h"
+#include <memory>
+
+namespace OCLRT {
+template <typename HWFamily>
+class BuiltInOp<HWFamily, EBuiltInOps::AuxTranslation> : public BuiltinDispatchInfoBuilder {
+  public:
+    BuiltInOp(BuiltIns &kernelsLib, Context &context, Device &device);
+    bool buildDispatchInfos(MultiDispatchInfo &multiDispatchInfo, const BuiltinOpParams &operationParams) const override;
+
+  protected:
+    void resizeKernelInstances(size_t size) const;
+
+    Kernel *baseKernel = nullptr;
+    mutable std::vector<std::unique_ptr<Kernel>> convertToNonAuxKernel;
+    mutable std::vector<std::unique_ptr<Kernel>> convertToAuxKernel;
+};
+
+} // namespace OCLRT

runtime/built_ins/built_ins.cpp

@@ -21,7 +21,9 @@
  */
 
 #include <cstdint>
+#include "runtime/built_ins/aux_translation_builtin.h"
 #include "runtime/built_ins/built_ins.h"
+#include "runtime/built_ins/built_ins.inl"
 #include "runtime/built_ins/vme_dispatch_builder.h"
 #include "runtime/built_ins/sip.h"
 #include "runtime/compiler_interface/compiler_interface.h"
@@ -236,14 +238,6 @@ void BuiltinDispatchInfoBuilder::releaseOwnership() {
     }
 }
 
-template <typename... KernelsDescArgsT>
-void BuiltinDispatchInfoBuilder::populate(Context &context, Device &device, EBuiltInOps op, const char *options, KernelsDescArgsT &&... desc) {
-    auto src = kernelsLib.getBuiltinsLib().getBuiltinCode(op, BuiltinCode::ECodeType::Any, device);
-    prog.reset(BuiltinsLib::createProgramFromCode(src, context, device).release());
-    prog->build(0, nullptr, options, nullptr, nullptr, kernelsLib.isCacheingEnabled());
-    grabKernels(std::forward<KernelsDescArgsT>(desc)...);
-}
-
 template <typename HWFamily>
 class BuiltInOp<HWFamily, EBuiltInOps::CopyBufferToBuffer> : public BuiltinDispatchInfoBuilder {
   public:
@@ -819,6 +813,9 @@ BuiltinDispatchInfoBuilder &BuiltIns::getBuiltinDispatchInfoBuilder(EBuiltInOps
     case EBuiltInOps::VmeBlockAdvancedMotionEstimateBidirectionalCheckIntel:
         std::call_once(operationBuilder.second, [&] { operationBuilder.first.reset(new BuiltInOp<HWFamily, EBuiltInOps::VmeBlockAdvancedMotionEstimateBidirectionalCheckIntel>(*this, context, device)); });
         break;
+    case EBuiltInOps::AuxTranslation:
+        std::call_once(operationBuilder.second, [&] { operationBuilder.first.reset(new BuiltInOp<HWFamily, EBuiltInOps::AuxTranslation>(*this, context, device)); });
+        break;
     }
     return *operationBuilder.first;
 }

runtime/built_ins/built_ins.h

@@ -49,7 +49,8 @@ class SchedulerKernel;
 extern const char *mediaKernelsBuildOptions;
 
 enum class EBuiltInOps : uint32_t {
-    CopyBufferToBuffer = 0,
+    AuxTranslation = 0,
+    CopyBufferToBuffer,
     CopyBufferRect,
     FillBuffer,
     CopyBufferToImage3d,

runtime/built_ins/built_ins.inl

@@ -0,0 +1,90 @@
+/*
+ * Copyright (c) 2018, Intel Corporation
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and associated documentation files (the "Software"),
+ * to deal in the Software without restriction, including without limitation
+ * the rights to use, copy, modify, merge, publish, distribute, sublicense,
+ * and/or sell copies of the Software, and to permit persons to whom the
+ * Software is furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS
+ * OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
+ * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR
+ * OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE,
+ * ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
+ * OTHER DEALINGS IN THE SOFTWARE.
+ */
+
+#include "runtime/built_ins/aux_translation_builtin.h"
+#include "runtime/mem_obj/buffer.h"
+#include "runtime/kernel/kernel.h"
+
+namespace OCLRT {
+template <typename... KernelsDescArgsT>
+void BuiltinDispatchInfoBuilder::populate(Context &context, Device &device, EBuiltInOps op, const char *options, KernelsDescArgsT &&... desc) {
+    auto src = kernelsLib.getBuiltinsLib().getBuiltinCode(op, BuiltinCode::ECodeType::Any, device);
+    prog.reset(BuiltinsLib::createProgramFromCode(src, context, device).release());
+    prog->build(0, nullptr, options, nullptr, nullptr, kernelsLib.isCacheingEnabled());
+    grabKernels(std::forward<KernelsDescArgsT>(desc)...);
+}
+
+template <typename HWFamily>
+BuiltInOp<HWFamily, EBuiltInOps::AuxTranslation>::BuiltInOp(BuiltIns &kernelsLib, Context &context, Device &device)
+    : BuiltinDispatchInfoBuilder(kernelsLib) {
+    BuiltinDispatchInfoBuilder::populate(context, device, EBuiltInOps::AuxTranslation, "", "fullCopy", baseKernel);
+    resizeKernelInstances(5);
+}
+
+template <typename HWFamily>
+bool BuiltInOp<HWFamily, EBuiltInOps::AuxTranslation>::buildDispatchInfos(MultiDispatchInfo &multiDispatchInfo, const BuiltinOpParams &operationParams) const {
+    size_t kernelInstanceNumber = 0;
+    resizeKernelInstances(operationParams.buffersForAuxTranslation->size());
+
+    for (auto &buffer : *operationParams.buffersForAuxTranslation) {
+        DispatchInfoBuilder<SplitDispatch::Dim::d1D, SplitDispatch::SplitMode::NoSplit> builder;
+        auto graphicsAllocation = buffer->getGraphicsAllocation();
+        size_t allocationSize = graphicsAllocation->getUnderlyingBufferSize();
+
+        if (operationParams.forceNonAuxMode) {
+            builder.setKernel(convertToNonAuxKernel.at(kernelInstanceNumber++).get());
+            builder.setArg(0, buffer);
+            builder.setArgSvm(1, allocationSize, reinterpret_cast<void *>(graphicsAllocation->getGpuAddress()));
+        } else {
+            builder.setKernel(convertToAuxKernel.at(kernelInstanceNumber++).get());
+            builder.setArgSvm(0, allocationSize, reinterpret_cast<void *>(graphicsAllocation->getGpuAddress()));
+            builder.setArg(1, buffer);
+        }
+
+        size_t elementSize = sizeof(uint32_t) * 4;
+        DEBUG_BREAK_IF(allocationSize < elementSize || !isAligned<4>(allocationSize));
+        size_t xGws = allocationSize / elementSize;
+
+        builder.setDispatchGeometry(Vec3<size_t>{xGws, 0, 0}, Vec3<size_t>{0, 0, 0}, Vec3<size_t>{0, 0, 0});
+        builder.bake(multiDispatchInfo);
+    }
+
+    return true;
+}
+
+template <typename HWFamily>
+void BuiltInOp<HWFamily, EBuiltInOps::AuxTranslation>::resizeKernelInstances(size_t size) const {
+    convertToNonAuxKernel.reserve(size);
+    convertToAuxKernel.reserve(size);
+
+    for (size_t i = convertToNonAuxKernel.size(); i < size; i++) {
+        auto clonedKernel1 = Kernel::create(baseKernel->getProgram(), baseKernel->getKernelInfo(), nullptr);
+        auto clonedKernel2 = Kernel::create(baseKernel->getProgram(), baseKernel->getKernelInfo(), nullptr);
+        clonedKernel1->cloneKernel(baseKernel);
+        clonedKernel2->cloneKernel(baseKernel);
+
+        convertToNonAuxKernel.emplace_back(clonedKernel1);
+        convertToAuxKernel.emplace_back(clonedKernel2);
+    }
+}
+
+} // namespace OCLRT

runtime/built_ins/built_ins_storage.cpp

@@ -32,6 +32,8 @@ const char *getBuiltinAsString(EBuiltInOps builtin) {
     switch (builtin) {
     default:
         return "unknown";
+    case EBuiltInOps::AuxTranslation:
+        return "aux_translation.igdrcl_built_in";
     case EBuiltInOps::CopyBufferToBuffer:
         return "copy_buffer_to_buffer.igdrcl_built_in";
     case EBuiltInOps::CopyBufferRect:

runtime/built_ins/builtins_binary.cmake

@@ -24,6 +24,7 @@ add_library(${BUILTINS_BINARIES_LIB_NAME} OBJECT builtins_binary.cmake)
 add_subdirectory(registry)
 
 set(GENERATED_BUILTINS
+  "aux_translation"
   "copy_buffer_rect"
   "copy_buffer_to_buffer"
   "copy_buffer_to_image3d"

runtime/built_ins/builtins_dispatch_builder.h

@@ -53,6 +53,8 @@ class BuiltinDispatchInfoBuilder {
         MemObj *dstMemObj = nullptr;
         GraphicsAllocation *srcSvmAlloc = nullptr;
         GraphicsAllocation *dstSvmAlloc = nullptr;
+        const BuffersForAuxTranslation *buffersForAuxTranslation = nullptr;
+        bool forceNonAuxMode = false;
         Vec3<size_t> srcOffset = {0, 0, 0};
         Vec3<size_t> dstOffset = {0, 0, 0};
         Vec3<size_t> size = {0, 0, 0};

runtime/built_ins/kernels/aux_translation.igdrcl_built_in

@@ -0,0 +1,29 @@
+/*
+ * Copyright (c) 2018, Intel Corporation
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and associated documentation files (the "Software"),
+ * to deal in the Software without restriction, including without limitation
+ * the rights to use, copy, modify, merge, publish, distribute, sublicense,
+ * and/or sell copies of the Software, and to permit persons to whom the
+ * Software is furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS
+ * OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
+ * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR
+ * OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE,
+ * ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
+ * OTHER DEALINGS IN THE SOFTWARE.
+ */
+
+R"===(
+__kernel void fullCopy(__global const uint* src, __global uint* dst) {
+    unsigned int gid = get_global_id(0);
+    uint4 loaded = vload4(gid, src);
+    vstore4(loaded, gid, dst);
+}
+)==="

runtime/built_ins/registry/register_copy_kernels_source.cpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2017, Intel Corporation
+ * Copyright (c) 2017 - 2018, Intel Corporation
  *
  * Permission is hereby granted, free of charge, to any person obtaining a
  * copy of this software and associated documentation files (the "Software"),
@@ -124,4 +124,13 @@ static RegisterEmbeddedResource registerFillImage3dSrc(
 #include "runtime/built_ins/kernels/fill_image3d.igdrcl_built_in"
         ));
 
+static RegisterEmbeddedResource registerAuxTranslationSrc(
+    createBuiltinResourceName(
+        EBuiltInOps::AuxTranslation,
+        BuiltinCode::getExtension(BuiltinCode::ECodeType::Source))
+        .c_str(),
+    std::string(
+#include "runtime/built_ins/kernels/aux_translation.igdrcl_built_in"
+        ));
+
 } // namespace OCLRT

runtime/kernel/kernel.h

@@ -129,7 +129,7 @@ class Kernel : public BaseObject<_cl_kernel> {
 
     cl_int initialize();
 
-    cl_int cloneKernel(Kernel *pSourceKernel);
+    MOCKABLE_VIRTUAL cl_int cloneKernel(Kernel *pSourceKernel);
 
     MOCKABLE_VIRTUAL bool canTransformImages() const;
     MOCKABLE_VIRTUAL bool isPatched() const;