lightseekorg--tokenspeed/tokenspeed-kernel/python/tokenspeed_kernel/ops/sampling/triton/probability.py

# Copyright (c) 2026 LightSeek Foundation
#
# Permission is hereby granted, free of charge, to any person obtaining a copy
# of this software and associated documentation files (the "Software"), to deal
# in the Software without restriction, including without limitation the rights
# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
# copies of the Software, and to permit persons to whom the Software is
# furnished to do so, subject to the following conditions:
#
# The above copyright notice and this permission notice shall be included in
# all copies or substantial portions of the Software.
#
# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
# SOFTWARE.

# Probability-route compatibility helper for the existing FlashInfer full backend.

from __future__ import annotations

import torch
from tokenspeed_kernel._triton import tl, triton


@triton.jit
def _min_p_renorm_prob_kernel(
    probs_ptr,
    min_p_ptr,
    out_ptr,
    vocab_size: tl.constexpr,
    probs_row_stride: tl.constexpr,
    out_row_stride: tl.constexpr,
    BLOCK_SIZE: tl.constexpr,
    ENABLE_PDL: tl.constexpr,
):
    if ENABLE_PDL:
        tl.extra.cuda.gdc_wait()

    row = tl.program_id(0)
    offs = tl.arange(0, BLOCK_SIZE)
    probs_row = probs_ptr + row * probs_row_stride
    out_row = out_ptr + row * out_row_stride

    max_prob = tl.full((), 0.0, tl.float32)
    for start in tl.range(0, vocab_size, BLOCK_SIZE, num_stages=3):
        cols = start + offs
        mask = cols < vocab_size
        vals = tl.load(probs_row + cols, mask=mask, other=0.0).to(tl.float32)
        max_prob = tl.maximum(max_prob, tl.max(tl.where(mask, vals, 0.0), axis=0))

    threshold = max_prob * tl.load(min_p_ptr + row).to(tl.float32)
    denom = tl.full((), 0.0, tl.float32)
    for start in tl.range(0, vocab_size, BLOCK_SIZE, num_stages=3):
        cols = start + offs
        mask = cols < vocab_size
        vals = tl.load(probs_row + cols, mask=mask, other=0.0).to(tl.float32)
        keep = mask & (vals >= threshold)
        denom += tl.sum(tl.where(keep, vals, 0.0), axis=0)

    inv_denom = 1.0 / tl.maximum(denom, 1.0e-20)
    for start in tl.range(0, vocab_size, BLOCK_SIZE, num_stages=3):
        cols = start + offs
        mask = cols < vocab_size
        vals = tl.load(probs_row + cols, mask=mask, other=0.0).to(tl.float32)
        keep = mask & (vals >= threshold)
        out = tl.where(keep, vals * inv_denom, 0.0)
        tl.store(out_row + cols, out, mask=mask)

    if ENABLE_PDL:
        tl.extra.cuda.gdc_launch_dependents()


def min_p_renorm_prob(
    probs: torch.Tensor,
    min_p: torch.Tensor,
    *,
    enable_pdl: bool = False,
) -> torch.Tensor:
    """Renormalize probabilities after applying a per-row min-p cutoff.

    For each row, this computes ``threshold = min_p[row] * max(probs[row])``,
    zeros probabilities below the threshold, and renormalizes the surviving
    probabilities so the row sums to one.
    """
    if probs.ndim != 2:
        raise ValueError(f"min_p_renorm_prob expects 2D probs, got {probs.ndim}D")
    if min_p.ndim != 1:
        raise ValueError(f"min_p_renorm_prob expects 1D min_p, got {min_p.ndim}D")
    if min_p.shape[0] != probs.shape[0]:
        raise ValueError(
            "min_p length must match probs rows, "
            f"got {min_p.shape[0]} and {probs.shape[0]}"
        )
    if probs.device.type != "cuda" or min_p.device.type != "cuda":
        raise ValueError("min_p_renorm_prob requires CUDA tensors")
    if probs.stride(-1) != 1:
        raise ValueError(
            f"min_p_renorm_prob requires stride-1 vocab dimension, got stride={probs.stride()}"
        )
    if not min_p.is_contiguous():
        min_p = min_p.contiguous()

    out = torch.empty_like(probs)
    rows, vocab_size = probs.shape
    if rows == 0:
        return out

    block_size = min(4096, triton.next_power_of_2(vocab_size))
    num_warps = 4 if block_size <= 1024 else 8
    extra_kwargs = {"launch_pdl": True} if enable_pdl else {}
    _min_p_renorm_prob_kernel[(rows,)](
        probs,
        min_p,
        out,
        vocab_size=vocab_size,
        probs_row_stride=probs.stride(0),
        out_row_stride=out.stride(0),
        BLOCK_SIZE=block_size,
        ENABLE_PDL=enable_pdl,
        num_warps=num_warps,
        num_stages=3,
        **extra_kwargs,
    )
    return out