Skip to content

Add Unique by Key Implementation for c.parallel - #3947

Merged
NaderAlAwar merged 55 commits into
NVIDIA:mainfrom
NaderAlAwar:unique-by-key-c-parallel
Mar 7, 2025
Merged

NaderAlAwar merged 55 commits into
NVIDIA:mainfrom
NaderAlAwar:unique-by-key-c-parallel

Conversation

@NaderAlAwar

Copy link
Copy Markdown
Contributor

Description

Closes #3807

This should only be merged after #3816

Checklist

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

…because this allows us to define it such that we can use the ActivePolicyT template parameter of invoke() as its policy, rather than the policy hub's MaxPolicy. ActivePolicyT will be the correct policy for the current ptx version, since calling invoke() from dispatch() causes us to move down the chained policies until we get to the correct one.
… needed because this allows us to define it such that we can use the ActivePolicyT template parameter of invoke() as its policy, rather than the policy hub's MaxPolicy. ActivePolicyT will be the correct policy for the current ptx version, since calling invoke() from dispatch() causes us to move down the chained policies until we get to the correct one."

This reverts commit 2a17133.
…s it from the c.parallel layer. Also we don't need it to have any run-time state
…und an nvcc 12.0 template instantiation issue with vsmem_helper_fallback_policy_t
…_key and fix issue with extra semicolon in kernel template instantiation
@NaderAlAwar
NaderAlAwar requested a review from shwina February 28, 2025 20:08

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

Looks great to me - nice work, @NaderAlAwar!

@github-actions

Copy link
Copy Markdown
Contributor
🟩 CI finished in 52m 15s: Pass: 100%/3 | Total: 1h 08m | Avg: 22m 56s | Max: 51m 53s | Hits: 97%/318
  • 🟩 cccl_c_parallel: Pass: 100%/2 | Total: 16m 57s | Avg: 8m 28s | Max: 14m 23s | Hits: 97%/318

    🟩 cpu
      🟩 amd64              Pass: 100%/2   | Total: 16m 57s | Avg:  8m 28s | Max: 14m 23s | Hits:  97%/318   
    🟩 ctk
      🟩 12.8               Pass: 100%/2   | Total: 16m 57s | Avg:  8m 28s | Max: 14m 23s | Hits:  97%/318   
    🟩 cudacxx
      🟩 nvcc12.8           Pass: 100%/2   | Total: 16m 57s | Avg:  8m 28s | Max: 14m 23s | Hits:  97%/318   
    🟩 cudacxx_family
      🟩 nvcc               Pass: 100%/2   | Total: 16m 57s | Avg:  8m 28s | Max: 14m 23s | Hits:  97%/318   
    🟩 cxx
      🟩 GCC13              Pass: 100%/2   | Total: 16m 57s | Avg:  8m 28s | Max: 14m 23s | Hits:  97%/318   
    🟩 cxx_family
      🟩 GCC                Pass: 100%/2   | Total: 16m 57s | Avg:  8m 28s | Max: 14m 23s | Hits:  97%/318   
    🟩 gpu
      🟩 rtx2080            Pass: 100%/2   | Total: 16m 57s | Avg:  8m 28s | Max: 14m 23s | Hits:  97%/318   
    🟩 jobs
      🟩 Build              Pass: 100%/1   | Total:  2m 34s | Avg:  2m 34s | Max:  2m 34s | Hits:  96%/159   
      🟩 Test               Pass: 100%/1   | Total: 14m 23s | Avg: 14m 23s | Max: 14m 23s | Hits:  98%/159   
    
  • 🟩 python: Pass: 100%/1 | Total: 51m 53s | Avg: 51m 53s | Max: 51m 53s

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

👃 Inspect Changes

Modifications in project?

Project
CCCL Infrastructure
libcu++
CUB
Thrust
CUDA Experimental
python
+/- CCCL C Parallel Library
Catch2Helper

Modifications in project or dependencies?

Project
CCCL Infrastructure
libcu++
CUB
Thrust
CUDA Experimental
+/- python
+/- CCCL C Parallel Library
Catch2Helper

🏃‍ Runner counts (total jobs: 3)

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

Comment thread c/parallel/test/test_unique_by_key.cpp Resolved
Comment thread c/parallel/test/test_unique_by_key.cpp Resolved
Comment thread c/parallel/test/test_unique_by_key.cpp Resolved
Comment thread c/parallel/src/unique_by_key.cu Resolved
Comment thread c/parallel/test/test_unique_by_key.cpp Resolved
Comment thread c/parallel/src/unique_by_key.cu Resolved
Comment thread c/parallel/src/unique_by_key.cu Outdated
Comment on lines +80 to +83
// TODO: we should update this once we figure out a way to reuse
// tuning logic from C++. Alternately, we should implement
// something better than a hardcoded default:
return {128, 10, cub::BLOCK_LOAD_DIRECT, cub::LOAD_DEFAULT};

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

important: we should wait with copy-pasting all the tuning policies we have on the C++ end until we have a machinery in place. That said, unique by key tuning uses items per thread scaling to better handle large data types. It shouldn't be a lot of effort to make the tuning policy depend on the key (similar to Nominal4BItemsToItems<KeyT>(10) logic) type along the lines of what we do in the reduce.cu. But if it's not the case, feel free to address this as a follow-up PR.

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 had original implemented Nominal4BItemsToItems but that triggered the memory op assertion in the tests, so I overrode that value it produced with a smaller one. I moved Nominal4BItemsToItems to a common file to use it here, but I ran into the same issue, so I fixed it the same way. I think I'm just going to fix both in a later PR

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.

Created #4013

Comment thread c/parallel/src/unique_by_key.cu Resolved
Comment thread c/parallel/src/unique_by_key.cu Resolved
Comment thread c/parallel/src/unique_by_key.cu Resolved
@github-actions

github-actions Bot commented Mar 3, 2025

Copy link
Copy Markdown
Contributor
🟩 CI finished in 54m 47s: Pass: 100%/3 | Total: 1h 08m | Avg: 22m 43s | Max: 51m 15s | Hits: 98%/318
  • 🟩 cccl_c_parallel: Pass: 100%/2 | Total: 16m 55s | Avg: 8m 27s | Max: 14m 26s | Hits: 98%/318

    🟩 cpu
      🟩 amd64              Pass: 100%/2   | Total: 16m 55s | Avg:  8m 27s | Max: 14m 26s | Hits:  98%/318   
    🟩 ctk
      🟩 12.8               Pass: 100%/2   | Total: 16m 55s | Avg:  8m 27s | Max: 14m 26s | Hits:  98%/318   
    🟩 cudacxx
      🟩 nvcc12.8           Pass: 100%/2   | Total: 16m 55s | Avg:  8m 27s | Max: 14m 26s | Hits:  98%/318   
    🟩 cudacxx_family
      🟩 nvcc               Pass: 100%/2   | Total: 16m 55s | Avg:  8m 27s | Max: 14m 26s | Hits:  98%/318   
    🟩 cxx
      🟩 GCC13              Pass: 100%/2   | Total: 16m 55s | Avg:  8m 27s | Max: 14m 26s | Hits:  98%/318   
    🟩 cxx_family
      🟩 GCC                Pass: 100%/2   | Total: 16m 55s | Avg:  8m 27s | Max: 14m 26s | Hits:  98%/318   
    🟩 gpu
      🟩 rtx2080            Pass: 100%/2   | Total: 16m 55s | Avg:  8m 27s | Max: 14m 26s | Hits:  98%/318   
    🟩 jobs
      🟩 Build              Pass: 100%/1   | Total:  2m 29s | Avg:  2m 29s | Max:  2m 29s | Hits:  98%/159   
      🟩 Test               Pass: 100%/1   | Total: 14m 26s | Avg: 14m 26s | Max: 14m 26s | Hits:  98%/159   
    
  • 🟩 python: Pass: 100%/1 | Total: 51m 15s | Avg: 51m 15s | Max: 51m 15s

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

👃 Inspect Changes

Modifications in project?

Project
CCCL Infrastructure
libcu++
CUB
Thrust
CUDA Experimental
python
+/- CCCL C Parallel Library
Catch2Helper

Modifications in project or dependencies?

Project
CCCL Infrastructure
libcu++
CUB
Thrust
CUDA Experimental
+/- python
+/- CCCL C Parallel Library
Catch2Helper

🏃‍ Runner counts (total jobs: 3)

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

@NaderAlAwar
NaderAlAwar requested a review from gevtushenko March 4, 2025 00:48
@github-actions

github-actions Bot commented Mar 4, 2025

Copy link
Copy Markdown
Contributor
🟩 CI finished in 50m 29s: Pass: 100%/3 | Total: 1h 07m | Avg: 22m 25s | Max: 50m 18s | Hits: 98%/320
  • 🟩 cccl_c_parallel: Pass: 100%/2 | Total: 16m 57s | Avg: 8m 28s | Max: 14m 26s | Hits: 98%/320

    🟩 cpu
      🟩 amd64              Pass: 100%/2   | Total: 16m 57s | Avg:  8m 28s | Max: 14m 26s | Hits:  98%/320   
    🟩 ctk
      🟩 12.8               Pass: 100%/2   | Total: 16m 57s | Avg:  8m 28s | Max: 14m 26s | Hits:  98%/320   
    🟩 cudacxx
      🟩 nvcc12.8           Pass: 100%/2   | Total: 16m 57s | Avg:  8m 28s | Max: 14m 26s | Hits:  98%/320   
    🟩 cudacxx_family
      🟩 nvcc               Pass: 100%/2   | Total: 16m 57s | Avg:  8m 28s | Max: 14m 26s | Hits:  98%/320   
    🟩 cxx
      🟩 GCC13              Pass: 100%/2   | Total: 16m 57s | Avg:  8m 28s | Max: 14m 26s | Hits:  98%/320   
    🟩 cxx_family
      🟩 GCC                Pass: 100%/2   | Total: 16m 57s | Avg:  8m 28s | Max: 14m 26s | Hits:  98%/320   
    🟩 gpu
      🟩 rtx2080            Pass: 100%/2   | Total: 16m 57s | Avg:  8m 28s | Max: 14m 26s | Hits:  98%/320   
    🟩 jobs
      🟩 Build              Pass: 100%/1   | Total:  2m 31s | Avg:  2m 31s | Max:  2m 31s | Hits:  98%/160   
      🟩 Test               Pass: 100%/1   | Total: 14m 26s | Avg: 14m 26s | Max: 14m 26s | Hits:  98%/160   
    
  • 🟩 python: Pass: 100%/1 | Total: 50m 18s | Avg: 50m 18s | Max: 50m 18s

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

👃 Inspect Changes

Modifications in project?

Project
CCCL Infrastructure
libcu++
CUB
Thrust
CUDA Experimental
python
+/- CCCL C Parallel Library
Catch2Helper

Modifications in project or dependencies?

Project
CCCL Infrastructure
libcu++
CUB
Thrust
CUDA Experimental
+/- python
+/- CCCL C Parallel Library
Catch2Helper

🏃‍ Runner counts (total jobs: 3)

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

@github-actions

github-actions Bot commented Mar 5, 2025

Copy link
Copy Markdown
Contributor
🟩 CI finished in 1h 01m: Pass: 100%/3 | Total: 1h 14m | Avg: 24m 47s | Max: 57m 26s | Hits: 98%/320
  • 🟩 cccl_c_parallel: Pass: 100%/2 | Total: 16m 56s | Avg: 8m 28s | Max: 14m 26s | Hits: 98%/320

    🟩 cpu
      🟩 amd64              Pass: 100%/2   | Total: 16m 56s | Avg:  8m 28s | Max: 14m 26s | Hits:  98%/320   
    🟩 ctk
      🟩 12.8               Pass: 100%/2   | Total: 16m 56s | Avg:  8m 28s | Max: 14m 26s | Hits:  98%/320   
    🟩 cudacxx
      🟩 nvcc12.8           Pass: 100%/2   | Total: 16m 56s | Avg:  8m 28s | Max: 14m 26s | Hits:  98%/320   
    🟩 cudacxx_family
      🟩 nvcc               Pass: 100%/2   | Total: 16m 56s | Avg:  8m 28s | Max: 14m 26s | Hits:  98%/320   
    🟩 cxx
      🟩 GCC13              Pass: 100%/2   | Total: 16m 56s | Avg:  8m 28s | Max: 14m 26s | Hits:  98%/320   
    🟩 cxx_family
      🟩 GCC                Pass: 100%/2   | Total: 16m 56s | Avg:  8m 28s | Max: 14m 26s | Hits:  98%/320   
    🟩 gpu
      🟩 rtx2080            Pass: 100%/2   | Total: 16m 56s | Avg:  8m 28s | Max: 14m 26s | Hits:  98%/320   
    🟩 jobs
      🟩 Build              Pass: 100%/1   | Total:  2m 30s | Avg:  2m 30s | Max:  2m 30s | Hits:  97%/160   
      🟩 Test               Pass: 100%/1   | Total: 14m 26s | Avg: 14m 26s | Max: 14m 26s | Hits:  98%/160   
    
  • 🟩 python: Pass: 100%/1 | Total: 57m 26s | Avg: 57m 26s | Max: 57m 26s

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

👃 Inspect Changes

Modifications in project?

Project
CCCL Infrastructure
libcu++
CUB
Thrust
CUDA Experimental
python
+/- CCCL C Parallel Library
Catch2Helper

Modifications in project or dependencies?

Project
CCCL Infrastructure
libcu++
CUB
Thrust
CUDA Experimental
+/- python
+/- CCCL C Parallel Library
Catch2Helper

🏃‍ Runner counts (total jobs: 3)

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

@@ -0,0 +1,13 @@
//===----------------------------------------------------------------------===//
//
// Part of CUDA Experimental in CUDA C++ Core Libraries,

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.

Is this correct? Isn't this the C parallel library?

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'm not really sure, this is the license used in the other files in the same directory

Comment thread c/parallel/src/util/tuning.cpp Resolved
@NaderAlAwar
NaderAlAwar enabled auto-merge (squash) March 7, 2025 15:35
@github-actions

github-actions Bot commented Mar 7, 2025

Copy link
Copy Markdown
Contributor
🟩 CI finished in 53m 59s: Pass: 100%/3 | Total: 1h 07m | Avg: 22m 31s | Max: 50m 49s | Hits: 98%/320
  • 🟩 cccl_c_parallel: Pass: 100%/2 | Total: 16m 45s | Avg: 8m 22s | Max: 14m 23s | Hits: 98%/320

    🟩 cpu
      🟩 amd64              Pass: 100%/2   | Total: 16m 45s | Avg:  8m 22s | Max: 14m 23s | Hits:  98%/320   
    🟩 ctk
      🟩 12.8               Pass: 100%/2   | Total: 16m 45s | Avg:  8m 22s | Max: 14m 23s | Hits:  98%/320   
    🟩 cudacxx
      🟩 nvcc12.8           Pass: 100%/2   | Total: 16m 45s | Avg:  8m 22s | Max: 14m 23s | Hits:  98%/320   
    🟩 cudacxx_family
      🟩 nvcc               Pass: 100%/2   | Total: 16m 45s | Avg:  8m 22s | Max: 14m 23s | Hits:  98%/320   
    🟩 cxx
      🟩 GCC13              Pass: 100%/2   | Total: 16m 45s | Avg:  8m 22s | Max: 14m 23s | Hits:  98%/320   
    🟩 cxx_family
      🟩 GCC                Pass: 100%/2   | Total: 16m 45s | Avg:  8m 22s | Max: 14m 23s | Hits:  98%/320   
    🟩 gpu
      🟩 rtx2080            Pass: 100%/2   | Total: 16m 45s | Avg:  8m 22s | Max: 14m 23s | Hits:  98%/320   
    🟩 jobs
      🟩 Build              Pass: 100%/1   | Total:  2m 22s | Avg:  2m 22s | Max:  2m 22s | Hits:  98%/160   
      🟩 Test               Pass: 100%/1   | Total: 14m 23s | Avg: 14m 23s | Max: 14m 23s | Hits:  98%/160   
    
  • 🟩 python: Pass: 100%/1 | Total: 50m 49s | Avg: 50m 49s | Max: 50m 49s

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

👃 Inspect Changes

Modifications in project?

Project
CCCL Infrastructure
libcu++
CUB
Thrust
CUDA Experimental
python
+/- CCCL C Parallel Library
Catch2Helper

Modifications in project or dependencies?

Project
CCCL Infrastructure
libcu++
CUB
Thrust
CUDA Experimental
+/- python
+/- CCCL C Parallel Library
Catch2Helper

🏃‍ Runner counts (total jobs: 3)

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

@NaderAlAwar
NaderAlAwar merged commit 84889c6 into NVIDIA:main Mar 7, 2025
@github-project-automation github-project-automation Bot moved this from In Review to Done in CCCL Mar 7, 2025
davebayer pushed a commit to davebayer/cccl that referenced this pull request Mar 12, 2025
* Move nominal_4b_items_to_items to separate file for reuse

* Use correct offset type for merge sort and unique by key
davebayer pushed a commit to davebayer/cccl that referenced this pull request Apr 7, 2025
* Move nominal_4b_items_to_items to separate file for reuse

* Use correct offset type for merge sort and unique by key
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.

[FEA]: Implement cccl.c.parallel version of unique by key

4 participants