Use native ("USM") pointers for backing buffer allocations #162
Conversation
This file contains 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
![github-actions[bot]](https://avatars.githubusercontent.com/in/15368?s=60&v=4){kind=small}
PeterTh
reviewed
Mar 8, 2023
fknorr
reviewed
Mar 8, 2023
facuMH
reviewed
Mar 9, 2023
|
Re the ComputeCpp CI failure which I'm also seeing in #163 : The backend library doesn't set the |
psalz
force-pushed
the
native-buffers
branch
2 times, most recently
from
4502e6f to
6de986d
Compare
This introduces the new SKIP macro for skipping tests at runtime.
Use native pointers (allocated using `sycl::malloc_device`) instead of relying on SYCL buffers for backing Celerity virtual buffers. This greatly simplifies various aspects of accessors and buffer management while enabling future optimizations. Futhermore, by using native pointers we completely circumvent any dataflow analysis performed by the SYCL runtime.
Since SYCL 2020 does not support multi-dimensional (rectangular) copies for USM pointers, we have to either do it in a loop (slow) or fall back to vendor-specific APIs. This introduces a new "backend" system that does the latter. Currently only "generic" (= SYCL, slow) and CUDA (when using OpenSYCL or DPC++) are supported. Since backends are configuration during compile time, this additionally introduces a new integration testing mechanism for testing backends. This requires Celerity to be built with different CMake options, so the test is implemented as a Python script.
...instead of range and offset.
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.
This transitions Celerity away from using SYCL buffers for backing buffer allocations to essentially device-native memory allocations through
sycl::malloc_deviceandsycl::free, which are part of SYCL 2020's USM APIs.This has several advantages:
Unfortunately there is also a major downside, which is that SYCL's USM capabilities are very much 1-dimensional at the moment.
This is reflected across all USM APIs, and particularly problematic for us is that there is no way of doing 2D/3D rectangular (strided) copies. Dispatching a series of 1D copies in a loop is not a viable solution, as it is extremely slow. Instead, for now (until SYCL Next, hopefully), we have to resort back to doing copies manually using the underlying vendor APIs. This means we need specialized code paths for each of the backends that we want to support efficiently.
Currently I have only implemented a specialized backend for CUDA as well as a generic (and slow) fallback, however adding new backends should be relatively straightforward. Backends are selected dynamically based on the device used. Since SYCL does not yet officially support CUDA as a backend, we also need an implementation-specific mechanism for detecting whether a device is a CUDA device; this is currently supported for OpenSYCL and DPC++.
I wasn't quite sure what the best way of structuring the backend system would be. Loading specialized libraries dynamically at runtime (similar to what OpenSYCL does) seemed overkill, so I've instead opted to make enabling/disabling backends a compile-time option (
-DCELERITY_ENABLE_CUDA_BACKEND).This makes testing somewhat more difficult though, as the behavior of test cases is now affected by how they are compiled, which SYCL implementation is used, and on the hardware available during runtime. While I've added some unit tests that just opportunistically test whatever is available, to cover more of these combinations in our CI setup, I've also created our first true integration test (we previously used our examples as integration tests, which was always suboptimal) in the form of a Python script that compiles and runs different backend configurations.