NVIDIA / NVIDIA/cuCollections

[ENHANCEMENT]: `__device__ for_each` that provides a mutable reference rather than a copy of the slot to the callback function

Open
#639 0 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

type: feature request
Dominant language
Cuda
Stars
667
Forks
120
Avg merge
7d 5h
Merged PRs (30d)
4

Description

Is your feature request related to a problem? Please describe.

Feature request on behalf of @LutzCle who is implementing a mark join for an internal project.

tl;dr The algorithm requires setting a "(matching) slot has been seen" bit during probing which is explicitly excluded from the hash value generation and key comparison, so it doesn't break the probing sequence if set or unset.

Our initial approach was to use the for_each device API and set the bit for each matching element but unfortunately this doesn't work since the current for_each implementation returns a copy of the matching slot rather than a mutable reference.

Describe the solution you'd like

Provide both copy and mutable reference options for for_each.

Naming ideas: for_each(_reference)/for_each_copy/for_each_mutable, ...

Describe alternatives you've considered

The current workaround consists of basically re-implementing for_each with the proposed feature in a custom kernel which involves using some of the internal cuco::detail machinery - not ideal.

Additional context

I'm not sure if this feature is a foot gun or not so I would like to spark a discussion around the potential dangers.

Contributor guide

Open the contributing guide

First steps

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. Open a pull request that references the issue number.

Research direction

Start by reviewing the current for_each device API and the cuco::detail machinery used by the custom-kernel workaround. Define the copy and mutable-reference options, including their safety implications, and consider how the existing API behavior should be covered.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
api
Issue type
Feature
Difficulty
5/5
Estimated time
Over a week
Activity status
Stale
Clarity
Mostly clear
Newbie friendliness
35/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.