diff --git a/CMakeLists.txt b/CMakeLists.txt index 1336da64bc..54587358eb 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -787,6 +787,7 @@ endif() # Please keep alphabetical order include_directories(${NEO_BUILD_DIR}) include_directories(${NEO_SOURCE_DIR}) +include_directories(${NEO_CORE_DIRECTORY}/built_ins/builtinops${BRANCH_DIR_SUFFIX}) include_directories(${NEO_CORE_DIRECTORY}/command_stream/definitions${BRANCH_DIR_SUFFIX}) include_directories(${NEO_CORE_DIRECTORY}/compiler_interface/compiler_options${BRANCH_DIR_SUFFIX}) include_directories(${NEO_CORE_DIRECTORY}/debug_settings/definitions${BRANCH_DIR_SUFFIX}) @@ -795,7 +796,6 @@ include_directories(${NEO_CORE_DIRECTORY}/gmm_helper/windows/gmm_memory${BRANCH_ include_directories(${NEO_CORE_DIRECTORY}/memory_manager/definitions${BRANCH_DIR_SUFFIX}) include_directories(${NEO_CORE_DIRECTORY}/memory_properties${BRANCH_DIR_SUFFIX}) include_directories(${NEO_CORE_DIRECTORY}/sku_info/definitions${BRANCH_DIR_SUFFIX}) -include_directories(${NEO_SOURCE_DIR}/opencl/source/built_ins/builtinops${BRANCH_DIR_SUFFIX}) include_directories(${NEO_SOURCE_DIR}/opencl/source/command_queue/definitions${BRANCH_DIR_SUFFIX}) include_directories(${NEO_SOURCE_DIR}/opencl/source/command_stream/definitions${BRANCH_DIR_SUFFIX}) include_directories(${NEO_SOURCE_DIR}/opencl/source/gen_common/reg_configs${BRANCH_DIR_SUFFIX}) @@ -832,6 +832,13 @@ else() set(NEO_CORE_MOCKABLE_LIB_NAME "neo_core_mockable") endif() endif() + +set(BIKSIM_LIB_NAME "biksim") +set(BUILTINS_SOURCES_LIB_NAME "builtins_sources") +set(BUILTINS_BINARIES_LIB_NAME "builtins_binaries") +set(BUILTINS_VME_LIB_NAME "builtins_vme") +set(SCHEDULER_BINARY_LIB_NAME "scheduler_binary") + add_subdirectory_unique(shared/source) add_subdirectory_unique(shared/test/unit_test) @@ -869,12 +876,6 @@ endmacro(generate_runtime_lib) set(NEO_DYNAMIC_LIB_NAME "igdrcl_dll") # single NEO dll set(NEO_DLL_NAME_BASE "igdrcl") -set(BIKSIM_LIB_NAME "biksim") -set(BUILTINS_SOURCES_LIB_NAME "builtins_sources") -set(BUILTINS_BINARIES_LIB_NAME "builtins_binaries") -set(BUILTINS_VME_LIB_NAME "builtins_vme") -set(SCHEDULER_BINARY_LIB_NAME "scheduler_binary") - if(DONT_CARE_OF_VIRTUALS) message(STATUS "All targets will use virtuals") diff --git a/offline_compiler/offline_compiler.cpp b/offline_compiler/offline_compiler.cpp index 14ce7166f9..1dbcdd068f 100644 --- a/offline_compiler/offline_compiler.cpp +++ b/offline_compiler/offline_compiler.cpp @@ -578,7 +578,7 @@ std::string OfflineCompiler::parseBinAsCharArray(uint8_t *binary, size_t size, s out << "};" << std::endl; out << std::endl - << "#include \"opencl/source/built_ins/registry/built_ins_registry.h\"\n" + << "#include \"shared/source/built_ins/registry/built_ins_registry.h\"\n" << std::endl; out << "namespace NEO {" << std::endl; out << "static RegisterEmbeddedResource register" << builtinName << "Bin(" << std::endl; diff --git a/opencl/source/api/api.cpp b/opencl/source/api/api.cpp index e9b9d802b8..a00c05069b 100644 --- a/opencl/source/api/api.cpp +++ b/opencl/source/api/api.cpp @@ -7,6 +7,7 @@ #include "api.h" +#include "shared/source/built_ins/built_ins.h" #include "shared/source/command_stream/command_stream_receiver.h" #include "shared/source/debug_settings/debug_settings_manager.h" #include "shared/source/execution_environment/root_device_environment.h" @@ -22,7 +23,6 @@ #include "opencl/source/accelerators/intel_motion_estimation.h" #include "opencl/source/api/additional_extensions.h" #include "opencl/source/aub/aub_center.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/built_ins/vme_builtin.h" #include "opencl/source/command_queue/command_queue.h" #include "opencl/source/context/context.h" diff --git a/opencl/source/built_ins/CMakeLists.txt b/opencl/source/built_ins/CMakeLists.txt index 000c533a23..d36fa56904 100644 --- a/opencl/source/built_ins/CMakeLists.txt +++ b/opencl/source/built_ins/CMakeLists.txt @@ -9,43 +9,18 @@ set(RUNTIME_SRCS_BUILT_INS ${CMAKE_CURRENT_SOURCE_DIR}/aux_translation_builtin.h ${CMAKE_CURRENT_SOURCE_DIR}/builtins_dispatch_builder.cpp ${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_in_ops_base.h ${CMAKE_CURRENT_SOURCE_DIR}/built_in_ops_vme.h ${CMAKE_CURRENT_SOURCE_DIR}/built_ins.inl - ${CMAKE_CURRENT_SOURCE_DIR}/sip.cpp - ${CMAKE_CURRENT_SOURCE_DIR}/sip.h ${CMAKE_CURRENT_SOURCE_DIR}${BRANCH_DIR_SUFFIX}/unknown_built_in_name.cpp ${CMAKE_CURRENT_SOURCE_DIR}/vme_builtin.cpp ${CMAKE_CURRENT_SOURCE_DIR}/vme_builtin.h ${CMAKE_CURRENT_SOURCE_DIR}/vme_dispatch_builder.h ) -add_subdirectory(builtinops) - 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.builtin_kernel - ${CMAKE_CURRENT_SOURCE_DIR}/kernels/copy_buffer_rect.builtin_kernel - ${CMAKE_CURRENT_SOURCE_DIR}/kernels/copy_buffer_rect_stateless.builtin_kernel - ${CMAKE_CURRENT_SOURCE_DIR}/kernels/copy_buffer_to_buffer.builtin_kernel - ${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_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 - ${CMAKE_CURRENT_SOURCE_DIR}/kernels/copy_image_to_image2d.builtin_kernel - ${CMAKE_CURRENT_SOURCE_DIR}/kernels/copy_image_to_image3d.builtin_kernel - ${CMAKE_CURRENT_SOURCE_DIR}/kernels/fill_buffer.builtin_kernel - ${CMAKE_CURRENT_SOURCE_DIR}/kernels/fill_buffer_stateless.builtin_kernel - ${CMAKE_CURRENT_SOURCE_DIR}/kernels/fill_image1d.builtin_kernel - ${CMAKE_CURRENT_SOURCE_DIR}/kernels/fill_image2d.builtin_kernel - ${CMAKE_CURRENT_SOURCE_DIR}/kernels/fill_image3d.builtin_kernel ${CMAKE_CURRENT_SOURCE_DIR}/kernels/vme_block_advanced_motion_estimate_bidirectional_check_intel.builtin_kernel ${CMAKE_CURRENT_SOURCE_DIR}/kernels/vme_block_advanced_motion_estimate_check_intel.builtin_kernel ${CMAKE_CURRENT_SOURCE_DIR}/kernels/vme_block_motion_estimate_intel.builtin_kernel @@ -53,6 +28,9 @@ set(RUNTIME_SRCS_BUILT_IN_KERNELS target_sources(${NEO_STATIC_LIB_NAME} PRIVATE ${RUNTIME_SRCS_BUILT_IN_KERNELS}) -if(NOT (TARGET ${BUILTINS_BINARIES_LIB_NAME})) - include(builtins_binary.cmake) +if(NOT (TARGET ${BUILTINS_VME_LIB_NAME})) + add_subdirectory(registry) + if(COMPILE_BUILT_INS) + add_subdirectory(kernels) + endif() endif() diff --git a/opencl/source/built_ins/aux_translation_builtin.h b/opencl/source/built_ins/aux_translation_builtin.h index 6562825dcc..114f98d7b7 100644 --- a/opencl/source/built_ins/aux_translation_builtin.h +++ b/opencl/source/built_ins/aux_translation_builtin.h @@ -6,9 +6,9 @@ */ #pragma once +#include "shared/source/built_ins/built_ins.h" #include "shared/source/helpers/hw_helper.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/built_ins/builtins_dispatch_builder.h" #include "opencl/source/helpers/dispatch_info_builder.h" diff --git a/opencl/source/built_ins/builtins_dispatch_builder.cpp b/opencl/source/built_ins/builtins_dispatch_builder.cpp index 4d594579a0..433e0078ac 100644 --- a/opencl/source/built_ins/builtins_dispatch_builder.cpp +++ b/opencl/source/built_ins/builtins_dispatch_builder.cpp @@ -7,14 +7,14 @@ #include "opencl/source/built_ins/builtins_dispatch_builder.h" +#include "shared/source/built_ins/built_ins.h" +#include "shared/source/built_ins/sip.h" #include "shared/source/compiler_interface/compiler_interface.h" #include "shared/source/helpers/basic_math.h" #include "shared/source/helpers/debug_helpers.h" #include "opencl/source/built_ins/aux_translation_builtin.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/built_ins/built_ins.inl" -#include "opencl/source/built_ins/sip.h" #include "opencl/source/built_ins/vme_dispatch_builder.h" #include "opencl/source/device/cl_device.h" #include "opencl/source/helpers/built_ins_helper.h" @@ -761,4 +761,25 @@ BuiltinDispatchInfoBuilder &BuiltInDispatchBuilderOp::getBuiltinDispatchInfoBuil } return *operationBuilder.first; } + +BuiltInOwnershipWrapper::BuiltInOwnershipWrapper(BuiltinDispatchInfoBuilder &inputBuilder, Context *context) { + takeOwnership(inputBuilder, context); +} +BuiltInOwnershipWrapper::~BuiltInOwnershipWrapper() { + if (builder) { + for (auto &kernel : builder->peekUsedKernels()) { + kernel->setContext(nullptr); + kernel->releaseOwnership(); + } + } +} +void BuiltInOwnershipWrapper::takeOwnership(BuiltinDispatchInfoBuilder &inputBuilder, Context *context) { + UNRECOVERABLE_IF(builder); + builder = &inputBuilder; + for (auto &kernel : builder->peekUsedKernels()) { + kernel->takeOwnership(); + kernel->setContext(context); + } +} + } // namespace NEO \ No newline at end of file diff --git a/opencl/source/built_ins/builtins_dispatch_builder.h b/opencl/source/built_ins/builtins_dispatch_builder.h index 7bb805a1fd..bb1bdc20b6 100644 --- a/opencl/source/built_ins/builtins_dispatch_builder.h +++ b/opencl/source/built_ins/builtins_dispatch_builder.h @@ -6,9 +6,9 @@ */ #pragma once +#include "shared/source/built_ins/built_ins.h" #include "shared/source/helpers/vec.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/kernel/kernel.h" #include "CL/cl.h" @@ -109,4 +109,16 @@ class BuiltInDispatchBuilderOp { std::unique_ptr newBuilder); }; +class BuiltInOwnershipWrapper : public NonCopyableOrMovableClass { + public: + BuiltInOwnershipWrapper() = default; + BuiltInOwnershipWrapper(BuiltinDispatchInfoBuilder &inputBuilder, Context *context); + ~BuiltInOwnershipWrapper(); + + void takeOwnership(BuiltinDispatchInfoBuilder &inputBuilder, Context *context); + + protected: + BuiltinDispatchInfoBuilder *builder = nullptr; +}; + } // namespace NEO diff --git a/opencl/source/built_ins/kernels/CMakeLists.txt b/opencl/source/built_ins/kernels/CMakeLists.txt index 41ac296e5f..bb9dcf5d34 100644 --- a/opencl/source/built_ins/kernels/CMakeLists.txt +++ b/opencl/source/built_ins/kernels/CMakeLists.txt @@ -4,10 +4,10 @@ # SPDX-License-Identifier: MIT # -add_custom_target(builtins) -set_target_properties(builtins PROPERTIES FOLDER "built_ins") +add_custom_target(builtins_vme_sources) +set_target_properties(builtins_vme_sources PROPERTIES FOLDER "built_ins") set(BUILTINS_OUTDIR_WITH_ARCH "${TargetDir}/built_ins/${NEO_ARCH}") -add_dependencies(${BUILTINS_BINARIES_LIB_NAME} builtins) +add_dependencies(${BUILTINS_BINARIES_LIB_NAME} builtins_vme_sources) add_subdirectories() set(GENERATED_BUILTINS ${GENERATED_BUILTINS} PARENT_SCOPE) set(GENERATED_BUILTINS_STATELESS ${GENERATED_BUILTINS_STATELESS} PARENT_SCOPE) @@ -95,7 +95,7 @@ macro(macro_for_each_gen) list(APPEND BUILTINS_COMMANDS ${TargetDir}/${BUILTIN_CPP}) set(RUNTIME_GENERATED_${GENERATED_BUILTIN_STATELESS}_${family_name_with_type} ${BUILTIN_CPP} PARENT_SCOPE) endforeach() - set(target_name builtins_${family_name_with_type}) + set(target_name builtins_${family_name_with_type}_vme) add_custom_target(${target_name} DEPENDS ${BUILTINS_COMMANDS}) add_dependencies(builtins ${target_name}) set_target_properties(${target_name} PROPERTIES FOLDER "opencl/source/built_ins/${family_name_with_type}") diff --git a/opencl/source/built_ins/kernels/copy_buffer_to_buffer.builtin_kernel b/opencl/source/built_ins/kernels/copy_buffer_to_buffer.builtin_kernel deleted file mode 100644 index d92ce2ded5..0000000000 --- a/opencl/source/built_ins/kernels/copy_buffer_to_buffer.builtin_kernel +++ /dev/null @@ -1,54 +0,0 @@ -/* - * Copyright (C) 2017-2020 Intel Corporation - * - * SPDX-License-Identifier: MIT - * - */ - -R"===( -__kernel void CopyBufferToBufferBytes( - const __global uchar* pSrc, - __global uchar* pDst, - uint srcOffsetInBytes, - uint dstOffsetInBytes, - uint bytesToRead ) -{ - pSrc += ( srcOffsetInBytes + get_global_id(0) ); - pDst += ( dstOffsetInBytes + get_global_id(0) ); - pDst[ 0 ] = pSrc[ 0 ]; -} - -__kernel void CopyBufferToBufferLeftLeftover( - const __global uchar* pSrc, - __global uchar* pDst, - uint srcOffsetInBytes, - uint dstOffsetInBytes) -{ - unsigned int gid = get_global_id(0); - pDst[ gid + dstOffsetInBytes ] = pSrc[ gid + srcOffsetInBytes ]; -} - -__kernel void CopyBufferToBufferMiddle( - const __global uint* pSrc, - __global uint* pDst, - uint srcOffsetInBytes, - uint dstOffsetInBytes) -{ - unsigned int gid = get_global_id(0); - pDst += dstOffsetInBytes >> 2; - pSrc += srcOffsetInBytes >> 2; - uint4 loaded = vload4(gid, pSrc); - vstore4(loaded, gid, pDst); -} - -__kernel void CopyBufferToBufferRightLeftover( - const __global uchar* pSrc, - __global uchar* pDst, - uint srcOffsetInBytes, - uint dstOffsetInBytes) -{ - unsigned int gid = get_global_id(0); - pDst[ gid + dstOffsetInBytes ] = pSrc[ gid + srcOffsetInBytes ]; -} - -)===" \ No newline at end of file diff --git a/opencl/source/built_ins/registry/CMakeLists.txt b/opencl/source/built_ins/registry/CMakeLists.txt index 35b195229d..672fd155c5 100644 --- a/opencl/source/built_ins/registry/CMakeLists.txt +++ b/opencl/source/built_ins/registry/CMakeLists.txt @@ -4,26 +4,6 @@ # SPDX-License-Identifier: MIT # -add_library(${BUILTINS_SOURCES_LIB_NAME} OBJECT EXCLUDE_FROM_ALL - CMakeLists.txt - built_ins_registry.h - register_copy_kernels_source.cpp -) - -set_target_properties(${BUILTINS_SOURCES_LIB_NAME} PROPERTIES POSITION_INDEPENDENT_CODE ON) -set_target_properties(${BUILTINS_SOURCES_LIB_NAME} PROPERTIES FOLDER "built_ins") - -target_include_directories(${BUILTINS_SOURCES_LIB_NAME} PRIVATE - ${ENGINE_NODE_DIR} - ${KHRONOS_HEADERS_DIR} - ${KHRONOS_GL_HEADERS_DIR} - ${NEO__GMM_INCLUDE_DIR} - ${NEO__IGC_INCLUDE_DIR} - ${THIRD_PARTY_DIR} -) - -add_subdirectories() - add_library(${BUILTINS_VME_LIB_NAME} OBJECT EXCLUDE_FROM_ALL CMakeLists.txt register_ext_vme_source.cpp diff --git a/opencl/source/built_ins/registry/register_ext_vme_source.cpp b/opencl/source/built_ins/registry/register_ext_vme_source.cpp index 19800397d2..e32062aa2c 100644 --- a/opencl/source/built_ins/registry/register_ext_vme_source.cpp +++ b/opencl/source/built_ins/registry/register_ext_vme_source.cpp @@ -5,8 +5,9 @@ * */ +#include "shared/source/built_ins/registry/built_ins_registry.h" + #include "opencl/source/built_ins/built_in_ops_vme.h" -#include "opencl/source/built_ins/registry/built_ins_registry.h" #include diff --git a/opencl/source/built_ins/unknown_built_in_name.cpp b/opencl/source/built_ins/unknown_built_in_name.cpp index 6d10fa2d55..3684f1abba 100644 --- a/opencl/source/built_ins/unknown_built_in_name.cpp +++ b/opencl/source/built_ins/unknown_built_in_name.cpp @@ -5,7 +5,8 @@ * */ -#include "opencl/source/built_ins/built_ins.h" +#include "shared/source/built_ins/built_ins.h" + #include "opencl/source/built_ins/builtins_dispatch_builder.h" namespace NEO { diff --git a/opencl/source/built_ins/vme_builtin.cpp b/opencl/source/built_ins/vme_builtin.cpp index 19ac1913ca..021f1b8b91 100644 --- a/opencl/source/built_ins/vme_builtin.cpp +++ b/opencl/source/built_ins/vme_builtin.cpp @@ -7,10 +7,10 @@ #include "opencl/source/built_ins/vme_builtin.h" +#include "shared/source/built_ins/built_ins.h" #include "shared/source/device/device.h" #include "opencl/source/built_ins/built_in_ops_vme.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/built_ins/builtins_dispatch_builder.h" #include "opencl/source/built_ins/populate_built_ins.inl" #include "opencl/source/built_ins/vme_dispatch_builder.h" diff --git a/opencl/source/built_ins/vme_dispatch_builder.h b/opencl/source/built_ins/vme_dispatch_builder.h index 86604f2b22..62b1742bfd 100644 --- a/opencl/source/built_ins/vme_dispatch_builder.h +++ b/opencl/source/built_ins/vme_dispatch_builder.h @@ -7,10 +7,11 @@ #pragma once +#include "shared/source/built_ins/built_ins.h" + #include "opencl/source/accelerators/intel_accelerator.h" #include "opencl/source/accelerators/intel_motion_estimation.h" #include "opencl/source/built_ins/built_in_ops_vme.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/built_ins/builtins_dispatch_builder.h" #include "opencl/source/helpers/dispatch_info_builder.h" #include "opencl/source/mem_obj/buffer.h" diff --git a/opencl/source/command_queue/enqueue_common.h b/opencl/source/command_queue/enqueue_common.h index 80fbb9c287..585d67cbea 100644 --- a/opencl/source/command_queue/enqueue_common.h +++ b/opencl/source/command_queue/enqueue_common.h @@ -6,6 +6,7 @@ */ #pragma once +#include "shared/source/built_ins/built_ins.h" #include "shared/source/command_stream/command_stream_receiver.h" #include "shared/source/helpers/array_count.h" #include "shared/source/helpers/engine_node_helper.h" @@ -17,7 +18,6 @@ #include "shared/source/utilities/range.h" #include "shared/source/utilities/tag_allocator.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/built_ins/builtins_dispatch_builder.h" #include "opencl/source/builtin_kernels_simulation/scheduler_simulation.h" #include "opencl/source/command_queue/command_queue_hw.h" diff --git a/opencl/source/command_queue/enqueue_copy_buffer.h b/opencl/source/command_queue/enqueue_copy_buffer.h index 67605591cc..e4863a434c 100644 --- a/opencl/source/command_queue/enqueue_copy_buffer.h +++ b/opencl/source/command_queue/enqueue_copy_buffer.h @@ -6,9 +6,9 @@ */ #pragma once +#include "shared/source/built_ins/built_ins.h" #include "shared/source/command_stream/command_stream_receiver.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/command_queue/command_queue_hw.h" #include "opencl/source/command_queue/enqueue_common.h" #include "opencl/source/helpers/hardware_commands_helper.h" diff --git a/opencl/source/command_queue/enqueue_copy_buffer_rect.h b/opencl/source/command_queue/enqueue_copy_buffer_rect.h index a16c5c7eb8..c7a89a2b44 100644 --- a/opencl/source/command_queue/enqueue_copy_buffer_rect.h +++ b/opencl/source/command_queue/enqueue_copy_buffer_rect.h @@ -6,9 +6,9 @@ */ #pragma once +#include "shared/source/built_ins/built_ins.h" #include "shared/source/command_stream/command_stream_receiver.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/command_queue/command_queue_hw.h" #include "opencl/source/helpers/hardware_commands_helper.h" #include "opencl/source/mem_obj/buffer.h" diff --git a/opencl/source/command_queue/enqueue_copy_buffer_to_image.h b/opencl/source/command_queue/enqueue_copy_buffer_to_image.h index ea83e3032b..ab93cee9fd 100644 --- a/opencl/source/command_queue/enqueue_copy_buffer_to_image.h +++ b/opencl/source/command_queue/enqueue_copy_buffer_to_image.h @@ -6,9 +6,9 @@ */ #pragma once +#include "shared/source/built_ins/built_ins.h" #include "shared/source/command_stream/command_stream_receiver.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/command_queue/command_queue_hw.h" #include "opencl/source/helpers/hardware_commands_helper.h" #include "opencl/source/helpers/mipmap.h" diff --git a/opencl/source/command_queue/enqueue_copy_image.h b/opencl/source/command_queue/enqueue_copy_image.h index 2ec7d21ee6..354dec9807 100644 --- a/opencl/source/command_queue/enqueue_copy_image.h +++ b/opencl/source/command_queue/enqueue_copy_image.h @@ -6,10 +6,10 @@ */ #pragma once +#include "shared/source/built_ins/built_ins.h" #include "shared/source/command_stream/command_stream_receiver.h" #include "shared/source/helpers/basic_math.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/command_queue/command_queue_hw.h" #include "opencl/source/helpers/hardware_commands_helper.h" #include "opencl/source/helpers/mipmap.h" diff --git a/opencl/source/command_queue/enqueue_copy_image_to_buffer.h b/opencl/source/command_queue/enqueue_copy_image_to_buffer.h index 8de4623e9c..0d13be76e3 100644 --- a/opencl/source/command_queue/enqueue_copy_image_to_buffer.h +++ b/opencl/source/command_queue/enqueue_copy_image_to_buffer.h @@ -6,9 +6,9 @@ */ #pragma once +#include "shared/source/built_ins/built_ins.h" #include "shared/source/command_stream/command_stream_receiver.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/command_queue/command_queue_hw.h" #include "opencl/source/helpers/hardware_commands_helper.h" #include "opencl/source/helpers/mipmap.h" diff --git a/opencl/source/command_queue/enqueue_fill_buffer.h b/opencl/source/command_queue/enqueue_fill_buffer.h index 8058d6e288..c3560be6f8 100644 --- a/opencl/source/command_queue/enqueue_fill_buffer.h +++ b/opencl/source/command_queue/enqueue_fill_buffer.h @@ -6,11 +6,11 @@ */ #pragma once +#include "shared/source/built_ins/built_ins.h" #include "shared/source/command_stream/command_stream_receiver.h" #include "shared/source/memory_manager/internal_allocation_storage.h" #include "shared/source/memory_manager/memory_manager.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/command_queue/command_queue_hw.h" #include "opencl/source/helpers/hardware_commands_helper.h" #include "opencl/source/mem_obj/buffer.h" diff --git a/opencl/source/command_queue/enqueue_fill_image.h b/opencl/source/command_queue/enqueue_fill_image.h index 8084f1ee12..5a810f149a 100644 --- a/opencl/source/command_queue/enqueue_fill_image.h +++ b/opencl/source/command_queue/enqueue_fill_image.h @@ -6,10 +6,10 @@ */ #pragma once +#include "shared/source/built_ins/built_ins.h" #include "shared/source/command_stream/command_stream_receiver.h" #include "shared/source/helpers/basic_math.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/command_queue/command_queue_hw.h" #include "opencl/source/helpers/hardware_commands_helper.h" #include "opencl/source/mem_obj/image.h" diff --git a/opencl/source/command_queue/enqueue_read_buffer.h b/opencl/source/command_queue/enqueue_read_buffer.h index ad26195c19..bc9ef2a4d4 100644 --- a/opencl/source/command_queue/enqueue_read_buffer.h +++ b/opencl/source/command_queue/enqueue_read_buffer.h @@ -6,11 +6,11 @@ */ #pragma once +#include "shared/source/built_ins/built_ins.h" #include "shared/source/command_stream/command_stream_receiver.h" #include "shared/source/helpers/cache_policy.h" #include "shared/source/memory_manager/unified_memory_manager.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/command_queue/command_queue_hw.h" #include "opencl/source/command_queue/enqueue_common.h" #include "opencl/source/helpers/hardware_commands_helper.h" diff --git a/opencl/source/command_queue/enqueue_read_buffer_rect.h b/opencl/source/command_queue/enqueue_read_buffer_rect.h index 08f3381ed5..54de8b8bfb 100644 --- a/opencl/source/command_queue/enqueue_read_buffer_rect.h +++ b/opencl/source/command_queue/enqueue_read_buffer_rect.h @@ -6,9 +6,9 @@ */ #pragma once +#include "shared/source/built_ins/built_ins.h" #include "shared/source/command_stream/command_stream_receiver.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/command_queue/command_queue_hw.h" #include "opencl/source/command_queue/enqueue_common.h" #include "opencl/source/helpers/hardware_commands_helper.h" diff --git a/opencl/source/command_queue/enqueue_read_image.h b/opencl/source/command_queue/enqueue_read_image.h index 2622c1274b..4d39c57d2c 100644 --- a/opencl/source/command_queue/enqueue_read_image.h +++ b/opencl/source/command_queue/enqueue_read_image.h @@ -6,12 +6,12 @@ */ #pragma once +#include "shared/source/built_ins/built_ins.h" #include "shared/source/command_stream/command_stream_receiver.h" #include "shared/source/helpers/basic_math.h" #include "shared/source/helpers/cache_policy.h" #include "shared/source/memory_manager/graphics_allocation.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/command_queue/command_queue_hw.h" #include "opencl/source/context/context.h" #include "opencl/source/event/event.h" diff --git a/opencl/source/command_queue/enqueue_svm.h b/opencl/source/command_queue/enqueue_svm.h index 766a7c42c1..f37db1e9c5 100644 --- a/opencl/source/command_queue/enqueue_svm.h +++ b/opencl/source/command_queue/enqueue_svm.h @@ -6,10 +6,10 @@ */ #pragma once +#include "shared/source/built_ins/built_ins.h" #include "shared/source/memory_manager/surface.h" #include "shared/source/memory_manager/unified_memory_manager.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/command_queue/command_queue_hw.h" #include "opencl/source/command_queue/enqueue_common.h" #include "opencl/source/event/event.h" diff --git a/opencl/source/command_queue/enqueue_write_buffer.h b/opencl/source/command_queue/enqueue_write_buffer.h index 99a3427323..5f0a03ec50 100644 --- a/opencl/source/command_queue/enqueue_write_buffer.h +++ b/opencl/source/command_queue/enqueue_write_buffer.h @@ -6,11 +6,11 @@ */ #pragma once +#include "shared/source/built_ins/built_ins.h" #include "shared/source/command_stream/command_stream_receiver.h" #include "shared/source/helpers/string.h" #include "shared/source/memory_manager/unified_memory_manager.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/command_queue/command_queue_hw.h" #include "opencl/source/helpers/hardware_commands_helper.h" #include "opencl/source/mem_obj/buffer.h" diff --git a/opencl/source/command_queue/enqueue_write_buffer_rect.h b/opencl/source/command_queue/enqueue_write_buffer_rect.h index 32c3c9f3fc..c0cab80d06 100644 --- a/opencl/source/command_queue/enqueue_write_buffer_rect.h +++ b/opencl/source/command_queue/enqueue_write_buffer_rect.h @@ -6,9 +6,9 @@ */ #pragma once +#include "shared/source/built_ins/built_ins.h" #include "shared/source/command_stream/command_stream_receiver.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/command_queue/command_queue_hw.h" #include "opencl/source/helpers/hardware_commands_helper.h" #include "opencl/source/mem_obj/buffer.h" diff --git a/opencl/source/command_queue/enqueue_write_image.h b/opencl/source/command_queue/enqueue_write_image.h index f4381b9172..469bd2aa69 100644 --- a/opencl/source/command_queue/enqueue_write_image.h +++ b/opencl/source/command_queue/enqueue_write_image.h @@ -6,11 +6,11 @@ */ #pragma once +#include "shared/source/built_ins/built_ins.h" #include "shared/source/command_stream/command_stream_receiver.h" #include "shared/source/helpers/basic_math.h" #include "shared/source/memory_manager/graphics_allocation.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/command_queue/command_queue_hw.h" #include "opencl/source/helpers/hardware_commands_helper.h" #include "opencl/source/helpers/mipmap.h" diff --git a/opencl/source/command_queue/gpgpu_walker.h b/opencl/source/command_queue/gpgpu_walker.h index 53516ebc79..4035dec9aa 100644 --- a/opencl/source/command_queue/gpgpu_walker.h +++ b/opencl/source/command_queue/gpgpu_walker.h @@ -7,6 +7,7 @@ #pragma once +#include "shared/source/built_ins/built_ins.h" #include "shared/source/command_stream/linear_stream.h" #include "shared/source/command_stream/preemption.h" #include "shared/source/helpers/register_offsets.h" @@ -15,7 +16,6 @@ #include "shared/source/indirect_heap/indirect_heap.h" #include "shared/source/utilities/tag_allocator.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/command_queue/command_queue.h" #include "opencl/source/context/context.h" #include "opencl/source/device_queue/device_queue_hw.h" diff --git a/opencl/source/context/context.cpp b/opencl/source/context/context.cpp index 2fb1ee2f84..b383788297 100644 --- a/opencl/source/context/context.cpp +++ b/opencl/source/context/context.cpp @@ -7,6 +7,7 @@ #include "opencl/source/context/context.h" +#include "shared/source/built_ins/built_ins.h" #include "shared/source/command_stream/command_stream_receiver.h" #include "shared/source/compiler_interface/compiler_interface.h" #include "shared/source/debug_settings/debug_settings_manager.h" @@ -17,7 +18,6 @@ #include "shared/source/memory_manager/memory_manager.h" #include "shared/source/memory_manager/unified_memory_manager.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/command_queue/command_queue.h" #include "opencl/source/device/cl_device.h" #include "opencl/source/device_queue/device_queue.h" diff --git a/opencl/source/dll/debugger.cpp b/opencl/source/dll/debugger.cpp index 62637181d6..51129a9483 100644 --- a/opencl/source/dll/debugger.cpp +++ b/opencl/source/dll/debugger.cpp @@ -7,10 +7,10 @@ #include "shared/source/debugger/debugger.h" +#include "shared/source/built_ins/sip.h" #include "shared/source/built_ins/sip_kernel_type.h" #include "shared/source/helpers/hw_info.h" -#include "opencl/source/built_ins/sip.h" #include "opencl/source/source_level_debugger/source_level_debugger.h" namespace NEO { diff --git a/opencl/source/helpers/built_ins_helper.h b/opencl/source/helpers/built_ins_helper.h index df4864c8bf..b94394b9d5 100644 --- a/opencl/source/helpers/built_ins_helper.h +++ b/opencl/source/helpers/built_ins_helper.h @@ -6,10 +6,9 @@ */ #pragma once +#include "shared/source/built_ins/built_ins.h" #include "shared/source/execution_environment/execution_environment.h" -#include "opencl/source/built_ins/built_ins.h" - namespace NEO { class Device; @@ -18,6 +17,6 @@ Program *createProgramForSip(ExecutionEnvironment &executionEnvironment, Context *context, std::vector &binary, size_t size, - cl_int *errcodeRet, + int *errcodeRet, Device *device); } // namespace NEO diff --git a/opencl/source/helpers/hardware_commands_helper.h b/opencl/source/helpers/hardware_commands_helper.h index 4dc2847cc5..365090c1dc 100644 --- a/opencl/source/helpers/hardware_commands_helper.h +++ b/opencl/source/helpers/hardware_commands_helper.h @@ -6,9 +6,9 @@ */ #pragma once +#include "shared/source/built_ins/built_ins.h" #include "shared/source/indirect_heap/indirect_heap.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/helpers/per_thread_data.h" #include "opencl/source/kernel/kernel.h" diff --git a/opencl/source/kernel/kernel.cpp b/opencl/source/kernel/kernel.cpp index ff7e3ba5cb..1f37bbe608 100644 --- a/opencl/source/kernel/kernel.cpp +++ b/opencl/source/kernel/kernel.cpp @@ -7,6 +7,7 @@ #include "opencl/source/kernel/kernel.h" +#include "shared/source/built_ins/built_ins.h" #include "shared/source/command_stream/command_stream_receiver.h" #include "shared/source/debug_settings/debug_settings_manager.h" #include "shared/source/gmm_helper/gmm_helper.h" @@ -22,7 +23,6 @@ #include "opencl/source/accelerators/intel_accelerator.h" #include "opencl/source/accelerators/intel_motion_estimation.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/built_ins/builtins_dispatch_builder.h" #include "opencl/source/command_queue/command_queue.h" #include "opencl/source/command_queue/gpgpu_walker.h" diff --git a/opencl/source/scheduler/scheduler_kernel.h b/opencl/source/scheduler/scheduler_kernel.h index deeb8d6871..7975646817 100644 --- a/opencl/source/scheduler/scheduler_kernel.h +++ b/opencl/source/scheduler/scheduler_kernel.h @@ -6,7 +6,8 @@ */ #pragma once -#include "opencl/source/built_ins/built_ins.h" +#include "shared/source/built_ins/built_ins.h" + #include "opencl/source/kernel/kernel.h" #include diff --git a/opencl/test/unit_test/CMakeLists.txt b/opencl/test/unit_test/CMakeLists.txt index 504f3e9173..8e454ad86b 100644 --- a/opencl/test/unit_test/CMakeLists.txt +++ b/opencl/test/unit_test/CMakeLists.txt @@ -161,7 +161,7 @@ target_include_directories(igdrcl_tests PRIVATE target_link_libraries(igdrcl_tests gmock-gtest ${IGDRCL_EXTRA_LIBS}) -set(BUILT_IN_KERNEL_DIR "${NEO_SOURCE_DIR}/opencl/source/built_ins") +set(BUILT_IN_KERNEL_DIR "${NEO_SOURCE_DIR}/shared/source/built_ins") function(neo_copy_test_files target product) set(outputdir "${TargetDir}/${product}") diff --git a/opencl/test/unit_test/api/cl_create_program_with_built_in_kernels_tests.cpp b/opencl/test/unit_test/api/cl_create_program_with_built_in_kernels_tests.cpp index 4363157f68..dec2877ce7 100644 --- a/opencl/test/unit_test/api/cl_create_program_with_built_in_kernels_tests.cpp +++ b/opencl/test/unit_test/api/cl_create_program_with_built_in_kernels_tests.cpp @@ -5,11 +5,11 @@ * */ +#include "shared/source/built_ins/built_ins.h" #include "shared/source/compiler_interface/compiler_interface.h" #include "shared/source/device/device.h" #include "opencl/source/built_ins/built_in_ops_vme.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/built_ins/vme_builtin.h" #include "opencl/source/context/context.h" #include "opencl/source/device/cl_device.h" diff --git a/opencl/test/unit_test/aub_tests/gen11/execution_model/enqueue_parent_kernel_tests_gen11.cpp b/opencl/test/unit_test/aub_tests/gen11/execution_model/enqueue_parent_kernel_tests_gen11.cpp index bd289fd4b8..44c60d71dc 100644 --- a/opencl/test/unit_test/aub_tests/gen11/execution_model/enqueue_parent_kernel_tests_gen11.cpp +++ b/opencl/test/unit_test/aub_tests/gen11/execution_model/enqueue_parent_kernel_tests_gen11.cpp @@ -5,7 +5,8 @@ * */ -#include "opencl/source/built_ins/built_ins.h" +#include "shared/source/built_ins/built_ins.h" + #include "opencl/source/mem_obj/image.h" #include "opencl/source/sampler/sampler.h" #include "opencl/test/unit_test/aub_tests/fixtures/aub_parent_kernel_fixture.h" diff --git a/opencl/test/unit_test/aub_tests/gen12lp/execution_model/enqueue_parent_kernel_tests_gen12lp.cpp b/opencl/test/unit_test/aub_tests/gen12lp/execution_model/enqueue_parent_kernel_tests_gen12lp.cpp index 7b3c1ed356..d67b882b4f 100644 --- a/opencl/test/unit_test/aub_tests/gen12lp/execution_model/enqueue_parent_kernel_tests_gen12lp.cpp +++ b/opencl/test/unit_test/aub_tests/gen12lp/execution_model/enqueue_parent_kernel_tests_gen12lp.cpp @@ -5,7 +5,8 @@ * */ -#include "opencl/source/built_ins/built_ins.h" +#include "shared/source/built_ins/built_ins.h" + #include "opencl/source/mem_obj/image.h" #include "opencl/source/sampler/sampler.h" #include "opencl/test/unit_test/aub_tests/fixtures/aub_parent_kernel_fixture.h" diff --git a/opencl/test/unit_test/aub_tests/gen8/execution_model/enqueue_parent_kernel_tests_gen8.cpp b/opencl/test/unit_test/aub_tests/gen8/execution_model/enqueue_parent_kernel_tests_gen8.cpp index 82502c07ab..0487a6aac3 100644 --- a/opencl/test/unit_test/aub_tests/gen8/execution_model/enqueue_parent_kernel_tests_gen8.cpp +++ b/opencl/test/unit_test/aub_tests/gen8/execution_model/enqueue_parent_kernel_tests_gen8.cpp @@ -5,7 +5,8 @@ * */ -#include "opencl/source/built_ins/built_ins.h" +#include "shared/source/built_ins/built_ins.h" + #include "opencl/source/mem_obj/image.h" #include "opencl/source/sampler/sampler.h" #include "opencl/test/unit_test/aub_tests/fixtures/aub_parent_kernel_fixture.h" diff --git a/opencl/test/unit_test/aub_tests/gen9/execution_model/enqueue_parent_kernel_tests_gen9.cpp b/opencl/test/unit_test/aub_tests/gen9/execution_model/enqueue_parent_kernel_tests_gen9.cpp index 28534fdc78..cb56c4c1f5 100644 --- a/opencl/test/unit_test/aub_tests/gen9/execution_model/enqueue_parent_kernel_tests_gen9.cpp +++ b/opencl/test/unit_test/aub_tests/gen9/execution_model/enqueue_parent_kernel_tests_gen9.cpp @@ -5,7 +5,8 @@ * */ -#include "opencl/source/built_ins/built_ins.h" +#include "shared/source/built_ins/built_ins.h" + #include "opencl/source/mem_obj/image.h" #include "opencl/source/sampler/sampler.h" #include "opencl/test/unit_test/aub_tests/fixtures/aub_parent_kernel_fixture.h" diff --git a/opencl/test/unit_test/built_ins/built_in_tests.cpp b/opencl/test/unit_test/built_ins/built_in_tests.cpp index 070fd3c1b9..e439bc0b7b 100644 --- a/opencl/test/unit_test/built_ins/built_in_tests.cpp +++ b/opencl/test/unit_test/built_ins/built_in_tests.cpp @@ -5,6 +5,7 @@ * */ +#include "shared/source/built_ins/built_ins.h" #include "shared/source/debug_settings/debug_settings_manager.h" #include "shared/source/gmm_helper/gmm.h" #include "shared/source/gmm_helper/gmm_helper.h" @@ -15,7 +16,6 @@ #include "shared/test/unit_test/utilities/base_object_utils.h" #include "opencl/source/built_ins/aux_translation_builtin.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/built_ins/builtins_dispatch_builder.h" #include "opencl/source/built_ins/vme_builtin.h" #include "opencl/source/built_ins/vme_dispatch_builder.h" diff --git a/opencl/test/unit_test/built_ins/sip_tests.cpp b/opencl/test/unit_test/built_ins/sip_tests.cpp index 5cf1f93af4..0ce2179273 100644 --- a/opencl/test/unit_test/built_ins/sip_tests.cpp +++ b/opencl/test/unit_test/built_ins/sip_tests.cpp @@ -5,8 +5,9 @@ * */ -#include "opencl/source/built_ins/built_ins.h" -#include "opencl/source/built_ins/sip.h" +#include "shared/source/built_ins/built_ins.h" +#include "shared/source/built_ins/sip.h" + #include "opencl/test/unit_test/global_environment.h" #include "opencl/test/unit_test/helpers/test_files.h" #include "opencl/test/unit_test/mocks/mock_device.h" diff --git a/opencl/test/unit_test/command_queue/enqueue_copy_buffer_rect_tests.cpp b/opencl/test/unit_test/command_queue/enqueue_copy_buffer_rect_tests.cpp index 5060e804d8..6a5b6dd583 100644 --- a/opencl/test/unit_test/command_queue/enqueue_copy_buffer_rect_tests.cpp +++ b/opencl/test/unit_test/command_queue/enqueue_copy_buffer_rect_tests.cpp @@ -5,9 +5,9 @@ * */ +#include "shared/source/built_ins/built_ins.h" #include "shared/source/memory_manager/memory_constants.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/built_ins/builtins_dispatch_builder.h" #include "opencl/source/helpers/dispatch_info.h" #include "opencl/test/unit_test/command_queue/enqueue_copy_buffer_rect_fixture.h" diff --git a/opencl/test/unit_test/command_queue/enqueue_copy_buffer_tests.cpp b/opencl/test/unit_test/command_queue/enqueue_copy_buffer_tests.cpp index aef4c253a9..4746ae8e5f 100644 --- a/opencl/test/unit_test/command_queue/enqueue_copy_buffer_tests.cpp +++ b/opencl/test/unit_test/command_queue/enqueue_copy_buffer_tests.cpp @@ -5,9 +5,9 @@ * */ +#include "shared/source/built_ins/built_ins.h" #include "shared/source/helpers/ptr_math.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/built_ins/builtins_dispatch_builder.h" #include "opencl/source/command_queue/command_queue_hw.h" #include "opencl/source/helpers/dispatch_info.h" diff --git a/opencl/test/unit_test/command_queue/enqueue_fill_buffer_tests.cpp b/opencl/test/unit_test/command_queue/enqueue_fill_buffer_tests.cpp index f7b2574fe5..859bfbe597 100644 --- a/opencl/test/unit_test/command_queue/enqueue_fill_buffer_tests.cpp +++ b/opencl/test/unit_test/command_queue/enqueue_fill_buffer_tests.cpp @@ -5,6 +5,7 @@ * */ +#include "shared/source/built_ins/built_ins.h" #include "shared/source/command_stream/command_stream_receiver.h" #include "shared/source/helpers/aligned_memory.h" #include "shared/source/helpers/ptr_math.h" @@ -12,7 +13,6 @@ #include "shared/source/memory_manager/memory_manager.h" #include "shared/source/os_interface/os_context.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/built_ins/builtins_dispatch_builder.h" #include "opencl/source/command_queue/command_queue.h" #include "opencl/source/helpers/dispatch_info.h" diff --git a/opencl/test/unit_test/command_queue/enqueue_read_buffer_rect_tests.cpp b/opencl/test/unit_test/command_queue/enqueue_read_buffer_rect_tests.cpp index 1c734cc880..6704fee821 100644 --- a/opencl/test/unit_test/command_queue/enqueue_read_buffer_rect_tests.cpp +++ b/opencl/test/unit_test/command_queue/enqueue_read_buffer_rect_tests.cpp @@ -5,9 +5,9 @@ * */ +#include "shared/source/built_ins/built_ins.h" #include "shared/source/memory_manager/memory_constants.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/built_ins/builtins_dispatch_builder.h" #include "opencl/source/event/event.h" #include "opencl/source/helpers/dispatch_info.h" diff --git a/opencl/test/unit_test/command_queue/enqueue_read_buffer_tests.cpp b/opencl/test/unit_test/command_queue/enqueue_read_buffer_tests.cpp index 6d25d78461..7f46491bab 100644 --- a/opencl/test/unit_test/command_queue/enqueue_read_buffer_tests.cpp +++ b/opencl/test/unit_test/command_queue/enqueue_read_buffer_tests.cpp @@ -5,12 +5,12 @@ * */ +#include "shared/source/built_ins/built_ins.h" #include "shared/source/gmm_helper/gmm_helper.h" #include "shared/source/helpers/cache_policy.h" #include "shared/source/memory_manager/allocations_list.h" #include "shared/test/unit_test/helpers/debug_manager_state_restore.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/built_ins/builtins_dispatch_builder.h" #include "opencl/source/helpers/dispatch_info.h" #include "opencl/test/unit_test/command_queue/enqueue_fixture.h" diff --git a/opencl/test/unit_test/command_queue/enqueue_write_buffer_rect_tests.cpp b/opencl/test/unit_test/command_queue/enqueue_write_buffer_rect_tests.cpp index fa47aaf3be..dbccc71a62 100644 --- a/opencl/test/unit_test/command_queue/enqueue_write_buffer_rect_tests.cpp +++ b/opencl/test/unit_test/command_queue/enqueue_write_buffer_rect_tests.cpp @@ -5,7 +5,8 @@ * */ -#include "opencl/source/built_ins/built_ins.h" +#include "shared/source/built_ins/built_ins.h" + #include "opencl/source/built_ins/builtins_dispatch_builder.h" #include "opencl/source/event/event.h" #include "opencl/source/helpers/dispatch_info.h" diff --git a/opencl/test/unit_test/command_queue/enqueue_write_buffer_tests.cpp b/opencl/test/unit_test/command_queue/enqueue_write_buffer_tests.cpp index 94dbcc70bd..89ea945ba3 100644 --- a/opencl/test/unit_test/command_queue/enqueue_write_buffer_tests.cpp +++ b/opencl/test/unit_test/command_queue/enqueue_write_buffer_tests.cpp @@ -5,10 +5,10 @@ * */ +#include "shared/source/built_ins/built_ins.h" #include "shared/source/memory_manager/allocations_list.h" #include "shared/test/unit_test/helpers/debug_manager_state_restore.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/built_ins/builtins_dispatch_builder.h" #include "opencl/source/helpers/dispatch_info.h" #include "opencl/test/unit_test/command_queue/buffer_operations_fixture.h" diff --git a/opencl/test/unit_test/command_queue/get_size_required_image_tests.cpp b/opencl/test/unit_test/command_queue/get_size_required_image_tests.cpp index 01d15f2992..9fac6b0762 100644 --- a/opencl/test/unit_test/command_queue/get_size_required_image_tests.cpp +++ b/opencl/test/unit_test/command_queue/get_size_required_image_tests.cpp @@ -5,7 +5,8 @@ * */ -#include "opencl/source/built_ins/built_ins.h" +#include "shared/source/built_ins/built_ins.h" + #include "opencl/source/built_ins/builtins_dispatch_builder.h" #include "opencl/source/command_queue/command_queue_hw.h" #include "opencl/source/command_queue/enqueue_copy_image.h" diff --git a/opencl/test/unit_test/command_queue/get_size_required_tests.cpp b/opencl/test/unit_test/command_queue/get_size_required_tests.cpp index 3dd25e4c95..8a66c78e01 100644 --- a/opencl/test/unit_test/command_queue/get_size_required_tests.cpp +++ b/opencl/test/unit_test/command_queue/get_size_required_tests.cpp @@ -5,7 +5,8 @@ * */ -#include "opencl/source/built_ins/built_ins.h" +#include "shared/source/built_ins/built_ins.h" + #include "opencl/source/command_queue/command_queue_hw.h" #include "opencl/source/command_queue/enqueue_barrier.h" #include "opencl/source/command_queue/enqueue_marker.h" diff --git a/opencl/test/unit_test/command_queue/gl/windows/enqueue_kernel_gl_tests_windows.cpp b/opencl/test/unit_test/command_queue/gl/windows/enqueue_kernel_gl_tests_windows.cpp index fce56806d3..61441288f8 100644 --- a/opencl/test/unit_test/command_queue/gl/windows/enqueue_kernel_gl_tests_windows.cpp +++ b/opencl/test/unit_test/command_queue/gl/windows/enqueue_kernel_gl_tests_windows.cpp @@ -5,11 +5,11 @@ * */ +#include "shared/source/built_ins/built_ins.h" #include "shared/source/helpers/preamble.h" #include "shared/source/memory_manager/graphics_allocation.h" #include "shared/source/memory_manager/memory_constants.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/sharings/gl/gl_buffer.h" #include "opencl/test/unit_test/command_queue/enqueue_fixture.h" #include "opencl/test/unit_test/fixtures/hello_world_fixture.h" diff --git a/opencl/test/unit_test/command_stream/command_stream_receiver_flush_task_gmock_tests.cpp b/opencl/test/unit_test/command_stream/command_stream_receiver_flush_task_gmock_tests.cpp index b3e45f4778..f227139473 100644 --- a/opencl/test/unit_test/command_stream/command_stream_receiver_flush_task_gmock_tests.cpp +++ b/opencl/test/unit_test/command_stream/command_stream_receiver_flush_task_gmock_tests.cpp @@ -5,6 +5,7 @@ * */ +#include "shared/source/built_ins/built_ins.h" #include "shared/source/command_stream/command_stream_receiver.h" #include "shared/source/command_stream/linear_stream.h" #include "shared/source/command_stream/preemption.h" @@ -20,7 +21,6 @@ #include "shared/test/unit_test/helpers/debug_manager_state_restore.h" #include "shared/test/unit_test/helpers/ult_hw_helper.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/command_queue/command_queue_hw.h" #include "opencl/source/command_queue/gpgpu_walker.h" #include "opencl/source/event/user_event.h" diff --git a/opencl/test/unit_test/command_stream/command_stream_receiver_hw_tests.cpp b/opencl/test/unit_test/command_stream/command_stream_receiver_hw_tests.cpp index f893193d5e..5c4549baa3 100644 --- a/opencl/test/unit_test/command_stream/command_stream_receiver_hw_tests.cpp +++ b/opencl/test/unit_test/command_stream/command_stream_receiver_hw_tests.cpp @@ -5,6 +5,7 @@ * */ +#include "shared/source/built_ins/built_ins.h" #include "shared/source/command_stream/command_stream_receiver.h" #include "shared/source/command_stream/linear_stream.h" #include "shared/source/command_stream/preemption.h" @@ -23,7 +24,6 @@ #include "shared/test/unit_test/helpers/debug_manager_state_restore.h" #include "shared/test/unit_test/utilities/base_object_utils.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/command_queue/command_queue_hw.h" #include "opencl/source/command_queue/gpgpu_walker.h" #include "opencl/source/event/user_event.h" diff --git a/opencl/test/unit_test/execution_environment/execution_environment_tests.cpp b/opencl/test/unit_test/execution_environment/execution_environment_tests.cpp index 2f1c5e4595..f6f6d54a60 100644 --- a/opencl/test/unit_test/execution_environment/execution_environment_tests.cpp +++ b/opencl/test/unit_test/execution_environment/execution_environment_tests.cpp @@ -5,6 +5,7 @@ * */ +#include "shared/source/built_ins/built_ins.h" #include "shared/source/command_stream/preemption.h" #include "shared/source/compiler_interface/compiler_interface.h" #include "shared/source/device/device.h" @@ -16,7 +17,6 @@ #include "shared/test/unit_test/utilities/destructor_counted.h" #include "opencl/source/aub/aub_center.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/memory_manager/os_agnostic_memory_manager.h" #include "opencl/source/platform/platform.h" #include "opencl/source/source_level_debugger/source_level_debugger.h" diff --git a/opencl/test/unit_test/execution_model/scheduler_dispatch_tests.cpp b/opencl/test/unit_test/execution_model/scheduler_dispatch_tests.cpp index 80fe992770..1a14a3f0f7 100644 --- a/opencl/test/unit_test/execution_model/scheduler_dispatch_tests.cpp +++ b/opencl/test/unit_test/execution_model/scheduler_dispatch_tests.cpp @@ -5,9 +5,9 @@ * */ +#include "shared/source/built_ins/built_ins.h" #include "shared/test/unit_test/helpers/debug_manager_state_restore.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/command_queue/enqueue_kernel.h" #include "opencl/source/device_queue/device_queue.h" #include "opencl/source/scheduler/scheduler_kernel.h" diff --git a/opencl/test/unit_test/fixtures/built_in_fixture.cpp b/opencl/test/unit_test/fixtures/built_in_fixture.cpp index 173380d78c..3857466be0 100644 --- a/opencl/test/unit_test/fixtures/built_in_fixture.cpp +++ b/opencl/test/unit_test/fixtures/built_in_fixture.cpp @@ -7,9 +7,9 @@ #include "opencl/test/unit_test/fixtures/built_in_fixture.h" +#include "shared/source/built_ins/built_ins.h" #include "shared/source/device/device.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/test/unit_test/global_environment.h" #include "opencl/test/unit_test/helpers/kernel_binary_helper.h" #include "opencl/test/unit_test/helpers/test_files.h" diff --git a/opencl/test/unit_test/gen8/scheduler_dispatch_tests_gen8.cpp b/opencl/test/unit_test/gen8/scheduler_dispatch_tests_gen8.cpp index 84e7f19675..b8d0767214 100644 --- a/opencl/test/unit_test/gen8/scheduler_dispatch_tests_gen8.cpp +++ b/opencl/test/unit_test/gen8/scheduler_dispatch_tests_gen8.cpp @@ -5,7 +5,8 @@ * */ -#include "opencl/source/built_ins/built_ins.h" +#include "shared/source/built_ins/built_ins.h" + #include "opencl/source/command_queue/enqueue_kernel.h" #include "opencl/source/command_queue/gpgpu_walker.h" #include "opencl/source/device_queue/device_queue.h" diff --git a/opencl/test/unit_test/gen9/command_stream_receiver_hw_tests_gen9.cpp b/opencl/test/unit_test/gen9/command_stream_receiver_hw_tests_gen9.cpp index 830ac3ed6d..25308a80c2 100644 --- a/opencl/test/unit_test/gen9/command_stream_receiver_hw_tests_gen9.cpp +++ b/opencl/test/unit_test/gen9/command_stream_receiver_hw_tests_gen9.cpp @@ -5,10 +5,10 @@ * */ +#include "shared/source/built_ins/built_ins.h" #include "shared/source/command_stream/command_stream_receiver.h" #include "shared/source/command_stream/linear_stream.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/command_queue/command_queue_hw.h" #include "opencl/source/event/user_event.h" #include "opencl/test/unit_test/fixtures/device_fixture.h" diff --git a/opencl/test/unit_test/gen9/sip_tests_gen9.cpp b/opencl/test/unit_test/gen9/sip_tests_gen9.cpp index 9dbcf01b0b..b272ce93d4 100644 --- a/opencl/test/unit_test/gen9/sip_tests_gen9.cpp +++ b/opencl/test/unit_test/gen9/sip_tests_gen9.cpp @@ -5,8 +5,9 @@ * */ -#include "opencl/source/built_ins/built_ins.h" -#include "opencl/source/built_ins/sip.h" +#include "shared/source/built_ins/built_ins.h" +#include "shared/source/built_ins/sip.h" + #include "opencl/test/unit_test/global_environment.h" #include "opencl/test/unit_test/helpers/test_files.h" #include "opencl/test/unit_test/mocks/mock_device.h" diff --git a/opencl/test/unit_test/helpers/hardware_commands_helper_tests.h b/opencl/test/unit_test/helpers/hardware_commands_helper_tests.h index f7d146798e..aeef2a3b2a 100644 --- a/opencl/test/unit_test/helpers/hardware_commands_helper_tests.h +++ b/opencl/test/unit_test/helpers/hardware_commands_helper_tests.h @@ -5,7 +5,8 @@ * */ -#include "opencl/source/built_ins/built_ins.h" +#include "shared/source/built_ins/built_ins.h" + #include "opencl/source/helpers/hardware_commands_helper.h" #include "opencl/source/kernel/kernel.h" #include "opencl/test/unit_test/fixtures/built_in_fixture.h" diff --git a/opencl/test/unit_test/helpers/kernel_binary_helper_hash_value.cpp b/opencl/test/unit_test/helpers/kernel_binary_helper_hash_value.cpp index 55cec409e3..5e9a180459 100644 --- a/opencl/test/unit_test/helpers/kernel_binary_helper_hash_value.cpp +++ b/opencl/test/unit_test/helpers/kernel_binary_helper_hash_value.cpp @@ -7,4 +7,4 @@ #include "opencl/test/unit_test/helpers/kernel_binary_helper.h" -const std::string KernelBinaryHelper::BUILT_INS("7030307152995455603"); \ No newline at end of file +const std::string KernelBinaryHelper::BUILT_INS("7206969092167061981"); \ No newline at end of file diff --git a/opencl/test/unit_test/mem_obj/image_tests.cpp b/opencl/test/unit_test/mem_obj/image_tests.cpp index 621c550955..6747bb806e 100644 --- a/opencl/test/unit_test/mem_obj/image_tests.cpp +++ b/opencl/test/unit_test/mem_obj/image_tests.cpp @@ -5,13 +5,13 @@ * */ +#include "shared/source/built_ins/built_ins.h" #include "shared/source/compiler_interface/compiler_interface.h" #include "shared/source/helpers/aligned_memory.h" #include "shared/source/image/image_surface_state.h" #include "shared/source/os_interface/os_context.h" #include "shared/test/unit_test/helpers/debug_manager_state_restore.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/helpers/mipmap.h" #include "opencl/source/mem_obj/image.h" #include "opencl/source/mem_obj/mem_obj_helper.h" diff --git a/opencl/test/unit_test/mocks/mock_builtin_dispatch_info_builder.h b/opencl/test/unit_test/mocks/mock_builtin_dispatch_info_builder.h index ddb03f4350..a9cff2154d 100644 --- a/opencl/test/unit_test/mocks/mock_builtin_dispatch_info_builder.h +++ b/opencl/test/unit_test/mocks/mock_builtin_dispatch_info_builder.h @@ -6,7 +6,8 @@ */ #pragma once -#include "opencl/source/built_ins/built_ins.h" +#include "shared/source/built_ins/built_ins.h" + #include "opencl/source/helpers/dispatch_info.h" using namespace NEO; diff --git a/opencl/test/unit_test/mocks/mock_builtins.h b/opencl/test/unit_test/mocks/mock_builtins.h index b5a6c111c4..ad3f68a9e5 100644 --- a/opencl/test/unit_test/mocks/mock_builtins.h +++ b/opencl/test/unit_test/mocks/mock_builtins.h @@ -7,9 +7,10 @@ #pragma once -#include "opencl/source/built_ins/built_ins.h" +#include "shared/source/built_ins/built_ins.h" +#include "shared/source/built_ins/sip.h" + #include "opencl/source/built_ins/builtins_dispatch_builder.h" -#include "opencl/source/built_ins/sip.h" #include "opencl/source/program/program.h" #include diff --git a/opencl/test/unit_test/mocks/mock_context.cpp b/opencl/test/unit_test/mocks/mock_context.cpp index bdcd546066..27d4a9abb3 100644 --- a/opencl/test/unit_test/mocks/mock_context.cpp +++ b/opencl/test/unit_test/mocks/mock_context.cpp @@ -7,11 +7,11 @@ #include "opencl/test/unit_test/mocks/mock_context.h" +#include "shared/source/built_ins/built_ins.h" #include "shared/source/compiler_interface/compiler_interface.h" #include "shared/source/memory_manager/deferred_deleter.h" #include "shared/source/memory_manager/unified_memory_manager.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/command_queue/command_queue.h" #include "opencl/source/memory_manager/os_agnostic_memory_manager.h" #include "opencl/source/sharings/sharing.h" diff --git a/opencl/test/unit_test/mocks/mock_sip.h b/opencl/test/unit_test/mocks/mock_sip.h index 2d241ef804..dbc37a2bb8 100644 --- a/opencl/test/unit_test/mocks/mock_sip.h +++ b/opencl/test/unit_test/mocks/mock_sip.h @@ -7,7 +7,8 @@ #pragma once -#include "opencl/source/built_ins/sip.h" +#include "shared/source/built_ins/sip.h" + #include "opencl/test/unit_test/mocks/mock_execution_environment.h" #include diff --git a/opencl/test/unit_test/offline_compiler/offline_compiler_tests.cpp b/opencl/test/unit_test/offline_compiler/offline_compiler_tests.cpp index 4dbf2e307e..7daf2accc7 100644 --- a/opencl/test/unit_test/offline_compiler/offline_compiler_tests.cpp +++ b/opencl/test/unit_test/offline_compiler/offline_compiler_tests.cpp @@ -349,7 +349,7 @@ TEST_F(OfflineCompilerTests, GoodParseBinToCharArray) { familyNameWithType + "[10] = {\n" " 0x40032302, 0x90800756, 0x05340301, 0x66097860, 0x101010ff, 0x40032302, 0x90800756, 0x05340301, \n" " 0x66097860, 0xff000000};\n\n" - "#include \"opencl/source/built_ins/registry/built_ins_registry.h\"\n\n" + "#include \"shared/source/built_ins/registry/built_ins_registry.h\"\n\n" "namespace NEO {\n" "static RegisterEmbeddedResource registerSchedulerBin(\n" " \"" + diff --git a/opencl/test/unit_test/program/program_from_binary.h b/opencl/test/unit_test/program/program_from_binary.h index f556801b01..a222c0b280 100644 --- a/opencl/test/unit_test/program/program_from_binary.h +++ b/opencl/test/unit_test/program/program_from_binary.h @@ -6,7 +6,8 @@ */ #pragma once -#include "opencl/source/built_ins/built_ins.h" +#include "shared/source/built_ins/built_ins.h" + #include "opencl/test/unit_test/fixtures/context_fixture.h" #include "opencl/test/unit_test/fixtures/device_fixture.h" #include "opencl/test/unit_test/fixtures/program_fixture.h" diff --git a/opencl/test/unit_test/test_files/7030307152995455603.cl b/opencl/test/unit_test/test_files/7206969092167061981.cl similarity index 57% rename from opencl/test/unit_test/test_files/7030307152995455603.cl rename to opencl/test/unit_test/test_files/7206969092167061981.cl index 2b38ca3c4e..871ee09199 100644 --- a/opencl/test/unit_test/test_files/7030307152995455603.cl +++ b/opencl/test/unit_test/test_files/7206969092167061981.cl @@ -5,40 +5,37 @@ * */ -__kernel void fullCopy(__global const uint* src, __global uint* dst) { +__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); } __kernel void CopyBufferToBufferBytes( - const __global uchar* pSrc, - __global uchar* pDst, + const __global uchar *pSrc, + __global uchar *pDst, uint srcOffsetInBytes, uint dstOffsetInBytes, - uint bytesToRead ) -{ - pSrc += ( srcOffsetInBytes + get_global_id(0) ); - pDst += ( dstOffsetInBytes + get_global_id(0) ); - pDst[ 0 ] = pSrc[ 0 ]; + uint bytesToRead) { + pSrc += (srcOffsetInBytes + get_global_id(0)); + pDst += (dstOffsetInBytes + get_global_id(0)); + pDst[0] = pSrc[0]; } __kernel void CopyBufferToBufferLeftLeftover( - const __global uchar* pSrc, - __global uchar* pDst, + const __global uchar *pSrc, + __global uchar *pDst, uint srcOffsetInBytes, - uint dstOffsetInBytes) -{ + uint dstOffsetInBytes) { unsigned int gid = get_global_id(0); - pDst[ gid + dstOffsetInBytes ] = pSrc[ gid + srcOffsetInBytes ]; + pDst[gid + dstOffsetInBytes] = pSrc[gid + srcOffsetInBytes]; } __kernel void CopyBufferToBufferMiddle( - const __global uint* pSrc, - __global uint* pDst, + const __global uint *pSrc, + __global uint *pDst, uint srcOffsetInBytes, - uint dstOffsetInBytes) -{ + uint dstOffsetInBytes) { unsigned int gid = get_global_id(0); pDst += dstOffsetInBytes >> 2; pSrc += srcOffsetInBytes >> 2; @@ -47,55 +44,49 @@ __kernel void CopyBufferToBufferMiddle( } __kernel void CopyBufferToBufferRightLeftover( - const __global uchar* pSrc, - __global uchar* pDst, + const __global uchar *pSrc, + __global uchar *pDst, uint srcOffsetInBytes, - uint dstOffsetInBytes) -{ + uint dstOffsetInBytes) { unsigned int gid = get_global_id(0); - pDst[ gid + dstOffsetInBytes ] = pSrc[ gid + srcOffsetInBytes ]; + pDst[gid + dstOffsetInBytes] = pSrc[gid + srcOffsetInBytes]; } - // assumption is local work size = pattern size __kernel void FillBufferBytes( - __global uchar* pDst, + __global uchar *pDst, uint dstOffsetInBytes, - const __global uchar* pPattern ) -{ + const __global uchar *pPattern) { uint dstIndex = get_global_id(0) + dstOffsetInBytes; uint srcIndex = get_local_id(0); pDst[dstIndex] = pPattern[srcIndex]; } __kernel void FillBufferLeftLeftover( - __global uchar* pDst, + __global uchar *pDst, uint dstOffsetInBytes, - const __global uchar* pPattern, - const uint patternSizeInEls ) -{ + const __global uchar *pPattern, + const uint patternSizeInEls) { uint gid = get_global_id(0); - pDst[ gid + dstOffsetInBytes ] = pPattern[ gid & (patternSizeInEls - 1) ]; + pDst[gid + dstOffsetInBytes] = pPattern[gid & (patternSizeInEls - 1)]; } __kernel void FillBufferMiddle( - __global uchar* pDst, + __global uchar *pDst, uint dstOffsetInBytes, - const __global uint* pPattern, - const uint patternSizeInEls ) -{ + const __global uint *pPattern, + const uint patternSizeInEls) { uint gid = get_global_id(0); - ((__global uint*)(pDst + dstOffsetInBytes))[gid] = pPattern[ gid & (patternSizeInEls - 1) ]; + ((__global uint *)(pDst + dstOffsetInBytes))[gid] = pPattern[gid & (patternSizeInEls - 1)]; } __kernel void FillBufferRightLeftover( - __global uchar* pDst, + __global uchar *pDst, uint dstOffsetInBytes, - const __global uchar* pPattern, - const uint patternSizeInEls ) -{ + const __global uchar *pPattern, + const uint patternSizeInEls) { uint gid = get_global_id(0); - pDst[ gid + dstOffsetInBytes ] = pPattern[ gid & (patternSizeInEls - 1) ]; + pDst[gid + dstOffsetInBytes] = pPattern[gid & (patternSizeInEls - 1)]; } __kernel void FillImage1d( @@ -179,42 +170,40 @@ __kernel void CopyImageToImage3d( ////////////////////////////////////////////////////////////////////////////// __kernel void CopyBufferRectBytes2d( - __global const char* src, - __global char* dst, + __global const char *src, + __global char *dst, uint4 SrcOrigin, uint4 DstOrigin, uint2 SrcPitch, - uint2 DstPitch ) + uint2 DstPitch) { int x = get_global_id(0); int y = get_global_id(1); - uint LSrcOffset = x + SrcOrigin.x + ( ( y + SrcOrigin.y ) * SrcPitch.x ); - uint LDstOffset = x + DstOrigin.x + ( ( y + DstOrigin.y ) * DstPitch.x ); - - *( dst + LDstOffset ) = *( src + LSrcOffset ); + uint LSrcOffset = x + SrcOrigin.x + ((y + SrcOrigin.y) * SrcPitch.x); + uint LDstOffset = x + DstOrigin.x + ((y + DstOrigin.y) * DstPitch.x); + *(dst + LDstOffset) = *(src + LSrcOffset); } ////////////////////////////////////////////////////////////////////////////// __kernel void CopyBufferRectBytes3d( - __global const char* src, - __global char* dst, - uint4 SrcOrigin, - uint4 DstOrigin, - uint2 SrcPitch, - uint2 DstPitch ) - -{ - int x = get_global_id(0); - int y = get_global_id(1); - int z = get_global_id(2); - - uint LSrcOffset = x + SrcOrigin.x + ( ( y + SrcOrigin.y ) * SrcPitch.x ) + ( ( z + SrcOrigin.z ) * SrcPitch.y ); - uint LDstOffset = x + DstOrigin.x + ( ( y + DstOrigin.y ) * DstPitch.x ) + ( ( z + DstOrigin.z ) * DstPitch.y ); - - *( dst + LDstOffset ) = *( src + LSrcOffset ); - + __global const char *src, + __global char *dst, + uint4 SrcOrigin, + uint4 DstOrigin, + uint2 SrcPitch, + uint2 DstPitch) + +{ + int x = get_global_id(0); + int y = get_global_id(1); + int z = get_global_id(2); + + uint LSrcOffset = x + SrcOrigin.x + ((y + SrcOrigin.y) * SrcPitch.x) + ((z + SrcOrigin.z) * SrcPitch.y); + uint LDstOffset = x + DstOrigin.x + ((y + DstOrigin.y) * DstPitch.x) + ((z + DstOrigin.z) * DstPitch.y); + + *(dst + LDstOffset) = *(src + LSrcOffset); } #pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable @@ -248,14 +237,13 @@ __kernel void CopyBufferToImage3d2Bytes(__global uchar *src, uint4 c = (uint4)(0, 0, 0, 1); - if(( ulong )(src + srcOffset) & 0x00000001){ - ushort upper = *((__global uchar*)(src + LOffset + x * 2 + 1)); - ushort lower = *((__global uchar*)(src + LOffset + x * 2)); + if ((ulong)(src + srcOffset) & 0x00000001) { + ushort upper = *((__global uchar *)(src + LOffset + x * 2 + 1)); + ushort lower = *((__global uchar *)(src + LOffset + x * 2)); ushort combined = (upper << 8) | lower; c.x = (uint)combined; - } - else{ - c.x = (uint)(*(__global ushort*)(src + LOffset + x * 2)); + } else { + c.x = (uint)(*(__global ushort *)(src + LOffset + x * 2)); } write_imageui(output, dstCoord, c); } @@ -274,16 +262,15 @@ __kernel void CopyBufferToImage3d4Bytes(__global uchar *src, uint4 c = (uint4)(0, 0, 0, 1); - if(( ulong )(src + srcOffset) & 0x00000003){ - uint upper2 = *((__global uchar*)(src + LOffset + x * 4 + 3)); - uint upper = *((__global uchar*)(src + LOffset + x * 4 + 2)); - uint lower2 = *((__global uchar*)(src + LOffset + x * 4 + 1)); - uint lower = *((__global uchar*)(src + LOffset + x * 4)); + if ((ulong)(src + srcOffset) & 0x00000003) { + uint upper2 = *((__global uchar *)(src + LOffset + x * 4 + 3)); + uint upper = *((__global uchar *)(src + LOffset + x * 4 + 2)); + uint lower2 = *((__global uchar *)(src + LOffset + x * 4 + 1)); + uint lower = *((__global uchar *)(src + LOffset + x * 4)); uint combined = (upper2 << 24) | (upper << 16) | (lower2 << 8) | lower; c.x = combined; - } - else{ - c.x = (*(__global uint*)(src + LOffset + x * 4)); + } else { + c.x = (*(__global uint *)(src + LOffset + x * 4)); } write_imageui(output, dstCoord, c); } @@ -300,24 +287,23 @@ __kernel void CopyBufferToImage3d8Bytes(__global uchar *src, int4 dstCoord = (int4)(x, y, z, 0) + dstOffset; uint LOffset = srcOffset + (y * Pitch.x) + (z * Pitch.y); - uint2 c = (uint2)(0, 0);//*((__global uint2*)(src + LOffset + x * 8)); + uint2 c = (uint2)(0, 0); //*((__global uint2*)(src + LOffset + x * 8)); - if(( ulong )(src + srcOffset) & 0x00000007){ - uint upper2 = *((__global uchar*)(src + LOffset + x * 8 + 3)); - uint upper = *((__global uchar*)(src + LOffset + x * 8 + 2)); - uint lower2 = *((__global uchar*)(src + LOffset + x * 8 + 1)); - uint lower = *((__global uchar*)(src + LOffset + x * 8)); + if ((ulong)(src + srcOffset) & 0x00000007) { + uint upper2 = *((__global uchar *)(src + LOffset + x * 8 + 3)); + uint upper = *((__global uchar *)(src + LOffset + x * 8 + 2)); + uint lower2 = *((__global uchar *)(src + LOffset + x * 8 + 1)); + uint lower = *((__global uchar *)(src + LOffset + x * 8)); uint combined = (upper2 << 24) | (upper << 16) | (lower2 << 8) | lower; c.x = combined; - upper2 = *((__global uchar*)(src + LOffset + x * 8 + 7)); - upper = *((__global uchar*)(src + LOffset + x * 8 + 6)); - lower2 = *((__global uchar*)(src + LOffset + x * 8 + 5)); - lower = *((__global uchar*)(src + LOffset + x * 8 + 4)); + upper2 = *((__global uchar *)(src + LOffset + x * 8 + 7)); + upper = *((__global uchar *)(src + LOffset + x * 8 + 6)); + lower2 = *((__global uchar *)(src + LOffset + x * 8 + 5)); + lower = *((__global uchar *)(src + LOffset + x * 8 + 4)); combined = ((uint)upper2 << 24) | ((uint)upper << 16) | ((uint)lower2 << 8) | lower; c.y = combined; - } - else{ - c = *((__global uint2*)(src + LOffset + x * 8)); + } else { + c = *((__global uint2 *)(src + LOffset + x * 8)); } write_imageui(output, dstCoord, (uint4)(c.x, c.y, 0, 1)); @@ -337,33 +323,32 @@ __kernel void CopyBufferToImage3d16Bytes(__global uchar *src, uint4 c = (uint4)(0, 0, 0, 0); - if(( ulong )(src + srcOffset) & 0x0000000f){ - uint upper2 = *((__global uchar*)(src + LOffset + x * 16 + 3)); - uint upper = *((__global uchar*)(src + LOffset + x * 16 + 2)); - uint lower2 = *((__global uchar*)(src + LOffset + x * 16 + 1)); - uint lower = *((__global uchar*)(src + LOffset + x * 16)); + if ((ulong)(src + srcOffset) & 0x0000000f) { + uint upper2 = *((__global uchar *)(src + LOffset + x * 16 + 3)); + uint upper = *((__global uchar *)(src + LOffset + x * 16 + 2)); + uint lower2 = *((__global uchar *)(src + LOffset + x * 16 + 1)); + uint lower = *((__global uchar *)(src + LOffset + x * 16)); uint combined = (upper2 << 24) | (upper << 16) | (lower2 << 8) | lower; c.x = combined; - upper2 = *((__global uchar*)(src + LOffset + x * 16 + 7)); - upper = *((__global uchar*)(src + LOffset + x * 16 + 6)); - lower2 = *((__global uchar*)(src + LOffset + x * 16 + 5)); - lower = *((__global uchar*)(src + LOffset + x * 16 + 4)); + upper2 = *((__global uchar *)(src + LOffset + x * 16 + 7)); + upper = *((__global uchar *)(src + LOffset + x * 16 + 6)); + lower2 = *((__global uchar *)(src + LOffset + x * 16 + 5)); + lower = *((__global uchar *)(src + LOffset + x * 16 + 4)); combined = (upper2 << 24) | (upper << 16) | (lower2 << 8) | lower; c.y = combined; - upper2 = *((__global uchar*)(src + LOffset + x * 16 + 11)); - upper = *((__global uchar*)(src + LOffset + x * 16 + 10)); - lower2 = *((__global uchar*)(src + LOffset + x * 16 + 9)); - lower = *((__global uchar*)(src + LOffset + x * 16 + 8)); + upper2 = *((__global uchar *)(src + LOffset + x * 16 + 11)); + upper = *((__global uchar *)(src + LOffset + x * 16 + 10)); + lower2 = *((__global uchar *)(src + LOffset + x * 16 + 9)); + lower = *((__global uchar *)(src + LOffset + x * 16 + 8)); combined = (upper2 << 24) | (upper << 16) | (lower2 << 8) | lower; c.z = combined; - upper2 = *((__global uchar*)(src + LOffset + x * 16 + 15)); - upper = *((__global uchar*)(src + LOffset + x * 16 + 14)); - lower2 = *((__global uchar*)(src + LOffset + x * 16 + 13)); - lower = *((__global uchar*)(src + LOffset + x * 16 + 12)); + upper2 = *((__global uchar *)(src + LOffset + x * 16 + 15)); + upper = *((__global uchar *)(src + LOffset + x * 16 + 14)); + lower2 = *((__global uchar *)(src + LOffset + x * 16 + 13)); + lower = *((__global uchar *)(src + LOffset + x * 16 + 12)); combined = (upper2 << 24) | (upper << 16) | (lower2 << 8) | lower; c.w = combined; - } - else{ + } else { c = *((__global uint4 *)(src + LOffset + x * 16)); } @@ -397,15 +382,14 @@ __kernel void CopyImage3dToBuffer2Bytes(__read_only image3d_t input, const int4 srcCoord = (int4)(x, y, z, 0) + srcOffset; uint DstOffset = dstOffset + (y * Pitch.x) + (z * Pitch.y); - + uint4 c = read_imageui(input, srcCoord); - if(( ulong )(dst + dstOffset) & 0x00000001){ - *((__global uchar*)(dst + DstOffset + x * 2 + 1)) = convert_uchar_sat((c.x >> 8 ) & 0xff); - *((__global uchar*)(dst + DstOffset + x * 2)) = convert_uchar_sat(c.x & 0xff); - } - else{ - *((__global ushort*)(dst + DstOffset + x * 2)) = convert_ushort_sat(c.x); + if ((ulong)(dst + dstOffset) & 0x00000001) { + *((__global uchar *)(dst + DstOffset + x * 2 + 1)) = convert_uchar_sat((c.x >> 8) & 0xff); + *((__global uchar *)(dst + DstOffset + x * 2)) = convert_uchar_sat(c.x & 0xff); + } else { + *((__global ushort *)(dst + DstOffset + x * 2)) = convert_ushort_sat(c.x); } } @@ -423,14 +407,13 @@ __kernel void CopyImage3dToBuffer4Bytes(__read_only image3d_t input, uint4 c = read_imageui(input, srcCoord); - if(( ulong )(dst + dstOffset) & 0x00000003){ - *((__global uchar*)(dst + DstOffset + x * 4 + 3)) = convert_uchar_sat((c.x >> 24 ) & 0xff); - *((__global uchar*)(dst + DstOffset + x * 4 + 2)) = convert_uchar_sat((c.x >> 16 ) & 0xff); - *((__global uchar*)(dst + DstOffset + x * 4 + 1)) = convert_uchar_sat((c.x >> 8 ) & 0xff); - *((__global uchar*)(dst + DstOffset + x * 4)) = convert_uchar_sat(c.x & 0xff); - } - else{ - *((__global uint*)(dst + DstOffset + x * 4)) = c.x; + if ((ulong)(dst + dstOffset) & 0x00000003) { + *((__global uchar *)(dst + DstOffset + x * 4 + 3)) = convert_uchar_sat((c.x >> 24) & 0xff); + *((__global uchar *)(dst + DstOffset + x * 4 + 2)) = convert_uchar_sat((c.x >> 16) & 0xff); + *((__global uchar *)(dst + DstOffset + x * 4 + 1)) = convert_uchar_sat((c.x >> 8) & 0xff); + *((__global uchar *)(dst + DstOffset + x * 4)) = convert_uchar_sat(c.x & 0xff); + } else { + *((__global uint *)(dst + DstOffset + x * 4)) = c.x; } } @@ -448,19 +431,18 @@ __kernel void CopyImage3dToBuffer8Bytes(__read_only image3d_t input, uint4 c = read_imageui(input, srcCoord); - if(( ulong )(dst + dstOffset) & 0x00000007){ - *((__global uchar*)(dst + DstOffset + x * 8 + 3)) = convert_uchar_sat((c.x >> 24 ) & 0xff); - *((__global uchar*)(dst + DstOffset + x * 8 + 2)) = convert_uchar_sat((c.x >> 16 ) & 0xff); - *((__global uchar*)(dst + DstOffset + x * 8 + 1)) = convert_uchar_sat((c.x >> 8 ) & 0xff); - *((__global uchar*)(dst + DstOffset + x * 8)) = convert_uchar_sat(c.x & 0xff); - *((__global uchar*)(dst + DstOffset + x * 8 + 7)) = convert_uchar_sat((c.y >> 24 ) & 0xff); - *((__global uchar*)(dst + DstOffset + x * 8 + 6)) = convert_uchar_sat((c.y >> 16 ) & 0xff); - *((__global uchar*)(dst + DstOffset + x * 8 + 5)) = convert_uchar_sat((c.y >> 8 ) & 0xff); - *((__global uchar*)(dst + DstOffset + x * 8 + 4)) = convert_uchar_sat(c.y & 0xff); - } - else{ - uint2 d = (uint2)(c.x,c.y); - *((__global uint2*)(dst + DstOffset + x * 8)) = d; + if ((ulong)(dst + dstOffset) & 0x00000007) { + *((__global uchar *)(dst + DstOffset + x * 8 + 3)) = convert_uchar_sat((c.x >> 24) & 0xff); + *((__global uchar *)(dst + DstOffset + x * 8 + 2)) = convert_uchar_sat((c.x >> 16) & 0xff); + *((__global uchar *)(dst + DstOffset + x * 8 + 1)) = convert_uchar_sat((c.x >> 8) & 0xff); + *((__global uchar *)(dst + DstOffset + x * 8)) = convert_uchar_sat(c.x & 0xff); + *((__global uchar *)(dst + DstOffset + x * 8 + 7)) = convert_uchar_sat((c.y >> 24) & 0xff); + *((__global uchar *)(dst + DstOffset + x * 8 + 6)) = convert_uchar_sat((c.y >> 16) & 0xff); + *((__global uchar *)(dst + DstOffset + x * 8 + 5)) = convert_uchar_sat((c.y >> 8) & 0xff); + *((__global uchar *)(dst + DstOffset + x * 8 + 4)) = convert_uchar_sat(c.y & 0xff); + } else { + uint2 d = (uint2)(c.x, c.y); + *((__global uint2 *)(dst + DstOffset + x * 8)) = d; } } @@ -478,25 +460,24 @@ __kernel void CopyImage3dToBuffer16Bytes(__read_only image3d_t input, const uint4 c = read_imageui(input, srcCoord); - if(( ulong )(dst + dstOffset) & 0x0000000f){ - *((__global uchar*)(dst + DstOffset + x * 16 + 3)) = convert_uchar_sat((c.x >> 24 ) & 0xff); - *((__global uchar*)(dst + DstOffset + x * 16 + 2)) = convert_uchar_sat((c.x >> 16 ) & 0xff); - *((__global uchar*)(dst + DstOffset + x * 16 + 1)) = convert_uchar_sat((c.x >> 8 ) & 0xff); - *((__global uchar*)(dst + DstOffset + x * 16)) = convert_uchar_sat(c.x & 0xff); - *((__global uchar*)(dst + DstOffset + x * 16 + 7)) = convert_uchar_sat((c.y >> 24 ) & 0xff); - *((__global uchar*)(dst + DstOffset + x * 16 + 6)) = convert_uchar_sat((c.y >> 16 ) & 0xff); - *((__global uchar*)(dst + DstOffset + x * 16 + 5)) = convert_uchar_sat((c.y >> 8 ) & 0xff); - *((__global uchar*)(dst + DstOffset + x * 16 + 4)) = convert_uchar_sat(c.y & 0xff); - *((__global uchar*)(dst + DstOffset + x * 16 + 11)) = convert_uchar_sat((c.z >> 24 ) & 0xff); - *((__global uchar*)(dst + DstOffset + x * 16 + 10)) = convert_uchar_sat((c.z >> 16 ) & 0xff); - *((__global uchar*)(dst + DstOffset + x * 16 + 9)) = convert_uchar_sat((c.z >> 8 ) & 0xff); - *((__global uchar*)(dst + DstOffset + x * 16 + 8)) = convert_uchar_sat(c.z & 0xff); - *((__global uchar*)(dst + DstOffset + x * 16 + 15)) = convert_uchar_sat((c.w >> 24 ) & 0xff); - *((__global uchar*)(dst + DstOffset + x * 16 + 14)) = convert_uchar_sat((c.w >> 16 ) & 0xff); - *((__global uchar*)(dst + DstOffset + x * 16 + 13)) = convert_uchar_sat((c.w >> 8 ) & 0xff); - *((__global uchar*)(dst + DstOffset + x * 16 + 12)) = convert_uchar_sat(c.w & 0xff); - } - else{ - *(__global uint4*)(dst + DstOffset + x * 16) = c; + if ((ulong)(dst + dstOffset) & 0x0000000f) { + *((__global uchar *)(dst + DstOffset + x * 16 + 3)) = convert_uchar_sat((c.x >> 24) & 0xff); + *((__global uchar *)(dst + DstOffset + x * 16 + 2)) = convert_uchar_sat((c.x >> 16) & 0xff); + *((__global uchar *)(dst + DstOffset + x * 16 + 1)) = convert_uchar_sat((c.x >> 8) & 0xff); + *((__global uchar *)(dst + DstOffset + x * 16)) = convert_uchar_sat(c.x & 0xff); + *((__global uchar *)(dst + DstOffset + x * 16 + 7)) = convert_uchar_sat((c.y >> 24) & 0xff); + *((__global uchar *)(dst + DstOffset + x * 16 + 6)) = convert_uchar_sat((c.y >> 16) & 0xff); + *((__global uchar *)(dst + DstOffset + x * 16 + 5)) = convert_uchar_sat((c.y >> 8) & 0xff); + *((__global uchar *)(dst + DstOffset + x * 16 + 4)) = convert_uchar_sat(c.y & 0xff); + *((__global uchar *)(dst + DstOffset + x * 16 + 11)) = convert_uchar_sat((c.z >> 24) & 0xff); + *((__global uchar *)(dst + DstOffset + x * 16 + 10)) = convert_uchar_sat((c.z >> 16) & 0xff); + *((__global uchar *)(dst + DstOffset + x * 16 + 9)) = convert_uchar_sat((c.z >> 8) & 0xff); + *((__global uchar *)(dst + DstOffset + x * 16 + 8)) = convert_uchar_sat(c.z & 0xff); + *((__global uchar *)(dst + DstOffset + x * 16 + 15)) = convert_uchar_sat((c.w >> 24) & 0xff); + *((__global uchar *)(dst + DstOffset + x * 16 + 14)) = convert_uchar_sat((c.w >> 16) & 0xff); + *((__global uchar *)(dst + DstOffset + x * 16 + 13)) = convert_uchar_sat((c.w >> 8) & 0xff); + *((__global uchar *)(dst + DstOffset + x * 16 + 12)) = convert_uchar_sat(c.w & 0xff); + } else { + *(__global uint4 *)(dst + DstOffset + x * 16) = c; } } diff --git a/opencl/test/unit_test/test_files/7030307152995455603_options.txt b/opencl/test/unit_test/test_files/7206969092167061981_options.txt similarity index 100% rename from opencl/test/unit_test/test_files/7030307152995455603_options.txt rename to opencl/test/unit_test/test_files/7206969092167061981_options.txt diff --git a/shared/source/CMakeLists.txt b/shared/source/CMakeLists.txt index 48738ef64e..3137b81c29 100644 --- a/shared/source/CMakeLists.txt +++ b/shared/source/CMakeLists.txt @@ -59,6 +59,8 @@ else() list(APPEND CORE_SOURCES ${CORE_SRCS_GENX_ALL_LINUX}) endif() +message("${NEO_CORE_SRCS_BUILT_INS}") + append_sources_from_properties(CORE_SOURCES NEO_CORE_COMMAND_CONTAINER NEO_CORE_DEBUG_SETTINGS @@ -69,6 +71,9 @@ append_sources_from_properties(CORE_SOURCES NEO_CORE_HELPERS NEO_CORE_INDIRECT_HEAP NEO_CORE_SKU_INFO_BASE + NEO_CORE_SRCS_BUILT_INS + NEO_CORE_SRCS_BUILT_IN_KERNELS + NEO_CORE_SRCS_BUILT_INS_OPS NEO_CORE_UTILITIES NEO_UNIFIED_MEMORY ) diff --git a/shared/source/built_ins/CMakeLists.txt b/shared/source/built_ins/CMakeLists.txt index c35699d2a8..b9da918ad8 100644 --- a/shared/source/built_ins/CMakeLists.txt +++ b/shared/source/built_ins/CMakeLists.txt @@ -1,12 +1,46 @@ # -# Copyright (C) 2020 Intel Corporation +# Copyright (C) 2017-2020 Intel Corporation # # SPDX-License-Identifier: MIT # -set(NEO_CORE_BUILT_INS +set(NEO_CORE_SRCS_BUILT_INS ${CMAKE_CURRENT_SOURCE_DIR}/CMakeLists.txt + ${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_in_ops_base.h + ${CMAKE_CURRENT_SOURCE_DIR}/sip.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/sip.h ${CMAKE_CURRENT_SOURCE_DIR}/sip_kernel_type.h ) -set_property(GLOBAL PROPERTY NEO_CORE_BUILT_INS ${NEO_CORE_BUILT_INS}) \ No newline at end of file +add_subdirectory(builtinops) + +set_property(GLOBAL PROPERTY NEO_CORE_SRCS_BUILT_INS ${NEO_CORE_SRCS_BUILT_INS}) + +set(NEO_CORE_SRCS_BUILT_IN_KERNELS + ${CMAKE_CURRENT_SOURCE_DIR}/kernels/aux_translation.builtin_kernel + ${CMAKE_CURRENT_SOURCE_DIR}/kernels/copy_buffer_rect.builtin_kernel + ${CMAKE_CURRENT_SOURCE_DIR}/kernels/copy_buffer_rect_stateless.builtin_kernel + ${CMAKE_CURRENT_SOURCE_DIR}/kernels/copy_buffer_to_buffer.builtin_kernel + ${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_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 + ${CMAKE_CURRENT_SOURCE_DIR}/kernels/copy_image_to_image2d.builtin_kernel + ${CMAKE_CURRENT_SOURCE_DIR}/kernels/copy_image_to_image3d.builtin_kernel + ${CMAKE_CURRENT_SOURCE_DIR}/kernels/fill_buffer.builtin_kernel + ${CMAKE_CURRENT_SOURCE_DIR}/kernels/fill_buffer_stateless.builtin_kernel + ${CMAKE_CURRENT_SOURCE_DIR}/kernels/fill_image1d.builtin_kernel + ${CMAKE_CURRENT_SOURCE_DIR}/kernels/fill_image2d.builtin_kernel + ${CMAKE_CURRENT_SOURCE_DIR}/kernels/fill_image3d.builtin_kernel +) + +set_property(GLOBAL PROPERTY NEO_CORE_SRCS_BUILT_IN_KERNELS ${NEO_CORE_SRCS_BUILT_IN_KERNELS}) + +if(NOT (TARGET ${BUILTINS_BINARIES_LIB_NAME})) + include(builtins_binary.cmake) +endif() diff --git a/opencl/source/built_ins/built_in_ops_base.h b/shared/source/built_ins/built_in_ops_base.h similarity index 100% rename from opencl/source/built_ins/built_in_ops_base.h rename to shared/source/built_ins/built_in_ops_base.h diff --git a/opencl/source/built_ins/built_ins.cpp b/shared/source/built_ins/built_ins.cpp similarity index 63% rename from opencl/source/built_ins/built_ins.cpp rename to shared/source/built_ins/built_ins.cpp index 0cf59d8433..877a8b2184 100644 --- a/opencl/source/built_ins/built_ins.cpp +++ b/shared/source/built_ins/built_ins.cpp @@ -5,21 +5,16 @@ * */ -#include "opencl/source/built_ins/built_ins.h" +#include "shared/source/built_ins/built_ins.h" +#include "shared/source/built_ins/sip.h" #include "shared/source/compiler_interface/compiler_interface.h" #include "shared/source/helpers/basic_math.h" #include "shared/source/helpers/debug_helpers.h" -#include "opencl/source/built_ins/aux_translation_builtin.h" -#include "opencl/source/built_ins/sip.h" -#include "opencl/source/device/cl_device.h" #include "opencl/source/helpers/built_ins_helper.h" #include "opencl/source/helpers/convert_color.h" #include "opencl/source/helpers/dispatch_info_builder.h" -#include "opencl/source/kernel/kernel.h" -#include "opencl/source/mem_obj/image.h" -#include "opencl/source/program/program.h" #include "compiler_options.h" @@ -40,7 +35,7 @@ const SipKernel &BuiltIns::getSipKernel(SipKernelType type, Device &device) { auto &sipBuiltIn = this->sipKernels[kernelId]; auto initializer = [&] { - cl_int retVal = CL_SUCCESS; + int retVal = 0; std::vector sipBinary; auto compilerInteface = device.getExecutionEnvironment()->getCompilerInterface(); @@ -56,13 +51,13 @@ const SipKernel &BuiltIns::getSipKernel(SipKernelType type, Device &device) { sipBinary.size(), &retVal, &device); - DEBUG_BREAK_IF(retVal != CL_SUCCESS); + DEBUG_BREAK_IF(retVal != 0); UNRECOVERABLE_IF(program == nullptr); program->setDevice(&device); retVal = program->processGenBinary(); - DEBUG_BREAK_IF(retVal != CL_SUCCESS); + DEBUG_BREAK_IF(retVal != 0); sipBuiltIn.first.reset(new SipKernel(type, program)); }; @@ -71,24 +66,4 @@ const SipKernel &BuiltIns::getSipKernel(SipKernelType type, Device &device) { return *sipBuiltIn.first; } -BuiltInOwnershipWrapper::BuiltInOwnershipWrapper(BuiltinDispatchInfoBuilder &inputBuilder, Context *context) { - takeOwnership(inputBuilder, context); -} -BuiltInOwnershipWrapper::~BuiltInOwnershipWrapper() { - if (builder) { - for (auto &kernel : builder->peekUsedKernels()) { - kernel->setContext(nullptr); - kernel->releaseOwnership(); - } - } -} -void BuiltInOwnershipWrapper::takeOwnership(BuiltinDispatchInfoBuilder &inputBuilder, Context *context) { - UNRECOVERABLE_IF(builder); - builder = &inputBuilder; - for (auto &kernel : builder->peekUsedKernels()) { - kernel->takeOwnership(); - kernel->setContext(context); - } -} - } // namespace NEO diff --git a/opencl/source/built_ins/built_ins.h b/shared/source/built_ins/built_ins.h similarity index 93% rename from opencl/source/built_ins/built_ins.h rename to shared/source/built_ins/built_ins.h index 39fee87c31..4c1c0db58c 100644 --- a/opencl/source/built_ins/built_ins.h +++ b/shared/source/built_ins/built_ins.h @@ -11,7 +11,6 @@ #include "shared/source/helpers/non_copyable_or_moveable.h" #include "shared/source/helpers/vec.h" -#include "CL/cl.h" #include "built_in_ops.h" #include "compiler_options.h" @@ -196,18 +195,6 @@ class BuiltIns { bool enableCacheing = true; }; -class BuiltInOwnershipWrapper : public NonCopyableOrMovableClass { - public: - BuiltInOwnershipWrapper() = default; - BuiltInOwnershipWrapper(BuiltinDispatchInfoBuilder &inputBuilder, Context *context); - ~BuiltInOwnershipWrapper(); - - void takeOwnership(BuiltinDispatchInfoBuilder &inputBuilder, Context *context); - - protected: - BuiltinDispatchInfoBuilder *builder = nullptr; -}; - template class BuiltInOp; diff --git a/opencl/source/built_ins/built_ins_storage.cpp b/shared/source/built_ins/built_ins_storage.cpp similarity index 99% rename from opencl/source/built_ins/built_ins_storage.cpp rename to shared/source/built_ins/built_ins_storage.cpp index 41e6f32c86..15db7bc934 100644 --- a/opencl/source/built_ins/built_ins_storage.cpp +++ b/shared/source/built_ins/built_ins_storage.cpp @@ -5,10 +5,10 @@ * */ +#include "shared/source/built_ins/built_ins.h" #include "shared/source/debug_settings/debug_settings_manager.h" #include "shared/source/device/device.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/built_ins/builtins_dispatch_builder.h" #include "os_inc.h" diff --git a/opencl/source/built_ins/builtinops/CMakeLists.txt b/shared/source/built_ins/builtinops/CMakeLists.txt similarity index 57% rename from opencl/source/built_ins/builtinops/CMakeLists.txt rename to shared/source/built_ins/builtinops/CMakeLists.txt index d0f0b220ae..c76e918427 100644 --- a/opencl/source/built_ins/builtinops/CMakeLists.txt +++ b/shared/source/built_ins/builtinops/CMakeLists.txt @@ -4,11 +4,10 @@ # SPDX-License-Identifier: MIT # -set(RUNTIME_SRCS_BUILT_INS_OPS +set(NEO_CORE_SRCS_BUILT_INS_OPS ${CMAKE_CURRENT_SOURCE_DIR}/CMakeLists.txt ${CMAKE_CURRENT_SOURCE_DIR}${BRANCH_DIR_SUFFIX}/built_in_ops.h ) +add_subdirectories() -target_sources(${NEO_STATIC_LIB_NAME} PRIVATE ${RUNTIME_SRCS_EMBARGO_BUILT_INS_OPS}) - -add_subdirectories() \ No newline at end of file +set_property(GLOBAL PROPERTY NEO_CORE_SRCS_BUILT_INS_OPS ${NEO_CORE_SRCS_BUILT_INS_OPS}) diff --git a/opencl/source/built_ins/builtinops/built_in_ops.h b/shared/source/built_ins/builtinops/built_in_ops.h similarity index 79% rename from opencl/source/built_ins/builtinops/built_in_ops.h rename to shared/source/built_ins/builtinops/built_in_ops.h index a83a02a841..46736d7ab8 100644 --- a/opencl/source/built_ins/builtinops/built_in_ops.h +++ b/shared/source/built_ins/builtinops/built_in_ops.h @@ -6,7 +6,7 @@ */ #pragma once -#include "opencl/source/built_ins/built_in_ops_base.h" +#include "shared/source/built_ins/built_in_ops_base.h" namespace NEO { namespace EBuiltInOps { diff --git a/opencl/source/built_ins/builtins_binary.cmake b/shared/source/built_ins/builtins_binary.cmake similarity index 100% rename from opencl/source/built_ins/builtins_binary.cmake rename to shared/source/built_ins/builtins_binary.cmake diff --git a/shared/source/built_ins/kernels/CMakeLists.txt b/shared/source/built_ins/kernels/CMakeLists.txt new file mode 100644 index 0000000000..41ac296e5f --- /dev/null +++ b/shared/source/built_ins/kernels/CMakeLists.txt @@ -0,0 +1,106 @@ +# +# Copyright (C) 2017-2020 Intel Corporation +# +# SPDX-License-Identifier: MIT +# + +add_custom_target(builtins) +set_target_properties(builtins PROPERTIES FOLDER "built_ins") +set(BUILTINS_OUTDIR_WITH_ARCH "${TargetDir}/built_ins/${NEO_ARCH}") +add_dependencies(${BUILTINS_BINARIES_LIB_NAME} builtins) +add_subdirectories() +set(GENERATED_BUILTINS ${GENERATED_BUILTINS} PARENT_SCOPE) +set(GENERATED_BUILTINS_STATELESS ${GENERATED_BUILTINS_STATELESS} PARENT_SCOPE) + +if("${NEO_ARCH}" STREQUAL "x32") + set(BUILTIN_OPTIONS "-cl-intel-greater-than-4GB-buffer-required") +else() + set(BUILTIN_OPTIONS "") +endif() + +set(BUILTIN_OPTIONS_STATELESS + "-cl-intel-greater-than-4GB-buffer-required" +) + +if("${CMAKE_BUILD_TYPE}" STREQUAL "Debug" ) + list(APPEND __cloc__options__ "-D DEBUG") +endif() + +set(BUILTINS_INCLUDE_DIR ${TargetDir} PARENT_SCOPE) +set(BUILTIN_CPP "") + +function(get_bits_for_stateless gen_type platform_type) + # Force 32bits compiling on gen9lp for stateless builtins + if( (${GEN_TYPE} STREQUAL "GEN9" ) AND ( ${PLATFORM_TYPE} STREQUAL "LP")) + set(BITS "32" PARENT_SCOPE) + else() + set(BITS ${NEO_BITS} PARENT_SCOPE) + endif() +endfunction() + +# Define function for compiling built-ins (with ocloc) +function(compile_builtin gen_type platform_type builtin bits builtin_options) + string(TOLOWER ${gen_type} gen_type_lower) + get_family_name_with_type(${gen_type} ${platform_type}) + set(OUTPUTDIR "${BUILTINS_OUTDIR_WITH_ARCH}/${gen_type_lower}") + # get filename + set(FILENAME ${builtin}) + # get name of the file w/o extension + get_filename_component(BASENAME ${builtin} NAME_WE) + + set(OUTPUTPATH_BASE "${OUTPUTDIR}/${BASENAME}_${family_name_with_type}") + set(OUTPUT_FILES + ${OUTPUTPATH_BASE}.spv + ${OUTPUTPATH_BASE}.bin + ${OUTPUTPATH_BASE}.cpp + ${OUTPUTPATH_BASE}.gen + ) + + # function returns builtin cpp filename + unset(BUILTIN_CPP) + # set variable outside function + set(BUILTIN_CPP built_ins/${NEO_ARCH}/${gen_type_lower}/${BASENAME}_${family_name_with_type}.cpp PARENT_SCOPE) + if(WIN32) + set(cloc_cmd_prefix ocloc) + else() + if(DEFINED NEO__IGC_LIBRARY_PATH) + set(cloc_cmd_prefix LD_LIBRARY_PATH=${NEO__IGC_LIBRARY_PATH} $) + else() + set(cloc_cmd_prefix LD_LIBRARY_PATH=$ $) + endif() + endif() + list(APPEND __cloc__options__ "-cl-kernel-arg-info") + add_custom_command( + OUTPUT ${OUTPUT_FILES} + COMMAND ${cloc_cmd_prefix} -q -file ${FILENAME} -device ${DEFAULT_SUPPORTED_${gen_type}_${platform_type}_PLATFORM} ${builtin_options} -${bits} -out_dir ${OUTPUTDIR} -cpp_file -options "$" + WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR} + DEPENDS ${builtin} ocloc copy_compiler_files + ) +endfunction() + +macro(macro_for_each_gen) + foreach(PLATFORM_TYPE ${PLATFORM_TYPES}) + if(${GEN_TYPE}_HAS_${PLATFORM_TYPE}) + get_family_name_with_type(${GEN_TYPE} ${PLATFORM_TYPE}) + string(TOLOWER ${PLATFORM_TYPE} PLATFORM_TYPE_LOWER) + unset(BUILTINS_COMMANDS) + foreach(GENERATED_BUILTIN ${GENERATED_BUILTINS}) + compile_builtin(${GEN_TYPE} ${PLATFORM_TYPE} ${GENERATED_BUILTIN}.builtin_kernel ${NEO_BITS} "${BUILTIN_OPTIONS}") + list(APPEND BUILTINS_COMMANDS ${TargetDir}/${BUILTIN_CPP}) + set(RUNTIME_GENERATED_${GENERATED_BUILTIN}_${family_name_with_type} ${BUILTIN_CPP} PARENT_SCOPE) + endforeach() + get_bits_for_stateless(${GEN_TYPE} ${PLATFORM_TYPE}) + foreach(GENERATED_BUILTIN_STATELESS ${GENERATED_BUILTINS_STATELESS}) + compile_builtin(${GEN_TYPE} ${PLATFORM_TYPE} ${GENERATED_BUILTIN_STATELESS}.builtin_kernel ${BITS} "${BUILTIN_OPTIONS_STATELESS}") + list(APPEND BUILTINS_COMMANDS ${TargetDir}/${BUILTIN_CPP}) + set(RUNTIME_GENERATED_${GENERATED_BUILTIN_STATELESS}_${family_name_with_type} ${BUILTIN_CPP} PARENT_SCOPE) + endforeach() + set(target_name builtins_${family_name_with_type}) + add_custom_target(${target_name} DEPENDS ${BUILTINS_COMMANDS}) + add_dependencies(builtins ${target_name}) + set_target_properties(${target_name} PROPERTIES FOLDER "opencl/source/built_ins/${family_name_with_type}") + endif() + endforeach() +endmacro() + +apply_macro_for_each_gen("SUPPORTED") diff --git a/opencl/source/built_ins/kernels/aux_translation.builtin_kernel b/shared/source/built_ins/kernels/aux_translation.builtin_kernel similarity index 100% rename from opencl/source/built_ins/kernels/aux_translation.builtin_kernel rename to shared/source/built_ins/kernels/aux_translation.builtin_kernel diff --git a/opencl/source/built_ins/kernels/copy_buffer_rect.builtin_kernel b/shared/source/built_ins/kernels/copy_buffer_rect.builtin_kernel similarity index 100% rename from opencl/source/built_ins/kernels/copy_buffer_rect.builtin_kernel rename to shared/source/built_ins/kernels/copy_buffer_rect.builtin_kernel diff --git a/opencl/source/built_ins/kernels/copy_buffer_rect_stateless.builtin_kernel b/shared/source/built_ins/kernels/copy_buffer_rect_stateless.builtin_kernel similarity index 100% rename from opencl/source/built_ins/kernels/copy_buffer_rect_stateless.builtin_kernel rename to shared/source/built_ins/kernels/copy_buffer_rect_stateless.builtin_kernel diff --git a/shared/source/built_ins/kernels/copy_buffer_to_buffer.builtin_kernel b/shared/source/built_ins/kernels/copy_buffer_to_buffer.builtin_kernel new file mode 100644 index 0000000000..7d47bb10d2 --- /dev/null +++ b/shared/source/built_ins/kernels/copy_buffer_to_buffer.builtin_kernel @@ -0,0 +1,90 @@ +/* + * Copyright (C) 2017-2020 Intel Corporation + * + * SPDX-License-Identifier: MIT + * + */ + +R"===( +__kernel void CopyBufferToBufferBytes( + const __global uchar* pSrc, + __global uchar* pDst, + uint srcOffsetInBytes, + uint dstOffsetInBytes, + uint bytesToRead ) +{ + pSrc += ( srcOffsetInBytes + get_global_id(0) ); + pDst += ( dstOffsetInBytes + get_global_id(0) ); + pDst[ 0 ] = pSrc[ 0 ]; +} + +__kernel void CopyBufferToBufferLeftLeftover( + const __global uchar* pSrc, + __global uchar* pDst, + uint srcOffsetInBytes, + uint dstOffsetInBytes) +{ + unsigned int gid = get_global_id(0); + pDst[ gid + dstOffsetInBytes ] = pSrc[ gid + srcOffsetInBytes ]; +} + +__kernel void CopyBufferToBufferMiddle( + const __global uint* pSrc, + __global uint* pDst, + uint srcOffsetInBytes, + uint dstOffsetInBytes) +{ + unsigned int gid = get_global_id(0); + pDst += dstOffsetInBytes >> 2; + pSrc += srcOffsetInBytes >> 2; + uint4 loaded = vload4(gid, pSrc); + vstore4(loaded, gid, pDst); +} + +__kernel void CopyBufferToBufferRightLeftover( + const __global uchar* pSrc, + __global uchar* pDst, + uint srcOffsetInBytes, + uint dstOffsetInBytes) +{ + unsigned int gid = get_global_id(0); + pDst[ gid + dstOffsetInBytes ] = pSrc[ gid + srcOffsetInBytes ]; +} + +__kernel void copyBufferToBufferBytesSingle(__global uchar *dst, const __global uchar *src) { + unsigned int gid = get_global_id(0); + dst[gid] = (uchar)(src[gid]); +} +__kernel void CopyBufferToBufferSideRegion( + __global uchar* pDst, + const __global uchar* pSrc, + unsigned int len, + uint dstSshOffset, // Offset needed in case ptr has been adjusted for SSH alignment + uint srcSshOffset // Offset needed in case ptr has been adjusted for SSH alignment + ) +{ + unsigned int gid = get_global_id(0); + __global uchar* pDstWithOffset = (__global uchar*)((__global uchar*)pDst + dstSshOffset); + __global uchar* pSrcWithOffset = (__global uchar*)((__global uchar*)pSrc + srcSshOffset); + if (gid < len) { + pDstWithOffset[ gid ] = pSrcWithOffset[ gid ]; + } +} + +__kernel void CopyBufferToBufferMiddleRegion( + __global uint* pDst, + const __global uint* pSrc, + unsigned int elems, + uint dstSshOffset, // Offset needed in case ptr has been adjusted for SSH alignment + uint srcSshOffset // Offset needed in case ptr has been adjusted for SSH alignment + ) +{ + unsigned int gid = get_global_id(0); + __global uint* pDstWithOffset = (__global uint*)((__global uchar*)pDst + dstSshOffset); + __global uint* pSrcWithOffset = (__global uint*)((__global uchar*)pSrc + srcSshOffset); + if (gid < elems) { + uint4 loaded = vload4(gid, pSrcWithOffset); + vstore4(loaded, gid, pDstWithOffset); + } +} +)===" \ No newline at end of file diff --git a/opencl/source/built_ins/kernels/copy_buffer_to_buffer_stateless.builtin_kernel b/shared/source/built_ins/kernels/copy_buffer_to_buffer_stateless.builtin_kernel similarity index 100% rename from opencl/source/built_ins/kernels/copy_buffer_to_buffer_stateless.builtin_kernel rename to shared/source/built_ins/kernels/copy_buffer_to_buffer_stateless.builtin_kernel diff --git a/opencl/source/built_ins/kernels/copy_buffer_to_image3d.builtin_kernel b/shared/source/built_ins/kernels/copy_buffer_to_image3d.builtin_kernel similarity index 100% rename from opencl/source/built_ins/kernels/copy_buffer_to_image3d.builtin_kernel rename to shared/source/built_ins/kernels/copy_buffer_to_image3d.builtin_kernel diff --git a/opencl/source/built_ins/kernels/copy_buffer_to_image3d_stateless.builtin_kernel b/shared/source/built_ins/kernels/copy_buffer_to_image3d_stateless.builtin_kernel similarity index 100% rename from opencl/source/built_ins/kernels/copy_buffer_to_image3d_stateless.builtin_kernel rename to shared/source/built_ins/kernels/copy_buffer_to_image3d_stateless.builtin_kernel diff --git a/opencl/source/built_ins/kernels/copy_image3d_to_buffer.builtin_kernel b/shared/source/built_ins/kernels/copy_image3d_to_buffer.builtin_kernel similarity index 100% rename from opencl/source/built_ins/kernels/copy_image3d_to_buffer.builtin_kernel rename to shared/source/built_ins/kernels/copy_image3d_to_buffer.builtin_kernel diff --git a/opencl/source/built_ins/kernels/copy_image3d_to_buffer_stateless.builtin_kernel b/shared/source/built_ins/kernels/copy_image3d_to_buffer_stateless.builtin_kernel similarity index 100% rename from opencl/source/built_ins/kernels/copy_image3d_to_buffer_stateless.builtin_kernel rename to shared/source/built_ins/kernels/copy_image3d_to_buffer_stateless.builtin_kernel diff --git a/opencl/source/built_ins/kernels/copy_image_to_image1d.builtin_kernel b/shared/source/built_ins/kernels/copy_image_to_image1d.builtin_kernel similarity index 100% rename from opencl/source/built_ins/kernels/copy_image_to_image1d.builtin_kernel rename to shared/source/built_ins/kernels/copy_image_to_image1d.builtin_kernel diff --git a/opencl/source/built_ins/kernels/copy_image_to_image2d.builtin_kernel b/shared/source/built_ins/kernels/copy_image_to_image2d.builtin_kernel similarity index 100% rename from opencl/source/built_ins/kernels/copy_image_to_image2d.builtin_kernel rename to shared/source/built_ins/kernels/copy_image_to_image2d.builtin_kernel diff --git a/opencl/source/built_ins/kernels/copy_image_to_image3d.builtin_kernel b/shared/source/built_ins/kernels/copy_image_to_image3d.builtin_kernel similarity index 100% rename from opencl/source/built_ins/kernels/copy_image_to_image3d.builtin_kernel rename to shared/source/built_ins/kernels/copy_image_to_image3d.builtin_kernel diff --git a/opencl/source/built_ins/kernels/fill_buffer.builtin_kernel b/shared/source/built_ins/kernels/fill_buffer.builtin_kernel similarity index 59% rename from opencl/source/built_ins/kernels/fill_buffer.builtin_kernel rename to shared/source/built_ins/kernels/fill_buffer.builtin_kernel index 1b09e46bdf..aa45973e14 100644 --- a/opencl/source/built_ins/kernels/fill_buffer.builtin_kernel +++ b/shared/source/built_ins/kernels/fill_buffer.builtin_kernel @@ -46,4 +46,28 @@ __kernel void FillBufferRightLeftover( uint gid = get_global_id(0); pDst[ gid + dstOffsetInBytes ] = pPattern[ gid & (patternSizeInEls - 1) ]; } + +__kernel void FillBufferImmediate( + __global uchar* ptr, + uint dstSshOffset, // Offset needed in case ptr has been adjusted for SSH alignment + const uint value) +{ + uint dstIndex = get_global_id(0); + __global uchar* pDst = (__global uchar*)ptr + dstSshOffset; + pDst[dstIndex] = value; +} + +__kernel void FillBufferSSHOffset( + __global uchar* ptr, + uint dstSshOffset, // Offset needed in case ptr has been adjusted for SSH alignment + const __global uchar* pPattern, + uint patternSshOffset // Offset needed in case pPattern has been adjusted for SSH alignment +) +{ + uint dstIndex = get_global_id(0); + uint srcIndex = get_local_id(0); + __global uchar* pDst = (__global uchar*)ptr + dstSshOffset; + __global uchar* pSrc = (__global uchar*)pPattern + patternSshOffset; + pDst[dstIndex] = pSrc[srcIndex]; +} )===" \ No newline at end of file diff --git a/opencl/source/built_ins/kernels/fill_buffer_stateless.builtin_kernel b/shared/source/built_ins/kernels/fill_buffer_stateless.builtin_kernel similarity index 100% rename from opencl/source/built_ins/kernels/fill_buffer_stateless.builtin_kernel rename to shared/source/built_ins/kernels/fill_buffer_stateless.builtin_kernel diff --git a/opencl/source/built_ins/kernels/fill_image1d.builtin_kernel b/shared/source/built_ins/kernels/fill_image1d.builtin_kernel similarity index 100% rename from opencl/source/built_ins/kernels/fill_image1d.builtin_kernel rename to shared/source/built_ins/kernels/fill_image1d.builtin_kernel diff --git a/opencl/source/built_ins/kernels/fill_image2d.builtin_kernel b/shared/source/built_ins/kernels/fill_image2d.builtin_kernel similarity index 100% rename from opencl/source/built_ins/kernels/fill_image2d.builtin_kernel rename to shared/source/built_ins/kernels/fill_image2d.builtin_kernel diff --git a/opencl/source/built_ins/kernels/fill_image3d.builtin_kernel b/shared/source/built_ins/kernels/fill_image3d.builtin_kernel similarity index 100% rename from opencl/source/built_ins/kernels/fill_image3d.builtin_kernel rename to shared/source/built_ins/kernels/fill_image3d.builtin_kernel diff --git a/shared/source/built_ins/registry/CMakeLists.txt b/shared/source/built_ins/registry/CMakeLists.txt new file mode 100644 index 0000000000..d48e5233df --- /dev/null +++ b/shared/source/built_ins/registry/CMakeLists.txt @@ -0,0 +1,25 @@ +# +# Copyright (C) 2017-2020 Intel Corporation +# +# SPDX-License-Identifier: MIT +# + +add_library(${BUILTINS_SOURCES_LIB_NAME} OBJECT EXCLUDE_FROM_ALL + CMakeLists.txt + built_ins_registry.h + register_copy_kernels_source.cpp +) + +set_target_properties(${BUILTINS_SOURCES_LIB_NAME} PROPERTIES POSITION_INDEPENDENT_CODE ON) +set_target_properties(${BUILTINS_SOURCES_LIB_NAME} PROPERTIES FOLDER "built_ins") + +target_include_directories(${BUILTINS_SOURCES_LIB_NAME} PRIVATE + ${ENGINE_NODE_DIR} + ${KHRONOS_HEADERS_DIR} + ${KHRONOS_GL_HEADERS_DIR} + ${NEO__GMM_INCLUDE_DIR} + ${NEO__IGC_INCLUDE_DIR} + ${THIRD_PARTY_DIR} +) + +add_subdirectories() diff --git a/opencl/source/built_ins/registry/built_ins_registry.h b/shared/source/built_ins/registry/built_ins_registry.h similarity index 93% rename from opencl/source/built_ins/registry/built_ins_registry.h rename to shared/source/built_ins/registry/built_ins_registry.h index 3d601b7a4c..4a7171d200 100644 --- a/opencl/source/built_ins/registry/built_ins_registry.h +++ b/shared/source/built_ins/registry/built_ins_registry.h @@ -7,7 +7,7 @@ #pragma once -#include "opencl/source/built_ins/built_ins.h" +#include "shared/source/built_ins/built_ins.h" #include #include diff --git a/opencl/source/built_ins/registry/register_copy_kernels_source.cpp b/shared/source/built_ins/registry/register_copy_kernels_source.cpp similarity index 80% rename from opencl/source/built_ins/registry/register_copy_kernels_source.cpp rename to shared/source/built_ins/registry/register_copy_kernels_source.cpp index f61e524e43..58c1f791fa 100644 --- a/opencl/source/built_ins/registry/register_copy_kernels_source.cpp +++ b/shared/source/built_ins/registry/register_copy_kernels_source.cpp @@ -5,7 +5,7 @@ * */ -#include "opencl/source/built_ins/registry/built_ins_registry.h" +#include "shared/source/built_ins/registry/built_ins_registry.h" #include @@ -17,7 +17,7 @@ static RegisterEmbeddedResource registerCopyBufferToBufferSrc( BuiltinCode::getExtension(BuiltinCode::ECodeType::Source)) .c_str(), std::string( -#include "opencl/source/built_ins/kernels/copy_buffer_to_buffer.builtin_kernel" +#include "shared/source/built_ins/kernels/copy_buffer_to_buffer.builtin_kernel" )); static RegisterEmbeddedResource registerCopyBufferToBufferStatelessSrc( @@ -26,7 +26,7 @@ static RegisterEmbeddedResource registerCopyBufferToBufferStatelessSrc( BuiltinCode::getExtension(BuiltinCode::ECodeType::Source)) .c_str(), std::string( -#include "opencl/source/built_ins/kernels/copy_buffer_to_buffer_stateless.builtin_kernel" +#include "shared/source/built_ins/kernels/copy_buffer_to_buffer_stateless.builtin_kernel" )); static RegisterEmbeddedResource registerCopyBufferRectSrc( @@ -35,7 +35,7 @@ static RegisterEmbeddedResource registerCopyBufferRectSrc( BuiltinCode::getExtension(BuiltinCode::ECodeType::Source)) .c_str(), std::string( -#include "opencl/source/built_ins/kernels/copy_buffer_rect.builtin_kernel" +#include "shared/source/built_ins/kernels/copy_buffer_rect.builtin_kernel" )); static RegisterEmbeddedResource registerCopyBufferRectStatelessSrc( @@ -44,7 +44,7 @@ static RegisterEmbeddedResource registerCopyBufferRectStatelessSrc( BuiltinCode::getExtension(BuiltinCode::ECodeType::Source)) .c_str(), std::string( -#include "opencl/source/built_ins/kernels/copy_buffer_rect_stateless.builtin_kernel" +#include "shared/source/built_ins/kernels/copy_buffer_rect_stateless.builtin_kernel" )); static RegisterEmbeddedResource registerFillBufferSrc( @@ -53,7 +53,7 @@ static RegisterEmbeddedResource registerFillBufferSrc( BuiltinCode::getExtension(BuiltinCode::ECodeType::Source)) .c_str(), std::string( -#include "opencl/source/built_ins/kernels/fill_buffer.builtin_kernel" +#include "shared/source/built_ins/kernels/fill_buffer.builtin_kernel" )); static RegisterEmbeddedResource registerFillBufferStatelessSrc( @@ -62,7 +62,7 @@ static RegisterEmbeddedResource registerFillBufferStatelessSrc( BuiltinCode::getExtension(BuiltinCode::ECodeType::Source)) .c_str(), std::string( -#include "opencl/source/built_ins/kernels/fill_buffer_stateless.builtin_kernel" +#include "shared/source/built_ins/kernels/fill_buffer_stateless.builtin_kernel" )); static RegisterEmbeddedResource registerCopyBufferToImage3dSrc( @@ -71,7 +71,7 @@ static RegisterEmbeddedResource registerCopyBufferToImage3dSrc( BuiltinCode::getExtension(BuiltinCode::ECodeType::Source)) .c_str(), std::string( -#include "opencl/source/built_ins/kernels/copy_buffer_to_image3d.builtin_kernel" +#include "shared/source/built_ins/kernels/copy_buffer_to_image3d.builtin_kernel" )); static RegisterEmbeddedResource registerCopyBufferToImage3dStatelessSrc( @@ -80,7 +80,7 @@ static RegisterEmbeddedResource registerCopyBufferToImage3dStatelessSrc( BuiltinCode::getExtension(BuiltinCode::ECodeType::Source)) .c_str(), std::string( -#include "opencl/source/built_ins/kernels/copy_buffer_to_image3d_stateless.builtin_kernel" +#include "shared/source/built_ins/kernels/copy_buffer_to_image3d_stateless.builtin_kernel" )); static RegisterEmbeddedResource registerCopyImage3dToBufferSrc( @@ -89,7 +89,7 @@ static RegisterEmbeddedResource registerCopyImage3dToBufferSrc( BuiltinCode::getExtension(BuiltinCode::ECodeType::Source)) .c_str(), std::string( -#include "opencl/source/built_ins/kernels/copy_image3d_to_buffer.builtin_kernel" +#include "shared/source/built_ins/kernels/copy_image3d_to_buffer.builtin_kernel" )); static RegisterEmbeddedResource registerCopyImage3dToBufferStatelessSrc( @@ -98,7 +98,7 @@ static RegisterEmbeddedResource registerCopyImage3dToBufferStatelessSrc( BuiltinCode::getExtension(BuiltinCode::ECodeType::Source)) .c_str(), std::string( -#include "opencl/source/built_ins/kernels/copy_image3d_to_buffer_stateless.builtin_kernel" +#include "shared/source/built_ins/kernels/copy_image3d_to_buffer_stateless.builtin_kernel" )); static RegisterEmbeddedResource registerCopyImageToImage1dSrc( @@ -107,7 +107,7 @@ static RegisterEmbeddedResource registerCopyImageToImage1dSrc( BuiltinCode::getExtension(BuiltinCode::ECodeType::Source)) .c_str(), std::string( -#include "opencl/source/built_ins/kernels/copy_image_to_image1d.builtin_kernel" +#include "shared/source/built_ins/kernels/copy_image_to_image1d.builtin_kernel" )); static RegisterEmbeddedResource registerCopyImageToImage2dSrc( @@ -116,7 +116,7 @@ static RegisterEmbeddedResource registerCopyImageToImage2dSrc( BuiltinCode::getExtension(BuiltinCode::ECodeType::Source)) .c_str(), std::string( -#include "opencl/source/built_ins/kernels/copy_image_to_image2d.builtin_kernel" +#include "shared/source/built_ins/kernels/copy_image_to_image2d.builtin_kernel" )); static RegisterEmbeddedResource registerCopyImageToImage3dSrc( @@ -125,7 +125,7 @@ static RegisterEmbeddedResource registerCopyImageToImage3dSrc( BuiltinCode::getExtension(BuiltinCode::ECodeType::Source)) .c_str(), std::string( -#include "opencl/source/built_ins/kernels/copy_image_to_image3d.builtin_kernel" +#include "shared/source/built_ins/kernels/copy_image_to_image3d.builtin_kernel" )); static RegisterEmbeddedResource registerFillImage1dSrc( @@ -134,7 +134,7 @@ static RegisterEmbeddedResource registerFillImage1dSrc( BuiltinCode::getExtension(BuiltinCode::ECodeType::Source)) .c_str(), std::string( -#include "opencl/source/built_ins/kernels/fill_image1d.builtin_kernel" +#include "shared/source/built_ins/kernels/fill_image1d.builtin_kernel" )); static RegisterEmbeddedResource registerFillImage2dSrc( @@ -143,7 +143,7 @@ static RegisterEmbeddedResource registerFillImage2dSrc( BuiltinCode::getExtension(BuiltinCode::ECodeType::Source)) .c_str(), std::string( -#include "opencl/source/built_ins/kernels/fill_image2d.builtin_kernel" +#include "shared/source/built_ins/kernels/fill_image2d.builtin_kernel" )); static RegisterEmbeddedResource registerFillImage3dSrc( @@ -152,7 +152,7 @@ static RegisterEmbeddedResource registerFillImage3dSrc( BuiltinCode::getExtension(BuiltinCode::ECodeType::Source)) .c_str(), std::string( -#include "opencl/source/built_ins/kernels/fill_image3d.builtin_kernel" +#include "shared/source/built_ins/kernels/fill_image3d.builtin_kernel" )); static RegisterEmbeddedResource registerAuxTranslationSrc( @@ -161,7 +161,7 @@ static RegisterEmbeddedResource registerAuxTranslationSrc( BuiltinCode::getExtension(BuiltinCode::ECodeType::Source)) .c_str(), std::string( -#include "opencl/source/built_ins/kernels/aux_translation.builtin_kernel" +#include "shared/source/built_ins/kernels/aux_translation.builtin_kernel" )); } // namespace NEO diff --git a/opencl/source/built_ins/sip.cpp b/shared/source/built_ins/sip.cpp similarity index 97% rename from opencl/source/built_ins/sip.cpp rename to shared/source/built_ins/sip.cpp index 9822b9d560..103fb4078a 100644 --- a/opencl/source/built_ins/sip.cpp +++ b/shared/source/built_ins/sip.cpp @@ -5,8 +5,9 @@ * */ -#include "opencl/source/built_ins/sip.h" +#include "shared/source/built_ins/sip.h" +#include "shared/source/built_ins/built_ins.h" #include "shared/source/device/device.h" #include "shared/source/execution_environment/execution_environment.h" #include "shared/source/helpers/debug_helpers.h" @@ -15,7 +16,6 @@ #include "shared/source/helpers/string.h" #include "shared/source/memory_manager/graphics_allocation.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/program/kernel_info.h" #include "opencl/source/program/program.h" diff --git a/opencl/source/built_ins/sip.h b/shared/source/built_ins/sip.h similarity index 100% rename from opencl/source/built_ins/sip.h rename to shared/source/built_ins/sip.h diff --git a/shared/source/command_stream/command_stream_receiver.cpp b/shared/source/command_stream/command_stream_receiver.cpp index a3e90c8780..408713a952 100644 --- a/shared/source/command_stream/command_stream_receiver.cpp +++ b/shared/source/command_stream/command_stream_receiver.cpp @@ -7,6 +7,7 @@ #include "shared/source/command_stream/command_stream_receiver.h" +#include "shared/source/built_ins/built_ins.h" #include "shared/source/command_stream/experimental_command_buffer.h" #include "shared/source/command_stream/preemption.h" #include "shared/source/command_stream/scratch_space_controller.h" @@ -26,8 +27,6 @@ #include "shared/source/utilities/cpuintrinsics.h" #include "shared/source/utilities/tag_allocator.h" -#include "opencl/source/built_ins/built_ins.h" - namespace NEO { // Global table of CommandStreamReceiver factories for HW and tests diff --git a/shared/source/command_stream/preemption.cpp b/shared/source/command_stream/preemption.cpp index 61848f5225..7c887f7d94 100644 --- a/shared/source/command_stream/preemption.cpp +++ b/shared/source/command_stream/preemption.cpp @@ -7,10 +7,10 @@ #include "shared/source/command_stream/preemption.h" +#include "shared/source/built_ins/built_ins.h" #include "shared/source/device/device.h" #include "shared/source/helpers/string.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/helpers/dispatch_info.h" #include "opencl/source/kernel/kernel.h" diff --git a/shared/source/command_stream/preemption.inl b/shared/source/command_stream/preemption.inl index 49530304f0..a58950bae2 100644 --- a/shared/source/command_stream/preemption.inl +++ b/shared/source/command_stream/preemption.inl @@ -5,12 +5,12 @@ * */ +#include "shared/source/built_ins/sip.h" #include "shared/source/command_stream/preemption.h" #include "shared/source/device/device.h" #include "shared/source/helpers/hw_helper.h" #include "shared/source/memory_manager/graphics_allocation.h" -#include "opencl/source/built_ins/sip.h" #include "opencl/source/command_queue/gpgpu_walker.h" namespace NEO { diff --git a/shared/source/compiler_interface/compiler_interface.h b/shared/source/compiler_interface/compiler_interface.h index 5d065ff423..6f721b628e 100644 --- a/shared/source/compiler_interface/compiler_interface.h +++ b/shared/source/compiler_interface/compiler_interface.h @@ -6,14 +6,13 @@ */ #pragma once +#include "shared/source/built_ins/sip.h" #include "shared/source/compiler_interface/compiler_cache.h" #include "shared/source/helpers/string.h" #include "shared/source/os_interface/os_library.h" #include "shared/source/utilities/arrayref.h" #include "shared/source/utilities/spinlock.h" -#include "opencl/source/built_ins/sip.h" - #include "cif/common/cif_main.h" #include "ocl_igc_interface/code_type.h" #include "ocl_igc_interface/fcl_ocl_device_ctx.h" diff --git a/shared/source/execution_environment/execution_environment.cpp b/shared/source/execution_environment/execution_environment.cpp index d3c0fa748c..ef574f403e 100644 --- a/shared/source/execution_environment/execution_environment.cpp +++ b/shared/source/execution_environment/execution_environment.cpp @@ -7,6 +7,7 @@ #include "shared/source/execution_environment/execution_environment.h" +#include "shared/source/built_ins/built_ins.h" #include "shared/source/compiler_interface/compiler_interface.h" #include "shared/source/compiler_interface/default_cache_config.h" #include "shared/source/debugger/debugger.h" @@ -14,7 +15,6 @@ #include "shared/source/gmm_helper/gmm_helper.h" #include "shared/source/helpers/hw_helper.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/source/memory_manager/os_agnostic_memory_manager.h" namespace NEO { diff --git a/shared/source/gen9/preemption_gen9.cpp b/shared/source/gen9/preemption_gen9.cpp index d2981d7bfe..f8ace1a206 100644 --- a/shared/source/gen9/preemption_gen9.cpp +++ b/shared/source/gen9/preemption_gen9.cpp @@ -5,12 +5,11 @@ * */ +#include "shared/source/built_ins/built_ins.h" #include "shared/source/command_stream/csr_definitions.h" #include "shared/source/command_stream/preemption.h" #include "shared/source/command_stream/preemption.inl" -#include "opencl/source/built_ins/built_ins.h" - #include namespace NEO { diff --git a/shared/source/helpers/hw_helper.h b/shared/source/helpers/hw_helper.h index f1095f2805..7890fbc35a 100644 --- a/shared/source/helpers/hw_helper.h +++ b/shared/source/helpers/hw_helper.h @@ -6,11 +6,11 @@ */ #pragma once +#include "shared/source/built_ins/sip.h" #include "shared/source/command_stream/linear_stream.h" #include "shared/source/helpers/aux_translation.h" #include "shared/source/helpers/hw_cmds.h" -#include "opencl/source/built_ins/sip.h" #include "opencl/source/gen_common/aub_mapper.h" #include "opencl/source/mem_obj/buffer.h" diff --git a/shared/test/unit_test/gen11/test_preemption_gen11.cpp b/shared/test/unit_test/gen11/test_preemption_gen11.cpp index d072dcc270..d6e9dcc4bd 100644 --- a/shared/test/unit_test/gen11/test_preemption_gen11.cpp +++ b/shared/test/unit_test/gen11/test_preemption_gen11.cpp @@ -5,10 +5,10 @@ * */ +#include "shared/source/built_ins/built_ins.h" #include "shared/source/helpers/hw_helper.h" #include "shared/test/unit_test/fixtures/preemption_fixture.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/test/unit_test/helpers/hw_parse.h" #include "opencl/test/unit_test/mocks/mock_buffer.h" #include "opencl/test/unit_test/mocks/mock_command_queue.h" diff --git a/shared/test/unit_test/gen9/test_preemption_gen9.cpp b/shared/test/unit_test/gen9/test_preemption_gen9.cpp index a8d724773b..f0d0732b87 100644 --- a/shared/test/unit_test/gen9/test_preemption_gen9.cpp +++ b/shared/test/unit_test/gen9/test_preemption_gen9.cpp @@ -5,11 +5,11 @@ * */ +#include "shared/source/built_ins/built_ins.h" #include "shared/source/command_stream/preemption.h" #include "shared/source/helpers/hw_helper.h" #include "shared/test/unit_test/fixtures/preemption_fixture.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/test/unit_test/command_queue/enqueue_fixture.h" #include "opencl/test/unit_test/helpers/hw_parse.h" #include "opencl/test/unit_test/mocks/mock_buffer.h" diff --git a/shared/test/unit_test/source_level_debugger/source_level_debugger_preamble_test.h b/shared/test/unit_test/source_level_debugger/source_level_debugger_preamble_test.h index d4258fd950..4859768051 100644 --- a/shared/test/unit_test/source_level_debugger/source_level_debugger_preamble_test.h +++ b/shared/test/unit_test/source_level_debugger/source_level_debugger_preamble_test.h @@ -6,10 +6,10 @@ */ #pragma once +#include "shared/source/built_ins/built_ins.h" #include "shared/source/command_stream/preemption.h" #include "shared/source/helpers/preamble.h" -#include "opencl/source/built_ins/built_ins.h" #include "opencl/test/unit_test/helpers/hw_parse.h" #include "opencl/test/unit_test/mocks/mock_device.h" #include "opencl/test/unit_test/mocks/mock_graphics_allocation.h"