Skip to content

Add Python wrappers for c.parallel radix_sort API - #4353

Merged
NaderAlAwar merged 77 commits into
NVIDIA:mainfrom
NaderAlAwar:radix-sort-python-wrappers
Apr 21, 2025
Merged

Add Python wrappers for c.parallel radix_sort API#4353
NaderAlAwar merged 77 commits into
NVIDIA:mainfrom
NaderAlAwar:radix-sort-python-wrappers

Conversation

@NaderAlAwar

@NaderAlAwar NaderAlAwar commented Apr 4, 2025

Copy link
Copy Markdown
Contributor

Description

closes #3799
closes #2548

This should only be merged after #4350

Checklist

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

Comment thread python/cuda_parallel/cuda/parallel/experimental/algorithms/_radix_sort.py Outdated
Comment thread python/cuda_parallel/tests/test_radix_sort_api.py Outdated
Comment thread c/parallel/src/radix_sort.cu Outdated
return error;
}

CUresult cccl_device_ascending_radix_sort(

@oleksandr-pavlyk oleksandr-pavlyk Apr 15, 2025

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.

Why not have single cccl_device_radix_sort which dispatches at run-time to cccl_device_radix_sort_impl<cub::SortOrder::Ascending> or cccl_device_radix_sort_impl<cub::SortOrder::Descending>?

Comment thread c/parallel/src/radix_sort.cu Outdated
Comment on lines +707 to +710
cub::DoubleBuffer<indirect_arg_t> d_keys_buffer(
static_cast<indirect_arg_t*>(d_keys_in.state), static_cast<indirect_arg_t*>(d_keys_out.state));
cub::DoubleBuffer<indirect_arg_t> d_values_buffer(
static_cast<indirect_arg_t*>(d_values_in.state), static_cast<indirect_arg_t*>(d_values_out.state));

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.

Suggested change
cub::DoubleBuffer<indirect_arg_t> d_keys_buffer(
static_cast<indirect_arg_t*>(d_keys_in.state), static_cast<indirect_arg_t*>(d_keys_out.state));
cub::DoubleBuffer<indirect_arg_t> d_values_buffer(
static_cast<indirect_arg_t*>(d_values_in.state), static_cast<indirect_arg_t*>(d_values_out.state));
indirect_arg_t key_arg1 = {d_keys_in};
indirect_arg_t key_arg2 = {d_keys_out};
cub::DoubleBuffer<indirect_arg_t> d_keys_buffer(&key_arg1, &key_arg2);
indirect_arg_t val_arg1 = {d_values_in};
indirect_arg_t val_arg2 = {d_values_out};
cub::DoubleBuffer<indirect_arg_t> d_values_buffer(&val_arg1, &val_arg2);

@NaderAlAwar

Copy link
Copy Markdown
Contributor Author

/ok to test 57bdd7d

@github-actions

Copy link
Copy Markdown
Contributor
🟨 CI finished in 1h 34m: Pass: 80%/5 | Total: 1h 06m | Avg: 13m 14s | Max: 34m 26s | Hits: 98%/324
  • 🟨 python: Pass: 66%/3 | Total: 29m 15s | Avg: 9m 45s | Max: 18m 47s

    🚨 jobs: cuda.parallel 🚨
      🟩 cuda.cccl          Pass: 100%/1   | Total:  3m 27s | Avg:  3m 27s | Max:  3m 27s
      🟩 cuda.cooperative   Pass: 100%/1   | Total: 18m 47s | Avg: 18m 47s | Max: 18m 47s
      🔥 cuda.parallel      Pass:   0%/1   | Total:  7m 01s | Avg:  7m 01s | Max:  7m 01s
    🟨 cpu
      🟨 amd64              Pass:  66%/3   | Total: 29m 15s | Avg:  9m 45s | Max: 18m 47s
    🟨 ctk
      🟨 12.8               Pass:  66%/3   | Total: 29m 15s | Avg:  9m 45s | Max: 18m 47s
    🟨 cudacxx
      🟨 nvcc12.8           Pass:  66%/3   | Total: 29m 15s | Avg:  9m 45s | Max: 18m 47s
    🟨 cudacxx_family
      🟨 nvcc               Pass:  66%/3   | Total: 29m 15s | Avg:  9m 45s | Max: 18m 47s
    🟨 cxx
      🟨 GCC13              Pass:  66%/3   | Total: 29m 15s | Avg:  9m 45s | Max: 18m 47s
    🟨 cxx_family
      🟨 GCC                Pass:  66%/3   | Total: 29m 15s | Avg:  9m 45s | Max: 18m 47s
    🟨 gpu
      🟨 rtx2080            Pass:  66%/3   | Total: 29m 15s | Avg:  9m 45s | Max: 18m 47s
    
  • 🟩 cccl_c_parallel: Pass: 100%/2 | Total: 36m 56s | Avg: 18m 28s | Max: 34m 26s | Hits: 98%/324

    🟩 cpu
      🟩 amd64              Pass: 100%/2   | Total: 36m 56s | Avg: 18m 28s | Max: 34m 26s | Hits:  98%/324   
    🟩 ctk
      🟩 12.8               Pass: 100%/2   | Total: 36m 56s | Avg: 18m 28s | Max: 34m 26s | Hits:  98%/324   
    🟩 cudacxx
      🟩 nvcc12.8           Pass: 100%/2   | Total: 36m 56s | Avg: 18m 28s | Max: 34m 26s | Hits:  98%/324   
    🟩 cudacxx_family
      🟩 nvcc               Pass: 100%/2   | Total: 36m 56s | Avg: 18m 28s | Max: 34m 26s | Hits:  98%/324   
    🟩 cxx
      🟩 GCC13              Pass: 100%/2   | Total: 36m 56s | Avg: 18m 28s | Max: 34m 26s | Hits:  98%/324   
    🟩 cxx_family
      🟩 GCC                Pass: 100%/2   | Total: 36m 56s | Avg: 18m 28s | Max: 34m 26s | Hits:  98%/324   
    🟩 gpu
      🟩 rtx2080            Pass: 100%/2   | Total: 36m 56s | Avg: 18m 28s | Max: 34m 26s | Hits:  98%/324   
    🟩 jobs
      🟩 Build              Pass: 100%/1   | Total:  2m 30s | Avg:  2m 30s | Max:  2m 30s | Hits:  98%/162   
      🟩 Test               Pass: 100%/1   | Total: 34m 26s | Avg: 34m 26s | Max: 34m 26s | Hits:  98%/162   
    

👃 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: 5)

# Runner
4 linux-amd64-gpu-rtx2080-latest-1
1 linux-amd64-cpu16

@NaderAlAwar

Copy link
Copy Markdown
Contributor Author

/ok to test e1f3250

@github-actions

Copy link
Copy Markdown
Contributor
🟨 CI finished in 35m 42s: Pass: 80%/5 | Total: 1h 05m | Avg: 13m 07s | Max: 32m 00s | Hits: 98%/324
  • 🟨 python: Pass: 66%/3 | Total: 31m 10s | Avg: 10m 23s | Max: 17m 26s

    🚨 jobs: cuda.parallel 🚨
      🟩 cuda.cccl          Pass: 100%/1   | Total:  7m 54s | Avg:  7m 54s | Max:  7m 54s
      🟩 cuda.cooperative   Pass: 100%/1   | Total: 17m 26s | Avg: 17m 26s | Max: 17m 26s
      🔥 cuda.parallel      Pass:   0%/1   | Total:  5m 50s | Avg:  5m 50s | Max:  5m 50s
    🟨 cpu
      🟨 amd64              Pass:  66%/3   | Total: 31m 10s | Avg: 10m 23s | Max: 17m 26s
    🟨 ctk
      🟨 12.8               Pass:  66%/3   | Total: 31m 10s | Avg: 10m 23s | Max: 17m 26s
    🟨 cudacxx
      🟨 nvcc12.8           Pass:  66%/3   | Total: 31m 10s | Avg: 10m 23s | Max: 17m 26s
    🟨 cudacxx_family
      🟨 nvcc               Pass:  66%/3   | Total: 31m 10s | Avg: 10m 23s | Max: 17m 26s
    🟨 cxx
      🟨 GCC13              Pass:  66%/3   | Total: 31m 10s | Avg: 10m 23s | Max: 17m 26s
    🟨 cxx_family
      🟨 GCC                Pass:  66%/3   | Total: 31m 10s | Avg: 10m 23s | Max: 17m 26s
    🟨 gpu
      🟨 rtx2080            Pass:  66%/3   | Total: 31m 10s | Avg: 10m 23s | Max: 17m 26s
    
  • 🟩 cccl_c_parallel: Pass: 100%/2 | Total: 34m 29s | Avg: 17m 14s | Max: 32m 00s | Hits: 98%/324

    🟩 cpu
      🟩 amd64              Pass: 100%/2   | Total: 34m 29s | Avg: 17m 14s | Max: 32m 00s | Hits:  98%/324   
    🟩 ctk
      🟩 12.8               Pass: 100%/2   | Total: 34m 29s | Avg: 17m 14s | Max: 32m 00s | Hits:  98%/324   
    🟩 cudacxx
      🟩 nvcc12.8           Pass: 100%/2   | Total: 34m 29s | Avg: 17m 14s | Max: 32m 00s | Hits:  98%/324   
    🟩 cudacxx_family
      🟩 nvcc               Pass: 100%/2   | Total: 34m 29s | Avg: 17m 14s | Max: 32m 00s | Hits:  98%/324   
    🟩 cxx
      🟩 GCC13              Pass: 100%/2   | Total: 34m 29s | Avg: 17m 14s | Max: 32m 00s | Hits:  98%/324   
    🟩 cxx_family
      🟩 GCC                Pass: 100%/2   | Total: 34m 29s | Avg: 17m 14s | Max: 32m 00s | Hits:  98%/324   
    🟩 gpu
      🟩 rtx2080            Pass: 100%/2   | Total: 34m 29s | Avg: 17m 14s | Max: 32m 00s | Hits:  98%/324   
    🟩 jobs
      🟩 Build              Pass: 100%/1   | Total:  2m 29s | Avg:  2m 29s | Max:  2m 29s | Hits:  98%/162   
      🟩 Test               Pass: 100%/1   | Total: 32m 00s | Avg: 32m 00s | Max: 32m 00s | Hits:  98%/162   
    

👃 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: 5)

# Runner
4 linux-amd64-gpu-rtx2080-latest-1
1 linux-amd64-cpu16

@NaderAlAwar
NaderAlAwar marked this pull request as ready for review April 18, 2025 13:27
@github-actions

Copy link
Copy Markdown
Contributor
🟨 CI finished in 21h 26m: Pass: 80%/5 | Total: 1h 05m | Avg: 13m 10s | Max: 32m 00s | Hits: 98%/324
  • 🟨 python: Pass: 66%/3 | Total: 31m 23s | Avg: 10m 27s | Max: 17m 26s

    🚨 jobs: cuda.parallel 🚨
      🟩 cuda.cccl          Pass: 100%/1   | Total:  7m 54s | Avg:  7m 54s | Max:  7m 54s
      🟩 cuda.cooperative   Pass: 100%/1   | Total: 17m 26s | Avg: 17m 26s | Max: 17m 26s
      🔥 cuda.parallel      Pass:   0%/1   | Total:  6m 03s | Avg:  6m 03s | Max:  6m 03s
    🟨 cpu
      🟨 amd64              Pass:  66%/3   | Total: 31m 23s | Avg: 10m 27s | Max: 17m 26s
    🟨 ctk
      🟨 12.8               Pass:  66%/3   | Total: 31m 23s | Avg: 10m 27s | Max: 17m 26s
    🟨 cudacxx
      🟨 nvcc12.8           Pass:  66%/3   | Total: 31m 23s | Avg: 10m 27s | Max: 17m 26s
    🟨 cudacxx_family
      🟨 nvcc               Pass:  66%/3   | Total: 31m 23s | Avg: 10m 27s | Max: 17m 26s
    🟨 cxx
      🟨 GCC13              Pass:  66%/3   | Total: 31m 23s | Avg: 10m 27s | Max: 17m 26s
    🟨 cxx_family
      🟨 GCC                Pass:  66%/3   | Total: 31m 23s | Avg: 10m 27s | Max: 17m 26s
    🟨 gpu
      🟨 rtx2080            Pass:  66%/3   | Total: 31m 23s | Avg: 10m 27s | Max: 17m 26s
    
  • 🟩 cccl_c_parallel: Pass: 100%/2 | Total: 34m 29s | Avg: 17m 14s | Max: 32m 00s | Hits: 98%/324

    🟩 cpu
      🟩 amd64              Pass: 100%/2   | Total: 34m 29s | Avg: 17m 14s | Max: 32m 00s | Hits:  98%/324   
    🟩 ctk
      🟩 12.8               Pass: 100%/2   | Total: 34m 29s | Avg: 17m 14s | Max: 32m 00s | Hits:  98%/324   
    🟩 cudacxx
      🟩 nvcc12.8           Pass: 100%/2   | Total: 34m 29s | Avg: 17m 14s | Max: 32m 00s | Hits:  98%/324   
    🟩 cudacxx_family
      🟩 nvcc               Pass: 100%/2   | Total: 34m 29s | Avg: 17m 14s | Max: 32m 00s | Hits:  98%/324   
    🟩 cxx
      🟩 GCC13              Pass: 100%/2   | Total: 34m 29s | Avg: 17m 14s | Max: 32m 00s | Hits:  98%/324   
    🟩 cxx_family
      🟩 GCC                Pass: 100%/2   | Total: 34m 29s | Avg: 17m 14s | Max: 32m 00s | Hits:  98%/324   
    🟩 gpu
      🟩 rtx2080            Pass: 100%/2   | Total: 34m 29s | Avg: 17m 14s | Max: 32m 00s | Hits:  98%/324   
    🟩 jobs
      🟩 Build              Pass: 100%/1   | Total:  2m 29s | Avg:  2m 29s | Max:  2m 29s | Hits:  98%/162   
      🟩 Test               Pass: 100%/1   | Total: 32m 00s | Avg: 32m 00s | Max: 32m 00s | Hits:  98%/162   
    

👃 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: 5)

# Runner
4 linux-amd64-gpu-rtx2080-latest-1
1 linux-amd64-cpu16

@cccl-authenticator-app cccl-authenticator-app Bot moved this from In Progress to In Review in CCCL Apr 18, 2025

@oleksandr-pavlyk oleksandr-pavlyk 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.

Thank you @NaderAlAwar , LGTM!

@github-actions

Copy link
Copy Markdown
Contributor
🟨 CI finished in 23m 07s: Pass: 66%/3 | Total: 33m 09s | Avg: 11m 03s | Max: 23m 07s
  • 🟨 python: Pass: 66%/3 | Total: 33m 09s | Avg: 11m 03s | Max: 23m 07s

    🚨 jobs: cuda.parallel 🚨
      🟩 cuda.cccl          Pass: 100%/1   | Total:  3m 16s | Avg:  3m 16s | Max:  3m 16s
      🟩 cuda.cooperative   Pass: 100%/1   | Total: 23m 07s | Avg: 23m 07s | Max: 23m 07s
      🔥 cuda.parallel      Pass:   0%/1   | Total:  6m 46s | Avg:  6m 46s | Max:  6m 46s
    🟨 cpu
      🟨 amd64              Pass:  66%/3   | Total: 33m 09s | Avg: 11m 03s | Max: 23m 07s
    🟨 ctk
      🟨 12.8               Pass:  66%/3   | Total: 33m 09s | Avg: 11m 03s | Max: 23m 07s
    🟨 cudacxx
      🟨 nvcc12.8           Pass:  66%/3   | Total: 33m 09s | Avg: 11m 03s | Max: 23m 07s
    🟨 cudacxx_family
      🟨 nvcc               Pass:  66%/3   | Total: 33m 09s | Avg: 11m 03s | Max: 23m 07s
    🟨 cxx
      🟨 GCC13              Pass:  66%/3   | Total: 33m 09s | Avg: 11m 03s | Max: 23m 07s
    🟨 cxx_family
      🟨 GCC                Pass:  66%/3   | Total: 33m 09s | Avg: 11m 03s | Max: 23m 07s
    🟨 gpu
      🟨 rtx2080            Pass:  66%/3   | Total: 33m 09s | Avg: 11m 03s | Max: 23m 07s
    

👃 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: 3)

# Runner
3 linux-amd64-gpu-rtx2080-latest-1

@github-actions

Copy link
Copy Markdown
Contributor
🟩 CI finished in 17m 08s: Pass: 100%/3 | Total: 29m 26s | Avg: 9m 48s | Max: 17m 07s
  • 🟩 python: Pass: 100%/3 | Total: 29m 26s | Avg: 9m 48s | Max: 17m 07s

    🟩 cpu
      🟩 amd64              Pass: 100%/3   | Total: 29m 26s | Avg:  9m 48s | Max: 17m 07s
    🟩 ctk
      🟩 12.8               Pass: 100%/3   | Total: 29m 26s | Avg:  9m 48s | Max: 17m 07s
    🟩 cudacxx
      🟩 nvcc12.8           Pass: 100%/3   | Total: 29m 26s | Avg:  9m 48s | Max: 17m 07s
    🟩 cudacxx_family
      🟩 nvcc               Pass: 100%/3   | Total: 29m 26s | Avg:  9m 48s | Max: 17m 07s
    🟩 cxx
      🟩 GCC13              Pass: 100%/3   | Total: 29m 26s | Avg:  9m 48s | Max: 17m 07s
    🟩 cxx_family
      🟩 GCC                Pass: 100%/3   | Total: 29m 26s | Avg:  9m 48s | Max: 17m 07s
    🟩 gpu
      🟩 rtx2080            Pass: 100%/3   | Total: 29m 26s | Avg:  9m 48s | Max: 17m 07s
    🟩 jobs
      🟩 cuda.cccl          Pass: 100%/1   | Total:  5m 56s | Avg:  5m 56s | Max:  5m 56s
      🟩 cuda.cooperative   Pass: 100%/1   | Total: 17m 07s | Avg: 17m 07s | Max: 17m 07s
      🟩 cuda.parallel      Pass: 100%/1   | Total:  6m 23s | Avg:  6m 23s | Max:  6m 23s
    

👃 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: 3)

# Runner
3 linux-amd64-gpu-rtx2080-latest-1

@NaderAlAwar
NaderAlAwar merged commit 8306d99 into NVIDIA:main Apr 21, 2025
@github-project-automation github-project-automation Bot moved this from In Review to Done in CCCL Apr 21, 2025
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.

Add Python wrappers for radix sort [FEA]: Implement cuda.parallel version of cub::DoubleBuffer

2 participants