---
id: pr-cccl-cub-8423
repo: NVIDIA/cccl
pr: 8423
title: Vectorize mbarrier initialization in warpspeed scan
author: bernhardmgruber
date: '2026-04-15'
url: https://github.com/NVIDIA/cccl/pull/8423
source_category: upstream-code
architectures:
- sm100
tags:
- arch_pipeline
- memory_primitives
- scan
- warpspeed scan
techniques:
- pipeline-stages
- vectorized-loads
- warp-specialization
hardware_features: []
kernel_types:
- memory
- runtime
- scan
languages:
- cuda-cpp
captured_at: '2026-05-16'
status: merged
merge_sha: 7b7a480ae3839ca9a3bacb49814b054b9b244947
inclusion_reason: 'Transfer primitive policy carefully: coalescing, shared-memory staging, partial tiles, and determinism are part of the contract.'
changed_paths:
- cub/cub/detail/warpspeed/sync_handler.cuh
- cub/cub/device/dispatch/kernels/kernel_scan_warpspeed.cuh
artifact_dir: evidence/pull-bundles/cccl-cub/gh-8423
head_ref_oid: ada5e2c285659bb797486be08a4d78d98f387bb0
---
# Vectorize mbarrier initialization in warpspeed scan

- Repository: `NVIDIA/cccl`
- PR: [8423](https://github.com/NVIDIA/cccl/pull/8423)
- Decision: `Transfer primitive policy carefully: coalescing, shared-memory staging, partial tiles, and determinism are part of the contract.`
- Reason: Transfer primitive policy carefully: coalescing, shared-memory staging, partial tiles, and determinism are part of the contract.

## Changed Paths

- `cub/cub/detail/warpspeed/sync_handler.cuh`
- `cub/cub/device/dispatch/kernels/kernel_scan_warpspeed.cuh`

## Evidence Bundle

`evidence/pull-bundles/cccl-cub/gh-8423`
