Skip to content

[FEA] Unify incrementer_t test functors#9896

Open
ahfoysal wants to merge 1 commit into
NVIDIA:mainfrom
ahfoysal:agent/unify-incrementer-t
Open

[FEA] Unify incrementer_t test functors#9896
ahfoysal wants to merge 1 commit into
NVIDIA:mainfrom
ahfoysal:agent/unify-incrementer-t

Conversation

@ahfoysal

@ahfoysal ahfoysal commented Jul 16, 2026

Copy link
Copy Markdown

Description

Closes #9872

Centralizes the five duplicate incrementer_t test functors in c2h/include/c2h/operator.cuh. The shared functor accepts both one-argument and two-argument DeviceFor callbacks, while retaining the compile-time offset-type check used by the DeviceFor::Bulk test.

Testing

  • pre-commit run --files c2h/include/c2h/operator.cuh cub/test/catch2_test_device_bulk.cu cub/test/catch2_test_device_for.cu cub/test/catch2_test_device_for_copy.cu cub/test/catch2_test_device_for_each_in_extents.cu cub/test/catch2_test_device_for_each_in_layout.cu
  • git diff --check
  • CUDA build/tests not run locally because this macOS environment does not have nvcc or an NVIDIA GPU.

Checklist

  • Existing tests cover these changes.
  • Documentation is up to date with these changes.

@copy-pr-bot

copy-pr-bot Bot commented Jul 16, 2026

Copy link
Copy Markdown
Contributor

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

@cccl-authenticator-app cccl-authenticator-app Bot moved this from Todo to In Progress in CCCL Jul 16, 2026
@ahfoysal
ahfoysal force-pushed the agent/unify-incrementer-t branch from d2af527 to c52d64d Compare July 16, 2026 09:00
@ahfoysal
ahfoysal marked this pull request as ready for review July 16, 2026 09:00
@ahfoysal
ahfoysal requested review from a team as code owners July 16, 2026 09:00
@ahfoysal
ahfoysal requested review from gevtushenko and wmaxey July 16, 2026 09:00
@cccl-authenticator-app cccl-authenticator-app Bot moved this from In Progress to In Review in CCCL Jul 16, 2026
@coderabbitai

coderabbitai Bot commented Jul 16, 2026

Copy link
Copy Markdown
Contributor

Review Change Stack

📝 Walkthrough

Summary by CodeRabbit

  • Tests

    • Expanded GPU test coverage for bulk and device iteration operations, including copying, extents, and layouts.
    • Added consistent validation for device-side counter increments and offset type handling across related test scenarios.
  • Refactor

    • Consolidated shared test functionality into a reusable device operation, reducing duplication while preserving existing test behavior.

Walkthrough

Changes

The shared incrementer_t device functor is added to c2h/operator.cuh. CUB device algorithm tests now include that header and remove duplicated local implementations while preserving their existing usages.

Incrementer functor unification

Layer / File(s) Summary
Shared incrementer contract
c2h/include/c2h/operator.cuh
Adds incrementer_t, offset-type validation, atomic counter updates, and a deduction guide.
Device test migration
cub/test/catch2_test_device_*.cu
Updates device tests to include the shared functor and removes duplicated local definitions.

Assessment against linked issues

Objective Addressed Explanation
Unify incrementer_t functors [#9872]

Suggested reviewers: miscco


Comment @coderabbitai help to get the list of available commands.

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

Actionable comments posted: 1


ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: b1f0e71d-51f3-40f5-b92e-2988c298c043

📥 Commits

Reviewing files that changed from the base of the PR and between 0d919ed and c52d64d.

📒 Files selected for processing (6)
  • c2h/include/c2h/operator.cuh
  • cub/test/catch2_test_device_bulk.cu
  • cub/test/catch2_test_device_for.cu
  • cub/test/catch2_test_device_for_copy.cu
  • cub/test/catch2_test_device_for_each_in_extents.cu
  • cub/test/catch2_test_device_for_each_in_layout.cu

Comment on lines +111 to +119
int* d_counts;

template <class OffsetT, class... Args>
__device__ void operator()(OffsetT i, Args...)
{
static_assert(cuda::std::is_void_v<ExpectedOffsetT> || cuda::std::is_same_v<ExpectedOffsetT, OffsetT>,
"ExpectedOffsetT and OffsetT must be the same type");
atomicAdd(d_counts + i, 1); // Check if `i` was served more than once
}

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.

📐 Maintainability & Code Quality | 🟡 Minor | ⚡ Quick win

important: Declare the immutable member const, and use the required device API annotation and noexcept on this non-throwing callable.

-  int* d_counts;
+  int* const d_counts;

-  __device__ void operator()(OffsetT i, Args...)
+  _CCCL_DEVICE_API void operator()(OffsetT i, Args...) const noexcept

As per coding guidelines, “All variables that are not modified must be declared const,” “Functions must be marked with _CCCL_*_API,” and non-throwing functions must be marked noexcept.

📝 Committable suggestion

‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.

Suggested change
int* d_counts;
template <class OffsetT, class... Args>
__device__ void operator()(OffsetT i, Args...)
{
static_assert(cuda::std::is_void_v<ExpectedOffsetT> || cuda::std::is_same_v<ExpectedOffsetT, OffsetT>,
"ExpectedOffsetT and OffsetT must be the same type");
atomicAdd(d_counts + i, 1); // Check if `i` was served more than once
}
int* const d_counts;
template <class OffsetT, class... Args>
_CCCL_DEVICE_API void operator()(OffsetT i, Args...) const noexcept
{
static_assert(cuda::std::is_void_v<ExpectedOffsetT> || cuda::std::is_same_v<ExpectedOffsetT, OffsetT>,
"ExpectedOffsetT and OffsetT must be the same type");
atomicAdd(d_counts + i, 1); // Check if `i` was served more than once
}

Source: Coding guidelines

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

Status: In Review

Development

Successfully merging this pull request may close these issues.

[FEA] Unify incrementer_t functors

1 participant