Skip to content

Restructure the RLE encode tuning policy#10028

Merged
bernhardmgruber merged 1 commit into
NVIDIA:mainfrom
nanan-nvidia:rle-encode-policy-container
Jul 21, 2026
Merged

Restructure the RLE encode tuning policy#10028
bernhardmgruber merged 1 commit into
NVIDIA:mainfrom
nanan-nvidia:rle-encode-policy-container

Conversation

@nanan-nvidia

Copy link
Copy Markdown
Contributor

This is taken out from #9699 so that we can ship the changes to the tuning APIs first.

The changes are:

// before
_CCCL_HOST_DEVICE_API constexpr auto operator()(cuda::compute_capability) const -> cub::RleEncodePolicy
{
  return {256, 14, cub::BLOCK_LOAD_DIRECT, cub::LOAD_CA, cub::BLOCK_SCAN_WARP_SCANS, {}};
}

// after
_CCCL_HOST_DEVICE_API constexpr auto operator()(cuda::compute_capability) const -> cub::RleEncodePolicy
{
  return {cub::RleAlgorithm::lookback,
          {256, 14, cub::BLOCK_LOAD_DIRECT, cub::LOAD_CA, cub::BLOCK_SCAN_WARP_SCANS, {}}};
}

After this is merged, #9699 needs to be rebased from here.

@nanan-nvidia
nanan-nvidia requested review from a team as code owners July 20, 2026 21:09
@nanan-nvidia
nanan-nvidia requested a review from gonidelis July 20, 2026 21:09
@nanan-nvidia
nanan-nvidia requested a review from pauleonix July 20, 2026 21:09
@github-project-automation github-project-automation Bot moved this to Todo in CCCL Jul 20, 2026
@cccl-authenticator-app cccl-authenticator-app Bot moved this from Todo to In Review in CCCL Jul 20, 2026
@coderabbitai

coderabbitai Bot commented Jul 20, 2026

Copy link
Copy Markdown
Contributor

Review Change Stack

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 51183fec-c718-4d48-90ac-3d261be7a96e

📥 Commits

Reviewing files that changed from the base of the PR and between c3a18e7 and cb660f6.

📒 Files selected for processing (4)
  • cub/benchmarks/bench/run_length_encode/encode.cu
  • cub/cub/device/device_run_length_encode.cuh
  • cub/cub/device/dispatch/tuning/tuning_rle_encode.cuh
  • cub/test/catch2_test_device_run_length_encode_env.cu

📝 Walkthrough

Summary by CodeRabbit

  • New Features

    • Added explicit algorithm selection for run-length encoding policies, currently supporting the lookback algorithm.
    • Added structured lookback tuning settings for clearer policy configuration and reporting.
  • Improvements

    • Updated device run-length encoding to use the new policy structure.
    • Preserved compute-capability-specific tuning behavior while improving policy organization and diagnostics.

Walkthrough

Changes

The RLE tuning policy now carries an explicit lookback algorithm and nested lookback configuration. Compute-capability selection, device adaptation, benchmarks, and tests use the updated structure.

RLE policy refactor

Layer / File(s) Summary
Policy types and serialization
cub/cub/device/dispatch/tuning/tuning_rle_encode.cuh
Adds RleAlgorithm and RleLookbackPolicy, nests lookback tuning in RleEncodePolicy, and updates comparisons and streaming.
Compute-capability tuning selection
cub/cub/device/dispatch/tuning/tuning_rle_encode.cuh
Returns lookback sub-policies for SM-specific cases and wraps them with RleAlgorithm::lookback.
Adapter and policy consumers
cub/cub/device/device_run_length_encode.cuh, cub/benchmarks/bench/run_length_encode/encode.cu, cub/test/catch2_test_device_run_length_encode_env.cu
Updates policy adaptation, benchmark construction, test fixtures, aggregate initialization, and expected serialized output.

Suggested reviewers: bernhardmgruber, gonidelis


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

@github-actions

Copy link
Copy Markdown
Contributor

🥳 CI Workflow Results

🟩 Finished in 1h 59m: Pass: 100%/287 | Total: 2d 09h | Max: 1h 54m | Hits: 78%/257856

See results here.

@bernhardmgruber
bernhardmgruber merged commit 426f9f8 into NVIDIA:main Jul 21, 2026
308 of 309 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

Archived in project

Development

Successfully merging this pull request may close these issues.

2 participants