diff --git a/buildbot/dependency.py b/buildbot/dependency.py index 8697dbfb0f991..e79eae2b62cc1 100644 --- a/buildbot/dependency.py +++ b/buildbot/dependency.py @@ -58,7 +58,7 @@ def do_dependency(args): # Checkout fixed version to avoid unexpected issues coming from upstream # Specific version can be uplifted as soon as such need arise - checkout_cmd = ["git", "checkout", "23710f1b99186065c1768fc3098ba681adc0f253"] + checkout_cmd = ["git", "checkout", "9ddb236e6eb3cf844f9e2f81677e1045f9bf838e"] subprocess.check_call(checkout_cmd, cwd=ocl_header_dir) # fetch and build OpenCL ICD loader diff --git a/llvm/include/llvm/SYCLLowerIR/HostPipes.h b/llvm/include/llvm/SYCLLowerIR/HostPipes.h new file mode 100644 index 0000000000000..9942a99187c3b --- /dev/null +++ b/llvm/include/llvm/SYCLLowerIR/HostPipes.h @@ -0,0 +1,59 @@ +//===------- HostPipes.h - get required info about FPGA Host Pipes --------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// The file contains a number of functions to extract corresponding attributes +// of the host pipe global variables and save them as a property set for the +// runtime. +//===----------------------------------------------------------------------===// + +#pragma once + +#include "llvm/ADT/MapVector.h" + +#include +#include + +namespace llvm { + +class GlobalVariable; +class Module; +class StringRef; + +// Represents a host pipe variable - at SYCL RT level host pipe +// variables are being represented as a byte-array. +struct HostPipeProperty { + HostPipeProperty(uint32_t Size) : Size(Size) {} + + // Encodes size of the underlying type T of the host pipe variable. + uint32_t Size; +}; + +using HostPipePropertyMapTy = + MapVector>; + +/// Return \c true if the variable @GV is a host pipe variable. +/// +/// The function checks whether the variable has the LLVM IR attribute \c +/// sycl-host-pipe +/// @param GV [in] A variable to test. +/// +/// @return \c true if the variable is a host pipe variable, \c false +/// otherwise. +bool isHostPipeVariable(const GlobalVariable &GV); + +/// Searches given module for occurrences of host pipe variable-specific +/// metadata and builds "host pipe variable name" -> +/// vector<"variable properties"> map. +/// +/// @param M [in] LLVM Module. +/// +/// @returns the "host pipe variable name" -> vector<"variable properties"> +/// map. +HostPipePropertyMapTy collectHostPipeProperties(const Module &M); + +} // end namespace llvm diff --git a/llvm/include/llvm/SYCLLowerIR/SYCLUtils.h b/llvm/include/llvm/SYCLLowerIR/SYCLUtils.h index 2d82e651b35fb..241859a47426d 100644 --- a/llvm/include/llvm/SYCLLowerIR/SYCLUtils.h +++ b/llvm/include/llvm/SYCLLowerIR/SYCLUtils.h @@ -22,7 +22,6 @@ namespace llvm { namespace sycl { namespace utils { constexpr char ATTR_SYCL_MODULE_ID[] = "sycl-module-id"; -constexpr StringRef SYCL_HOST_PIPE_ATTR = "sycl-host-pipe"; using CallGraphNodeAction = ::std::function; using CallGraphFunctionFilter = @@ -117,10 +116,6 @@ inline bool isSYCLExternalFunction(const Function *F) { return F->hasFnAttribute(ATTR_SYCL_MODULE_ID); } -inline bool isHostPipeVariable(const GlobalVariable &GV) { - return GV.hasAttribute(SYCL_HOST_PIPE_ATTR); -} - } // namespace utils } // namespace sycl } // namespace llvm diff --git a/llvm/include/llvm/Support/PropertySetIO.h b/llvm/include/llvm/Support/PropertySetIO.h index 95dfa0190caf9..d878065809692 100644 --- a/llvm/include/llvm/Support/PropertySetIO.h +++ b/llvm/include/llvm/Support/PropertySetIO.h @@ -197,6 +197,7 @@ class PropertySetRegistry { static constexpr char SYCL_EXPORTED_SYMBOLS[] = "SYCL/exported symbols"; static constexpr char SYCL_DEVICE_GLOBALS[] = "SYCL/device globals"; static constexpr char SYCL_DEVICE_REQUIREMENTS[] = "SYCL/device requirements"; + static constexpr char SYCL_HOST_PIPES[] = "SYCL/host pipes"; // Function for bulk addition of an entire property set under given category // (property set name). diff --git a/llvm/lib/SYCLLowerIR/CMakeLists.txt b/llvm/lib/SYCLLowerIR/CMakeLists.txt index b8f6ab50d15b5..3fe3bb06cc69d 100644 --- a/llvm/lib/SYCLLowerIR/CMakeLists.txt +++ b/llvm/lib/SYCLLowerIR/CMakeLists.txt @@ -57,6 +57,7 @@ add_llvm_component_library(LLVMSYCLLowerIR ESIMD/LowerESIMDVecArg.cpp ESIMD/LowerESIMDVLoadVStore.cpp ESIMD/LowerESIMDSlmReservation.cpp + HostPipes.cpp LowerInvokeSimd.cpp LowerKernelProps.cpp LowerWGLocalMemory.cpp diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp index e24d9ef236390..a652f3cfb0608 100644 --- a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp +++ b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp @@ -10,7 +10,7 @@ #include "llvm/SYCLLowerIR/CompileTimePropertiesPass.h" #include "llvm/SYCLLowerIR/DeviceGlobals.h" -#include "llvm/SYCLLowerIR/SYCLUtils.h" +#include "llvm/SYCLLowerIR/HostPipes.h" #include "llvm/ADT/APInt.h" #include "llvm/ADT/StringMap.h" @@ -343,7 +343,7 @@ PreservedAnalyses CompileTimePropertiesPass::run(Module &M, HostAccessDecorValue, VarName)); } - if (sycl::utils::isHostPipeVariable(GV)) { + if (isHostPipeVariable(GV)) { auto VarName = getGlobalVariableUniqueId(GV); MDOps.push_back(buildSpirvDecorMetadata(Ctx, SPIRV_HOST_ACCESS_DECOR, SPIRV_HOST_ACCESS_DEFAULT_VALUE, diff --git a/llvm/lib/SYCLLowerIR/DeviceGlobals.cpp b/llvm/lib/SYCLLowerIR/DeviceGlobals.cpp index b28f0d63f8cbd..2dd2bb0560552 100644 --- a/llvm/lib/SYCLLowerIR/DeviceGlobals.cpp +++ b/llvm/lib/SYCLLowerIR/DeviceGlobals.cpp @@ -72,19 +72,19 @@ bool hasDeviceImageScopeProperty(const GlobalVariable &GV) { return hasProperty(GV, SYCL_DEVICE_IMAGE_SCOPE_ATTR); } -/// Returns the unique id for the device global variable. +/// Returns the unique id for the device global or host pipe variable. /// /// The function gets this value from the LLVM IR attribute \c /// sycl-unique-id. /// -/// @param GV [in] Device Global variable. +/// @param GV [in] Device Global or Hostpipe variable. /// -/// @returns the unique id of the device global variable represented -/// in the LLVM IR by \c GV. +/// @returns the unique id of the device global or hostpipe variable +/// represented in the LLVM IR by \c GV. StringRef getGlobalVariableUniqueId(const GlobalVariable &GV) { assert(GV.hasAttribute(SYCL_UNIQUE_ID_ATTR) && "a 'sycl-unique-id' string must be associated with every device " - "global variable"); + "global or hostpipe variable"); return GV.getAttribute(SYCL_UNIQUE_ID_ATTR).getValueAsString(); } diff --git a/llvm/lib/SYCLLowerIR/HostPipes.cpp b/llvm/lib/SYCLLowerIR/HostPipes.cpp new file mode 100644 index 0000000000000..ca116c702c1d4 --- /dev/null +++ b/llvm/lib/SYCLLowerIR/HostPipes.cpp @@ -0,0 +1,80 @@ +//===------------- HostPipes.cpp - SYCL Host Pipes Pass -------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// See comments in the header. +//===----------------------------------------------------------------------===// + +#include "llvm/SYCLLowerIR/HostPipes.h" +#include "llvm/SYCLLowerIR/CompileTimePropertiesPass.h" +#include "llvm/SYCLLowerIR/DeviceGlobals.h" + +#include "llvm/ADT/STLExtras.h" +#include "llvm/ADT/StringRef.h" +#include "llvm/IR/Module.h" + +#include + +using namespace llvm; + +namespace { + +constexpr StringRef SYCL_HOST_PIPE_ATTR = "sycl-host-pipe"; +constexpr StringRef SYCL_HOST_PIPE_SIZE_ATTR = "sycl-host-pipe-size"; + +/// Returns the size (in bytes) of the type \c T of the host +/// pipe variable. +/// +/// The function gets this value from the LLVM IR attribute \c +/// sycl-host-pipe-size. +/// +/// @param GV [in] Host Pipe variable. +/// +/// @returns the size (int bytes) of the underlying type \c T of the +/// host pipe variable represented in the LLVM IR by @GV. +uint32_t getHostPipeTypeSize(const GlobalVariable &GV) { + assert(GV.hasAttribute(SYCL_HOST_PIPE_SIZE_ATTR) && + "The host pipe variable must have the 'sycl-host-pipe-size' " + "attribute that must contain a number representing the size of the " + "underlying type T of the host pipe variable"); + return getAttributeAsInteger(GV, SYCL_HOST_PIPE_SIZE_ATTR); +} + +} // anonymous namespace + +namespace llvm { + +/// Return \c true if the variable @GV is a host pipe variable. +/// +/// The function checks whether the variable has the LLVM IR attribute \c +/// sycl-host-pipe. +/// @param GV [in] A variable to test. +/// +/// @return \c true if the variable is a host pipe variable, \c false +/// otherwise. +bool isHostPipeVariable(const GlobalVariable &GV) { + return GV.hasAttribute(SYCL_HOST_PIPE_ATTR); +} + +HostPipePropertyMapTy collectHostPipeProperties(const Module &M) { + HostPipePropertyMapTy HPM; + auto HostPipeNum = count_if(M.globals(), isHostPipeVariable); + if (HostPipeNum == 0) + return HPM; + + HPM.reserve(HostPipeNum); + + for (auto &GV : M.globals()) { + if (!isHostPipeVariable(GV)) + continue; + + HPM[getGlobalVariableUniqueId(GV)] = {getHostPipeTypeSize(GV)}; + } + + return HPM; +} + +} // namespace llvm diff --git a/llvm/lib/Support/PropertySetIO.cpp b/llvm/lib/Support/PropertySetIO.cpp index 3639879214e0e..796769956401f 100644 --- a/llvm/lib/Support/PropertySetIO.cpp +++ b/llvm/lib/Support/PropertySetIO.cpp @@ -203,6 +203,7 @@ constexpr char PropertySetRegistry::SYCL_ASSERT_USED[]; constexpr char PropertySetRegistry::SYCL_EXPORTED_SYMBOLS[]; constexpr char PropertySetRegistry::SYCL_DEVICE_GLOBALS[]; constexpr char PropertySetRegistry::SYCL_DEVICE_REQUIREMENTS[]; +constexpr char PropertySetRegistry::SYCL_HOST_PIPES[]; } // namespace util } // namespace llvm diff --git a/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/host-pipes/basic.ll b/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/host-pipes/basic.ll index f6009ad0b4583..8619156b79e7d 100644 --- a/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/host-pipes/basic.ll +++ b/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/host-pipes/basic.ll @@ -14,7 +14,7 @@ $_ZN4sycl3_V13ext5intel12experimental9host_pipeI9D2HPipeIDiNS1_6oneapi12experime @_ZN4sycl3_V13ext5intel12experimental9host_pipeI9D2HPipeIDiNS1_6oneapi12experimental10propertiesISt5tupleIJEEEEE6__pipeE = linkonce_odr dso_local addrspace(1) constant %struct.BasicKernel zeroinitializer, comdat, align 1 #0 ; CHECK-IR: @_ZN4sycl3_V13ext5intel12experimental9host_pipeI9D2HPipeIDiNS1_6oneapi12experimental10propertiesISt5tupleIJEEEEE6__pipeE = linkonce_odr dso_local addrspace(1) constant %struct.BasicKernel zeroinitializer, comdat, align 1, !spirv.Decorations ![[#MN0:]] -attributes #0 = { "sycl-host-pipe" "sycl-unique-id"="_ZN4sycl3_V13ext5intel12experimental9host_pipeI9H2DPipeIDiNS1_6oneapi12experimental10propertiesISt5tupleIJEEEEE6__pipeE" } +attributes #0 = { "sycl-host-pipe" "sycl-host-pipe-size"="4" "sycl-unique-id"="_ZN4sycl3_V13ext5intel12experimental9host_pipeI9H2DPipeIDiNS1_6oneapi12experimental10propertiesISt5tupleIJEEEEE6__pipeE" } ; Ensure that the generated metadata nodes are correct ; CHECK-IR-DAG: ![[#MN0]] = !{![[#MN1:]]} diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index f9110752bc331..13468c241da0a 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -37,6 +37,7 @@ #include "llvm/Passes/PassBuilder.h" #include "llvm/SYCLLowerIR/DeviceGlobals.h" #include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.h" +#include "llvm/SYCLLowerIR/HostPipes.h" #include "llvm/SYCLLowerIR/LowerInvokeSimd.h" #include "llvm/SYCLLowerIR/LowerKernelProps.h" #include "llvm/Support/CommandLine.h" @@ -466,6 +467,11 @@ std::string saveModuleProperties(module_split::ModuleDesc &MD, PropSet.add(PropSetRegTy::SYCL_DEVICE_GLOBALS, DevGlobalPropertyMap); } + auto HostPipePropertyMap = collectHostPipeProperties(M); + if (!HostPipePropertyMap.empty()) { + PropSet.add(PropSetRegTy::SYCL_HOST_PIPES, HostPipePropertyMap); + } + std::error_code EC; std::string SCFile = makeResultFileName(".prop", I, Suff); raw_fd_ostream SCOut(SCFile, EC); diff --git a/opencl/CMakeLists.txt b/opencl/CMakeLists.txt index 4ed26bd5e719d..a8da95ea6f66c 100644 --- a/opencl/CMakeLists.txt +++ b/opencl/CMakeLists.txt @@ -20,7 +20,7 @@ set(OCL_LOADER_REPO # Repo tags/hashes -set(OCL_HEADERS_TAG dcd5bede6859d26833cd85f0d6bbcee7382dc9b3) +set(OCL_HEADERS_TAG 9ddb236e6eb3cf844f9e2f81677e1045f9bf838e) set(OCL_LOADER_TAG 9a3e962f16f5097d2054233ad8b6dad51b6f41b7) # OpenCL Headers diff --git a/sycl/include/sycl/detail/cg.hpp b/sycl/include/sycl/detail/cg.hpp index 3bad45de43080..aa94ea6fed335 100644 --- a/sycl/include/sycl/detail/cg.hpp +++ b/sycl/include/sycl/detail/cg.hpp @@ -74,6 +74,7 @@ class CG { Memset2DUSM = 18, CopyToDeviceGlobal = 19, CopyFromDeviceGlobal = 20, + ReadWriteHostPipe = 21, }; CG(CGTYPE Type, std::vector> ArgsStorage, @@ -495,6 +496,36 @@ class CGMemset2DUSM : public CG { char getValue() const { return MValue; } }; +/// "ReadWriteHostPipe" command group class. +class CGReadWriteHostPipe : public CG { + std::string PipeName; + bool Blocking; + void *HostPtr; + size_t TypeSize; + bool IsReadOp; + +public: + CGReadWriteHostPipe(const std::string &Name, bool Block, void *Ptr, + size_t Size, bool Read, + std::vector> ArgsStorage, + std::vector AccStorage, + std::vector> SharedPtrStorage, + std::vector Requirements, + std::vector Events, + detail::code_location loc = {}) + : CG(ReadWriteHostPipe, std::move(ArgsStorage), std::move(AccStorage), + std::move(SharedPtrStorage), std::move(Requirements), + std::move(Events), std::move(loc)), + PipeName(Name), Blocking(Block), HostPtr(Ptr), TypeSize(Size), + IsReadOp(Read) {} + + std::string getPipeName() { return PipeName; } + void *getHostPtr() { return HostPtr; } + size_t getTypeSize() { return TypeSize; } + bool isBlocking() { return Blocking; } + bool isReadHostPipe() { return IsReadOp; } +}; + /// "Copy to device_global" command group class. class CGCopyToDeviceGlobal : public CG { void *MSrc; diff --git a/sycl/include/sycl/detail/pi.def b/sycl/include/sycl/detail/pi.def index eda09035c883e..2e71783720def 100644 --- a/sycl/include/sycl/detail/pi.def +++ b/sycl/include/sycl/detail/pi.def @@ -131,6 +131,9 @@ _PI_API(piextUSMEnqueueMemcpy) _PI_API(piextUSMEnqueuePrefetch) _PI_API(piextUSMEnqueueMemAdvise) _PI_API(piextUSMGetMemAllocInfo) +// Host pipes +_PI_API(piextEnqueueReadHostPipe) +_PI_API(piextEnqueueWriteHostPipe) _PI_API(piextKernelSetArgMemObj) _PI_API(piextKernelSetArgSampler) diff --git a/sycl/include/sycl/detail/pi.h b/sycl/include/sycl/detail/pi.h index e89ea947c28e5..01ece111742d8 100644 --- a/sycl/include/sycl/detail/pi.h +++ b/sycl/include/sycl/detail/pi.h @@ -82,9 +82,10 @@ // the new PI_EXT_KERNEL_EXEC_INFO_CACHE_CONFIG property. // 12.25 Added PI_EXT_DEVICE_INFO_ATOMIC_FENCE_ORDER_CAPABILITIES and // PI_EXT_DEVICE_INFO_ATOMIC_FENCE_SCOPE_CAPABILITIES for piDeviceGetInfo. +// 12.26 Added piextEnqueueReadHostPipe and piextEnqueueWriteHostPipe functions. #define _PI_H_VERSION_MAJOR 12 -#define _PI_H_VERSION_MINOR 25 +#define _PI_H_VERSION_MINOR 26 #define _PI_STRING_HELPER(a) #a #define _PI_CONCAT(a, b) _PI_STRING_HELPER(a.b) @@ -827,6 +828,8 @@ static const uint8_t PI_DEVICE_BINARY_OFFLOAD_KIND_SYCL = 4; /// PropertySetRegistry::SYCL_DEVICE_REQUIREMENTS defined in PropertySetIO.h #define __SYCL_PI_PROPERTY_SET_SYCL_DEVICE_REQUIREMENTS \ "SYCL/device requirements" +/// PropertySetRegistry::SYCL_HOST_PIPES defined in PropertySetIO.h +#define __SYCL_PI_PROPERTY_SET_SYCL_HOST_PIPES "SYCL/host pipes" /// Program metadata tags recognized by the PI backends. For kernels the tag /// must appear after the kernel name. @@ -1947,6 +1950,55 @@ pi_result piextEnqueueDeviceGlobalVariableRead( /// /// Plugin /// +/// +// Host Pipes +/// + +/// Read from pipe of a given name +/// +/// @param queue a valid host command-queue in which the read / write command +/// will be queued. command_queue and program must be created with the same +/// OpenCL context. +/// @param program a program object with a successfully built executable. +/// @param pipe_symbol the name of the program scope pipe global variable. +/// @param blocking indicate if the read and write operations are blocking or +/// non-blocking +/// @param ptr a pointer to buffer in host memory that will hold resulting data +/// from pipe +/// @param size size of the memory region to read or write, in bytes. +/// @param num_events_in_waitlist number of events in the wait list. +/// @param events_waitlist specify events that need to complete before this +/// particular command can be executed. +/// @param event returns an event object that identifies this read / write +/// command and can be used to query or queue a wait for this command to +/// complete. +__SYCL_EXPORT pi_result piextEnqueueReadHostPipe( + pi_queue queue, pi_program program, const char *pipe_symbol, + pi_bool blocking, void *ptr, size_t size, pi_uint32 num_events_in_waitlist, + const pi_event *events_waitlist, pi_event *event); + +/// Write to pipe of a given name +/// +/// @param queue a valid host command-queue in which the read / write command +/// will be queued. command_queue and program must be created with the same +/// OpenCL context. +/// @param program a program object with a successfully built executable. +/// @param pipe_symbol the name of the program scope pipe global variable. +/// @param blocking indicate if the read and write operations are blocking or +/// non-blocking +/// @param ptr a pointer to buffer in host memory that holds data to be written +/// to host pipe. +/// @param size size of the memory region to read or write, in bytes. +/// @param num_events_in_waitlist number of events in the wait list. +/// @param events_waitlist specify events that need to complete before this +/// particular command can be executed. +/// @param event returns an event object that identifies this read / write +/// command and can be used to query or queue a wait for this command to +/// complete. +__SYCL_EXPORT pi_result piextEnqueueWriteHostPipe( + pi_queue queue, pi_program program, const char *pipe_symbol, + pi_bool blocking, void *ptr, size_t size, pi_uint32 num_events_in_waitlist, + const pi_event *events_waitlist, pi_event *event); /// API to get Plugin internal data, opaque to SYCL RT. Some devices whose /// device code is compiled by the host compiler (e.g. CPU emulators) may use it diff --git a/sycl/include/sycl/ext/intel/experimental/pipe_properties.hpp b/sycl/include/sycl/ext/intel/experimental/pipe_properties.hpp new file mode 100644 index 0000000000000..747c6359c7012 --- /dev/null +++ b/sycl/include/sycl/ext/intel/experimental/pipe_properties.hpp @@ -0,0 +1,190 @@ +//==----- pipe_properties.hpp - SYCL properties associated with data flow pipe +//---==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include +#include + +namespace sycl { +__SYCL_INLINE_VER_NAMESPACE(_V1) { +namespace ext { +namespace intel { +namespace experimental { + +struct min_capacity_key { + template + using value_t = oneapi::experimental::property_value< + min_capacity_key, std::integral_constant>; +}; + +struct ready_latency_key { + template + using value_t = oneapi::experimental::property_value< + ready_latency_key, std::integral_constant>; +}; + +struct bits_per_symbol_key { + template + using value_t = + oneapi::experimental::property_value>; +}; + +struct uses_valid_key { + template + using value_t = + oneapi::experimental::property_value>; +}; + +struct in_csr_key { + template + using value_t = + oneapi::experimental::property_value>; +}; + +struct first_symbol_in_high_order_bits_key { + template + using value_t = + oneapi::experimental::property_value>; +}; + +enum class protocol_name : std::uint16_t { + AVALON_STREAMING = 0, + AVALON_STREAMING_USES_READY = 1, + AVALON_MM = 2, + AVALON_MM_USES_READY = 3 +}; + +struct protocol_key { + template + using value_t = oneapi::experimental::property_value< + protocol_key, std::integral_constant>; +}; + +template +inline constexpr min_capacity_key::value_t min_capacity; + +template +inline constexpr ready_latency_key::value_t ready_latency; + +template +inline constexpr bits_per_symbol_key::value_t bits_per_symbol; + +template +inline constexpr uses_valid_key::value_t uses_valid; +inline constexpr uses_valid_key::value_t uses_valid_on; +inline constexpr uses_valid_key::value_t uses_valid_off; + +template inline constexpr in_csr_key::value_t in_csr; +inline constexpr in_csr_key::value_t in_csr_on; +inline constexpr in_csr_key::value_t in_csr_off; + +template +inline constexpr first_symbol_in_high_order_bits_key::value_t + first_symbol_in_high_order_bits; +inline constexpr first_symbol_in_high_order_bits_key::value_t + first_symbol_in_high_order_bits_on; +inline constexpr first_symbol_in_high_order_bits_key::value_t + first_symbol_in_high_order_bits_off; + +template +inline constexpr protocol_key::value_t protocol; +inline constexpr protocol_key::value_t + protocol_avalon_streaming; +inline constexpr protocol_key::value_t< + protocol_name::AVALON_STREAMING_USES_READY> + protocol_avalon_streaming_uses_ready; +inline constexpr protocol_key::value_t + protocol_avalon_mm; +inline constexpr protocol_key::value_t + protocol_avalon_mm_uses_ready; + +} // namespace experimental +} // namespace intel + +namespace oneapi { +namespace experimental { + +template <> +struct is_property_key : std::true_type { +}; +template <> +struct is_property_key + : std::true_type {}; +template <> +struct is_property_key + : std::true_type {}; +template <> +struct is_property_key : std::true_type {}; +template <> +struct is_property_key : std::true_type {}; +template <> +struct is_property_key + : std::true_type {}; +template <> +struct is_property_key : std::true_type {}; + +namespace detail { +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::MinCapacity; +}; +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::ReadyLatency; +}; +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::BitsPerSymbol; +}; +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::UsesValid; +}; +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::ImplementInCSR; +}; +template <> +struct PropertyToKind< + intel::experimental::first_symbol_in_high_order_bits_key> { + static constexpr PropKind Kind = PropKind::FirstSymbolInHigherOrderBit; +}; +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::PipeProtocol; +}; + +template <> +struct IsCompileTimeProperty + : std::true_type {}; +template <> +struct IsCompileTimeProperty + : std::true_type {}; +template <> +struct IsCompileTimeProperty + : std::true_type {}; +template <> +struct IsCompileTimeProperty + : std::true_type {}; +template <> +struct IsCompileTimeProperty : std::true_type { +}; +template <> +struct IsCompileTimeProperty< + intel::experimental::first_symbol_in_high_order_bits_key> : std::true_type { +}; +template <> +struct IsCompileTimeProperty + : std::true_type {}; + +} // namespace detail +} // namespace experimental +} // namespace oneapi +} // namespace ext +} // __SYCL_INLINE_VER_NAMESPACE(_V1) +} // namespace sycl diff --git a/sycl/include/sycl/ext/intel/experimental/pipes.hpp b/sycl/include/sycl/ext/intel/experimental/pipes.hpp index b327e1fdd167d..7dedc1d4cbc26 100644 --- a/sycl/include/sycl/ext/intel/experimental/pipes.hpp +++ b/sycl/include/sycl/ext/intel/experimental/pipes.hpp @@ -11,29 +11,110 @@ #include "fpga_utils.hpp" #include #include +#include +#include +#include #include +#include #include #include +#ifdef XPTI_ENABLE_INSTRUMENTATION +#include +#include +#endif + namespace sycl { __SYCL_INLINE_VER_NAMESPACE(_V1) { -namespace ext::intel::experimental { +namespace ext { +namespace intel { +namespace experimental { + +// A helper templateless base class to get the host_pipe name. +class pipe_base { + +protected: + pipe_base(); + ~pipe_base(); + + __SYCL_EXPORT static std::string get_pipe_name(const void *HostPipePtr); +}; template -class pipe { - static_assert(std::is_same_v<_propertiesT, - decltype(oneapi::experimental::properties{})>, - "experimental pipe properties are not yet implemented"); -}; - -template -class pipe<_name, _dataT, _min_capacity, _propertiesT, - std::enable_if_t>> { +class pipe : public pipe_base { public: + struct +#ifdef __SYCL_DEVICE_ONLY__ + [[__sycl_detail__::add_ir_attributes_global_variable( + "sycl-host-pipe", "sycl-host-pipe-size", nullptr, + sizeof(_dataT))]] [[__sycl_detail__::sycl_type(host_pipe)]] +#endif // __SYCL_DEVICE_ONLY___ + ConstantPipeStorageExp +#ifdef __SYCL_DEVICE_ONLY__ + : ConstantPipeStorage +#endif // __SYCL_DEVICE_ONLY___ + { + int32_t _ReadyLatency; + int32_t _BitsPerSymbol; + bool _UsesValid; + bool _FirstSymInHighOrderBits; + protocol_name _Protocol; + }; + // Non-blocking pipes + + // Host API + static _dataT read(queue &Q, bool &Success, + memory_order Order = memory_order::seq_cst) { + const device Dev = Q.get_device(); + bool IsPipeSupported = + Dev.has_extension("cl_intel_program_scope_host_pipe"); + if (!IsPipeSupported) { + return _dataT(); + } + _dataT Data; + void *DataPtr = &Data; + const void *HostPipePtr = &m_Storage; + const std::string PipeName = pipe_base::get_pipe_name(HostPipePtr); + + event E = Q.submit([=](handler &CGH) { + CGH.ext_intel_read_host_pipe(PipeName, DataPtr, sizeof(_dataT) /* non-blocking */); + }); + E.wait(); + if (E.get_info() == + sycl::info::event_command_status::complete) { + Success = true; + return *(_dataT *)DataPtr; + } else { + Success = false; + return _dataT(); + } + } + + static void write(queue &Q, const _dataT &Data, bool &Success, + memory_order Order = memory_order::seq_cst) { + const device Dev = Q.get_device(); + bool IsPipeSupported = + Dev.has_extension("cl_intel_program_scope_host_pipe"); + if (!IsPipeSupported) { + return; + } + + const void *HostPipePtr = &m_Storage; + const std::string PipeName = pipe_base::get_pipe_name(HostPipePtr); + void *DataPtr = const_cast<_dataT *>(&Data); + + event E = Q.submit([=](handler &CGH) { + CGH.ext_intel_write_host_pipe( + PipeName, DataPtr, sizeof(_dataT) /* non-blocking */); + }); + E.wait(); + Success = E.get_info() == + sycl::info::event_command_status::complete; + } + // Reading from pipe is lowered to SPIR-V instruction OpReadPipe via SPIR-V // friendly LLVM IR. template @@ -75,7 +156,8 @@ class pipe<_name, _dataT, _min_capacity, _propertiesT, (void)Properties; throw sycl::exception( sycl::make_error_code(sycl::errc::feature_not_supported), - "Pipes are not supported on a host device."); + "Device-side API are not supported on a host device. Please use " + "host-side API instead."); #endif // __SYCL_DEVICE_ONLY__ } @@ -124,7 +206,8 @@ class pipe<_name, _dataT, _min_capacity, _propertiesT, (void)Properties; throw sycl::exception( sycl::make_error_code(sycl::errc::feature_not_supported), - "Pipes are not supported on a host device."); + "Device-side API are not supported on a host device. Please use " + "host-side API instead."); #endif // __SYCL_DEVICE_ONLY__ } @@ -132,7 +215,48 @@ class pipe<_name, _dataT, _min_capacity, _propertiesT, write(Data, Success, oneapi::experimental::properties{}); } + static const void *get_host_ptr() { return &m_Storage; } + // Blocking pipes + + // Host API + static _dataT read(queue &Q, memory_order Order = memory_order::seq_cst) { + const device Dev = Q.get_device(); + bool IsPipeSupported = + Dev.has_extension("cl_intel_program_scope_host_pipe"); + if (!IsPipeSupported) { + return _dataT(); + } + _dataT Data; + void *DataPtr = &Data; + const void *HostPipePtr = &m_Storage; + const std::string PipeName = pipe_base::get_pipe_name(HostPipePtr); + event E = Q.submit([=](handler &CGH) { + CGH.ext_intel_read_host_pipe(PipeName, DataPtr, sizeof(_dataT), + true /*blocking*/); + }); + E.wait(); + return *(_dataT *)DataPtr; + } + + static void write(queue &Q, const _dataT &Data, + memory_order Order = memory_order::seq_cst) { + const device Dev = Q.get_device(); + bool IsPipeSupported = + Dev.has_extension("cl_intel_program_scope_host_pipe"); + if (!IsPipeSupported) { + return; + } + const void *HostPipePtr = &m_Storage; + const std::string PipeName = pipe_base::get_pipe_name(HostPipePtr); + void *DataPtr = const_cast<_dataT *>(&Data); + event E = Q.submit([=](handler &CGH) { + CGH.ext_intel_write_host_pipe(PipeName, DataPtr, + sizeof(_dataT), true /*blocking */); + }); + E.wait(); + } + // Reading from pipe is lowered to SPIR-V instruction OpReadPipe via SPIR-V // friendly LLVM IR. template @@ -173,7 +297,8 @@ class pipe<_name, _dataT, _min_capacity, _propertiesT, (void)Properties; throw sycl::exception( sycl::make_error_code(sycl::errc::feature_not_supported), - "Pipes are not supported on a host device."); + "Device-side API are not supported on a host device. Please use " + "host-side API instead."); #endif // __SYCL_DEVICE_ONLY__ } @@ -218,7 +343,8 @@ class pipe<_name, _dataT, _min_capacity, _propertiesT, (void)Properties; throw sycl::exception( sycl::make_error_code(sycl::errc::feature_not_supported), - "Pipes are not supported on a host device."); + "Device-side API are not supported on a host device. Please use " + "host-side API instead."); #endif // __SYCL_DEVICE_ONLY__ } @@ -230,10 +356,37 @@ class pipe<_name, _dataT, _min_capacity, _propertiesT, static constexpr int32_t m_Size = sizeof(_dataT); static constexpr int32_t m_Alignment = alignof(_dataT); static constexpr int32_t m_Capacity = _min_capacity; + + static constexpr int32_t m_ready_latency = + oneapi::experimental::detail::ValueOrDefault< + _propertiesT, ready_latency_key>::template get(0); + static constexpr int32_t m_bits_per_symbol = + oneapi::experimental::detail::ValueOrDefault< + _propertiesT, bits_per_symbol_key>::template get(8); + static constexpr bool m_uses_valid = + oneapi::experimental::detail::ValueOrDefault< + _propertiesT, uses_valid_key>::template get(true); + static constexpr bool m_first_symbol_in_high_order_bits = + oneapi::experimental::detail::ValueOrDefault< + _propertiesT, + first_symbol_in_high_order_bits_key>::template get(0); + static constexpr protocol_name m_protocol = oneapi::experimental::detail:: + ValueOrDefault<_propertiesT, protocol_key>::template get( + protocol_name::AVALON_STREAMING_USES_READY); + +public: + static constexpr struct ConstantPipeStorageExp m_Storage = { #ifdef __SYCL_DEVICE_ONLY__ - static constexpr struct ConstantPipeStorage m_Storage = {m_Size, m_Alignment, - m_Capacity}; + {m_Size, m_Alignment, m_Capacity}, +#endif // __SYCL_DEVICE_ONLY___ + m_ready_latency, + m_bits_per_symbol, + m_uses_valid, + m_first_symbol_in_high_order_bits, + m_protocol}; +#ifdef __SYCL_DEVICE_ONLY__ +private: // FPGA BE will recognize this function and extract its arguments. // TODO: Pass latency control parameters via the __spirv_* builtin when ready. template @@ -276,6 +429,8 @@ class pipe<_name, _dataT, _min_capacity, _propertiesT, #endif // __SYCL_DEVICE_ONLY__ }; -} // namespace ext::intel::experimental +} // namespace experimental +} // namespace intel +} // namespace ext } // __SYCL_INLINE_VER_NAMESPACE(_V1) } // namespace sycl diff --git a/sycl/include/sycl/ext/intel/pipes.hpp b/sycl/include/sycl/ext/intel/pipes.hpp index 61a9e742e5cb1..8b8dffeba04f4 100644 --- a/sycl/include/sycl/ext/intel/pipes.hpp +++ b/sycl/include/sycl/ext/intel/pipes.hpp @@ -35,7 +35,8 @@ template class pipe { (void)_Success; throw sycl::exception( sycl::make_error_code(sycl::errc::feature_not_supported), - "Pipes are not supported on a host device."); + "Device-side API are not supported on a host device. Please use " + "host-side API instead."); #endif // __SYCL_DEVICE_ONLY__ } @@ -52,7 +53,8 @@ template class pipe { (void)_Data; throw sycl::exception( sycl::make_error_code(sycl::errc::feature_not_supported), - "Pipes are not supported on a host device."); + "Device-side API are not supported on a host device. Please use " + "host-side API instead."); #endif // __SYCL_DEVICE_ONLY__ } @@ -69,7 +71,8 @@ template class pipe { #else throw sycl::exception( sycl::make_error_code(sycl::errc::feature_not_supported), - "Pipes are not supported on a host device."); + "Device-side API are not supported on a host device. Please use " + "host-side API instead.."); #endif // __SYCL_DEVICE_ONLY__ } @@ -84,7 +87,8 @@ template class pipe { (void)_Data; throw sycl::exception( sycl::make_error_code(sycl::errc::feature_not_supported), - "Pipes are not supported on a host device."); + "Device-side API are not supported on a host device. Please use " + "host-side API instead."); #endif // __SYCL_DEVICE_ONLY__ } @@ -137,7 +141,8 @@ class kernel_readable_io_pipe { (void)_Success; throw sycl::exception( sycl::make_error_code(sycl::errc::feature_not_supported), - "Pipes are not supported on a host device."); + "Device-side API are not supported on a host device. Please use " + "host-side API instead."); #endif // __SYCL_DEVICE_ONLY__ } @@ -154,7 +159,8 @@ class kernel_readable_io_pipe { #else throw sycl::exception( sycl::make_error_code(sycl::errc::feature_not_supported), - "Pipes are not supported on a host device."); + "Device-side API are not supported on a host device. Please use " + "host-side API instead."); #endif // __SYCL_DEVICE_ONLY__ } @@ -187,7 +193,8 @@ class kernel_writeable_io_pipe { (void)_Success; throw sycl::exception( sycl::make_error_code(sycl::errc::feature_not_supported), - "Pipes are not supported on a host device."); + "Device-side API are not supported on a host device. Please use " + "host-side API instead."); #endif // __SYCL_DEVICE_ONLY__ } @@ -203,7 +210,8 @@ class kernel_writeable_io_pipe { (void)_Data; throw sycl::exception( sycl::make_error_code(sycl::errc::feature_not_supported), - "Pipes are not supported on a host device."); + "Device-side API are not supported on a host device. Please use " + "host-side API instead."); #endif // __SYCL_DEVICE_ONLY__ } diff --git a/sycl/include/sycl/ext/oneapi/properties/properties.hpp b/sycl/include/sycl/ext/oneapi/properties/properties.hpp index 5b83abf6cef5a..9cb851a639a21 100644 --- a/sycl/include/sycl/ext/oneapi/properties/properties.hpp +++ b/sycl/include/sycl/ext/oneapi/properties/properties.hpp @@ -221,6 +221,23 @@ template using merged_properties_t = typename merged_properties::type; +template +struct ValueOrDefault { + template static constexpr ValT get(ValT Default) { + return Default; + } +}; + +template +struct ValueOrDefault< + Properties, PropertyKey, + std::enable_if_t && + Properties::template has_property()>> { + template static constexpr ValT get(ValT) { + return Properties::template get_property().value; + } +}; + } // namespace detail } // namespace ext::oneapi::experimental diff --git a/sycl/include/sycl/ext/oneapi/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp index 4421c5d66de4d..7c2e3063ace13 100644 --- a/sycl/include/sycl/ext/oneapi/properties/property.hpp +++ b/sycl/include/sycl/ext/oneapi/properties/property.hpp @@ -188,8 +188,15 @@ enum PropKind : uint32_t { WaitRequest = 22, Alignment = 23, CacheConfig = 24, + BitsPerSymbol = 25, + FirstSymbolInHigherOrderBit = 26, + MinCapacity = 27, + PipeProtocol = 28, + ReadyLatency = 29, + UsesReady = 30, + UsesValid = 31, // PropKindSize must always be the last value. - PropKindSize = 25, + PropKindSize = 32, }; // This trait must be specialized for all properties and must have a unique diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 0985e60d8eeb9..4d234ab9c06d3 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -95,6 +95,14 @@ __SYCL_INLINE_VER_NAMESPACE(_V1) { class handler; template class buffer; + +namespace ext::intel::experimental +{ + template + class pipe; +} + namespace detail { class handler_impl; @@ -2888,6 +2896,29 @@ class __SYCL_EXPORT handler { friend class ::MockHandler; friend class detail::queue_impl; + // Make pipe class friend to be able to call ext_intel_read/write_host_pipe method. + template + friend class ext::intel::experimental::pipe; + + /// Read from a host pipe given a host address and + /// \param Name name of the host pipe to be passed into lower level runtime + /// \param Ptr host pointer of host pipe as identified by address of its const + /// expr m_Storage member \param Size the size of data getting read back / to. + /// /// \param Size the size of data getting read back / to. \param Block + /// if read opeartion is blocking, default to false. + void ext_intel_read_host_pipe(const std::string &Name, void *Ptr, + size_t Size, bool Block=false); + + /// Write to host pipes given a host address and + /// \param Name name of the host pipe to be passed into lower level runtime + /// \param Ptr host pointer of host pipe as identified by address of its const + /// expr m_Storage member \param Size the size of data getting read back / to. + /// /// \param Size the size of data write / to. \param Block + /// if write opeartion is blocking, default to false. + void ext_intel_write_host_pipe(const std::string &Name, void *Ptr, + size_t Size, bool Block=false); + bool DisableRangeRounding(); bool RangeRoundingTrace(); diff --git a/sycl/include/sycl/sycl.hpp b/sycl/include/sycl/sycl.hpp index dd062cc2d2fb8..03be402053f81 100644 --- a/sycl/include/sycl/sycl.hpp +++ b/sycl/include/sycl/sycl.hpp @@ -63,6 +63,8 @@ #include #endif #include +#include +#include #include #include #include diff --git a/sycl/plugins/cuda/pi_cuda.cpp b/sycl/plugins/cuda/pi_cuda.cpp index 6f7a31c3af8ef..83a986f193d2f 100644 --- a/sycl/plugins/cuda/pi_cuda.cpp +++ b/sycl/plugins/cuda/pi_cuda.cpp @@ -5542,6 +5542,43 @@ pi_result cuda_piextEnqueueDeviceGlobalVariableRead( return result; } +/// Host Pipes +pi_result cuda_piextEnqueueReadHostPipe( + pi_queue queue, pi_program program, const char *pipe_symbol, + pi_bool blocking, void *ptr, size_t size, pi_uint32 num_events_in_waitlist, + const pi_event *events_waitlist, pi_event *event) { + (void)queue; + (void)program; + (void)pipe_symbol; + (void)blocking; + (void)ptr; + (void)size; + (void)num_events_in_waitlist; + (void)events_waitlist; + (void)event; + + sycl::detail::pi::die("cuda_piextEnqueueReadHostPipe not implemented"); + return {}; +} + +pi_result cuda_piextEnqueueWriteHostPipe( + pi_queue queue, pi_program program, const char *pipe_symbol, + pi_bool blocking, void *ptr, size_t size, pi_uint32 num_events_in_waitlist, + const pi_event *events_waitlist, pi_event *event) { + (void)queue; + (void)program; + (void)pipe_symbol; + (void)blocking; + (void)ptr; + (void)size; + (void)num_events_in_waitlist; + (void)events_waitlist; + (void)event; + + sycl::detail::pi::die("cuda_piextEnqueueWriteHostPipe not implemented"); + return {}; +} + // This API is called by Sycl RT to notify the end of the plugin lifetime. // Windows: dynamically loaded plugins might have been unloaded already // when this is called. Sycl RT holds onto the PI plugin so it can be @@ -5728,6 +5765,10 @@ pi_result piPluginInit(pi_plugin *PluginInit) { _PI_CL(piextEnqueueDeviceGlobalVariableRead, cuda_piextEnqueueDeviceGlobalVariableRead) + // Host Pipe + _PI_CL(piextEnqueueReadHostPipe, cuda_piextEnqueueReadHostPipe) + _PI_CL(piextEnqueueWriteHostPipe, cuda_piextEnqueueWriteHostPipe) + _PI_CL(piextKernelSetArgMemObj, cuda_piextKernelSetArgMemObj) _PI_CL(piextKernelSetArgSampler, cuda_piextKernelSetArgSampler) _PI_CL(piPluginGetLastError, cuda_piPluginGetLastError) diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp index 873abdc64c26b..8eeb2432bcf08 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -2005,6 +2005,25 @@ pi_result piextUSMGetMemAllocInfo(pi_context, const void *, pi_mem_alloc_info, DIE_NO_IMPLEMENTATION; } +/// Host Pipes +pi_result piextEnqueueReadHostPipe(pi_queue queue, pi_program program, + const char *pipe_symbol, pi_bool blocking, + void *ptr, size_t size, + pi_uint32 num_events_in_waitlist, + const pi_event *events_waitlist, + pi_event *event) { + DIE_NO_IMPLEMENTATION; +} + +pi_result piextEnqueueWriteHostPipe(pi_queue queue, pi_program program, + const char *pipe_symbol, pi_bool blocking, + void *ptr, size_t size, + pi_uint32 num_events_in_waitlist, + const pi_event *events_waitlist, + pi_event *event) { + DIE_NO_IMPLEMENTATION; +} + pi_result piKernelSetExecInfo(pi_kernel, pi_kernel_exec_info, size_t, const void *) { DIE_NO_IMPLEMENTATION; diff --git a/sycl/plugins/hip/pi_hip.cpp b/sycl/plugins/hip/pi_hip.cpp index 110fa8452ade6..986b85102a3cb 100644 --- a/sycl/plugins/hip/pi_hip.cpp +++ b/sycl/plugins/hip/pi_hip.cpp @@ -5373,6 +5373,44 @@ pi_result hip_piextEnqueueDeviceGlobalVariableRead( sycl::detail::pi::die( "hip_piextEnqueueDeviceGlobalVariableRead not implemented"); +} + +/// Host Pipes +pi_result hip_piextEnqueueReadHostPipe(pi_queue queue, pi_program program, + const char *pipe_symbol, + pi_bool blocking, void *ptr, size_t size, + pi_uint32 num_events_in_waitlist, + const pi_event *events_waitlist, + pi_event *event) { + (void)queue; + (void)program; + (void)pipe_symbol; + (void)blocking; + (void)ptr; + (void)size; + (void)num_events_in_waitlist; + (void)events_waitlist; + (void)event; + + sycl::detail::pi::die("hip_piextEnqueueReadHostPipe not implemented"); + return {}; +} + +pi_result hip_piextEnqueueWriteHostPipe( + pi_queue queue, pi_program program, const char *pipe_symbol, + pi_bool blocking, void *ptr, size_t size, pi_uint32 num_events_in_waitlist, + const pi_event *events_waitlist, pi_event *event) { + (void)queue; + (void)program; + (void)pipe_symbol; + (void)blocking; + (void)ptr; + (void)size; + (void)num_events_in_waitlist; + (void)events_waitlist; + (void)event; + + sycl::detail::pi::die("hip_piextEnqueueWriteHostPipe not implemented"); return {}; } @@ -5562,6 +5600,10 @@ pi_result piPluginInit(pi_plugin *PluginInit) { _PI_CL(piextEnqueueDeviceGlobalVariableRead, hip_piextEnqueueDeviceGlobalVariableRead) + // Host Pipe + _PI_CL(piextEnqueueReadHostPipe, hip_piextEnqueueReadHostPipe) + _PI_CL(piextEnqueueWriteHostPipe, hip_piextEnqueueWriteHostPipe) + _PI_CL(piextKernelSetArgMemObj, hip_piextKernelSetArgMemObj) _PI_CL(piextKernelSetArgSampler, hip_piextKernelSetArgSampler) _PI_CL(piPluginGetLastError, hip_piPluginGetLastError) diff --git a/sycl/plugins/level_zero/pi_level_zero.cpp b/sycl/plugins/level_zero/pi_level_zero.cpp index 01a111fbb79de..2793c76798981 100644 --- a/sycl/plugins/level_zero/pi_level_zero.cpp +++ b/sycl/plugins/level_zero/pi_level_zero.cpp @@ -7960,6 +7960,71 @@ pi_result piextEnqueueDeviceGlobalVariableRead( Count, pi_cast(GlobalVarPtr) + Offset, NumEventsInWaitList, EventsWaitList, Event, PreferCopyEngine); } +/// API for Read from host pipe. +/// +/// \param Queue is the queue +/// \param Program is the program containing the device variable +/// \param PipeSymbol is the unique identifier for the device variable +/// \param Blocking is true if the write should block +/// \param Ptr is a pointer to where the data will be copied to +/// \param Size is size of the data that is read/written from/to pipe +/// \param NumEventsInWaitList is a number of events in the wait list +/// \param EventWaitList is the wait list +/// \param Event is the resulting event +pi_result piextEnqueueReadHostPipe(pi_queue Queue, pi_program Program, + const char *PipeSymbol, pi_bool Blocking, + void *Ptr, size_t Size, + pi_uint32 NumEventsInWaitList, + const pi_event *EventsWaitList, + pi_event *Event) { + (void)Queue; + (void)Program; + (void)PipeSymbol; + (void)Blocking; + (void)Ptr; + (void)Size; + (void)NumEventsInWaitList; + (void)EventsWaitList; + (void)Event; + + PI_ASSERT(Queue, PI_ERROR_INVALID_QUEUE); + + die("piextEnqueueReadHostPipe: not implemented"); + return {}; +} + +/// API for write to pipe of a given name. +/// +/// \param Queue is the queue +/// \param Program is the program containing the device variable +/// \param PipeSymbol is the unique identifier for the device variable +/// \param Blocking is true if the write should block +/// \param Ptr is a pointer to where the data must be copied from +/// \param Size is size of the data that is read/written from/to pipe +/// \param NumEventsInWaitList is a number of events in the wait list +/// \param EventWaitList is the wait list +/// \param Event is the resulting event +pi_result piextEnqueueWriteHostPipe(pi_queue Queue, pi_program Program, + const char *PipeSymbol, pi_bool Blocking, + void *Ptr, size_t Size, + pi_uint32 NumEventsInWaitList, + const pi_event *EventsWaitList, + pi_event *Event) { + (void)Queue; + (void)Program; + (void)PipeSymbol; + (void)Blocking; + (void)Ptr; + (void)Size; + (void)NumEventsInWaitList; + (void)EventsWaitList; + (void)Event; + + PI_ASSERT(Queue, PI_ERROR_INVALID_QUEUE); + + die("piextEnqueueWriteHostPipe: not implemented"); + return {}; +} pi_result piKernelSetExecInfo(pi_kernel Kernel, pi_kernel_exec_info ParamName, size_t ParamValueSize, const void *ParamValue) { diff --git a/sycl/plugins/opencl/pi_opencl.cpp b/sycl/plugins/opencl/pi_opencl.cpp index 61111af7f40c9..5c5c24ffc130d 100644 --- a/sycl/plugins/opencl/pi_opencl.cpp +++ b/sycl/plugins/opencl/pi_opencl.cpp @@ -71,6 +71,9 @@ CONSTFIX char clEnqueueWriteGlobalVariableName[] = "clEnqueueWriteGlobalVariableINTEL"; CONSTFIX char clEnqueueReadGlobalVariableName[] = "clEnqueueReadGlobalVariableINTEL"; +// Names of host pipe functions queried from OpenCL +CONSTFIX char clEnqueueReadHostPipeName[] = "clEnqueueReadHostPipeINTEL"; +CONSTFIX char clEnqueueWriteHostPipeName[] = "clEnqueueWriteHostPipeINTEL"; #undef CONSTFIX @@ -1927,6 +1930,64 @@ pi_result piextEnqueueDeviceGlobalVariableRead( return cast(Res); } +pi_result piextEnqueueReadHostPipe(pi_queue queue, pi_program program, + const char *pipe_symbol, pi_bool blocking, + void *ptr, size_t size, + pi_uint32 num_events_in_waitlist, + const pi_event *events_waitlist, + pi_event *event) { + cl_context CLContext; + cl_int CLErr = + clGetCommandQueueInfo(cast(queue), CL_QUEUE_CONTEXT, + sizeof(cl_context), &CLContext, nullptr); + if (CLErr != CL_SUCCESS) { + return cast(CLErr); + } + + clEnqueueReadHostPipeINTEL_fn FuncPtr = nullptr; + pi_result RetVal = getExtFuncFromContext( + cast(CLContext), &FuncPtr); + + if (FuncPtr) { + RetVal = cast(FuncPtr( + cast(queue), cast(program), pipe_symbol, + blocking, ptr, size, num_events_in_waitlist, + cast(events_waitlist), cast(event))); + } + + return RetVal; +} + +pi_result piextEnqueueWriteHostPipe(pi_queue queue, pi_program program, + const char *pipe_symbol, pi_bool blocking, + void *ptr, size_t size, + pi_uint32 num_events_in_waitlist, + const pi_event *events_waitlist, + pi_event *event) { + cl_context CLContext; + cl_int CLErr = + clGetCommandQueueInfo(cast(queue), CL_QUEUE_CONTEXT, + sizeof(cl_context), &CLContext, nullptr); + if (CLErr != CL_SUCCESS) { + return cast(CLErr); + } + + clEnqueueWriteHostPipeINTEL_fn FuncPtr = nullptr; + pi_result RetVal = getExtFuncFromContext( + cast(CLContext), &FuncPtr); + + if (FuncPtr) { + RetVal = cast(FuncPtr( + cast(queue), cast(program), pipe_symbol, + blocking, ptr, size, num_events_in_waitlist, + cast(events_waitlist), cast(event))); + } + + return RetVal; +} + /// API to set attributes controlling kernel execution /// /// \param kernel is the pi kernel to execute @@ -2213,6 +2274,9 @@ pi_result piPluginInit(pi_plugin *PluginInit) { piextEnqueueDeviceGlobalVariableWrite) _PI_CL(piextEnqueueDeviceGlobalVariableRead, piextEnqueueDeviceGlobalVariableRead) + // Host Pipe + _PI_CL(piextEnqueueReadHostPipe, piextEnqueueReadHostPipe) + _PI_CL(piextEnqueueWriteHostPipe, piextEnqueueWriteHostPipe) _PI_CL(piextKernelSetArgMemObj, piextKernelSetArgMemObj) _PI_CL(piextKernelSetArgSampler, piextKernelSetArgSampler) diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt index 1423d6d4880ad..99112641aeda2 100644 --- a/sycl/source/CMakeLists.txt +++ b/sycl/source/CMakeLists.txt @@ -165,6 +165,7 @@ set(SYCL_SOURCES "detail/context_impl.cpp" "detail/device_binary_image.cpp" "detail/device_filter.cpp" + "detail/host_pipe_map.cpp" "detail/device_global_map.cpp" "detail/device_global_map_entry.cpp" "detail/device_impl.cpp" @@ -183,6 +184,7 @@ set(SYCL_SOURCES "detail/kernel_impl.cpp" "detail/kernel_program_cache.cpp" "detail/memory_manager.cpp" + "detail/pipes.cpp" "detail/platform_impl.cpp" "detail/program_impl.cpp" "detail/program_manager/program_manager.cpp" diff --git a/sycl/source/detail/device_binary_image.cpp b/sycl/source/detail/device_binary_image.cpp index 38aab48e0229e..210516675277c 100644 --- a/sycl/source/detail/device_binary_image.cpp +++ b/sycl/source/detail/device_binary_image.cpp @@ -176,6 +176,7 @@ void RTDeviceBinaryImage::init(pi_device_binary Bin) { ExportedSymbols.init(Bin, __SYCL_PI_PROPERTY_SET_SYCL_EXPORTED_SYMBOLS); DeviceGlobals.init(Bin, __SYCL_PI_PROPERTY_SET_SYCL_DEVICE_GLOBALS); DeviceRequirements.init(Bin, __SYCL_PI_PROPERTY_SET_SYCL_DEVICE_REQUIREMENTS); + HostPipes.init(Bin, __SYCL_PI_PROPERTY_SET_SYCL_HOST_PIPES); } DynRTDeviceBinaryImage::DynRTDeviceBinaryImage( diff --git a/sycl/source/detail/device_binary_image.hpp b/sycl/source/detail/device_binary_image.hpp index 6e92a288778ba..3cbd24e63642e 100644 --- a/sycl/source/detail/device_binary_image.hpp +++ b/sycl/source/detail/device_binary_image.hpp @@ -222,6 +222,7 @@ class RTDeviceBinaryImage { const PropertyRange &getDeviceRequirements() const { return DeviceRequirements; } + const PropertyRange &getHostPipes() const { return HostPipes; } std::uintptr_t getImageID() const { assert(Bin && "Image ID is not available without a binary image."); @@ -245,6 +246,7 @@ class RTDeviceBinaryImage { RTDeviceBinaryImage::PropertyRange ExportedSymbols; RTDeviceBinaryImage::PropertyRange DeviceGlobals; RTDeviceBinaryImage::PropertyRange DeviceRequirements; + RTDeviceBinaryImage::PropertyRange HostPipes; }; // Dynamically allocated device binary image, which de-allocates its binary diff --git a/sycl/source/detail/handler_impl.hpp b/sycl/source/detail/handler_impl.hpp index 80618e139404f..3cb23c1df8d54 100644 --- a/sycl/source/detail/handler_impl.hpp +++ b/sycl/source/detail/handler_impl.hpp @@ -85,6 +85,19 @@ class handler_impl { /// property. bool MIsDeviceImageScoped = false; + // Program scope pipe information. + + // Pipe name that uniquely identifies a pipe. + std::string HostPipeName; + // Pipe host pointer, the address of its constexpr __pipe member. + void *HostPipePtr = nullptr; + // Host pipe read write operation is blocking. + bool HostPipeBlocking = false; + // The size of returned type for each read. + size_t HostPipeTypeSize = 0; + // If the pipe operation is read or write, 1 for read 0 for write. + bool HostPipeRead = true; + RT::PiKernelCacheConfig MKernelCacheConfig = PI_EXT_KERNEL_EXEC_INFO_CACHE_DEFAULT; }; diff --git a/sycl/source/detail/host_pipe_map.cpp b/sycl/source/detail/host_pipe_map.cpp new file mode 100644 index 0000000000000..3b0315e7f00b2 --- /dev/null +++ b/sycl/source/detail/host_pipe_map.cpp @@ -0,0 +1,24 @@ +//==-------------------- host_pipe_map.cpp -----------------------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include + +namespace sycl { +__SYCL_INLINE_VER_NAMESPACE(_V1) { +namespace detail { +namespace host_pipe_map { + +__SYCL_EXPORT void add(const void *HostPipePtr, const char *UniqueId) { + detail::ProgramManager::getInstance().addOrInitHostPipeEntry(HostPipePtr, + UniqueId); +} + +} // namespace host_pipe_map +} // namespace detail +} // __SYCL_INLINE_VER_NAMESPACE(_V1) +} // namespace sycl diff --git a/sycl/source/detail/host_pipe_map_entry.hpp b/sycl/source/detail/host_pipe_map_entry.hpp new file mode 100644 index 0000000000000..2a7412e167990 --- /dev/null +++ b/sycl/source/detail/host_pipe_map_entry.hpp @@ -0,0 +1,57 @@ +//==----------------- host_pipe_map_entry.hpp --------------------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include +#include +#include + +namespace sycl { +__SYCL_INLINE_VER_NAMESPACE(_V1) { +namespace detail { + +struct HostPipeMapEntry { + std::string MUniqueId; + // Pointer to the host_pipe on host. + const void *MHostPipePtr; + // Size of the underlying type in the host_pipe. + std::uint32_t MHostPipeTSize; + // The device image that pipe is associated with + const RTDeviceBinaryImage *mDeviceImage; + + // Constructor only initializes with the pointer and ID. + // Other members will be initialized later + HostPipeMapEntry(std::string UniqueId, const void *HostPipePtr) + : MUniqueId(UniqueId), MHostPipePtr(HostPipePtr), MHostPipeTSize(0) {} + + // Constructor only initializes with the size and ID. + // Other members will be initialized later + HostPipeMapEntry(std::string UniqueId, std::uint32_t HostPipeTSize) + : MUniqueId(UniqueId), MHostPipePtr(nullptr), + MHostPipeTSize(HostPipeTSize) {} + + void initialize(std::uint32_t HostPipeTSize) { + assert(HostPipeTSize != 0 && "Host pipe initialized with 0 size."); + assert(MHostPipeTSize == 0 && "Host pipe has already been initialized."); + MHostPipeTSize = HostPipeTSize; + } + + void initialize(const void *HostPipePtr) { + assert(!MHostPipePtr && "Host pipe pointer has already been initialized."); + MHostPipePtr = HostPipePtr; + } + + void initialize(const RTDeviceBinaryImage *DeviceImage) { + mDeviceImage = DeviceImage; + } +}; + +} // namespace detail +} // __SYCL_INLINE_VER_NAMESPACE(_V1) +} // namespace sycl diff --git a/sycl/source/detail/pipes.cpp b/sycl/source/detail/pipes.cpp new file mode 100755 index 0000000000000..9b8b805a07662 --- /dev/null +++ b/sycl/source/detail/pipes.cpp @@ -0,0 +1,29 @@ +//==-------------------- pipes.cpp -----------------------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include +#include + +namespace sycl { +__SYCL_INLINE_VER_NAMESPACE(_V1) { +namespace ext { +namespace intel { +namespace experimental { + +__SYCL_EXPORT std::string pipe_base::get_pipe_name(const void *HostPipePtr) { + return sycl::_V1::detail::ProgramManager::getInstance() + .getHostPipeEntry(HostPipePtr) + ->MUniqueId; +} + +} // namespace experimental +} // namespace intel +} // namespace ext +} // __SYCL_INLINE_VER_NAMESPACE(_V1) +} // namespace sycl diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index d6bc161428708..0c829abbd1629 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -1177,7 +1177,6 @@ bool ProgramManager::kernelUsesAssert(OSModuleHandle M, void ProgramManager::addImages(pi_device_binaries DeviceBinary) { std::lock_guard Guard(Sync::getGlobalLock()); const bool DumpImages = std::getenv("SYCL_DUMP_IMAGES") && !m_UseSpvFile; - for (int I = 0; I < DeviceBinary->NumDeviceBinaries; I++) { pi_device_binary RawImg = &(DeviceBinary->DeviceBinaries[I]); OSModuleHandle M = OSUtil::getOSModuleHandle(RawImg); @@ -1320,6 +1319,38 @@ void ProgramManager::addImages(pi_device_binaries DeviceBinary) { } } } + // ... and initialize associated host_pipe information + { + std::lock_guard HostPipesGuard(m_HostPipesMutex); + auto HostPipes = Img->getHostPipes(); + for (const pi_device_binary_property &HostPipe : HostPipes) { + ByteArray HostPipeInfo = DeviceBinaryProperty(HostPipe).asByteArray(); + + // The supplied host_pipe info property is expected to contain: + // * 8 bytes - Size of the property. + // * 4 bytes - Size of the underlying type in the host_pipe. + // Note: Property may be padded. + + HostPipeInfo.dropBytes(8); + auto TypeSize = HostPipeInfo.consume(); + assert(HostPipeInfo.empty() && "Extra data left!"); + + auto ExistingHostPipe = m_HostPipes.find(HostPipe->Name); + if (ExistingHostPipe != m_HostPipes.end()) { + // If it has already been registered we update the information. + ExistingHostPipe->second->initialize(TypeSize); + ExistingHostPipe->second->initialize(Img.get()); + } else { + // If it has not already been registered we create a new entry. + // Note: Pointer to the host pipe is not available here, so it + // cannot be set until registration happens. + auto EntryUPtr = + std::make_unique(HostPipe->Name, TypeSize); + EntryUPtr->initialize(Img.get()); + m_HostPipes.emplace(HostPipe->Name, std::move(EntryUPtr)); + } + } + } m_DeviceImages[KSId].reset(new std::vector()); cacheKernelUsesAssertInfo(M, *Img); @@ -1643,6 +1674,37 @@ std::vector ProgramManager::getDeviceGlobalEntries( return FoundEntries; } +void ProgramManager::addOrInitHostPipeEntry(const void *HostPipePtr, + const char *UniqueId) { + std::lock_guard HostPipesGuard(m_HostPipesMutex); + + auto ExistingHostPipe = m_HostPipes.find(UniqueId); + if (ExistingHostPipe != m_HostPipes.end()) { + ExistingHostPipe->second->initialize(HostPipePtr); + m_Ptr2HostPipe.insert({HostPipePtr, ExistingHostPipe->second.get()}); + return; + } + + auto EntryUPtr = std::make_unique(UniqueId, HostPipePtr); + auto NewEntry = m_HostPipes.emplace(UniqueId, std::move(EntryUPtr)); + m_Ptr2HostPipe.insert({HostPipePtr, NewEntry.first->second.get()}); +} + +HostPipeMapEntry * +ProgramManager::getHostPipeEntry(const std::string &UniqueId) { + std::lock_guard HostPipesGuard(m_HostPipesMutex); + auto Entry = m_HostPipes.find(UniqueId); + assert(Entry != m_HostPipes.end() && "Host pipe entry not found"); + return Entry->second.get(); +} + +HostPipeMapEntry *ProgramManager::getHostPipeEntry(const void *HostPipePtr) { + std::lock_guard HostPipesGuard(m_HostPipesMutex); + auto Entry = m_Ptr2HostPipe.find(HostPipePtr); + assert(Entry != m_Ptr2HostPipe.end() && "Host pipe entry not found"); + return Entry->second; +} + device_image_plain ProgramManager::getDeviceImageFromBinaryImage( RTDeviceBinaryImage *BinImage, const context &Ctx, const device &Dev) { const bundle_state ImgState = getBinImageState(BinImage); diff --git a/sycl/source/detail/program_manager/program_manager.hpp b/sycl/source/detail/program_manager/program_manager.hpp index 1bd81491462fa..3dba0a499ab61 100644 --- a/sycl/source/detail/program_manager/program_manager.hpp +++ b/sycl/source/detail/program_manager/program_manager.hpp @@ -9,10 +9,12 @@ #pragma once #include #include +#include #include #include #include #include +#include #include #include #include @@ -214,6 +216,17 @@ class ProgramManager { std::vector getDeviceGlobalEntries(const std::vector &UniqueIds, bool ExcludeDeviceImageScopeDecorated = false); + // The function inserts or initializes a host_pipe entry into the + // host_pipe map. + void addOrInitHostPipeEntry(const void *HostPipePtr, const char *UniqueId); + + // The function gets a host_pipe entry identified by the unique ID from + // the host_pipe map. + HostPipeMapEntry *getHostPipeEntry(const std::string &UniqueId); + + // The function gets a host_pipe entry identified by the pointer to the + // host_pipe object from the host_pipe map. + HostPipeMapEntry *getHostPipeEntry(const void *HostPipePtr); device_image_plain getDeviceImageFromBinaryImage(RTDeviceBinaryImage *BinImage, @@ -425,6 +438,14 @@ class ProgramManager { /// Protects m_DeviceGlobals and m_Ptr2DeviceGlobal. std::mutex m_DeviceGlobalsMutex; + + // Maps between host_pipe identifiers and associated information. + std::unordered_map> + m_HostPipes; + std::unordered_map m_Ptr2HostPipe; + + /// Protects m_HostPipes and m_Ptr2HostPipe. + std::mutex m_HostPipesMutex; }; } // namespace detail } // __SYCL_INLINE_VER_NAMESPACE(_V1) diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index 9283e415aee01..4ab7c6bdce954 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -10,6 +10,7 @@ #include #include +#include #include #include #include @@ -2343,6 +2344,39 @@ pi_int32 enqueueImpKernel( return PI_SUCCESS; } +pi_int32 enqueueReadWriteHostPipe(const QueueImplPtr &Queue, + const std::string &PipeName, bool blocking, + void *ptr, size_t size, + std::vector &RawEvents, + RT::PiEvent *OutEvent, bool read) { + detail::HostPipeMapEntry *hostPipeEntry = + ProgramManager::getInstance().getHostPipeEntry(PipeName); + + RT::PiProgram Program = ProgramManager::getInstance().createPIProgram( + *(hostPipeEntry->mDeviceImage), Queue->get_context(), + Queue->get_device()); + + // Get plugin for calling opencl functions + const detail::plugin &Plugin = Queue->getPlugin(); + + pi_queue pi_q = Queue->getHandleRef(); + pi_result Error; + if (read) { + Error = + Plugin.call_nocheck( + pi_q, Program, PipeName.c_str(), blocking, ptr, size, + RawEvents.size(), RawEvents.empty() ? nullptr : &RawEvents[0], + OutEvent); + } else { + Error = + Plugin.call_nocheck( + pi_q, Program, PipeName.c_str(), blocking, ptr, size, + RawEvents.size(), RawEvents.empty() ? nullptr : &RawEvents[0], + OutEvent); + } + return Error; +} + pi_int32 ExecCGCommand::enqueueImp() { if (getCG().getType() != CG::CGTYPE::CodeplayHostTask) waitForPreparedHostEvents(); @@ -2752,6 +2786,21 @@ pi_int32 ExecCGCommand::enqueueImp() { return CL_SUCCESS; } + case CG::CGTYPE::ReadWriteHostPipe: { + CGReadWriteHostPipe *ExecReadWriteHostPipe = + (CGReadWriteHostPipe *)MCommandGroup.get(); + std::string pipeName = ExecReadWriteHostPipe->getPipeName(); + void *hostPtr = ExecReadWriteHostPipe->getHostPtr(); + size_t typeSize = ExecReadWriteHostPipe->getTypeSize(); + bool blocking = ExecReadWriteHostPipe->isBlocking(); + bool read = ExecReadWriteHostPipe->isReadHostPipe(); + + if (!Event) { + Event = &MEvent->getHandleRef(); + } + return enqueueReadWriteHostPipe(MQueue, pipeName, blocking, hostPtr, + typeSize, RawEvents, Event, read); + } case CG::CGTYPE::None: throw runtime_error("CG type not implemented.", PI_ERROR_INVALID_OPERATION); } diff --git a/sycl/source/detail/scheduler/commands.hpp b/sycl/source/detail/scheduler/commands.hpp index 86944ac4dcc1e..35df6ef614cc7 100644 --- a/sycl/source/detail/scheduler/commands.hpp +++ b/sycl/source/detail/scheduler/commands.hpp @@ -585,6 +585,12 @@ class MemCpyCommandHost : public Command { void **MDstPtr = nullptr; }; +pi_int32 enqueueReadWriteHostPipe(const QueueImplPtr &Queue, + const std::string &PipeName, bool blocking, + void *ptr, size_t size, + std::vector &RawEvents, + RT::PiEvent *OutEvent, bool read); + pi_int32 enqueueImpKernel( const QueueImplPtr &Queue, NDRDescT &NDRDesc, std::vector &Args, const std::shared_ptr &KernelBundleImplPtr, diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 8c565e4b5fa61..3a44c5e810011 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -359,6 +359,14 @@ event handler::finalize() { std::move(MEvents), MOSModuleHandle, MCodeLoc)); break; } + case detail::CG::ReadWriteHostPipe: { + CommandGroup.reset(new detail::CGReadWriteHostPipe( + MImpl->HostPipeName, MImpl->HostPipeBlocking, MImpl->HostPipePtr, + MImpl->HostPipeTypeSize, MImpl->HostPipeRead, std::move(MArgsStorage), + std::move(MAccStorage), std::move(MSharedPtrStorage), + std::move(MRequirements), std::move(MEvents), MCodeLoc)); + break; + } case detail::CG::None: if (detail::pi::trace(detail::pi::TraceLevel::PI_TRACE_ALL)) { std::cout << "WARNING: An empty command group is submitted." << std::endl; @@ -855,6 +863,26 @@ id<2> handler::computeFallbackKernelBounds(size_t Width, size_t Height) { return id<2>{std::min(ItemLimit[0], Height), std::min(ItemLimit[1], Width)}; } +void handler::ext_intel_read_host_pipe(const std::string &Name, void *Ptr, + size_t Size, bool Block) { + MImpl->HostPipeName = Name; + MImpl->HostPipePtr = Ptr; + MImpl->HostPipeTypeSize = Size; + MImpl->HostPipeBlocking = Block; + MImpl->HostPipeRead = 1; + setType(detail::CG::ReadWriteHostPipe); +} + +void handler::ext_intel_write_host_pipe(const std::string &Name, void *Ptr, + size_t Size, bool Block) { + MImpl->HostPipeName = Name; + MImpl->HostPipePtr = Ptr; + MImpl->HostPipeTypeSize = Size; + MImpl->HostPipeBlocking = Block; + MImpl->HostPipeRead = 0; + setType(detail::CG::ReadWriteHostPipe); +} + void handler::memcpyToDeviceGlobal(const void *DeviceGlobalPtr, const void *Src, bool IsDeviceImageScoped, size_t NumBytes, size_t Offset) { @@ -884,8 +912,7 @@ handler::getContextImplPtr() const { return MQueue->getContextImplPtr(); } -void handler::setKernelCacheConfig( - detail::RT::PiKernelCacheConfig Config) { +void handler::setKernelCacheConfig(detail::RT::PiKernelCacheConfig Config) { MImpl->MKernelCacheConfig = Config; } diff --git a/sycl/test/abi/pi_level_zero_symbol_check.dump b/sycl/test/abi/pi_level_zero_symbol_check.dump index fbefe601f3675..4363024cda8a5 100644 --- a/sycl/test/abi/pi_level_zero_symbol_check.dump +++ b/sycl/test/abi/pi_level_zero_symbol_check.dump @@ -89,6 +89,8 @@ piextContextSetExtendedDeleter piextDeviceCreateWithNativeHandle piextDeviceGetNativeHandle piextDeviceSelectBinary +piextEnqueueReadHostPipe +piextEnqueueWriteHostPipe piextEventCreateWithNativeHandle piextEventGetNativeHandle piextGetDeviceFunctionPointer diff --git a/sycl/test/abi/pi_opencl_symbol_check.dump b/sycl/test/abi/pi_opencl_symbol_check.dump index 7925dfcbc6b53..9707c21163b95 100644 --- a/sycl/test/abi/pi_opencl_symbol_check.dump +++ b/sycl/test/abi/pi_opencl_symbol_check.dump @@ -38,6 +38,8 @@ piextContextGetNativeHandle piextDeviceCreateWithNativeHandle piextDeviceGetNativeHandle piextDeviceSelectBinary +piextEnqueueReadHostPipe +piextEnqueueWriteHostPipe piextEventCreateWithNativeHandle piextGetDeviceFunctionPointer piextKernelCreateWithNativeHandle diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 480938399db63..d96b589b9a846 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3655,6 +3655,7 @@ _ZN4sycl3_V120aligned_alloc_sharedEmmRKNS0_6deviceERKNS0_7contextERKNS0_6detail1 _ZN4sycl3_V122accelerator_selector_vERKNS0_6deviceE _ZN4sycl3_V13ext5intel12experimental15online_compilerILNS3_15source_languageE0EE7compileIJSt6vectorINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESaISE_EEEEES8_IhSaIhEERKSE_DpRKT_ _ZN4sycl3_V13ext5intel12experimental15online_compilerILNS3_15source_languageE1EE7compileIJSt6vectorINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESaISE_EEEEES8_IhSaIhEERKSE_DpRKT_ +_ZN4sycl3_V13ext5intel12experimental9pipe_base13get_pipe_nameB5cxx11EPKv _ZN4sycl3_V13ext6oneapi10level_zero10make_eventERKNS0_7contextEmb _ZN4sycl3_V13ext6oneapi10level_zero10make_queueERKNS0_7contextERKNS0_6deviceEmb _ZN4sycl3_V13ext6oneapi10level_zero10make_queueERKNS0_7contextEmb @@ -3823,6 +3824,7 @@ _ZN4sycl3_V16detail13MemoryManager7releaseESt10shared_ptrINS1_12context_implEEPN _ZN4sycl3_V16detail13MemoryManager8allocateESt10shared_ptrINS1_12context_implEEPNS1_11SYCLMemObjIEbPvSt6vectorIS3_INS1_10event_implEESaISB_EERP9_pi_event _ZN4sycl3_V16detail13MemoryManager8copy_usmEPKvSt10shared_ptrINS1_10queue_implEEmPvSt6vectorIP9_pi_eventSaISB_EEPSB_ _ZN4sycl3_V16detail13MemoryManager8fill_usmEPvSt10shared_ptrINS1_10queue_implEEmiSt6vectorIP9_pi_eventSaIS9_EEPS9_ +_ZN4sycl3_V16detail13host_pipe_map3addEPKvPKc _ZN4sycl3_V16detail13make_platformEmNS0_7backendE _ZN4sycl3_V16detail13select_deviceERKSt8functionIFiRKNS0_6deviceEEE _ZN4sycl3_V16detail13select_deviceERKSt8functionIFiRKNS0_6deviceEEERKNS0_7contextE @@ -3964,16 +3966,18 @@ _ZN4sycl3_V17handler19supportsUSMMemset2DEv _ZN4sycl3_V17handler20DisableRangeRoundingEv _ZN4sycl3_V17handler20associateWithHandlerEPNS0_6detail16AccessorBaseHostENS0_6access6targetE _ZN4sycl3_V17handler20memcpyToDeviceGlobalEPKvS3_bmm +_ZN4sycl3_V17handler20setKernelCacheConfigE23_pi_kernel_cache_config _ZN4sycl3_V17handler20setStateSpecConstSetEv _ZN4sycl3_V17handler22ext_oneapi_fill2d_implEPvmPKvmmm _ZN4sycl3_V17handler22memcpyFromDeviceGlobalEPvPKvbmm _ZN4sycl3_V17handler22setHandlerKernelBundleENS0_6kernelE _ZN4sycl3_V17handler22setHandlerKernelBundleERKSt10shared_ptrINS0_6detail18kernel_bundle_implEE _ZN4sycl3_V17handler22verifyUsedKernelBundleERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE -_ZN4sycl3_V17handler20setKernelCacheConfigE23_pi_kernel_cache_config _ZN4sycl3_V17handler24GetRangeRoundingSettingsERmS2_S2_ +_ZN4sycl3_V17handler24ext_intel_read_host_pipeERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPvmb _ZN4sycl3_V17handler24ext_oneapi_memcpy2d_implEPvmPKvmmm _ZN4sycl3_V17handler24ext_oneapi_memset2d_implEPvmimm +_ZN4sycl3_V17handler25ext_intel_write_host_pipeERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPvmb _ZN4sycl3_V17handler27computeFallbackKernelBoundsEmm _ZN4sycl3_V17handler28extractArgsAndReqsFromLambdaEPcmPKNS0_6detail19kernel_param_desc_tEb _ZN4sycl3_V17handler28setStateExplicitKernelBundleEv diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 8168cc42e5b9d..56559c8fdba87 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -755,6 +755,7 @@ ?accessGlobalFlushBuf@stream_impl@detail@_V1@sycl@@QEAA?AV?$accessor@D$00$0EAC@$0HNO@$0A@V?$accessor_property_list@$$V@oneapi@ext@_V1@sycl@@@34@AEAVhandler@34@@Z ?accessGlobalOffset@stream_impl@detail@_V1@sycl@@QEAA?AV?$accessor@I$00$0EAF@$0HNO@$0A@V?$accessor_property_list@$$V@oneapi@ext@_V1@sycl@@@34@AEAVhandler@34@@Z ?add@device_global_map@detail@_V1@sycl@@YAXPEBXPEBD@Z +?add@host_pipe_map@detail@_V1@sycl@@YAXPEBXPEBD@Z ?addHostAccessorAndWait@detail@_V1@sycl@@YAXPEAVAccessorImplHost@123@@Z ?addInteropObject@buffer_impl@detail@_V1@sycl@@QEBAXAEAV?$vector@_KV?$allocator@_K@std@@@std@@@Z ?addOrReplaceAccessorProperties@SYCLMemObjT@detail@_V1@sycl@@QEAAXAEBVproperty_list@34@@Z @@ -841,6 +842,8 @@ ?end@exception_list@_V1@sycl@@QEBA?AV?$_Vector_const_iterator@V?$_Vector_val@U?$_Simple_types@Vexception_ptr@std@@@std@@@std@@@std@@XZ ?end@kernel_bundle_plain@detail@_V1@sycl@@IEBAPEBVdevice_image_plain@234@XZ ?ext_codeplay_supports_fusion@queue@_V1@sycl@@QEBA_NXZ +?ext_intel_read_host_pipe@handler@_V1@sycl@@AEAAXAEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@PEAX_K_N@Z +?ext_intel_write_host_pipe@handler@_V1@sycl@@AEAAXAEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@PEAX_K_N@Z ?ext_oneapi_barrier@handler@_V1@sycl@@QEAAXAEBV?$vector@Vevent@_V1@sycl@@V?$allocator@Vevent@_V1@sycl@@@std@@@std@@@Z ?ext_oneapi_barrier@handler@_V1@sycl@@QEAAXXZ ?ext_oneapi_empty@queue@_V1@sycl@@QEBA_NXZ @@ -1002,6 +1005,7 @@ ?get_max_statement_size@stream@_V1@sycl@@QEBA_KXZ ?get_max_statement_size@stream_impl@detail@_V1@sycl@@QEBA_KXZ ?get_name@kernel_id@_V1@sycl@@QEBAPEBDXZ +?get_pipe_name@pipe_base@experimental@intel@ext@_V1@sycl@@KA?AV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@PEBX@Z ?get_pitch@image_impl@detail@_V1@sycl@@QEBA?AV?$range@$01@34@XZ ?get_pitch@image_plain@detail@_V1@sycl@@IEBA?AV?$range@$01@34@XZ ?get_platform@context@_V1@sycl@@QEBA?AVplatform@23@XZ diff --git a/sycl/test/extensions/properties/properties_pipe.cpp b/sycl/test/extensions/properties/properties_pipe.cpp new file mode 100644 index 0000000000000..70e2c30078db3 --- /dev/null +++ b/sycl/test/extensions/properties/properties_pipe.cpp @@ -0,0 +1,167 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s +// expected-no-diagnostics + +#include + +#include + +using namespace sycl::ext; + +constexpr sycl::ext::intel::experimental::protocol_name TestProtocol = + sycl::ext::intel::experimental::protocol_name::AVALON_STREAMING; + +int main() { + // Check that is_property_key is correctly specialized. + static_assert(sycl::ext::oneapi::experimental::is_property_key< + sycl::ext::intel::experimental::min_capacity_key>::value); + static_assert(sycl::ext::oneapi::experimental::is_property_key< + sycl::ext::intel::experimental::ready_latency_key>::value); + static_assert(sycl::ext::oneapi::experimental::is_property_key< + sycl::ext::intel::experimental::bits_per_symbol_key>::value); + static_assert(sycl::ext::oneapi::experimental::is_property_key< + sycl::ext::intel::experimental::uses_valid_key>::value); + static_assert(sycl::ext::oneapi::experimental::is_property_key< + sycl::ext::intel::experimental::in_csr_key>::value); + static_assert( + sycl::ext::oneapi::experimental::is_property_key< + sycl::ext::intel::experimental::first_symbol_in_high_order_bits_key>:: + value); + static_assert(sycl::ext::oneapi::experimental::is_property_key< + sycl::ext::intel::experimental::protocol_key>::value); + + // Check that is_property_value is correctly specialized. + static_assert( + sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental::min_capacity<3>)>::value); + static_assert( + sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental::ready_latency<3>)>::value); + static_assert( + sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental::bits_per_symbol<3>)>::value); + + static_assert( + sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental::uses_valid)>::value); + static_assert( + sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental::uses_valid_on)>::value); + static_assert( + sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental::uses_valid_off)>::value); + static_assert(sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental::in_csr)>::value); + static_assert(sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental::in_csr_on)>::value); + static_assert(sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental::in_csr_off)>::value); + + static_assert(sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental:: + first_symbol_in_high_order_bits)>::value); + static_assert(sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental:: + first_symbol_in_high_order_bits_on)>::value); + static_assert(sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental:: + first_symbol_in_high_order_bits_off)>::value); + + static_assert( + sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental::protocol)>:: + value); + static_assert( + sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental::protocol_avalon_streaming)>:: + value); + static_assert(sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental:: + protocol_avalon_streaming_uses_ready)>::value); + static_assert( + sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental::protocol_avalon_mm)>::value); + static_assert(sycl::ext::oneapi::experimental::is_property_value< + decltype(sycl::ext::intel::experimental:: + protocol_avalon_mm_uses_ready)>::value); + + // Checks that fully specialized properties are the same as the templated + // variants. + static_assert(std::is_same_v< + decltype(sycl::ext::intel::experimental::uses_valid_on), + decltype(sycl::ext::intel::experimental::uses_valid)>); + static_assert( + std::is_same_v)>); + static_assert( + std::is_same_v)>); + static_assert( + std::is_same_v< + decltype(sycl::ext::intel::experimental::protocol_avalon_streaming), + decltype(sycl::ext::intel::experimental::protocol)>); + static_assert( + std::is_same_v)>); + static_assert( + std::is_same_v< + decltype(sycl::ext::intel::experimental::protocol_avalon_mm), + decltype(sycl::ext::intel::experimental::protocol< + sycl::ext::intel::experimental::protocol_name::AVALON_MM>)>); + static_assert( + std::is_same_v)>); + + // Check that property lists will accept the new properties. + using P = decltype(sycl::ext::oneapi::experimental::properties( + sycl::ext::intel::experimental::min_capacity<0>, + sycl::ext::intel::experimental::ready_latency<1>, + sycl::ext::intel::experimental::bits_per_symbol<2>, + sycl::ext::intel::experimental::uses_valid, + sycl::ext::intel::experimental::in_csr, + sycl::ext::intel::experimental::first_symbol_in_high_order_bits_off, + sycl::ext::intel::experimental::protocol_avalon_streaming)); + static_assert(sycl::ext::oneapi::experimental::is_property_list_v

); + static_assert( + P::has_property()); + static_assert( + P::has_property()); + static_assert( + P::has_property()); + static_assert( + P::has_property()); + static_assert(P::has_property()); + static_assert(P::has_property()); + static_assert( + P::has_property()); + + static_assert( + P::get_property() == + sycl::ext::intel::experimental::min_capacity<0>); + static_assert( + P::get_property() == + sycl::ext::intel::experimental::ready_latency<1>); + static_assert( + P::get_property() == + sycl::ext::intel::experimental::bits_per_symbol<2>); + static_assert( + P::get_property() == + sycl::ext::intel::experimental::uses_valid); + static_assert(P::get_property() == + sycl::ext::intel::experimental::in_csr); + static_assert( + P::get_property() == + sycl::ext::intel::experimental::first_symbol_in_high_order_bits_off); + static_assert( + P::get_property() == + sycl::ext::intel::experimental::protocol_avalon_streaming); +} diff --git a/sycl/unittests/CMakeLists.txt b/sycl/unittests/CMakeLists.txt index 9ec74d34b8d2e..55c8470e7d133 100644 --- a/sycl/unittests/CMakeLists.txt +++ b/sycl/unittests/CMakeLists.txt @@ -40,6 +40,7 @@ add_subdirectory(scheduler) add_subdirectory(stream) add_subdirectory(SYCL2020) add_subdirectory(thread_safety) +add_subdirectory(pipes) add_subdirectory(program_manager) add_subdirectory(assert) add_subdirectory(Extensions) diff --git a/sycl/unittests/helpers/PiImage.hpp b/sycl/unittests/helpers/PiImage.hpp index 5a536c96a569c..dd61ed48ffe3d 100644 --- a/sycl/unittests/helpers/PiImage.hpp +++ b/sycl/unittests/helpers/PiImage.hpp @@ -479,6 +479,25 @@ inline PiProperty makeDeviceGlobalInfo(const std::string &Name, return Prop; } +/// Utility function to create a host pipe info property. +/// +/// \param Name is the name of the hostpipe name. +/// \param TypeSize is the size of the underlying type in the hostpipe. +/// decorated. +inline PiProperty makeHostPipeInfo(const std::string &Name, + const uint32_t TypeSize) { + constexpr size_t BYTES_FOR_SIZE = 8; + const std::uint64_t BytesForArgs = sizeof(std::uint32_t); + std::vector DescData; + DescData.resize(BYTES_FOR_SIZE + BytesForArgs); + std::memcpy(DescData.data(), &BytesForArgs, sizeof(BytesForArgs)); + std::memcpy(DescData.data() + BYTES_FOR_SIZE, &TypeSize, sizeof(TypeSize)); + + PiProperty Prop{Name, DescData, PI_PROPERTY_TYPE_BYTE_ARRAY}; + + return Prop; +} + /// Utility function to add aspects to property set. inline PiProperty makeAspectsProp(const std::vector &Aspects) { const size_t BYTES_FOR_SIZE = 8; diff --git a/sycl/unittests/helpers/PiMockPlugin.hpp b/sycl/unittests/helpers/PiMockPlugin.hpp index c2ac5e6863b8b..0faa30e9b7407 100644 --- a/sycl/unittests/helpers/PiMockPlugin.hpp +++ b/sycl/unittests/helpers/PiMockPlugin.hpp @@ -1138,3 +1138,19 @@ inline pi_result mock_piGetDeviceAndHostTimer(pi_device device, } return PI_SUCCESS; } + +inline pi_result mock_piextEnqueueReadHostPipe( + pi_queue queue, pi_program program, const char *pipe_symbol, + pi_bool blocking, void *ptr, size_t size, pi_uint32 num_events_in_waitlist, + const pi_event *events_waitlist, pi_event *event) { + *event = createDummyHandle(); + return PI_SUCCESS; +} + +inline pi_result mock_piextEnqueueWriteHostPipe( + pi_queue queue, pi_program program, const char *pipe_symbol, + pi_bool blocking, void *ptr, size_t size, pi_uint32 num_events_in_waitlist, + const pi_event *events_waitlist, pi_event *event) { + *event = createDummyHandle(); + return PI_SUCCESS; +} diff --git a/sycl/unittests/pipes/CMakeLists.txt b/sycl/unittests/pipes/CMakeLists.txt new file mode 100644 index 0000000000000..58069920f5cb4 --- /dev/null +++ b/sycl/unittests/pipes/CMakeLists.txt @@ -0,0 +1,8 @@ +set(CMAKE_CXX_EXTENSIONS OFF) + +add_sycl_unittest(PipeTests OBJECT + host_pipe_registration.cpp +) + +add_dependencies(PipeTests sycl) +target_include_directories(PipeTests PRIVATE SYSTEM ${sycl_inc_dir}) diff --git a/sycl/unittests/pipes/host_pipe_registration.cpp b/sycl/unittests/pipes/host_pipe_registration.cpp new file mode 100644 index 0000000000000..dd18d98d864d1 --- /dev/null +++ b/sycl/unittests/pipes/host_pipe_registration.cpp @@ -0,0 +1,162 @@ +//==-------------- host_pipe_registration.cpp - Host pipe tests------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include + +#include +#include +#include +#include +#include +#include + +template class TestKernel; + +namespace sycl { +__SYCL_INLINE_VER_NAMESPACE(_V1) { +namespace detail { +template struct KernelInfo> { + static constexpr unsigned getNumParams() { return 0; } + static const kernel_param_desc_t &getParamDesc(int) { + static kernel_param_desc_t Dummy; + return Dummy; + } + static constexpr const char *getName() { return "TestKernel"; } + static constexpr bool isESIMD() { return false; } + static constexpr bool callsThisItem() { return false; } + static constexpr bool callsAnyThisFreeFunction() { return false; } + static constexpr int64_t getKernelSize() { return KernelSize; } +}; + +} // namespace detail +} // __SYCL_INLINE_VER_NAMESPACE(_V1) +} // namespace sycl + +using namespace sycl; +using default_pipe_properties = + decltype(sycl::ext::oneapi::experimental::properties( + sycl::ext::intel::experimental::uses_valid)); + +class PipeID; +using Pipe = sycl::ext::intel::experimental::pipe; + +static sycl::unittest::PiImage generateDefaultImage() { + using namespace sycl::unittest; + + sycl::detail::host_pipe_map::add(Pipe::get_host_ptr(), + "test_host_pipe_unique_id"); + + PiPropertySet PropSet; + PiProperty HostPipeInfo = + makeHostPipeInfo("test_host_pipe_unique_id", sizeof(int)); + PropSet.insert(__SYCL_PI_PROPERTY_SET_SYCL_HOST_PIPES, + PiArray{std::move(HostPipeInfo)}); + + std::vector Bin{0, 1, 2, 3, 4, 5}; // Random data + + PiArray Entries = makeEmptyKernels({"TestKernel"}); + + PiImage Img{PI_DEVICE_BINARY_TYPE_SPIRV, // Format + __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64, // DeviceTargetSpec + "", // Compile options + "", // Link options + std::move(Bin), + std::move(Entries), + std::move(PropSet)}; + + return Img; +} + +pi_event READ = reinterpret_cast(0); +pi_event WRITE = reinterpret_cast(1); +static constexpr int PipeReadVal = 8; +static int PipeWriteVal = 0; +pi_result redefinedEnqueueReadHostPipe(pi_queue, pi_program, const char *, + pi_bool, void *ptr, size_t, pi_uint32, + const pi_event *, pi_event *event) { + *(((int *)ptr)) = PipeReadVal; + return PI_SUCCESS; +} +pi_result redefinedEnqueueWriteHostPipe(pi_queue, pi_program, const char *, + pi_bool, void *ptr, size_t, pi_uint32, + const pi_event *, pi_event *event) { + PipeWriteVal = 9; + return PI_SUCCESS; +} + +pi_result after_piDeviceGetInfo(pi_device device, pi_device_info param_name, + size_t param_value_size, void *param_value, + size_t *param_value_size_ret) { + constexpr char MockSupportedExtensions[] = + "cl_khr_fp64 cl_khr_fp16 cl_khr_il_program " + "cl_intel_program_scope_host_pipe"; + switch (param_name) { + case PI_DEVICE_INFO_EXTENSIONS: { + if (param_value) { + assert(param_value_size >= sizeof(MockSupportedExtensions)); + std::memcpy(param_value, MockSupportedExtensions, + sizeof(MockSupportedExtensions)); + } + if (param_value_size_ret) + *param_value_size_ret = sizeof(MockSupportedExtensions); + return PI_SUCCESS; + } + default:; + } + return PI_SUCCESS; +} + +void preparePiMock(unittest::PiMock &Mock) { + Mock.redefine( + redefinedEnqueueReadHostPipe); + Mock.redefine( + redefinedEnqueueWriteHostPipe); +} + +class PipeTest : public ::testing::Test { +public: + PipeTest() : Mock{}, Plt{Mock.getPlatform()} {} + +protected: + void SetUp() override { + preparePiMock(Mock); + const sycl::device Dev = Plt.get_devices()[0]; + sycl::context Ctx{Dev}; + sycl::queue Q{Ctx, Dev}; + ctx = Ctx; + q = Q; + } + +protected: + unittest::PiMock Mock; + sycl::platform Plt; + context ctx; + queue q; +}; + +TEST_F(PipeTest, Basic) { + // Fake extension + Mock.redefineAfter( + after_piDeviceGetInfo); + + // Device registration + static sycl::unittest::PiImage Img = generateDefaultImage(); + static sycl::unittest::PiImageArray<1> ImgArray{&Img}; + + // Testing read + int HostPipeReadData; + HostPipeReadData = Pipe::read(q); + assert(HostPipeReadData == PipeReadVal); + + // Testing write + int HostPipeWriteData = 9; + Pipe::write(q, HostPipeWriteData); + assert(PipeWriteVal == 9); +}