Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this request and that I agree to the Code of Conduct

Area

CUB

Is your feature request related to a problem? Please describe.

Vectorized global loads in CUB

This note lists every place where a CUB device algorithm loads global memory with vector instructions. For each one
it names the implementation, and it traces how the loads reach algorithms that are built on other algorithms. It is
meant as the starting point for centralizing the vectorized-load code.

Source: a static read of cub/ at upstream main (4984d75517), with default tuning policies and no user-provided
policies.

Terms used below:

• Striped: thread t loads vectors t, t + N, t + 2N, and so on, where N is the number of threads in the
  block or warp.
• Blocked: each thread loads consecutive vectors.
• Async copies to shared memory (LDGSTS, TMA) are a separate mechanism and are listed at the end.

Summary

• Six separate implementations. CUB has four striped implementations of vectorized register loads: AgentReduce,
  AgentFind, the DeviceTransform vectorized kernel, and the DeviceFor::ForEachCopy wrapper. It has two blocked
  ones: BlockLoad/WarpLoad VECTORIZE, and AgentHistogram's BlockLoad of vector-typed items. A helper in
  agent_batch_memcpy.cuh also builds vectors out of 32-bit loads so that it can issue vector stores.
• No two agree. They differ in vector type, eligibility check, how and where alignment is checked, partial-tile
  handling, and cache-modifier support (see Where they differ).
• VECTORIZE is rare. Only DeviceTopK uses BLOCK_LOAD_VECTORIZE in its default tunings, and no tuning uses
  WARP_LOAD_VECTORIZE. Every other BlockLoad user defaults to DIRECT, STRIPED, TRANSPOSE or WARP_TRANSPOSE.
• Reuse spreads the striped loads. DeviceSegmentedReduce gets them through AgentReduce and AgentWarpReduce.
  DeviceFind::LowerBound/UpperBound and DeviceCopy::Copy get them through DeviceTransform.
• DeviceTransform mostly uses async copies. On sm_80 and newer it defaults to LDGSTS (sm_80–89) or TMA (sm_90+).
  Its vectorized kernel is the default only below sm_80 and a fallback above it.

Implementations

Where they differ

• Vector type. There are three conventions.
  • CubVector<T, N> over the value type (reduce, find, for, histogram). This is a native CUDA vector only for the
    built-in arithmetic types. For any other T it's a plain struct aligned like T, so its "vector" load is issued
    as narrower loads.
  • CubVector over a machine word (BlockLoad). It uses the largest word that divides T
    (UnitWord<T>::DeviceWord), so it works for any T.
  • Raw integer words (DeviceTransform). This is the only implementation that can emit 32-byte loads.
• Eligibility. All of them gate on thrust::is_contiguous_iterator_v, but the type gates differ:
  • reduce: primitive or trivially copyable, and at most 8 B
  • find: trivially copyable
  • for: none
  • BlockLoad: the value type must equal T and must not be volatile
  • transform: power-of-two value sizes
• Alignment.
  • Device-side alignof checks: reduce, BlockLoad, histogram.
  • A device-side sizeof check on every tile: find.
  • Host-side checks: transform (sizeof(T) * vec_size) and for (alignof).
• Vector width. It comes from the policy (reduce, find, histogram, transform), is hard-coded to 4 (for), or is
  derived from the bytes per thread (BlockLoad).
• Partial tiles. Every implementation loads them scalar, except DeviceFor, which finishes the remainder in the
  last vector slot.
• Cache modifiers. Reduce, find and histogram apply them through CacheModifiedInputIterator<VectorT>, and
  BlockLoad through ThreadLoad<MODIFIER>. Transform and for have none.
• AgentReduce sizing (from reading the code, not run). vec_size is clamped to ITEMS_PER_THREAD, but
  VectorT is built from the unclamped VECTOR_LOAD_LENGTH. The default sm_60+ policy uses a vector length of 4 and
  clamp(64 / sizeof(AccumT), 1, 32) items per thread. A 17–32 B accumulator over a small input type therefore gives
  2 or 3 items per thread, yet each thread loads a 4-item vector. Those vectors start vec_size items apart, so most
  lanes' loads are misaligned, and the loads write past the per-thread array.

Device algorithms

Propagation through other algorithms

No device algorithm reaches DeviceTopK, so blocked VECTORIZE loads do not propagate.

Async copies to shared memory

These copy global memory straight into shared memory, so they're a different facility from register vector loads.
BlockLoadToShared is the only shared implementation.

Describe the solution you'd like

We should unify all loads to registers into a single facility probably under BlockLoad.

There should be two top level entry points

• BlockLoadDirectVectorized
• BlockLoadStripedVectorized

Every algorithm in CUB that utilizes vectorized loads should go through one of the two

Describe alternatives you've considered

No response

Additional context

No response