[SYCL] Host pipe runtime implementation #7468
Merged
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
zibaiwan
force-pushed
the
sycl-hostpipe-runtime
branch
5 times, most recently
from
00c823c to
a754b91
Compare
zibaiwan
changed the title
zibaiwan
force-pushed
the
sycl-hostpipe-runtime
branch
from
a754b91 to
13e24ee
Compare
zibaiwan
force-pushed
the
sycl-hostpipe-runtime
branch
4 times, most recently
from
e590e72 to
9123db6
Compare
zibaiwan
changed the title
asudarsa
reviewed
Jan 11, 2023
|
@steffenlarsen, can you please take a look at this PR when you have a chance? Any feedback is much appreciated! |
|
Converted to draft to avoid accidental merge. |
zibaiwan
force-pushed
the
sycl-hostpipe-runtime
branch
from
ff1912a to
e6e3e99
Compare
bader
pushed a commit
that referenced
this pull request
Feb 7, 2023
…8009) Implementation of host pipes outlined in the design document in this PR: #5850 1. Generation a unique pipe id for GVs marked with the new "sycl-host-pipe" attribute. Id generation utilizes the same method as used for name generation for device global. 2. Added a host pipe map to map the addresses of marked GVs with the unique id. This host pipe map is generated by a constructor and method calls added to the header and footer. 3. Modified the sycl-post-link tool to generate compile time properties metadata for these GVs. This metadata contains the unique id generated for the GV to be consumed by the device backend compiler. PR for accompanying runtime changes: #7468 --------- Co-authored-by: Alexey Sachkov <alexey.sachkov@intel.com>
|
/testwin |
|
/testwin |
|
/testwin |
|
/testwin |
|
The CI for the latest code commit passed: https://github.com/intel/llvm/actions/runs/4556840235/jobs/8043028160, except for the known failures. @steffenlarsen , this PR is ready to be merged now. Thank you very much for your help. |
steffenlarsen
added a commit
that referenced
this pull request
Mar 30, 2023
#8789) A continuation of #5838. Accompanying runtime change is #7468. Add a memory order parameter to device-side read/write members and default to sycl::memory_order::seq_cst. This parameter is in place but not being used at this moment, it's intended for the future work. Add host pipe read/write members with additional sycl::queue parameter. --------- Co-authored-by: Steffen Larsen <steffen.larsen@intel.com>
steffenlarsen
added a commit
to steffenlarsen/llvm
that referenced
this pull request
Mar 30, 2023
This commit fixes unusued variable warnings after intel#7468. Additionally it replaces asserts in the unittests with EXPECT_EQ. Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
steffenlarsen
added a commit
that referenced
this pull request
Mar 30, 2023
This commit fixes unusued variable warnings after #7468. Additionally it replaces asserts in the unittests with EXPECT_EQ. Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
steffenlarsen
pushed a commit
that referenced
this pull request
Mar 30, 2023
This patch fixes post-commit failure introduced in #7468: ``` /Users/runner/work/llvm/llvm/src/sycl/source/detail/program_manager/program_manager.cpp:1324:9: error: 'lock_guard' may not intend to support class template argument deduction [-Werror,-Wctad-maybe-unsupported] std::lock_guard HostPipesGuard(m_HostPipesMutex); ^ ```
steffenlarsen
added a commit
that referenced
this pull request
May 2, 2023
Pipe properties support was added in this PR: #7468 (sycl/ext/intel/experimental/pipe_properties.hpp). This is the accompanying spec. --------- Co-authored-by: Steffen Larsen <steffen.larsen@intel.com>
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Disclaimer:
This work is a continuation of a previous approved Sherry's PR: #5766 and her draft work #5894.
We are implementing Hostpipes based on the Spec here. OpenCL spec is here.
The following is the outline of the design.
2.The frontend calls the registration of device image, maybe similar to this code: which is where the host pipe information is available. This is where the mapping from host pipe name to host pipe pointer is extracted.
The frontend also calls
host_pipe_map::addto register/initialize the host pipes. This is the new function. We are not sure about the ordering of registration of device image / registration of host pipe. for which ever one that comes later, it need to initialize the remaining attribute of host pipe (such as its properties).The opencl runtime will need to get a cl_program object, which is typically not available until the first kernel launch. To get a program object early on, the host pipe name/pointer to device image mapping is cached during registration. And when the specific host pipe is needed, build the program and get its ocl runtime representation. This is done in the first couple commits.
Since a host pipe read/write need to depend on other write operation finishing before it (including the inter kernel write). This means the pipe needs to know the dependency of kernel execution. For this reason, the host pipe read and write ocl function cannot be called with no dep event. therefore, it is implemented with handler , which is aware of the event that it is supposed to wait upon. This is done in the "Register new command group .." commit.
Unit test