Skip to content

Add pi extension API for host pipes #5885

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Closed
wants to merge 1 commit into from
Closed
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 3 additions & 3 deletions buildbot/dependency.py
Original file line number Diff line number Diff line change
Expand Up @@ -49,16 +49,16 @@ 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"]
subprocess.check_call(fetch_cmd, cwd=ocl_header_dir)

# 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
Expand Down
4 changes: 2 additions & 2 deletions opencl/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
3 changes: 3 additions & 0 deletions sycl/include/CL/sycl/detail/pi.def
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down
50 changes: 50 additions & 0 deletions sycl/include/CL/sycl/detail/pi.h
Original file line number Diff line number Diff line change
Expand Up @@ -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.
Expand Down
41 changes: 41 additions & 0 deletions sycl/plugins/cuda/pi_cuda.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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.
Expand Down Expand Up @@ -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)
Expand Down
19 changes: 19 additions & 0 deletions sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
43 changes: 43 additions & 0 deletions sycl/plugins/hip/pi_hip.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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.
Expand Down Expand Up @@ -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)
Expand Down
78 changes: 72 additions & 6 deletions sycl/plugins/level_zero/pi_level_zero.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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<const char *[]> 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;
}
Expand Down Expand Up @@ -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;
Expand Down
Loading