Clusters, Data Types, Inline PTX, Pointers

Overview

Modern CUDA performance work often requires thinking across scheduling, memory topology, compiler behavior, and instruction-level control at the same time. Lesson 2 is where those layers begin to collapse into one picture. Thread block clusters change how blocks cooperate. Distributed shared memory changes what can be shared without bouncing through global memory. PTX and inline PTX expose instruction forms that CUDA C++ does not always surface directly. Pointer conversion instructions tie all of that back to the actual hardware address spaces.

Thread block clusters

A thread block cluster is a group of thread blocks guaranteed to be co-scheduled concurrently on adjacent SMs within a GPC. The important point is that this extends the hierarchy from threads -> thread blocks -> clusters -> grid without pretending one block is now spread across many SMs.

Key properties

• Clusters are explicit software constructs, not an automatic scheduling side effect.
• All CTAs in a cluster are guaranteed concurrent residency on nearby SMs.
• They are useful when cross-block data sharing is frequent enough to make global memory too expensive.
• They trade some coordination overhead for a larger effective working set.

Cluster sizing tradeoffs

Declaring a cluster

__global__ void __cluster_dims__(2, 1, 1)
cluster_kernel(float* input, float* output) {
  // kernel body
}

// launch shape must remain compatible with cluster dimensions

Cluster special registers

Distributed shared memory

Distributed shared memory is the cluster-era extension of shared memory. It allows direct memory access between SMs inside the same thread block cluster. Ownership still stays per block, but the programmer can address a neighboring CTA's shared-memory allocation without routing everything through global memory.

The practical win is a larger effective working set at a much cheaper memory tier than global memory. That is especially useful for top-k, matrix pipelines, reductions, or tile exchange patterns where neighboring CTAs need each other's intermediate state.

PTX, inline PTX, operands, and state spaces

PTX is NVIDIA's virtual ISA inside the CUDA toolchain. CUDA C++ lowers into PTX, which is then lowered again into machine-specific code. Inline PTX matters when you want instruction-level control without rewriting an entire kernel in PTX.

Why use inline PTX

• To access hardware features that CUDA C++ does not expose cleanly.
• To hand-optimize small performance-critical fragments.
• To request exact operand types, state spaces, or synchronization forms.
• To inject constants or instruction parameters through templates.

asm volatile(
  "ptx instruction here;"
  : /* outputs */
  : /* inputs */
  : /* clobbers */
);

Constraint modifiers and common operand classes

A memory clobber is often required when the assembly changes memory in ways the compiler cannot infer from the operand list alone. Without it, the compiler may legally reorder surrounding code in a way that breaks the intended semantics.

PTX state spaces

Data types that show up in Hopper code

• Standard integer and floating-point scalar families still matter for control flow and addressing.
• BF16, FP16, TF32, and FP8 matter because they change throughput, packing density, and tensor-core issue shape.
• Packed types matter because multiple low-precision values are often carried inside one 32-bit register.

Addressing, banking, and swizzling

Shared memory is still banked, and the bank structure still matters even on Hopper. Poor access patterns can serialize otherwise parallel traffic, which is why low-level layout choices continue to be part of performance engineering instead of a side detail.

Shared-memory banks

• Threads in a warp ideally spread their requests across banks instead of piling onto the same bank.
• Bank conflicts increase replay and reduce effective bandwidth.
• The more structured the access pattern, the easier it is to reason about conflict behavior.

128-byte transaction granularity

Many high-performance paths reason in 128-byte chunks because that aligns well with the granularity of efficient movement and layout transformations. This is one reason tile shapes and swizzles show up so often in Hopper code.

Generic pointers, state-space pointers, cvta, and mapa

CUDA C++ usually gives you generic pointers, but PTX instructions often need a pointer that is already understood in the correct state space. That is why conversion steps matter. The compiler cannot always assume which physical memory region a generic pointer should target inside low-level assembly.

cvta and mapa

// Convert a generic pointer into a specific address space
asm volatile("cvta.to.shared.u64 %0, %1;" : "=l"(smem_addr) : "l"(ptr));

// Map a shared-memory address into another CTA's shared-memory region
asm volatile("mapa.shared::cluster.u64 %0, %1, %2;" : "=l"(remote_addr) : "l"(addr), "r"(rank));

cvta turns a generic address into the appropriate state-space address. mapa takes that one step further for cluster programming by mapping a shared-memory address into the address space of another CTA inside the cluster.

Integration takeaways

1. Do not flatten clusters into “shared memory, but bigger.” Residency guarantees and ownership still matter.
2. Keep state-space terminology precise. .global, .shared, and generic pointers are not interchangeable.
3. Treat inline PTX as a surgical tool. Use it when exact semantics are the point, not as a stylistic preference.
4. Respect bank behavior and 128-byte layout assumptions. Swizzles and packed layouts exist for a reason.
5. Be exact about rank and mapping. Many cluster bugs are really address-space or rank-mapping bugs.

Glossary

Continue the course

Lesson 2 builds the address-space and instruction-level vocabulary Hopper kernels need. The next big step is synchronization under overlap: barriers, wait patterns, and correctness when loads and compute are intentionally in flight at the same time.