Skip to content

Native extension to bind to cccl c parallel library - #4325

Merged
oleksandr-pavlyk merged 58 commits into
NVIDIA:mainfrom
oleksandr-pavlyk:native-extension-to-bind-to-cccl-c-parallel-library
Apr 9, 2025
Merged

oleksandr-pavlyk merged 58 commits into
NVIDIA:mainfrom
oleksandr-pavlyk:native-extension-to-bind-to-cccl-c-parallel-library

Conversation

@oleksandr-pavlyk

Copy link
Copy Markdown
Contributor

Description

closes gh-3854

This PR replaces use of ctypes to call cccl.c.parallel library 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

  • New or existing tests cover these changes.
  • The documentation is up to date with these changes.

@copy-pr-bot

copy-pr-bot Bot commented Apr 2, 2025

Copy link
Copy Markdown
Contributor

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.

@cccl-authenticator-app cccl-authenticator-app Bot moved this from Todo to In Progress in CCCL Apr 2, 2025
@oleksandr-pavlyk

Copy link
Copy Markdown
Contributor Author

/ok to test

@leofang leofang left a comment

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

(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.)

@github-actions

github-actions Bot commented Apr 2, 2025

Copy link
Copy Markdown
Contributor
🟩 CI finished in 1h 06m: Pass: 100%/1 | Total: 1h 06m | Avg: 1h 06m | Max: 1h 06m
  • 🟩 python: Pass: 100%/1 | Total: 1h 06m | Avg: 1h 06m | Max: 1h 06m

    🟩 cpu
      🟩 amd64              Pass: 100%/1   | Total:  1h 06m | Avg:  1h 06m | Max:  1h 06m
    🟩 ctk
      🟩 12.8               Pass: 100%/1   | Total:  1h 06m | Avg:  1h 06m | Max:  1h 06m
    🟩 cudacxx
      🟩 nvcc12.8           Pass: 100%/1   | Total:  1h 06m | Avg:  1h 06m | Max:  1h 06m
    🟩 cudacxx_family
      🟩 nvcc               Pass: 100%/1   | Total:  1h 06m | Avg:  1h 06m | Max:  1h 06m
    🟩 cxx
      🟩 GCC13              Pass: 100%/1   | Total:  1h 06m | Avg:  1h 06m | Max:  1h 06m
    🟩 cxx_family
      🟩 GCC                Pass: 100%/1   | Total:  1h 06m | Avg:  1h 06m | Max:  1h 06m
    🟩 gpu
      🟩 rtx2080            Pass: 100%/1   | Total:  1h 06m | Avg:  1h 06m | Max:  1h 06m
    🟩 jobs
      🟩 Test               Pass: 100%/1   | Total:  1h 06m | Avg:  1h 06m | Max:  1h 06m
    

👃 Inspect Changes

Modifications in project?

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

@oleksandr-pavlyk

Copy link
Copy Markdown
Contributor Author

Here is the performance data:

combined-plot

It is comparing the time of execution of __call__ operator for two algorithms (reduce_into and merge_sort) for two usages: inputs are iterators of kind POINTER, and inputs are iterators of kind ITERATOR.

With ctypes the latter was noticeably slower than the former. Using native extension they virtually identical.

@oleksandr-pavlyk

Copy link
Copy Markdown
Contributor Author

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.

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Given that we have dropped the legacy bindings, can we also drop the _cy_ prefix in these files?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

@oleksandr-pavlyk oleksandr-pavlyk Apr 3, 2025 •

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

@oleksandr-pavlyk

Copy link
Copy Markdown
Contributor Author

/ok to test

@github-actions

github-actions Bot commented Apr 4, 2025

Copy link
Copy Markdown
Contributor
🟩 CI finished in 1h 05m: Pass: 100%/1 | Total: 1h 05m | Avg: 1h 05m | Max: 1h 05m
  • 🟩 python: Pass: 100%/1 | Total: 1h 05m | Avg: 1h 05m | Max: 1h 05m

    🟩 cpu
      🟩 amd64              Pass: 100%/1   | Total:  1h 05m | Avg:  1h 05m | Max:  1h 05m
    🟩 ctk
      🟩 12.8               Pass: 100%/1   | Total:  1h 05m | Avg:  1h 05m | Max:  1h 05m
    🟩 cudacxx
      🟩 nvcc12.8           Pass: 100%/1   | Total:  1h 05m | Avg:  1h 05m | Max:  1h 05m
    🟩 cudacxx_family
      🟩 nvcc               Pass: 100%/1   | Total:  1h 05m | Avg:  1h 05m | Max:  1h 05m
    🟩 cxx
      🟩 GCC13              Pass: 100%/1   | Total:  1h 05m | Avg:  1h 05m | Max:  1h 05m
    🟩 cxx_family
      🟩 GCC                Pass: 100%/1   | Total:  1h 05m | Avg:  1h 05m | Max:  1h 05m
    🟩 gpu
      🟩 rtx2080            Pass: 100%/1   | Total:  1h 05m | Avg:  1h 05m | Max:  1h 05m
    🟩 jobs
      🟩 Test               Pass: 100%/1   | Total:  1h 05m | Avg:  1h 05m | Max:  1h 05m
    

👃 Inspect Changes

Modifications in project?

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

@oleksandr-pavlyk
oleksandr-pavlyk force-pushed the native-extension-to-bind-to-cccl-c-parallel-library branch from 1c142e5 to fcc5ed7 Compare April 7, 2025 15:43

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Can we put all the cdef extern declarations into _bindings.pxd, and reserve _bindings.pyx for the wrapper definitions?

@leofang
leofang self-requested a review April 7, 2025 17:43

@leofang leofang left a comment

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.)

@oleksandr-pavlyk
oleksandr-pavlyk marked this pull request as ready for review April 7, 2025 18:50
@oleksandr-pavlyk
oleksandr-pavlyk requested review from a team as code owners April 7, 2025 18:50
@cccl-authenticator-app cccl-authenticator-app Bot moved this from In Progress to In Review in CCCL Apr 7, 2025

@shwina shwina left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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:

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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?

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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:

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I don't think I fully understand the need for the following:

  • IteratorStateView
  • IteratorState,
  • 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?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

@github-actions

github-actions Bot commented Apr 7, 2025

Copy link
Copy Markdown
Contributor
🟩 CI finished in 1h 27m: Pass: 100%/1 | Total: 1h 27m | Avg: 1h 27m | Max: 1h 27m
  • 🟩 python: Pass: 100%/1 | Total: 1h 27m | Avg: 1h 27m | Max: 1h 27m

    🟩 cpu
      🟩 amd64              Pass: 100%/1   | Total:  1h 27m | Avg:  1h 27m | Max:  1h 27m
    🟩 ctk
      🟩 12.8               Pass: 100%/1   | Total:  1h 27m | Avg:  1h 27m | Max:  1h 27m
    🟩 cudacxx
      🟩 nvcc12.8           Pass: 100%/1   | Total:  1h 27m | Avg:  1h 27m | Max:  1h 27m
    🟩 cudacxx_family
      🟩 nvcc               Pass: 100%/1   | Total:  1h 27m | Avg:  1h 27m | Max:  1h 27m
    🟩 cxx
      🟩 GCC13              Pass: 100%/1   | Total:  1h 27m | Avg:  1h 27m | Max:  1h 27m
    🟩 cxx_family
      🟩 GCC                Pass: 100%/1   | Total:  1h 27m | Avg:  1h 27m | Max:  1h 27m
    🟩 gpu
      🟩 rtx2080            Pass: 100%/1   | Total:  1h 27m | Avg:  1h 27m | Max:  1h 27m
    🟩 jobs
      🟩 Test               Pass: 100%/1   | Total:  1h 27m | Avg:  1h 27m | Max:  1h 27m
    

👃 Inspect Changes

Modifications in project?

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".

@alliepiper alliepiper left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Just looked at CMake stuff.

Comment thread python/cuda_parallel/CMakeLists.txt Resolved
Comment thread python/cuda_parallel/CMakeLists.txt Resolved
Comment thread python/cuda_parallel/CMakeLists.txt Resolved
@github-project-automation github-project-automation Bot moved this from In Review to In Progress in CCCL Apr 8, 2025
Comment thread python/cuda_parallel/CMakeLists.txt Resolved
Comment thread python/cuda_parallel/CMakeLists.txt Resolved
Comment thread python/cuda_parallel/CMakeLists.txt Resolved
Comment thread python/cuda_parallel/CMakeLists.txt Resolved
Comment thread python/cuda_parallel/CMakeLists.txt Resolved
Comment thread python/cuda_parallel/CMakeLists.txt Resolved
Comment thread python/cuda_parallel/CMakeLists.txt Resolved
1. Missing quotes
2. Removed FIND_PACKAGE_HANDLE_STANDARD_ARGS and mark_as_advanced
@github-project-automation github-project-automation Bot moved this from In Progress to In Review in CCCL Apr 8, 2025
@rwgk

rwgk commented Apr 8, 2025

Copy link
Copy Markdown
Contributor

For completeness, when reviewing this PR, I also looked at this cloc overview:

$ cloc --diff ~/ref/cccl/python/cuda_parallel python/cuda_parallel
      34 text files.
      36 text files.
       6 files ignored.                             

github.com/AlDanial/cloc v 1.98  T=0.08 s (533.5 files/s, 99524.8 lines/s)
-------------------------------------------------------------------------------
Language                     files          blank        comment           code
-------------------------------------------------------------------------------
Python
 same                           15            351            382           2208
 modified                        7              0              0             42
 added                           8            231            165           1219
 removed                         7            192            170            913
Cython
 same                            0              0              0              0
 modified                        0              0              0              0
 added                           1            244             28           1348
 removed                         0              0              0              0
CMake
 same                            0              0              1             17
 modified                        1              0              0              1
 added                           0              9              4             52
 removed                         0              0              0              0
TOML
 same                            0              0              4             54
 modified                        1              0              0              1
 added                           0              0              0              0
 removed                         0              0              0              0
Markdown
 same                            1              4              0              9
 modified                        0              0              0              0
 added                           0              0              0              0
 removed                         0              0              0              0
-------------------------------------------------------------------------------
SUM:
 same                           16            355            387           2288
 modified                        9              0              0             44
 added                           9            484            197           2619
 removed                         7            192            170            913
-------------------------------------------------------------------------------

Comment thread python/cuda_parallel/CMakeLists.txt Resolved
Comment thread python/cuda_parallel/CMakeLists.txt Resolved
@oleksandr-pavlyk

Copy link
Copy Markdown
Contributor Author

Thank you for good review everybody. I think this PR can be merged once CI's badge of approval arrives.

@github-actions

github-actions Bot commented Apr 8, 2025

Copy link
Copy Markdown
Contributor
🟩 CI finished in 1h 33m: Pass: 100%/1 | Total: 1h 33m | Avg: 1h 33m | Max: 1h 33m
  • 🟩 python: Pass: 100%/1 | Total: 1h 33m | Avg: 1h 33m | Max: 1h 33m

    🟩 cpu
      🟩 amd64              Pass: 100%/1   | Total:  1h 33m | Avg:  1h 33m | Max:  1h 33m
    🟩 ctk
      🟩 12.8               Pass: 100%/1   | Total:  1h 33m | Avg:  1h 33m | Max:  1h 33m
    🟩 cudacxx
      🟩 nvcc12.8           Pass: 100%/1   | Total:  1h 33m | Avg:  1h 33m | Max:  1h 33m
    🟩 cudacxx_family
      🟩 nvcc               Pass: 100%/1   | Total:  1h 33m | Avg:  1h 33m | Max:  1h 33m
    🟩 cxx
      🟩 GCC13              Pass: 100%/1   | Total:  1h 33m | Avg:  1h 33m | Max:  1h 33m
    🟩 cxx_family
      🟩 GCC                Pass: 100%/1   | Total:  1h 33m | Avg:  1h 33m | Max:  1h 33m
    🟩 gpu
      🟩 rtx2080            Pass: 100%/1   | Total:  1h 33m | Avg:  1h 33m | Max:  1h 33m
    🟩 jobs
      🟩 Test               Pass: 100%/1   | Total:  1h 33m | Avg:  1h 33m | Max:  1h 33m
    

👃 Inspect Changes

Modifications in project?

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

@github-actions

github-actions Bot commented Apr 9, 2025

Copy link
Copy Markdown
Contributor
🟩 CI finished in 1h 29m: Pass: 100%/1 | Total: 1h 29m | Avg: 1h 29m | Max: 1h 29m
  • 🟩 python: Pass: 100%/1 | Total: 1h 29m | Avg: 1h 29m | Max: 1h 29m

    🟩 cpu
      🟩 amd64              Pass: 100%/1   | Total:  1h 29m | Avg:  1h 29m | Max:  1h 29m
    🟩 ctk
      🟩 12.8               Pass: 100%/1   | Total:  1h 29m | Avg:  1h 29m | Max:  1h 29m
    🟩 cudacxx
      🟩 nvcc12.8           Pass: 100%/1   | Total:  1h 29m | Avg:  1h 29m | Max:  1h 29m
    🟩 cudacxx_family
      🟩 nvcc               Pass: 100%/1   | Total:  1h 29m | Avg:  1h 29m | Max:  1h 29m
    🟩 cxx
      🟩 GCC13              Pass: 100%/1   | Total:  1h 29m | Avg:  1h 29m | Max:  1h 29m
    🟩 cxx_family
      🟩 GCC                Pass: 100%/1   | Total:  1h 29m | Avg:  1h 29m | Max:  1h 29m
    🟩 gpu
      🟩 rtx2080            Pass: 100%/1   | Total:  1h 29m | Avg:  1h 29m | Max:  1h 29m
    🟩 jobs
      🟩 Test               Pass: 100%/1   | Total:  1h 29m | Avg:  1h 29m | Max:  1h 29m
    

👃 Inspect Changes

Modifications in project?

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

@oleksandr-pavlyk
oleksandr-pavlyk merged commit 7dea383 into NVIDIA:main Apr 9, 2025
@github-project-automation github-project-automation Bot moved this from In Review to Done in CCCL Apr 9, 2025
@oleksandr-pavlyk
oleksandr-pavlyk deleted the native-extension-to-bind-to-cccl-c-parallel-library branch April 9, 2025 02:50
oleksandr-pavlyk added a commit to oleksandr-pavlyk/cccl that referenced this pull request Apr 9, 2025
davebayer pushed a commit to davebayer/cccl that referenced this pull request Apr 11, 2025
* 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]>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

Archived in project

Development

Successfully merging this pull request may close these issues.

cuda.parallel: move away from using ctypes

7 participants