Upload folder using huggingface_hub

build/torch-universal/triton_kernels/__init__.py

@@ -1,3 +0,0 @@
-from . import matmul_ogs
-
-__all__ = ["matmul_ogs"]

build/torch-universal/triton_kernels/_ops.py

@@ -1,8 +1,8 @@
 import torch
-ops = torch.ops._triton_kernels_8830f14_dirty
+ops = torch.ops._triton_kernels_10e8091_dirty
 
 def add_op_namespace_prefix(op_name: str):
     """
     Prefix op by namespace.
     """
-    return f"_triton_kernels_8830f14_dirty::{op_name}"
+    return f"_triton_kernels_10e8091_dirty::{op_name}"

build/torch-universal/triton_kernels/matmul_ogs.py

@@ -7,10 +7,10 @@ import torch
 import triton
 from enum import Enum, auto
 # utilities
-from triton_kernels import target_info
-from triton_kernels.numerics import InFlexData, OutFlexData
-from triton_kernels.routing import GatherIndx, RoutingData, ScatterIndx
-from triton_kernels.target_info import is_cuda
+from . import target_info
+from .numerics import InFlexData, OutFlexData
+from .routing import GatherIndx, RoutingData, ScatterIndx
+from .target_info import is_cuda
 # details
 from .matmul_ogs_details._matmul_ogs import _compute_writeback_idx
 from .matmul_ogs_details._matmul_ogs import _matmul_ogs

build/torch-universal/triton_kernels/matmul_ogs_details/_finalize_matmul.py

@@ -1,9 +1,9 @@
 import triton
 import triton.language as tl
-from triton_kernels.numerics_details.flexpoint import float_to_flex, load_scale, update_scale
-from triton_kernels.numerics_details.mxfp_details._downcast_to_mxfp import MXFP_BLOCK_SIZE
-from triton_kernels.target_info import cuda_capability_geq as _cuda_capability_geq
-from triton_kernels.target_info import is_hip as _is_hip
+from ..numerics_details.flexpoint import float_to_flex, load_scale, update_scale
+from ..numerics_details.mxfp_details._downcast_to_mxfp import MXFP_BLOCK_SIZE
+from ..target_info import cuda_capability_geq as _cuda_capability_geq
+from ..target_info import is_hip as _is_hip
 
 
 # fmt: off

build/torch-universal/triton_kernels/matmul_ogs_details/_matmul_ogs.py

@@ -2,11 +2,11 @@
 # fmt: off
 import triton
 import triton.language as tl
-from triton_kernels.tensor_details.layout_details.blackwell_scale import unswizzle_mx_scale_bw
-from triton_kernels.tensor_details.layout_details.hopper_scale import unswizzle_mxfp4_scale_hopper
-from triton_kernels.tensor_details.layout_details.hopper_value import mxfp4_to_bf16_triton
-from triton_kernels.numerics_details.flexpoint import float_to_flex, load_scale
-from triton_kernels.numerics_details.mxfp_details._downcast_to_mxfp import MXFP_BLOCK_SIZE
+from ..tensor_details.layout_details.blackwell_scale import unswizzle_mx_scale_bw
+from ..tensor_details.layout_details.hopper_scale import unswizzle_mxfp4_scale_hopper
+from ..tensor_details.layout_details.hopper_value import mxfp4_to_bf16_triton
+from ..numerics_details.flexpoint import float_to_flex, load_scale
+from ..numerics_details.mxfp_details._downcast_to_mxfp import MXFP_BLOCK_SIZE
 from ._common import make_matmul_repr, matmul_launch_metadata, swizzle2d, xcd_swizzle, get_scaled_dot_format_string
 
 

build/torch-universal/triton_kernels/matmul_ogs_details/_p_matmul_ogs.py

@@ -3,15 +3,15 @@
 import torch
 import triton
 import triton.language as tl
-from triton_kernels import target_info
-from triton_kernels.tensor_details.layout_details.blackwell_scale import unswizzle_mx_scale_bw
-from triton_kernels.numerics_details.flexpoint import (
+from . import target_info
+from ..tensor_details.layout_details.blackwell_scale import unswizzle_mx_scale_bw
+from ..numerics_details.flexpoint import (
     float_to_flex,
     load_scale,
     nan_propagating_absmax_reduce,
     compute_scale,
 )
-from triton_kernels.numerics_details.mxfp_details._downcast_to_mxfp import MXFP_BLOCK_SIZE
+from ..numerics_details.mxfp_details._downcast_to_mxfp import MXFP_BLOCK_SIZE
 from ._common import make_matmul_repr, matmul_launch_metadata, swizzle2d, xcd_swizzle, get_scaled_dot_format_string
 
 

build/torch-universal/triton_kernels/matmul_ogs_details/opt_flags.py

@@ -2,7 +2,7 @@
 # fmt: off
 from dataclasses import dataclass
 import triton
-from triton_kernels.target_info import get_cdna_version
+from ..target_info import get_cdna_version
 import torch
 from .opt_flags_details import opt_flags_amd, opt_flags_nvidia
 

build/torch-universal/triton_kernels/matmul_ogs_details/opt_flags_details/opt_flags_amd.py

@@ -1,7 +1,7 @@
 import torch
 import triton
-from triton_kernels.target_info import get_cdna_version
-from triton_kernels.tensor import bitwidth
+from ...target_info import get_cdna_version
+from ...tensor import bitwidth
 
 
 def compute_block_nk(n, block_m, grid_m, num_xcds, lhs_dtype, rhs_dtype, precision_config):

build/torch-universal/triton_kernels/matmul_ogs_details/opt_flags_details/opt_flags_nvidia.py

@@ -1,9 +1,9 @@
 import torch
 import triton
-from triton_kernels import target_info
-from triton_kernels.tensor import get_layout, bitwidth, FP4
-from triton_kernels.tensor_details.layout import HopperMXScaleLayout
-from triton_kernels.numerics_details.mxfp_details._downcast_to_mxfp import MXFP_BLOCK_SIZE
+from ... import target_info
+from ...tensor import get_layout, bitwidth, FP4
+from ...tensor_details.layout import HopperMXScaleLayout
+from ...numerics_details.mxfp_details._downcast_to_mxfp import MXFP_BLOCK_SIZE
 
 
 def compute_grid_size(routing_data, m, n, block_m, block_n):

build/torch-universal/triton_kernels/numerics_details/flexpoint.py

@@ -1,5 +1,5 @@
 from ..numerics import MAX_FINITE_FLOAT8E4B8, MAX_FINITE_FLOAT8E4NV, MAX_FINITE_FLOAT8E5
-from triton_kernels import target_info
+from .. import target_info
 import triton
 import triton.language as tl
 

build/torch-universal/triton_kernels/swiglu.py

@@ -1,9 +1,9 @@
 from dataclasses import dataclass
-from triton_kernels.numerics import InFlexData, OutFlexData
+from .numerics import InFlexData, OutFlexData
 import torch
 import triton
 from .swiglu_details._swiglu import _swiglu, _swiglu_fn
-from triton_kernels import target_info
+from . import target_info
 
 
 @dataclass(frozen=True)

build/torch-universal/triton_kernels/swiglu_details/_swiglu.py

@@ -1,4 +1,4 @@
-from triton_kernels.numerics_details.flexpoint import load_scale, float_to_flex, update_scale
+from ..numerics_details.flexpoint import load_scale, float_to_flex, update_scale
 import triton
 import triton.language as tl
 

build/torch-universal/triton_kernels/testing.py

@@ -4,7 +4,7 @@ import os
 import subprocess
 import sys
 import torch
-from triton_kernels.numerics import MAX_FINITE_FLOAT8E4B8, MAX_FINITE_FLOAT8E4NV, MAX_FINITE_FLOAT8E5
+from .numerics import MAX_FINITE_FLOAT8E4B8, MAX_FINITE_FLOAT8E4NV, MAX_FINITE_FLOAT8E5
 
 
 def assert_equal(ref, tri):

build/torch-universal/triton_kernels/topk.py

@@ -1,8 +1,8 @@
 import torch
 import triton
-from triton_kernels.topk_details._topk_forward import _topk_forward
-from triton_kernels.topk_details._topk_backward import _topk_backward
-from triton_kernels.tensor import Tensor, Bitmatrix
+from .topk_details._topk_forward import _topk_forward
+from .topk_details._topk_backward import _topk_backward
+from .tensor import Tensor, Bitmatrix
 
 
 def topk_forward(x, k, apply_softmax=True, dim=1, return_bitmatrix=True, y_indx=None, n_rows=None):

torch-ext/triton_kernels/matmul_ogs.py

@@ -7,10 +7,10 @@ import torch
 import triton
 from enum import Enum, auto
 # utilities
-from triton_kernels import target_info
-from triton_kernels.numerics import InFlexData, OutFlexData
-from triton_kernels.routing import GatherIndx, RoutingData, ScatterIndx
-from triton_kernels.target_info import is_cuda
+from . import target_info
+from .numerics import InFlexData, OutFlexData
+from .routing import GatherIndx, RoutingData, ScatterIndx
+from .target_info import is_cuda
 # details
 from .matmul_ogs_details._matmul_ogs import _compute_writeback_idx
 from .matmul_ogs_details._matmul_ogs import _matmul_ogs

torch-ext/triton_kernels/matmul_ogs_details/_finalize_matmul.py

@@ -1,9 +1,9 @@
 import triton
 import triton.language as tl
-from triton_kernels.numerics_details.flexpoint import float_to_flex, load_scale, update_scale
-from triton_kernels.numerics_details.mxfp_details._downcast_to_mxfp import MXFP_BLOCK_SIZE
-from triton_kernels.target_info import cuda_capability_geq as _cuda_capability_geq
-from triton_kernels.target_info import is_hip as _is_hip
+from ..numerics_details.flexpoint import float_to_flex, load_scale, update_scale
+from ..numerics_details.mxfp_details._downcast_to_mxfp import MXFP_BLOCK_SIZE
+from ..target_info import cuda_capability_geq as _cuda_capability_geq
+from ..target_info import is_hip as _is_hip
 
 
 # fmt: off

torch-ext/triton_kernels/matmul_ogs_details/_matmul_ogs.py

@@ -2,11 +2,11 @@
 # fmt: off
 import triton
 import triton.language as tl
-from triton_kernels.tensor_details.layout_details.blackwell_scale import unswizzle_mx_scale_bw
-from triton_kernels.tensor_details.layout_details.hopper_scale import unswizzle_mxfp4_scale_hopper
-from triton_kernels.tensor_details.layout_details.hopper_value import mxfp4_to_bf16_triton
-from triton_kernels.numerics_details.flexpoint import float_to_flex, load_scale
-from triton_kernels.numerics_details.mxfp_details._downcast_to_mxfp import MXFP_BLOCK_SIZE
+from ..tensor_details.layout_details.blackwell_scale import unswizzle_mx_scale_bw
+from ..tensor_details.layout_details.hopper_scale import unswizzle_mxfp4_scale_hopper
+from ..tensor_details.layout_details.hopper_value import mxfp4_to_bf16_triton
+from ..numerics_details.flexpoint import float_to_flex, load_scale
+from ..numerics_details.mxfp_details._downcast_to_mxfp import MXFP_BLOCK_SIZE
 from ._common import make_matmul_repr, matmul_launch_metadata, swizzle2d, xcd_swizzle, get_scaled_dot_format_string
 
 

torch-ext/triton_kernels/matmul_ogs_details/_p_matmul_ogs.py

@@ -3,15 +3,15 @@
 import torch
 import triton
 import triton.language as tl
-from triton_kernels import target_info
-from triton_kernels.tensor_details.layout_details.blackwell_scale import unswizzle_mx_scale_bw
-from triton_kernels.numerics_details.flexpoint import (
+from . import target_info
+from ..tensor_details.layout_details.blackwell_scale import unswizzle_mx_scale_bw
+from ..numerics_details.flexpoint import (
     float_to_flex,
     load_scale,
     nan_propagating_absmax_reduce,
     compute_scale,
 )
-from triton_kernels.numerics_details.mxfp_details._downcast_to_mxfp import MXFP_BLOCK_SIZE
+from ..numerics_details.mxfp_details._downcast_to_mxfp import MXFP_BLOCK_SIZE
 from ._common import make_matmul_repr, matmul_launch_metadata, swizzle2d, xcd_swizzle, get_scaled_dot_format_string
 
 

torch-ext/triton_kernels/matmul_ogs_details/opt_flags.py

@@ -2,7 +2,7 @@
 # fmt: off
 from dataclasses import dataclass
 import triton
-from triton_kernels.target_info import get_cdna_version
+from ..target_info import get_cdna_version
 import torch
 from .opt_flags_details import opt_flags_amd, opt_flags_nvidia
 

torch-ext/triton_kernels/matmul_ogs_details/opt_flags_details/opt_flags_amd.py

@@ -1,7 +1,7 @@
 import torch
 import triton
-from triton_kernels.target_info import get_cdna_version
-from triton_kernels.tensor import bitwidth
+from ...target_info import get_cdna_version
+from ...tensor import bitwidth
 
 
 def compute_block_nk(n, block_m, grid_m, num_xcds, lhs_dtype, rhs_dtype, precision_config):

torch-ext/triton_kernels/matmul_ogs_details/opt_flags_details/opt_flags_nvidia.py

@@ -1,9 +1,9 @@
 import torch
 import triton
-from triton_kernels import target_info
-from triton_kernels.tensor import get_layout, bitwidth, FP4
-from triton_kernels.tensor_details.layout import HopperMXScaleLayout
-from triton_kernels.numerics_details.mxfp_details._downcast_to_mxfp import MXFP_BLOCK_SIZE
+from ... import target_info
+from ...tensor import get_layout, bitwidth, FP4
+from ...tensor_details.layout import HopperMXScaleLayout
+from ...numerics_details.mxfp_details._downcast_to_mxfp import MXFP_BLOCK_SIZE
 
 
 def compute_grid_size(routing_data, m, n, block_m, block_n):

torch-ext/triton_kernels/numerics_details/flexpoint.py

@@ -1,5 +1,5 @@
 from ..numerics import MAX_FINITE_FLOAT8E4B8, MAX_FINITE_FLOAT8E4NV, MAX_FINITE_FLOAT8E5
-from triton_kernels import target_info
+from .. import target_info
 import triton
 import triton.language as tl
 

torch-ext/triton_kernels/swiglu.py

@@ -1,9 +1,9 @@
 from dataclasses import dataclass
-from triton_kernels.numerics import InFlexData, OutFlexData
+from .numerics import InFlexData, OutFlexData
 import torch
 import triton
 from .swiglu_details._swiglu import _swiglu, _swiglu_fn
-from triton_kernels import target_info
+from . import target_info
 
 
 @dataclass(frozen=True)

torch-ext/triton_kernels/swiglu_details/_swiglu.py

@@ -1,4 +1,4 @@
-from triton_kernels.numerics_details.flexpoint import load_scale, float_to_flex, update_scale
+from ..numerics_details.flexpoint import load_scale, float_to_flex, update_scale
 import triton
 import triton.language as tl
 

torch-ext/triton_kernels/testing.py

@@ -4,7 +4,7 @@ import os
 import subprocess
 import sys
 import torch
-from triton_kernels.numerics import MAX_FINITE_FLOAT8E4B8, MAX_FINITE_FLOAT8E4NV, MAX_FINITE_FLOAT8E5
+from .numerics import MAX_FINITE_FLOAT8E4B8, MAX_FINITE_FLOAT8E4NV, MAX_FINITE_FLOAT8E5
 
 
 def assert_equal(ref, tri):

torch-ext/triton_kernels/topk.py

@@ -1,8 +1,8 @@
 import torch
 import triton
-from triton_kernels.topk_details._topk_forward import _topk_forward
-from triton_kernels.topk_details._topk_backward import _topk_backward
-from triton_kernels.tensor import Tensor, Bitmatrix
+from .topk_details._topk_forward import _topk_forward
+from .topk_details._topk_backward import _topk_backward
+from .tensor import Tensor, Bitmatrix
 
 
 def topk_forward(x, k, apply_softmax=True, dim=1, return_bitmatrix=True, y_indx=None, n_rows=None):