j@dZddlZddlZddlZddlZddlZddlZddlZddlmZddl m Z ddl m Z ddlmZdZd ZeeeeedejejfZdadagZgZdaeZ gZ!gZ"iZ#eZ$eZ%gZ&d a'd a(ej)Z*d Z+ej,d;d Z-dZ.dZ/dZ0ddd ddZ1dZ2dd!d"DZ?Gd#d$e Z@d%ZAd&ZBd'ZCGd(d)eDZEd*ZFd+ZGeHfd,ZId-ZJd.ZKd/ZLd0ZMd1ZNd2ZOGd3d4e ZPGd5d6ZQd7ZRd8ZSGd9d:ZTdS)>a Automatic CUDA graphs of the blocks of the towers (the repeated layers of a transformer): the host work of a block call (Python, kernel launches) is recorded once and replayed, for models whose steps are short enough for the host to be the bottleneck. Inside `scope()` (the denoising loop of a model; elsewhere blocks run as usual), each call of a block is keyed by the tensor objects of its weights, its active LoRAs and their scalings, the shapes of its tensor arguments and extra inputs, its other arguments and the registered global keys: the first call of a key runs eagerly (it compiles and tunes the kernels, makes the caches of the weights), then is captured into a CUDA graph that the next calls replay. Blocks run eagerly, without captures, while the placement of their weights is not settled (`ready()` False: e.g. an automatic preload measuring a new workload, which keeps other blocks in VRAM at the next step). A block thus has a graph per location of its weights (kept in VRAM, or each slot of the swapped blocks), and a model loaded again captures its graphs again. The tensor arguments, extra inputs (tensors that come and go with a block, e.g. its LoRA copies) and outputs go through static buffers shared by the graphs, the arguments a block modifies in place are copied back when needed, and an output that is an argument is that argument. A block keeps its `MAX_KEYS` most recently used graphs (a new resolution replaces them); all the graphs go with the model or `invalidate()`. With MMGP's VMM allocator, eager preparation and capture share a pool. Suitable live arguments and first-call outputs supply the static storage; retained caller outputs are protected before another call can overwrite them. A large argument that only its caller holds (a tower's hidden state, new at each step) moves into its static buffer instead of being copied there: the caller's tensor takes the buffer's storage and frees its own, so the graphs keep no second copy of it and the next blocks copy nothing; while the caller holds it, other calls use another buffer. Eager stages beside the blocks (input preparation, output head) can run on the pool's idle memory with `pooled()`. Graph warmup is scoped to its manager's pool. A consumed single-tensor list can transfer its input's execution lifetime: after its last execution view dies, later temporaries may use its physical pages while every graph address stays fixed. Caller references, storage aliases and recorded stream consumers prevent this transfer. Large outputs can reuse their warmup storage at the corresponding captured allocation; a changed allocation trace is recaptured without that output plan. These plans add no replay-time mapping or synchronization, and do not participate in ordinary, non-graph allocation. Contract of a block call in a scope: the block reads only its weights, its tensor arguments (in lists, tuples, dicts and dataclasses), the state registered for its model (register_state) and tensors it creates; its effects are its outputs and in-place updates of its tensor arguments or of that state (Python side effects are not replayed, and the capture runs the Python of the block once more); no synchronization or CPU round trip. The capture checks the reads and raises an error naming any other tensor. A model's own loop over a tower (run_blocks) replays each run of consecutive blocks kept in VRAM as one graph, while the blocks streamed in between go through the hooks that transfer them. N)tqdm)_pytree)TorchDispatchMode)runtimeFcxda ddlm}n#t$rYdSwxYw|jjfd}||j_dS)aTriton's autotuners (@triton.autotune) benchmark their configurations on a cache miss: inside a capture, the benchmark launches would be recorded in the graph and CUDA refuses its timing synchronizations. The eager call that precedes each capture tunes the keys a block uses (sizes and dtypes); a miss that still happens in a capture (a key that differs there) raises an error naming the kernel.Tr) autotunerNctr8t|d|j}tdt|d|d|g|Ri|S)Nbase_fnzthe Triton kernel __name__zz would tune itself (@triton.autotune) during a CUDA graph capture, for a key the eager call before the capture did not use) _capturinggetattrfn RuntimeError)selfargskwargsrbenchs 8C:\Users\johnn\OneDrive\Documents\AI\mmgp\cuda_graphs.py guarded_benchz-_guard_triton_autotune..guarded_benchMsz  [y$'22B ZGB B4O4O Z Z Z[[ [uT+D+++F+++)_autotune_guardedtriton.runtimer ImportError Autotuner_bench)r rrs @r_guard_triton_autotunerAs~ ,,,,,,,    &E,,,,,"/Is  Tc#K|rtdkr~|r)tddtDDdndatJttvrzscope..]s'TTFVTTTTTTTrc3*K|]}|VdSr"r#)r$gates rr&zscope..]s(.I.I$ttvv.I.I.I.I.I.IrNz[MMGP] CUDA graphs are off: .) _depthnext_gates_blocked _reportedaddrwritebool)enabledgateds rscoper4Ws C6Q;;_dn4TT.I.I&.I.I.ITTTVZ[[[jn  HI$=$= MM( # # # JAhAAA B B B d7mmF  $w--$w--s &CCctdkS)aTrue while a block of a graphed model runs to be captured (its first, eager call and its capture): a kernel launched outside PyTorch's current stream (e.g. on the legacy default stream) would be missing from its graph, so code with such kernels uses others.r) _recordingr#rr recordingr7hs >rcptj|}|dS|\}}|j|Sr")r_owner cuda_graphsget)moduleownermanagermodel_ids r_graphsr@ns< N6 " "E }tGX   " "8 , ,,rctrt|nd}|duo%tduo|jduo|S)a Whether a managed module can use graphs in the current scope. Includes mode selection, registered gates and weight-residency readiness. Use only to skip preparation specific to a graph stage; run() also enforces these conditions. Unmanaged/released modules return False. N)r*r@r-disabledready)r<graphss r is_enabledrEvsK!' 0WV___DF   a(d"2 av$7N aSYS_S_SaSaarforward)r=methodr2ct||}|rtrt||n|nd}|||i|n|||||S)aExecute an extra stage using its owner's existing graph cache and lifetime. Call inside scope(). An unregistered stage may supply a managed owner whose weights it uses; its parameters need not be registered a second time. The selected method must satisfy the same graph contract as a tower block. Outside a scope, with enabled=False, or without an active graph manager, execute the method normally. Gates/readiness and cache invalidation/release remain MMGP-owned. PyTorch compilation does not enable these CUDA graphs. A stage may span several blocks of its owner (e.g. a whole decoding step) only while runtime.weights_resident(owner): graphed blocks it calls are then captured inline. N)rr*r@call)r<r=rGr2rrrFrDs rrunrJsjff%%G