Characterizing Warp Divergence from Pascal to Blackwell

Hugging Face Daily Papers Papers
gpu warp-divergence nvidia architecture microbenchmarks

Summary

This paper characterizes warp divergence behavior across NVIDIA GPU generations from Pascal to Blackwell, finding stable linear serialization costs despite evolving reconvergence mechanisms.

Since Volta introduced Independent Thread Scheduling (ITS), NVIDIA GPUs have been widely assumed to handle warp divergence in a fixed manner. We test this assumption across Ampere, Hopper, and datacenter and consumer Blackwell GPUs, using pre-ITS Pascal as a baseline. Combining cycle-accurate microbenchmarks, hardware counters, and static analysis of compiler-generated SASS, we separate stable behavior from architectural change. Across all tested generations, divergent paths serialize linearly with the number of paths k, following T(k) approx sk with no super-linear reconvergence penalty. Warp execution efficiency falls as 32/k, the penalty is independent of occupancy, and predication removes the serialization cost. The same behavior appears on Pascal, showing that this programmer-visible cost model predates ITS. The compiler-emitted reconvergence machinery, however, has changed substantially. Pascal uses a per-warp SSY/SYNC instruction stack, whereas later generations use barrier-register instructions. Deferred reconvergence beyond the immediate post-dominator falls from 29 cases on Ampere to 2 on Blackwell. Blackwell also introduces a two-tier convergence-barrier classification, uniform-branch instructions, and explicit partial-mask warp synchronization, none of which appear on Ampere or Hopper. Controlled bit-flip experiments indicate that the new barrier class is a static compiler classification with no observable runtime effect in our tests. Thus, divergence retains a stable and predictable performance cost even as NVIDIA's control-flow ISA and reconvergence mechanisms continue to evolve.
Original Article
View Cached Full Text

Cached at: 07/28/26, 10:24 AM

Paper page - Characterizing Warp Divergence from Pascal to Blackwell

Source: https://huggingface.co/papers/2607.23402 Published on Jul 26

·

Submitted byhttps://huggingface.co/alpindale

Alpinon Jul 28

Abstract

SinceVoltaintroducedIndependentThreadScheduling(ITS),NVIDIAGPUshavebeenwidelyassumedtohandlewarpdivergenceinafixedmanner.WetestthisassumptionacrossAmpere,Hopper,anddatacenterandconsumerBlackwellGPUs,usingpre-ITSPascalasabaseline.Combiningcycle-accuratemicrobenchmarks,hardwarecounters,andstaticanalysisofcompiler-generatedSASS,weseparatestablebehaviorfromarchitecturalchange.Acrossalltestedgenerations,divergentpathsserializelinearlywiththenumberofpathsk,followingT(k)approxskwithnosuper-linearreconvergencepenalty.Warpexecutionefficiencyfallsas32/k,thepenaltyisindependentofoccupancy,andpredicationremovestheserializationcost.ThesamebehaviorappearsonPascal,showingthatthisprogrammer-visiblecostmodelpredatesITS.Thecompiler-emittedreconvergencemachinery,however,haschangedsubstantially.Pascalusesaper-warpSSY/SYNCinstructionstack,whereaslatergenerationsusebarrier-registerinstructions.Deferredreconvergencebeyondtheimmediatepost-dominatorfallsfrom29casesonAmpereto2onBlackwell.Blackwellalsointroducesatwo-tierconvergence-barrierclassification,uniform-branchinstructions,andexplicitpartial-maskwarpsynchronization,noneofwhichappearonAmpereorHopper.Controlledbit-flipexperimentsindicatethatthenewbarrierclassisastaticcompilerclassificationwithnoobservableruntimeeffectinourtests.Thus,divergenceretainsastableandpredictableperformancecostevenasNVIDIA’scontrol-flowISAandreconvergencemechanismscontinuetoevolve.

View arXiv pageView PDFAdd to collection

Get this paper in your agent:

hf papers read 2607\.23402

Don’t have the latest CLI?curl \-LsSf https://hf\.co/cli/install\.sh \| bash

Models citing this paper0

No model linking this paper

Cite arxiv.org/abs/2607.23402 in a model README.md to link it from this page.

Datasets citing this paper0

No dataset linking this paper

Cite arxiv.org/abs/2607.23402 in a dataset README.md to link it from this page.

Spaces citing this paper0

No Space linking this paper

Cite arxiv.org/abs/2607.23402 in a Space README.md to link it from this page.

Collections including this paper0

No Collection including this paper

Add this paper to acollectionto link it from this page.

Similar Articles

Blackwell and PDL performance increase

Reddit r/LocalLLaMA

Llama.cpp now supports Nvidia's Programmatic Dependent Launch (PDL) for Blackwell GPUs, offering a 5-10% performance boost on token generation. The feature is not enabled by default and requires a build flag.

@ZhihuFrontier: GPU programming changed because Tensor Cores became too fast to feed Zhihu contributor THU-PACMAN实验室 shared a sharp bre…

X AI KOLs Timeline

A detailed analysis of how NVIDIA GPU programming evolved from Volta to Blackwell, highlighting the shift from synchronous thread models to asynchronous dataflow and the challenges of feeding Tensor Cores. The article discusses new hardware features like TMA, TMEM, and tcgen05 MMA, and shows how modern kernels like FlashAttention-3 and FlashMLA exploit these changes for higher utilization.

@charles_irl: dflash go brr

X AI KOLs Timeline

NVIDIA announces DFlash, an open source block diffusion model for speculative decoding that achieves up to 15x higher inference throughput on Blackwell GPUs while maintaining interactivity.