Detects the LSE format if we're in a varlen case. Returns `None` if the format is not relevant (eg not varlen) Raises an exception if the

Computes the best operator for forward Raises: NotImplementedError: if not operator was found Returns: AttentionOp: The best

Given sequence lengths, returns the min/max value and the sequence start positions (offsets), with first element being 0 (returned in

Kernel _fwd_kernel_splitK needs to be post-processed by unroll_varargs to specialize it for a given number of quantization groups N_GROUPS

Warning: grad/ctx.out is potentially in BMK format

Converts dense 2d matrix to a csr sparse matrix.

Given a list of operator names, their corresponding runtimes, and the maximum amount of memory available, find the subset of operators that c

More flexible unused version of _slow_rope - allows varying dtypes.

Applies fn to the attention output, as if we were a tensor. fn must expect tensors of shape BMGHK or BMHK, but cannot actually

Construct a tree based on given branching factor for each non-root level.

Create CLI argument parser.

Decodes a list of token IDs into a string. Args: t (list[int]): The list of token IDs to be decoded. Returns:

Benchmark the runtime of the provided function. Args: fn: Function to benchmark rep: Repetition time (in ms) grad_to_

Encodes a string into a list of token IDs. Args: s (str): The input string to be encoded. bos (bool): Whethe

Terminate the current phase.

Marks the transition to either comms or compute in the FW pass

* kernels expect 4D bias/bias.grad with shape * (batch_sz, n_heads, n_queries, n_keys). common bias shapes users may pass * are: * - (n_queries, n_

Given a function, its arguments, and the maximum amount of memory available, find the subset of operators that can be optimized to reduce run

rsqrt = libdevice_find("rsqrt") is a triton-version-friendly way to say from triton.language.extra.libdevice import rsqrt

Makes each block causal

Experimental: Makes each block causal with local attention

Returns a list of reasons why this is not supported. The kernel can run these inputs only if the returned list is empty

Create block tables and pages K/V cache for testing paged attention. Args: cache_k, cache_v: K/V caches, each of shape [B, MAX_T, H_k