Repository navigation
Native extension to bind to cccl c parallel library - #4325
oleksandr-pavlyk merged 58 commits into
Conversation
|
Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually. Contributors can view more details about this message here. |
|
/ok to test |
leofang
left a comment
There was a problem hiding this comment.
(As discussed offline, I'd still like to see how much of the Cython code can be autogenerated. As it stands today this is not scalable as we expand the library.)
🟩 CI finished in 1h 06m: Pass: 100%/1 | Total: 1h 06m | Avg: 1h 06m | Max: 1h 06m
|
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| CUB | |
| Thrust | |
| CUDA Experimental | |
| stdpar | |
| +/- | python |
| CCCL C Parallel Library | |
| Catch2Helper |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| CUB | |
| Thrust | |
| CUDA Experimental | |
| stdpar | |
| +/- | python |
| CCCL C Parallel Library | |
| Catch2Helper |
🏃 Runner counts (total jobs: 1)
| # | Runner |
|---|---|
| 1 | linux-amd64-gpu-rtx2080-latest-1 |
|
Here is the performance data: It is comparing the time of execution of With |
|
I created branch https://github.com/oleksandr-pavlyk/cccl/tree/coexisting-ctypes-and-native-bindings-to-cccl-c-parallel for ease of reference to state where both bindings coexisted and deleted code for legacy binding from this branch. |
There was a problem hiding this comment.
Given that we have dropped the legacy bindings, can we also drop the _cy_ prefix in these files?
There was a problem hiding this comment.
Yes, but I would prefer to keep the leading underscore to avoid the merge_sort identifier retrieved using from cuda.parallel.experimental.algorithms import merge_sort from having a type module, rather than a type function.
There was a problem hiding this comment.
I renamed _cy_${algo_name}.py to ${algo_name}.py.
| from typing import Callable | ||
|
|
||
| import numba | ||
| from numba.cuda.cudadrv import enums | ||
|
|
||
| from .. import _cccl as cccl |
There was a problem hiding this comment.
Drive-by comment, I only took a quick peek here: My vote would be to not change the filenames in this directory. At this level, from a user perspective, it's a technical detail what bindings tool is used.
Changing the names under cuda/parallel/experimental seems fine.
|
/ok to test |
🟩 CI finished in 1h 05m: Pass: 100%/1 | Total: 1h 05m | Avg: 1h 05m | Max: 1h 05m
|
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| CUB | |
| Thrust | |
| CUDA Experimental | |
| stdpar | |
| +/- | python |
| CCCL C Parallel Library | |
| Catch2Helper |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| CUB | |
| Thrust | |
| CUDA Experimental | |
| stdpar | |
| +/- | python |
| CCCL C Parallel Library | |
| Catch2Helper |
🏃 Runner counts (total jobs: 1)
| # | Runner |
|---|---|
| 1 | linux-amd64-gpu-rtx2080-latest-1 |
1c142e5 to
fcc5ed7
Compare
There was a problem hiding this comment.
Can we put all the cdef extern declarations into _bindings.pxd, and reserve _bindings.pyx for the wrapper definitions?
leofang
left a comment
There was a problem hiding this comment.
FWIW I've revoked my blocking review. As noted offline, a big chunk of the bindings can be autogenerated, and the hand-written code will be discarded by then, so I am convinced we should move on as soon as other team members are happy. (Leaving a note here in case I can't respond timely before I return.)
shwina
left a comment
There was a problem hiding this comment.
Looking good. Mainly I'm wondering if it's possible to simplify our handling of iterator state objects by assuming they're convertible to (typed) memoryviews.
| ) | ||
|
|
||
|
|
||
| cdef class IntEnum: |
There was a problem hiding this comment.
For future readers, can we add a note in the code here explaining why Python's native Enum isn't enough for our use case?
There was a problem hiding this comment.
I was thinking it'd be best to give this type a name that's different from the standard library IntEnum, so that it is always obvious which is which. E.g. FastIntEnum?
| return <void *>(ptr_val) | ||
|
|
||
|
|
||
| cdef class IteratorStateView: |
There was a problem hiding this comment.
I don't think I fully understand the need for the following:
IteratorStateViewIteratorState,Pointer
I'm trying to understand: can we simplify by assuming that an iterator state is some Python object that is coercible either to a memoryview or typedmemoryview?
There was a problem hiding this comment.
IteratorStateView has been removed.
IteratorState essentially amounts to memory_view associated with contiguous bytes. The content of that memory represents the state of an iterator. IteratorState is to be used to set the state of CCCL iterators of kind ITERATOR. The size of memory allocation represented by IteratorState is checked against expected size computed from state_type stored in IteratorBase.
Pointer is python object that stores a pointer, but also records reference to the python object that may own the memory the pointer points to. This type is used to set the state of CCCL iterators of kind Pointer. For CCCL iterator of kind POINTER, the size and alignment struct members are not meaningful.
Use of these distinct classes better reflects different semantic of void *state in CCCL iterators of different kinds.
🟩 CI finished in 1h 27m: Pass: 100%/1 | Total: 1h 27m | Avg: 1h 27m | Max: 1h 27m
|
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| CUB | |
| Thrust | |
| CUDA Experimental | |
| stdpar | |
| +/- | python |
| CCCL C Parallel Library | |
| Catch2Helper |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| CUB | |
| Thrust | |
| CUDA Experimental | |
| stdpar | |
| +/- | python |
| CCCL C Parallel Library | |
| Catch2Helper |
🏃 Runner counts (total jobs: 1)
| # | Runner |
|---|---|
| 1 | linux-amd64-gpu-rtx2080-latest-1 |
The NdArrayIterator is an input iterator which traverses the array elements of a strided nd-array in the same order of corresponding flat array, but without making copies. This iterator enables two `test_segmented_reduce_api.py` examples: "segmented-reduce-columnwise-maximum" and "segmented-reduce-multiaxis-sum".
…segmented_reduce_api
alliepiper
left a comment
There was a problem hiding this comment.
Just looked at CMake stuff.
Also directly require Python3 and use Python3_add_library
1. Missing quotes 2. Removed FIND_PACKAGE_HANDLE_STANDARD_ARGS and mark_as_advanced
|
For completeness, when reviewing this PR, I also looked at this |
|
Thank you for good review everybody. I think this PR can be merged once CI's badge of approval arrives. |
🟩 CI finished in 1h 33m: Pass: 100%/1 | Total: 1h 33m | Avg: 1h 33m | Max: 1h 33m
|
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| CUB | |
| Thrust | |
| CUDA Experimental | |
| stdpar | |
| +/- | python |
| CCCL C Parallel Library | |
| Catch2Helper |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| CUB | |
| Thrust | |
| CUDA Experimental | |
| stdpar | |
| +/- | python |
| CCCL C Parallel Library | |
| Catch2Helper |
🏃 Runner counts (total jobs: 1)
| # | Runner |
|---|---|
| 1 | linux-amd64-gpu-rtx2080-latest-1 |
🟩 CI finished in 1h 29m: Pass: 100%/1 | Total: 1h 29m | Avg: 1h 29m | Max: 1h 29m
|
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| CUB | |
| Thrust | |
| CUDA Experimental | |
| stdpar | |
| +/- | python |
| CCCL C Parallel Library | |
| Catch2Helper |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| CUB | |
| Thrust | |
| CUDA Experimental | |
| stdpar | |
| +/- | python |
| CCCL C Parallel Library | |
| Catch2Helper |
🏃 Runner counts (total jobs: 1)
| # | Runner |
|---|---|
| 1 | linux-amd64-gpu-rtx2080-latest-1 |
The state_type was added in NVIDIAgh-4325
* Add Cython extension that implements bindings
* Native extension for launching reduce_into algorithm
It is as fast as ctypes, also needs to redeclare structs,
but in Cython.
Using the following script to compare time of submission
using ctypes-based and using Cython-based approaches
```python
import cuda.parallel.experimental.algorithms.reduce as impl_base
import cuda.parallel.experimental.algorithms._cy_reduce as impl_new
import numpy as np
import cupy as cp
import time
def time1(reps):
global pr
d = cp.ones(10, dtype="i4")
res = cp.empty(tuple(), dtype=d.dtype)
h_init = np.zeros(tuple(), dtype=d.dtype)
def my_add(a, b):
return a + b
alg = impl_new.reduce_into(d, res, my_add, h_init)
temp_bytes = alg(None, d, res, d.size, h_init)
scratch = cp.empty(temp_bytes, dtype=cp.uint8)
t0 = time.perf_counter_ns()
for i in range(reps):
alg(scratch, d, res, d.size, h_init)
t1 = time.perf_counter_ns()
return t1 - t0
def time2(reps):
global pr
d = cp.ones(10, dtype="i4")
res = cp.empty(tuple(), dtype=d.dtype)
h_init = np.zeros(tuple(), dtype=d.dtype)
def my_add(a, b):
return a + b
alg = impl_base.reduce_into(d, res, my_add, h_init)
temp_bytes = alg(None, d, res, d.size, h_init)
scratch = cp.empty(temp_bytes, dtype=cp.uint8)
t0 = time.perf_counter_ns()
for i in range(reps):
alg(scratch, d, res, d.size, h_init)
t1 = time.perf_counter_ns()
return t1 - t0
if __name__ == "__main__":
print("Base: ", time2(100000))
print("New: ", time1(100000))
```
I see
```
(cuda_parallel_venv) (cccl) coder ➜ ~/cccl/python/cuda_parallel $ LD_LIBRARY_PATH=${LD_LIBRARY_PATH}:${VIRTUAL_ENV}/lib/python3.10/site-packages/cuda/parallel/experimental/cccl python bench.py
Base: 454121541
New: 448204824
```
The timings are roughly the same.
* Add bench script and test_cyb.py
* Add steps to build Cython extension to replace ctypes-based bindings
* Introduce static attribute _impl = _cy_bindings
This is to prevent conditions under which the model
`_cy_bindings` had been unloaded before `_Reduce.__del__`
operator has been called resulting in attribute errors
what NoneType does not have device_reduce_cleanup
attribute.
* Add stub file for native extension for mypy linter to pass
* Add cython as build-time dependency of cuda.parallel
* WIP
* Represent state of iterators using IteratorState
State of cuda.parallel.experimental.iterators is stored
as IteratorState class instance.
* Reorder runtime type-check in Pointer constructor to optimize for PointerProxy argument handling
* Bench can time pointer-base workflow and iterator-base workflow
* Add comment about why self._impl is needed
Replaced use of `cyb` module with self._impl
* Add cython-based Scan algorithm
* Add Cython bindings for segmented reduce
* Reduce problem size for uint64 data type reduction test
This avoid out-of-memory exception arising on entry level GPUs
* Add tests for cython-binding of segmented reduce
* Create tests_for_cything_bindings
* Added cython bindings for all existing algorithms
Tests added to exercise cython bindings
* Corrected import
* Introduce _cy_bindings.make_pointer
Replace use of array.__cuda_array_interface__["data"][0] with
protocols.get_pointer
* Remove PointerProxy class
* Initialize enum values in constructor, instead of in property body
This makes the cost of `IteratorKind.POINTER` much lower (no need
to recreate new copy of IntEnum on every call).
Also create two new query methods in Iterator class, `.is_kind_pointer()`
( it is a 3x faster than `it.type == IteratorKind.POINTER` )
and `.is_kind_iterator()`
* Commend out all get() and get_ptr() cdef methods
* Save repeated computation in locals, use slots of _MergeSort class
* Bench scripts
* Delete commented out code, eliminate build result struct member redeclaration
* Make Cython-based implementation default
Legacy implementation is still accessible via cuda.parallel.experimental.algorithms.legacy,
and cuda.parallel.experimental.iterator_legacy
Legacy test suite is still accessible via tests_legacy/
Both `pytest tests` and `pytest tests_legacy` pass.
* Update benchmark scripts for file reorg
* Add device synchronization in benchmark scripts after each bench
* Optimize attribute away lookup out of hot loop
Instead of looking up function from module in __call__,
import it into the implementations' namespace directly.
This shaves a hundrend nanoseconds per instance changed.
* Changed API exposed by Cython
Instead of exposing cccl_device_reduce_build, cccl_device_reduce
and cccl_device_reduce_cleanup which act on build_result struct (C-style),
we expose these actions as methods of Python class that stores the
build_result variable.
The cleanup method is just called by `__dealloc__` special methods
of that class and is never exposed to the user.
* Delete benchmarking scripts and legacy interface
* Release GIL when calling cccl.c.parallel functions
This should be particularly beneficial for build step.
* Drop use of '_cy' prefix in filenames
* Reverse iterator to set state_type parameter of its parent IteratorBase class
* Strenghten test_iterator test checking kind comparison
Test checks whether iteratator's kind compares unequal
for iterators with different value types even though
the values themselves compare equal.
* Add state_value to IteratorKind constructor
IteratorKind's hash values are used to cache compiled programs.
Iterators generating different LTOIR should never hash to the same
value. To ensure that parametric iterators with different parameter
values (for example, 1D strided iterator with different step sizes)
has to different values we must mix state_type into the hash.
For this reason, state_value is now passed to IteratorKind's
constructor along with value_type.
* Address PR feedback
1. Replace use of Cython deprecated syntax for property getter/setter definition
with conventional @Property and @state.setter decorators
2. Removed IteratorStateView and its uses
3. Added docstring documenting that `as_bytes` are debug utilities to
obtain memory view of native library structs.
4. Removed unused pointer_as_bytes function in Cython extension
* Removed cy_ prefix from Cython types and redeclarations
This makes it much less verbose and easier to maintain.
* Remove _cccl.py as it creeped back in after rebasing
* add _bindings to mock'd imports
* Remove resources/cmake
Implement cythonization step via add_custom_command instead.
Cython command generates DEPFILE. Custom target is created
to link to the custom commandm. The custom target is added
as dependency of `_bindings` target that generates the extension.
* Addressed feedback from @rwgk review
1. Add comment at the top of _bindings.pyx advising reader of existence
of _bindings.pyi stub file and the need to keep Python API in sync
to clear checks of Python's static typing tools, e.g., `mypy`.
2. Renamed native class member in `IteratorState` type from
`size` (which coincided with Python's property) and misled
AI agents in believing that infinite recursion is possible.
The native struct member is now called state_nbytes instead.
* Use cuda_parallel_SOURCE_DIR instead of CMAKE_SOURCE_DIR, use quotes
Also directly require Python3 and use Python3_add_library
* IntEnum -> IntEnumerationMember
* Add missing quote around cmake variable
* Changes to PR review
1. Missing quotes
2. Removed FIND_PACKAGE_HANDLE_STANDARD_ARGS and mark_as_advanced
* Set GENERATED property of Cython-generated source file to TRUE
* Add quotes around Python3_EXECUTABLE
* Use cuda_parallel_BINARY_DIR instead of CMAKE_BINARY_DIR
* Change signature of arg_type_check to have void return type with except *
* Corrected docstring in predicates
* Change signature of alignment validation function to have return type void and qualifier except *
---------
Co-authored-by: Ashwin Srinath <[email protected]>

Description
closes gh-3854
This PR replaces use of
ctypesto callcccl.c.parallellibrary with use of Cython native extension.In addition to improved type-safety of using Cython, we also get a nice reduction in the overhead in
__call__special method implementation for each algorithms, particularly for iterator inputs.Checklist