Vectorized Memory Access: float4, Alignment, and Tails

Core Concept

float2, float4, or 128-bit load/store operations move multiple adjacent elements per thread. They mainly reduce memory-instruction count, address arithmetic, and loop iterations; they do not reduce required data bytes.

Why It Affects Performance

Vectorization helps when instruction or issue overhead matters, but natural alignment, contiguous layout, and correct tail handling are mandatory. Misaligned reinterpretation can be invalid or slower.

Example or Pseudocode

int j = (blockIdx.x * blockDim.x + threadIdx.x) * 4;
if (j + 3 < n) {
  float4 v = reinterpret_cast<const float4*>(x)[j/4];
  reinterpret_cast<float4*>(y)[j/4] = v;
}
// Handle n%4 remaining elements with a scalar tail.

Profiling and Performance Judgment

Inspect SASS for wide loads/stores and total instruction count. Compare scalar, float2, and float4 latency, DRAM bytes, sectors, and LG throttle. Test nonmultiples of four and offset tensor views.

Five-Minute Exercise

For N=1,000,003, how many elements use the float4 path and how many remain?