Hacktoberfest 2026：メンテナが10月に向けて印を付けた、オープンで初心者向けの issue。 Hacktoberfest の issue を見る

[ENHANCEMENT]: Implement cooperative device bulk operations

オープン
  • 何か問題がありますか？
#674 コメント 1 件 リアクション 1 件 担当者 0 名 GitHub で見る

メンテナーはふだん 2 日以内に返信

@0z5a がすでに取り組んでいます。

2026年9月18日 から。

評価

難易度
5/5
見積もり時間
1週間以上
初心者へのやさしさ
35/100
issue の種類
機能追加
明瞭さ
おおむね明確
活発さ
活発
技術スタック
cpp
領域
data, performance

調査の方向性

まず、この issue で説明されている cooperative device API と、#672 で参照されている例を読んでください。次に、既存の単一キー挿入パスを追跡し、そのうえで、ハッシュをブロードキャストしながら、範囲ベースの操作で cooperative groups 間に作業を分配する方法を判断してください。bulk API が入力範囲を受け付け、冗長なハッシュ計算なしで worker groups の適応的な細分化をサポートすれば完了です。

索引モデルが issue の本文から書いたものです。

説明

type: improvement

Pitch

Device bulk operations such as

template <class CG, class InputIt>
__device__ void insert(CG const& group, InputIt first, InputIt last)

, i.e., assigning a cooperative group to process N input elements cooperatively, allow for adaptive and optimized work distribution.

Background

Our current device CG APIs look as follows:

template <class CG, class Key>
__device__ void insert(CG const& group, Key key)

, i.e., they take a single input datum and cooperatively insert it into the container. Each thread in the CG is expected to get the same key as input.

The problem with this approach is that it leads to redundant computation:
Each thread in the CG needs to compute the hash value of the (exact same) input key, before proceeding to the probing/insertion part. This imposes unnecessary pressure on the compute pipelines, which reduces overall throughput mainly in cases where the kernel is compute bound (which likely happens if the data structure fits (partially) into the L2$). However, experiments show that even for large data structures (>2GB) the performance improvement is still visible (>5%).

The bulk device operation eliminates this drawback as follows:
Each thread loads one input key at a time, resulting in a coalesced memory acces pattern, and then computes its respective hash value. From here we go into a cooperative loop, iterating over all ranks in the group. At each iteration, we shuffle the computed hash value from the target rank to all threads in the group (broadcast). From here we can proceed with the standard cooperative insertion/probing approach, as each thread holds the same hash value.

Another benefit is in the API:
The user may give us any CG size as input, which we can then adaptively subdivide into multiple worker groups that insert the given range of input keys as efficient as possible. In that sense, this API becomes adaptive in relation to the given input CG size.

One example of this approach is given in #672

主要言語
Cuda
スター
671
フォーク
122
平均マージ
4日 19時間
マージ済み PR（30日）
10

環境構築

Codespaces で開く

このプロジェクトの開発コンテナを、あなたの GitHub アカウントでブラウザ上に起動します。

はじめの一歩

  1. issue を最後まで読み、次にプロジェクトのコントリビューションガイドを読みます。
  2. 着手することを issue にコメントします — 二人が同じ作業をするのを防げます。
  3. リポジトリをフォークし、ブランチを切って変更します。
  4. issue 番号を参照したプルリクエストを送ります。

NVIDIA/cuCollections のほかの issue

NVIDIA/cuCollections の issue をすべて見る

似ている issue

Data Engineering の issue をもっと見る

新しい issue をメールで受け取る

初心者向けの GitHub issue を短くまとめたダイジェスト。