Hacktoberfest 2026:维护者为十月标记出来的 issue,仍然开放、适合新手。 浏览 Hacktoberfest issue

[ENHANCEMENT]: Implement cooperative device bulk operations

未关闭
#674 1 条评论 1 个 reaction 已指派 0 人 在 GitHub 查看

维护者通常 2 天内回复

@0z5a 已经在做这个了。

开始于 2026年9月18日。

评估

难度
5/5
预计耗时
一周以上
新手友好度
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 小时
30 天内合并 PR
10

环境准备

在 Codespaces 中打开

在浏览器里用你自己的 GitHub 账号启动这个项目的开发容器。

从这里开始

  1. 先读完整个 Issue,再读项目的贡献指南。
  2. 在 Issue 下留言说明你要接手 —— 这能避免两个人做同样的事。
  3. Fork 仓库,在一个分支上完成修改。
  4. 提交 Pull Request,并在描述里引用这个 Issue 编号。

NVIDIA/cuCollections 的其他 Issue

查看 NVIDIA/cuCollections 的全部 Issue

相似的 Issue

更多 Data Engineering Issue

把新 issue 发到你的邮箱

精选适合新手参与的 GitHub issue 摘要。