The 99KB Barrier: Probing Blackwell's Shared Memory Limits in the GLM-5-NVFP4 Optimization

# Message 912 — Assistant executing a Python probe on the remote server
ssh root@10.1.230.174 'source /root/ml-env/bin/activate && python3 -c "
import torch

props = torch.cuda.get_device_properties(0)
# List all attributes
attrs = [a for a in dir(props) if \"shared\" in a.lower() or \"smem\" in a.lower() or \"memory\" in a.lower() and not a.startswith(\"_\")]
for a in attrs:
    print(f\"  {a}: {getattr(props, a)}\")

import ctypes
cuda = ctypes.cdll.LoadLibrary(\"libcuda.so\")
for attr_id, name in [(8, \"MaxSharedMemPerBlock\"), (81, \"MaxSharedMemPerSM\"), (97, \"MaxSharedMemPerBlockOptin\")]:
    val = ctypes.c_int()
    cuda.cuDeviceGetAttribute(ctypes.byref(val), attr_id, 0)
    print(f\"  CUDA attr {attr_id} ({name}): {val.value} bytes ({val.value/1024:.1f} KB)\")
"'
  shared_memory_per_block: 49152
  shared_memory_per_block_optin: 101376
  shared_memory_per_multiprocessor: 102400
  total_memory: 101976440832
  CUDA attr 8 (MaxSharedMemPerBlock): 49152 bytes (48.0 KB)
  CUDA attr 81 (MaxSharedMemPerSM): 102400 bytes (100.0 KB)
  CUDA attr 97 (MaxSharedMemPerBlockOptin): 101376 bytes (99.0 KB)

The Moment of Measurement

At first glance, message 912 appears to be a routine diagnostic command—a Python script that queries CUDA device properties on a remote server. But in the context of the broader optimization effort for the GLM-5-NVFP4 model on eight NVIDIA RTX PRO 6000 Blackwell GPUs, this message represents a critical inflection point. It is the moment when a hypothesis about why certain CUTLASS tile configurations fail is tested against hardware reality, and the results fundamentally reshape the optimization landscape.

The message is deceptively simple: a single bash tool call that runs a Python script on the remote machine at 10.1.230.174. The script does two things: it enumerates all PyTorch CUDA device properties containing "shared," "smem," or "memory" in their names, and it independently queries the CUDA driver API via ctypes to cross-validate three specific shared memory attributes. The output reveals the precise shared memory budget of the Blackwell SM120 architecture: 100 KB per multiprocessor, 99 KB available per block with opt-in, and a default of only 48 KB per block.## The Context: A Mystery in CUTLASS Tile Initialization

To understand why this message matters, we must step back. In the preceding messages, the assistant had been systematically investigating why the GLM-5-NVFP4 model—a Mixture-of-Experts architecture with 256 experts, 8 activated per token—was achieving far below its theoretical peak throughput on the Blackwell GPUs. A critical clue emerged during micro-benchmarking of the CUTLASS FP4 GEMM kernels: two tile configurations, M128×N256×K128 and M256×N128×K128, crashed during initialization with the cryptic error "Error Internal" from initializeMoeGroupedGemm. Meanwhile, the smaller M128×N128×K128 and M128×N128×K256 tiles worked fine.

This was not a minor curiosity. The failing tiles—particularly the wider N256 configuration—are precisely the ones that would allow better work distribution across the GPU's 188 Streaming Multiprocessors (SMs) during the grouped GEMM that implements the MoE computation. The MoE layer processes tokens by routing them through a subset of experts, and each expert's GEMM operates on a small batch of tokens. With wider tiles, more SMs can participate in parallel, reducing idle time and improving throughput. The failure of these tiles was therefore a direct performance bottleneck.

The assistant's hypothesis was that the larger tiles required more shared memory than the SM120 architecture could provide. CUTLASS uses shared memory as a software-managed cache for tile data and pipeline stages. The StageCountAutoCarveout mechanism automatically computes how many pipeline stages fit within the available shared memory budget. If even a single stage requires more shared memory than available, the initialization fails. The assistant needed to confirm the exact shared memory limits to validate this hypothesis—and that is precisely what message 912 delivers.

The Numbers That Matter

The output from the Python probe reveals three critical numbers:

shared_memory_per_block (default): 49,152 bytes (48 KB) — This is the default shared memory allocation per thread block, without any opt-in request. Many standard CUDA kernels operate within this budget.
shared_memory_per_block_optin: 101,376 bytes (99 KB) — This is the maximum shared memory a kernel can request by calling cudaFuncSetAttribute with cudaFuncAttributeMaxDynamicSharedMemorySize. This is the relevant number for CUTLASS kernels that use dynamic shared memory.
shared_memory_per_multiprocessor: 102,400 bytes (100 KB) — This is the total shared memory physically available per SM. The 99 KB opt-in limit is slightly less than the 100 KB physical capacity, likely because some shared memory is reserved for the CUDA runtime's own use (e.g., the constant cache or register file spill areas). These numbers are consistent with NVIDIA's documented shared memory configuration for the Blackwell architecture. Compared to Hopper (SM90), which offers up to 228 KB of shared memory per block with opt-in, Blackwell's 99 KB represents a significant reduction. This is a deliberate design choice: Blackwell trades shared memory capacity for other architectural improvements, such as increased register file size and enhanced tensor core capabilities. For the GLM-5-NVFP4 optimization effort, this reduction has direct and painful consequences.## Connecting the Dots: Why 99 KB Kills the Larger Tiles The assistant's hypothesis now has concrete evidence. The M128×N256×K128 tile configuration requires more than 99 KB of shared memory for its CUTLASS collective pipeline. The StageCountAutoCarveout algorithm, when evaluating this tile, determines that even a single pipeline stage exceeds the available shared memory, and the CUTLASS initialize() call returns kErrorInternal. The can_implement() check passes—meaning the tile is theoretically supported by the hardware—but the actual resource allocation fails at runtime. This is a classic tension in GPU kernel design: larger tiles improve arithmetic intensity and parallelism but consume more shared memory. On Hopper GPUs with 228 KB of shared memory, these tile configurations would likely succeed. On Blackwell's 99 KB budget, they are impossible without fundamental changes to the kernel implementation—such as reducing the pipeline depth, using a different epilogue strategy, or partitioning the computation across multiple kernel launches. The cross-validation using both PyTorch's device properties API and the raw CUDA driver API is noteworthy. The assistant deliberately checks both sources because PyTorch sometimes reports different values than the underlying CUDA runtime, and the ctypes approach guarantees access to the actual hardware attributes. The fact that both APIs agree (99 KB opt-in, 100 KB per SM) gives high confidence in the measurement.

Assumptions and Knowledge Boundaries

This message operates on several implicit assumptions. First, it assumes that the CUTLASS kernel initialization failure is indeed caused by shared memory overflow rather than some other resource constraint (such as register pressure or warp allocation limits). This is a reasonable inference given the error message ("Error Internal" from initialize after can_implement passes), but it is not definitively proven by this diagnostic alone. A more rigorous approach would involve instrumenting the CUTLASS kernel builder to print the computed shared memory requirements for each tile configuration.

Second, the assistant assumes that the shared memory limit is a hard architectural constraint that cannot be circumvented through software configuration. This is correct for the default kernel launch path, but there are potential workarounds: using persistent kernel designs that reuse shared memory across multiple tile iterations, or splitting the grouped GEMM into multiple kernel launches that each operate within the 99 KB budget. The assistant implicitly acknowledges this by listing "fix the failing larger tiles" as a priority item, but the diagnostic in message 912 suggests that the fix may require changes to the CUTLASS kernel template itself rather than simple parameter tuning.

Third, the probe assumes that the shared memory values reported by CUDA are accurate for the actual kernel execution environment. In a virtualized environment (the GPUs are passed through to a Proxmox VM), there is a risk that the hypervisor reports different device properties than the bare metal. However, the assistant has already verified in earlier segments that the GPUs are operating with their full compute capability (SM120), so this concern is minimal.

The input knowledge required to interpret this message is substantial. The reader must understand: the role of shared memory in CUTLASS GEMM kernels, the concept of tile configurations (M×N×K dimensions), the difference between default and opt-in shared memory limits, the MoE grouped GEMM architecture, and the specific challenges of FP4 quantization on Blackwell GPUs. Without this background, the numbers are meaningless—they are just three lines of output.## The Output Knowledge Created

Message 912 produces three distinct pieces of knowledge that directly inform subsequent optimization decisions:

Knowledge 1: The shared memory ceiling is 99 KB. This is the single most important number. It explains why the M128×N256 and M256×N128 tiles fail, and it sets a hard upper bound on what CUTLASS tile configurations can be used on Blackwell for this workload. Any optimization strategy that relies on larger tiles must either find a way to reduce per-tile shared memory consumption or abandon the approach entirely.

Knowledge 2: The default shared memory per block (48 KB) is insufficient for any non-trivial CUTLASS tile. This means that every CUTLASS kernel used in the MoE grouped GEMM must explicitly request the opt-in shared memory allocation. If the FlashInfer or TensorRT-LLM kernel launcher code does not call cudaFuncSetAttribute to request the full 99 KB, the kernel will silently operate with only 48 KB, potentially causing suboptimal tile selection or runtime failures. This is a potential bug surface that the assistant could investigate.

Knowledge 3: The total memory per GPU is approximately 97.3 GB (101,976,440,832 bytes / 1024³ ≈ 97.3 GB, or about 101,976,440,832 / 1e9 ≈ 102 GB in decimal). While not directly relevant to the shared memory investigation, this confirms that the GPUs are the 96 GB variant of the RTX PRO 6000 Blackwell, consistent with the earlier device identification.

The Thinking Process

The reasoning visible in this message reveals a methodical diagnostic approach. The assistant has already traced the error to the initializeMoeGroupedGemm function in the CUTLASS launcher template, examined the can_implement vs. initialize distinction, and identified shared memory as the likely culprit. The Python probe is the natural next step: measure the actual hardware limit to confirm the hypothesis.

The choice to cross-validate with both PyTorch and raw CUDA APIs demonstrates careful methodology. PyTorch's device properties API is convenient but may not expose all attributes correctly (as evidenced by the initial AttributeError in message 910, where max_shared_memory_per_block was renamed to shared_memory_per_block in the PyTorch version being used). The ctypes approach bypasses PyTorch's abstraction layer entirely, querying the CUDA driver directly. This dual-path validation is a best practice when dealing with low-level hardware properties.

The assistant also shows awareness of the broader optimization context. The shared memory probe is not an isolated diagnostic—it is one piece of a larger puzzle that includes GPU power draw measurements (235W out of 600W TDP), micro-benchmarking of CUTLASS kernel efficiency (0.8–55 TFLOPS during actual decode vs. 1,300 TFLOPS peak for large matrices), and server parameter tuning. The shared memory limit is the "why" behind the failing tiles, and the failing tiles are one of several bottlenecks limiting throughput to approximately 1,640 tok/s at 2048 concurrency.

Implications for the Optimization Roadmap

The discovery that Blackwell's shared memory is capped at 99 KB has profound implications for the GLM-5-NVFP4 deployment. It means that the CUTLASS kernel templates written for Hopper (with 228 KB shared memory) cannot be directly ported to Blackwell without modification. The KernelScheduleAuto mechanism that CUTLASS uses for SM120 may be selecting suboptimal schedules precisely because it is constrained by the tighter shared memory budget.

For the optimization roadmap, this knowledge redirects effort away from "fixing the tiles" (which may be impossible without upstream changes to CUTLASS or FlashInfer) and toward alternative strategies: increasing the effective batch size per expert through higher num_continuous_decode_steps, exploring the cuBLASLt FP4 path which may have different shared memory characteristics, or investigating persistent kernel designs that amortize shared memory overhead across multiple tile iterations.

In the broader narrative of the coding session, message 912 is the moment when a hardware constraint is empirically confirmed, transforming a mysterious crash into a well-understood architectural limitation. It exemplifies the kind of low-level systems investigation that distinguishes effective ML infrastructure work: tracing an error from its symptom (a crash during initialization) through intermediate hypotheses (shared memory overflow) to a definitive measurement (99 KB opt-in limit), and then using that knowledge to guide the next iteration of optimization decisions.