diff --git a/buildbot/dependency.py b/buildbot/dependency.py index ec8dc9d423fe4..8ab0e08bb9c16 100644 --- a/buildbot/dependency.py +++ b/buildbot/dependency.py @@ -49,8 +49,8 @@ def do_dependency(args): # fetch OpenCL headers ocl_header_dir = os.path.join(args.obj_dir, "OpenCL-Headers") if not os.path.isdir(ocl_header_dir): - clone_cmd = ["git", "clone", "https://github.com/KhronosGroup/OpenCL-Headers", - "OpenCL-Headers", "-b", "main"] + clone_cmd = ["git", "clone", "https://github.com/sherry-yuan/OpenCL-Headers", + "OpenCL-Headers", "-b", "host_pipe"] # TODO: Remove change once upstream header changed subprocess.check_call(clone_cmd, cwd=args.obj_dir) else: fetch_cmd = ["git", "pull", "--ff", "--ff-only", "origin"] @@ -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", "1f2cb76195fb77be7c0b4d811ecff244c864d2e2"] # TODO: Remove change once upstream header changed subprocess.check_call(checkout_cmd, cwd=ocl_header_dir) # fetch and build OpenCL ICD loader diff --git a/opencl/CMakeLists.txt b/opencl/CMakeLists.txt index ee6e76829537a..00bc0c590de5b 100644 --- a/opencl/CMakeLists.txt +++ b/opencl/CMakeLists.txt @@ -14,13 +14,13 @@ endif() # Repo URLs set(OCL_HEADERS_REPO - "https://github.com/KhronosGroup/OpenCL-Headers.git") + "https://github.com/sherry-yuan/OpenCL-Headers.git") set(OCL_LOADER_REPO "https://github.com/KhronosGroup/OpenCL-ICD-Loader.git") # Repo tags/hashes -set(OCL_HEADERS_TAG dcd5bede6859d26833cd85f0d6bbcee7382dc9b3) +set(OCL_HEADERS_TAG 1f2cb76195fb77be7c0b4d811ecff244c864d2e2) set(OCL_LOADER_TAG 5f8249691ec8c25775789498951f8e9eb62c201d) # OpenCL Headers diff --git a/sycl/include/CL/sycl/detail/pi.def b/sycl/include/CL/sycl/detail/pi.def index c9a68c6cadec3..a0ee268af01e5 100644 --- a/sycl/include/CL/sycl/detail/pi.def +++ b/sycl/include/CL/sycl/detail/pi.def @@ -130,6 +130,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/CL/sycl/detail/pi.h b/sycl/include/CL/sycl/detail/pi.h index 16482c07e1fa5..a9070dba9bc40 100644 --- a/sycl/include/CL/sycl/detail/pi.h +++ b/sycl/include/CL/sycl/detail/pi.h @@ -1769,6 +1769,56 @@ __SYCL_EXPORT pi_result piextUSMGetMemAllocInfo( pi_context context, const void *ptr, pi_mem_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret); +/// +// 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 /// to access some device code functionality implemented in/behind the plugin. diff --git a/sycl/plugins/cuda/pi_cuda.cpp b/sycl/plugins/cuda/pi_cuda.cpp index a29d4152ec4c9..7096350c0ea37 100644 --- a/sycl/plugins/cuda/pi_cuda.cpp +++ b/sycl/plugins/cuda/pi_cuda.cpp @@ -4943,6 +4943,43 @@ pi_result cuda_piextUSMGetMemAllocInfo(pi_context context, const void *ptr, 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; + + cl::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; + + cl::sycl::detail::pi::die("cuda_piextEnqueueWriteHostPipe not implemented"); + return {}; +} + // This API is called by Sycl RT to notify the end of the plugin lifetime. // TODO: add a global variable lifetime management code here (see // pi_level_zero.cpp for reference) Currently this is just a NOOP. @@ -5085,6 +5122,10 @@ pi_result piPluginInit(pi_plugin *PluginInit) { _PI_CL(piextUSMEnqueueMemAdvise, cuda_piextUSMEnqueueMemAdvise) _PI_CL(piextUSMGetMemAllocInfo, cuda_piextUSMGetMemAllocInfo) + // Host Pipe + _PI_CL(piextEnqueueReadHostPipe, cuda_piextEnqueueReadHostPipe) + _PI_CL(piextEnqueueWriteHostPipe, cuda_piextEnqueueWriteHostPipe) + _PI_CL(piextKernelSetArgMemObj, cuda_piextKernelSetArgMemObj) _PI_CL(piextKernelSetArgSampler, cuda_piextKernelSetArgSampler) _PI_CL(piTearDown, cuda_piTearDown) diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp index 03f0061f8a8f0..cbee730cb3897 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -1768,6 +1768,25 @@ pi_result piextUSMGetMemAllocInfo(pi_context, const void *, pi_mem_info, size_t, 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 d565fbb87c838..b45501d49a2c4 100644 --- a/sycl/plugins/hip/pi_hip.cpp +++ b/sycl/plugins/hip/pi_hip.cpp @@ -4833,6 +4833,45 @@ pi_result hip_piextUSMGetMemAllocInfo(pi_context context, const void *ptr, return result; } +/// 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; + + cl::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; + + cl::sycl::detail::pi::die("hip_piextEnqueueWriteHostPipe not implemented"); + return {}; +} + // This API is called by Sycl RT to notify the end of the plugin lifetime. // TODO: add a global variable lifetime management code here (see // pi_level_zero.cpp for reference) Currently this is just a NOOP. @@ -4974,6 +5013,10 @@ pi_result piPluginInit(pi_plugin *PluginInit) { _PI_CL(piextUSMEnqueueMemAdvise, hip_piextUSMEnqueueMemAdvise) _PI_CL(piextUSMGetMemAllocInfo, hip_piextUSMGetMemAllocInfo) + // Host Pipe + _PI_CL(piextEnqueueReadHostPipe, hip_piextEnqueueReadHostPipe) + _PI_CL(piextEnqueueWriteHostPipe, hip_piextEnqueueWriteHostPipe) + _PI_CL(piextKernelSetArgMemObj, hip_piextKernelSetArgMemObj) _PI_CL(piextKernelSetArgSampler, hip_piextKernelSetArgSampler) _PI_CL(piTearDown, hip_piTearDown) diff --git a/sycl/plugins/level_zero/pi_level_zero.cpp b/sycl/plugins/level_zero/pi_level_zero.cpp index e0cb87dbde8a1..7ee022a1201d5 100644 --- a/sycl/plugins/level_zero/pi_level_zero.cpp +++ b/sycl/plugins/level_zero/pi_level_zero.cpp @@ -3999,12 +3999,12 @@ pi_result piProgramGetInfo(pi_program Program, pi_program_info ParamName, uint32_t Count = 0; ZE_CALL(zeModuleGetKernelNames, (Program->ZeModule, &Count, nullptr)); std::unique_ptr PNames(new const char *[Count]); - ZE_CALL(zeModuleGetKernelNames, - (Program->ZeModule, &Count, PNames.get())); - for (uint32_t I = 0; I < Count; ++I) { - PINames += (I > 0 ? ";" : ""); - PINames += PNames[I]; - } + ZE_CALL(zeModuleGetKernelNames, + (Program->ZeModule, &Count, PNames.get())); + for (uint32_t I = 0; I < Count; ++I) { + PINames += (I > 0 ? ";" : ""); + PINames += PNames[I]; + } } else { return PI_INVALID_PROGRAM; } @@ -7607,6 +7607,72 @@ pi_result piextUSMGetMemAllocInfo(pi_context Context, const void *Ptr, return PI_SUCCESS; } +/// 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_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_INVALID_QUEUE); + + die("piextEnqueueWriteHostPipe: not implemented"); + return {}; +} + pi_result piKernelSetExecInfo(pi_kernel Kernel, pi_kernel_exec_info ParamName, size_t ParamValueSize, const void *ParamValue) { (void)ParamValueSize; diff --git a/sycl/plugins/opencl/pi_opencl.cpp b/sycl/plugins/opencl/pi_opencl.cpp index 8075f169347e4..a314ce0e0c546 100644 --- a/sycl/plugins/opencl/pi_opencl.cpp +++ b/sycl/plugins/opencl/pi_opencl.cpp @@ -68,6 +68,9 @@ CONSTFIX char clSetProgramSpecializationConstantName[] = "clSetProgramSpecializationConstant"; CONSTFIX char clGetDeviceFunctionPointerName[] = "clGetDeviceFunctionPointerINTEL"; +// Names of host pipe functions queried from OpenCL +CONSTFIX char clEnqueueReadHostPipeName[] = "clEnqueueReadHostPipeIntelFPGA"; +CONSTFIX char clEnqueueWriteHostPipeName[] = "clEnqueueWriteHostPipeIntelFPGA"; #undef CONSTFIX @@ -1305,6 +1308,64 @@ pi_result piextUSMGetMemAllocInfo(pi_context context, const void *ptr, return RetVal; } +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); + } + + clEnqueueReadHostPipeIntelFPGA_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); + } + + clEnqueueWriteHostPipeIntelFPGA_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 @@ -1537,6 +1598,9 @@ pi_result piPluginInit(pi_plugin *PluginInit) { _PI_CL(piextUSMEnqueuePrefetch, piextUSMEnqueuePrefetch) _PI_CL(piextUSMEnqueueMemAdvise, piextUSMEnqueueMemAdvise) _PI_CL(piextUSMGetMemAllocInfo, piextUSMGetMemAllocInfo) + // Host Pipe + _PI_CL(piextEnqueueReadHostPipe, piextEnqueueReadHostPipe) + _PI_CL(piextEnqueueWriteHostPipe, piextEnqueueWriteHostPipe) _PI_CL(piextKernelSetArgMemObj, piextKernelSetArgMemObj) _PI_CL(piextKernelSetArgSampler, piextKernelSetArgSampler) diff --git a/sycl/test/abi/pi_level_zero_symbol_check.dump b/sycl/test/abi/pi_level_zero_symbol_check.dump index 491eb9d87f155..0443faf3b401e 100644 --- a/sycl/test/abi/pi_level_zero_symbol_check.dump +++ b/sycl/test/abi/pi_level_zero_symbol_check.dump @@ -86,6 +86,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 0e6aaa9dfa7a5..db890125035cb 100644 --- a/sycl/test/abi/pi_opencl_symbol_check.dump +++ b/sycl/test/abi/pi_opencl_symbol_check.dump @@ -33,6 +33,8 @@ piextContextGetNativeHandle piextDeviceCreateWithNativeHandle piextDeviceGetNativeHandle piextDeviceSelectBinary +piextEnqueueReadHostPipe +piextEnqueueWriteHostPipe piextEventCreateWithNativeHandle piextGetDeviceFunctionPointer piextKernelCreateWithNativeHandle