Skip to content

feat: unify is_even#9120

Draft
thedevmystic wants to merge 5 commits into
NVIDIA:mainfrom
thedevmystic:feat/is_even
Draft

feat: unify is_even#9120
thedevmystic wants to merge 5 commits into
NVIDIA:mainfrom
thedevmystic:feat/is_even

Conversation

@thedevmystic
Copy link
Copy Markdown
Contributor

@thedevmystic thedevmystic commented May 23, 2026

This is a continuation of #7455.

This PR focuses on unifying is_even functor across CCCL testing and benchmarking environment.

Quick Summary of Changed

  • Moved is_even logic to c2h/include/c2h/operator.cuh and enclosed it in c2h namespace.
  • changed scattered functors to utilize unified one.

Note

This does not unify cub/* files, because they uses custom_type. If we have to unify them, then I'll.

(Comment from @bernhardmgruber in #7455 specified duplicates in cub/*)

Also, this is an issue from my end, my system has integrated M2 chip and testing CCCL on it, is either impossible or painfully slow. So, changes can contain few errors.

@thedevmystic thedevmystic requested review from a team as code owners May 23, 2026 18:44
@github-project-automation github-project-automation Bot moved this to Todo in CCCL May 23, 2026
@copy-pr-bot
Copy link
Copy Markdown
Contributor

copy-pr-bot Bot commented May 23, 2026

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 Review in CCCL May 23, 2026
@thedevmystic thedevmystic marked this pull request as draft May 23, 2026 18:45
@cccl-authenticator-app cccl-authenticator-app Bot moved this from In Review to In Progress in CCCL May 23, 2026
@coderabbitai
Copy link
Copy Markdown
Contributor

coderabbitai Bot commented May 23, 2026

Review Change Stack

📝 Walkthrough

Summary by CodeRabbit

  • New Features

    • Added shared c2h::is_even predicate and complex32/complex64 type aliases to operator utilities for consistent use across tests and benchmarks.
  • Refactor

    • Consolidated duplicate is_even predicate definitions in test files and benchmarks to use the centralized utility, reducing code duplication.

Walkthrough

Adds a centralized c2h::is_even predicate functor to replace ~40 scattered local implementations across CUB, libcudacxx, and Thrust. The functor checks bit-level parity for integral types and real-component parity for complex types. All algorithm tests (partition, remove, copy, select) and benchmarks are updated to import and use the shared predicate.

Changes

Centralized is_even predicate

Layer / File(s) Summary
Core predicate definition and export
c2h/include/c2h/operator.cuh
Adds CUDA complex support and defines is_even_fn functor with operator() overloads for integral parity and cuda::std::complex real-component parity. Exports inline constexpr c2h::is_even instance and type aliases for complex32/complex64.
Partition and stable_partition algorithm tests
libcudacxx/test/.../alg.partitions/pstl_partition*.cu, thrust/testing/partition*.cu, thrust/testing/cuda/partition*.cu, thrust/testing/is_partitioned.cu
Updates partition/partition_copy/stable_partition test invocations across empty, contiguous, random-access, and converting-predicate cases to use c2h::is_even instead of local is_even<T> functors.
Remove and remove_if algorithm tests
libcudacxx/benchmarks/bench/remove_*.cu, thrust/testing/remove.cu, thrust/testing/cuda/remove.cu
Updates remove_if and remove_copy_if test invocations covering basic, stream, and stencil-based paths to use c2h::is_even; removes local functor definitions.
Copy_if and device select tests
cub/test/catch2_test_device_select_*.cu, libcudacxx/benchmarks/bench/copy_if/basic.cu, thrust/testing/copy.cu, thrust/testing/cuda/copy_if.cu
Updates copy_if and DeviceSelect::FlaggedIf tests for both in-place and out-of-place invocations; removes local predicate definitions.
SIMD mask, loadstore, and reduction tests
libcudacxx/test/.../simd/simd_test_utils.h, libcudacxx/test/.../simd/simd.mask.class/*.cpp, libcudacxx/test/.../simd/simd.loadstore/*.cpp, libcudacxx/test/.../simd/simd.reductions/*.cpp, libcudacxx/test/.../simd/simd.vec.class/ctor.pass.cpp
Removes is_even from test utilities and updates ~30 SIMD test files for mask construction, compound assignments, conversions, unary operators, load/store operations, and reduction tests.
Algorithm benchmarks
libcudacxx/benchmarks/bench/{copy_if,remove_if,remove_copy_if}/basic.cu, thrust/benchmarks/bench/{copy_if,remove_if,remove_copy_if}/basic.cu
Updates copy_if, remove_if, and remove_copy_if benchmarks to use c2h::is_even; removes local functor definitions including complex-type overloads.

Suggested labels

libcu++

Suggested reviewers

  • griwes
  • davebayer
  • pauleonix

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

Copy link
Copy Markdown
Contributor

@coderabbitai coderabbitai Bot left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actionable comments posted: 14


ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 5aed0315-6ddb-43dd-8a4f-12869d69cf33

📥 Commits

Reviewing files that changed from the base of the PR and between c47f140 and db3eb57.

📒 Files selected for processing (42)
  • c2h/include/c2h/operator.cuh
  • cub/test/catch2_test_device_select_api.cu
  • cub/test/catch2_test_device_select_flagged_if.cu
  • libcudacxx/benchmarks/bench/copy_if/basic.cu
  • libcudacxx/benchmarks/bench/remove_copy_if/basic.cu
  • libcudacxx/benchmarks/bench/remove_if/basic.cu
  • libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.partitions/pstl_partition.cu
  • libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.partitions/pstl_partition_copy.cu
  • libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.partitions/pstl_stable_partition.cu
  • libcudacxx/test/libcudacxx/std/numerics/simd/simd.loadstore/partial_load.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/simd/simd.loadstore/partial_store.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/simd/simd.loadstore/unchecked_load.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/simd/simd.loadstore/unchecked_store.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/simd/simd.mask.class/compound_assign.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/simd/simd.mask.class/conversion.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/simd/simd.mask.class/ctor.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/simd/simd.mask.class/subscript.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/simd/simd.mask.class/unary.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/simd/simd.reductions/all_of.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/simd/simd.reductions/any_of.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/simd/simd.reductions/none_of.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/simd/simd.reductions/reduce.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/simd/simd.reductions/reduce_count.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/simd/simd.reductions/reduce_max.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/simd/simd.reductions/reduce_max_index.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/simd/simd.reductions/reduce_min.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/simd/simd.reductions/reduce_min_index.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/simd/simd.vec.class/ctor.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/simd/simd_test_utils.h
  • thrust/benchmarks/bench/copy_if/basic.cu
  • thrust/benchmarks/bench/remove_copy_if/basic.cu
  • thrust/benchmarks/bench/remove_if/basic.cu
  • thrust/testing/copy.cu
  • thrust/testing/cuda/copy_if.cu
  • thrust/testing/cuda/is_partitioned.cu
  • thrust/testing/cuda/partition.cu
  • thrust/testing/cuda/partition_point.cu
  • thrust/testing/cuda/remove.cu
  • thrust/testing/is_partitioned.cu
  • thrust/testing/partition.cu
  • thrust/testing/partition_point.cu
  • thrust/testing/remove.cu

Comment thread cub/test/catch2_test_device_select_api.cu Outdated
Comment thread libcudacxx/test/libcudacxx/std/numerics/simd/simd.mask.class/ctor.pass.cpp Outdated
Mask all_true(true);
Mask all_false(false);
Mask alternating(is_even{});
Mask alternating(c2h::is_even);
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.

⚠️ Potential issue | 🟡 Minor | ⚡ Quick win

important: Line 43 declares alternating as mutable but never modifies it; make it const to satisfy the project constness rule.

As per coding guidelines "All non-const variables must use const if they are not modified".

Mask all_true(true);
Mask all_false(false);
Mask mixed(is_even{});
Mask mixed(c2h::is_even);
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.

⚠️ Potential issue | 🟡 Minor | ⚡ Quick win

important: Lines 45, 63, 90, and 117 create mask variables that are never mutated; declare each as const Mask ... to comply with constness requirements.

As per coding guidelines "All non-const variables must use const if they are not modified".

Also applies to: 63-63, 90-90, 117-117

if constexpr (N > 1)
{
Mask even(is_even{});
Mask even(c2h::is_even);
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.

⚠️ Potential issue | 🟡 Minor | ⚡ Quick win

important: Line 43 should use const Mask even(c2h::is_even); since even is not modified.

As per coding guidelines "All non-const variables must use const if they are not modified".

using Mask = typename Vec::mask_type;
Vec vec = make_iota_vec<T, N>();
Mask even(is_even{});
Mask even(c2h::is_even);
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.

⚠️ Potential issue | 🟡 Minor | ⚡ Quick win

important: Line 135 creates even as non-const but never mutates it; declare it const.

As per coding guidelines "All non-const variables must use const if they are not modified".

state.exec(nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch | nvbench::exec_tag::sync,
[&](nvbench::launch& launch) {
thrust::remove_if(policy(alloc, launch), in.begin(), in.end(), is_even{});
thrust::remove_if(policy(alloc, launch), in.begin(), in.end(), c2h::is_even);
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.

⚠️ Potential issue | 🟡 Minor | ⚡ Quick win

🧩 Analysis chain

🏁 Script executed:

#!/bin/bash
set -euo pipefail
# Inspect the exact snippet around the reported line
FILE="thrust/benchmarks/bench/remove_if/basic.cu"
if [ ! -f "$FILE" ]; then
  echo "Missing file: $FILE"
  exit 1
fi

echo "== File: $FILE =="
# Show enough context around line 30
nl -ba "$FILE" | sed -n '1,120p' | sed -n '20,60p'
echo
echo "== Search for remove_if call patterns in this file =="
rg -n "thrust::remove_if\(|::thrust::remove_if\(" "$FILE" || true
echo
echo "== Check for any using directives that might affect qualification =="
rg -n "using\s+namespace\s+thrust|using\s+thrust::|namespace\s+thrust" "$FILE" || true
echo
echo "== Check local aliases for remove_if =="
rg -n "remove_if" "$FILE" || true

Repository: NVIDIA/cccl

Length of output: 172


🏁 Script executed:

#!/bin/bash
set -euo pipefail
FILE="thrust/benchmarks/bench/remove_if/basic.cu"
if [ ! -f "$FILE" ]; then
  echo "Missing file: $FILE"
  exit 1
fi

echo "== File: $FILE =="
# Print lines 20-60 with line numbers using awk (nl not available)
awk 'NR>=20 && NR<=60 {printf "%5d\t%s\n", NR, $0}' "$FILE"

echo
echo "== Search for remove_if call patterns in this file =="
rg -n "thrust::remove_if\(|::thrust::remove_if\(" "$FILE" || true

echo
echo "== Check for any using directives that might affect qualification =="
rg -n "using\s+namespace\s+thrust|using\s+thrust::|namespace\s+thrust" "$FILE" || true

echo
echo "== Check for local aliases for remove_if =="
rg -n "remove_if" "$FILE" || true

Repository: NVIDIA/cccl

Length of output: 1274


suggestion: Qualify the free function call from the global namespace. File thrust/benchmarks/bench/remove_if/basic.cu:30 uses thrust::remove_if(...); update to ::thrust::remove_if(...) per the global qualification requirement.

Diff
-               thrust::remove_if(policy(alloc, launch), in.begin(), in.end(), c2h::is_even);
+               ::thrust::remove_if(policy(alloc, launch), in.begin(), in.end(), c2h::is_even);
📝 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
thrust::remove_if(policy(alloc, launch), in.begin(), in.end(), c2h::is_even);
::thrust::remove_if(policy(alloc, launch), in.begin(), in.end(), c2h::is_even);

Comment thread thrust/testing/cuda/copy_if.cu Outdated
Comment on lines +192 to +195
h_new_end = thrust::copy_if(h_data.begin(), h_data.end(), h_result.begin(), c2h::is_even);

copy_if_kernel<<<1, 1>>>(
exec, d_data.begin(), d_data.end(), d_result.begin(), is_even<int>(), d_new_end_vec.begin());
exec, d_data.begin(), d_data.end(), d_result.begin(), c2h::is_even, d_new_end_vec.begin());
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.

⚠️ Potential issue | 🟠 Major | ⚡ Quick win

important: Line 192 and Line 195 use the non-stencil copy_if overload inside TestCopyIfStencilDevice, so this test does not validate stencil behavior even though stencil inputs are prepared; wire these calls to the stencil overload (and apply the same fix to the non-bool block) to avoid missing stencil regressions.

Comment on lines +22 to +25
iterator ref = thrust::stable_partition(v.begin(), v.end(), c2h::is_even);

thrust::device_vector<iterator> result(1);
partition_point_kernel<<<1, 1>>>(exec, v.begin(), v.end(), is_even<int>(), result.begin());
partition_point_kernel<<<1, 1>>>(exec, v.begin(), v.end(), c2h::is_even, result.begin());
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.

⚠️ Potential issue | 🟠 Major | ⚡ Quick win

important: add a direct include for c2h::is_even in this file (#include <c2h/operator.cuh>). Line 22 and Line 25 currently depend on transitive includes, which is fragile and can fail under include-order changes. As per coding guidelines: "Files must include all headers related to the symbols they are using".

Comment on lines +312 to 313
std::ptrdiff_t n_true = thrust::count_if(h_data.begin(), h_data.end(), c2h::is_even);
std::ptrdiff_t n_false = n - n_true;
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.

⚠️ Potential issue | 🟠 Major | ⚡ Quick win

important: compute the true-count from the stencil range, not the data range. Line 312 uses h_data, but this test’s partition_copy classification uses h_stencil (Line 323 and Line 327). That can mis-size h_true_results / h_false_results and cause out-of-bounds writes when stencil parity differs from data parity.

@thedevmystic
Copy link
Copy Markdown
Contributor Author

Most issues that the AI pointed out was already there, but if we gotta fix them, tell me.

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

Labels

None yet

Projects

Status: In Progress

Development

Successfully merging this pull request may close these issues.

1 participant