Hacktoberfest 2026: những issue maintainer đã đánh dấu cho tháng Mười, đang mở và phù hợp người mới. Xem issue Hacktoberfest

[ENHANCEMENT]: Implement cooperative device bulk operations

Đang mở
  • Có gì không ổn?
#674 1 bình luận 1 reaction 0 người được giao Xem trên GitHub

Maintainer thường phản hồi trong vòng 2 ngày

@0z5a đang làm issue này rồi.

Từ ngày 18/9/2026.

Đánh giá

Độ khó
5/5
Thời gian dự kiến
Hơn một tuần
Mức phù hợp với người mới
35/100
Loại issue
Tính năng
Độ rõ ràng
Khá rõ ràng
Mức độ hoạt động
Sôi nổi
Công nghệ
cpp
Lĩnh vực
data, performance

Hướng nghiên cứu

Bắt đầu bằng cách đọc cooperative device API được mô tả trong issue này và ví dụ được tham chiếu trong #672. Truy theo đường đi hiện có của thao tác chèn một khóa, sau đó xác định cách một thao tác dựa trên phạm vi có thể phân phối công việc giữa các cooperative groups trong khi broadcast các hash. Hoàn tất có nghĩa là bulk API chấp nhận một phạm vi đầu vào và hỗ trợ việc phân chia thích ứng các worker groups mà không tính toán hash dư thừa.

Do mô hình lập chỉ mục viết ra từ nội dung của issue.

Mô tả

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

Ngôn ngữ chính
Cuda
Star
671
Fork
122
Merge trung bình
4 ngày 19 giờ
Pull request đã merge (30 ngày)
10

Chuẩn bị môi trường

Mở trong Codespaces

Khởi chạy dev container của dự án ngay trên trình duyệt, bằng tài khoản GitHub của bạn.

Bắt đầu từ đâu

  1. Đọc hết issue, rồi đọc hướng dẫn đóng góp của dự án.
  2. Bình luận trên issue rằng bạn sẽ nhận — tránh hai người làm cùng một việc.
  3. Fork repository và làm thay đổi trên một nhánh.
  4. Mở pull request có tham chiếu số hiệu của issue.

Issue khác của NVIDIA/cuCollections

Tất cả issue của NVIDIA/cuCollections

Issue tương tự

Thêm issue về Data Engineering

Nhận issue mới trong hộp thư của bạn

Bản tóm tắt ngắn những issue GitHub phù hợp với người mới.