MCPcopy Create free account

hub / github.com/ByteDance-Seed/Triton-distributed / functions

Functions5,103 in github.com/ByteDance-Seed/Triton-distributed

↓ 2 callersMethodregister_codegen
Register a codegen handler for a loop modifier. Args: modifier_name: Name of the modifier (e.g., "unroll", "para
python/little_kernel/codegen/registries/loop_modifier_codegen.py:44
↓ 2 callersMethodregister_modifier
Register a loop modifier. Args: name: Name of the modifier (e.g., "unroll", "parallel") modifier: Lo
python/little_kernel/core/passes/utils/registries/loop_modifier_registry.py:79
↓ 2 callersMethodregister_pass
Register a pass for a specific backend. Args: backend: Backend name (e.g., "cuda") pass_func: Pass f
python/little_kernel/core/passes/pass_registry.py:57
↓ 2 callersMethodreset_cusum_seq_lens
(self, local_seqlen, seq_lens_cpu=None)
python/triton_dist/kernels/nvidia/sp_ulysess_o_all2all_gemm.py:634
↓ 2 callersFunctionrotate_half
Rotates half the hidden dims of the input.
python/triton_dist/test/nvidia/test_patch_torch_compile.py:30
↓ 2 callersFunctionrun_allgather
()
python/triton_dist/test/amd/test_bandwidth.py:166
↓ 2 callersFunctionrun_command
Execute command with retries, show full output on failure
python/triton_dist/benchmark/bench_tp_attn.py:142
↓ 2 callersFunctionrun_moe_reduce_ar_triton_non_overlap
(x: torch.Tensor, weights: torch.Tensor, chosen_experts: torch.Tensor,
python/triton_dist/kernels/nvidia/moe_reduce_ar.py:221
↓ 2 callersFunctionrun_moe_reduce_rs_triton_non_overlap
(x: torch.Tensor, weights: torch.Tensor, chosen_experts: torch.Tensor,
python/triton_dist/kernels/nvidia/moe_reduce_rs.py:825
↓ 2 callersFunctionrun_p2p
()
python/triton_dist/test/amd/test_bandwidth.py:305
↓ 2 callersFunctionrun_reduce
(t: torch.Tensor, reduce_order: List[int], acc_dtype: torch.dtype = None)
python/triton_dist/test/nvidia/test_multimem_ld_reduce.py:80
↓ 2 callersFunctionsend
(ctx, ts, rank, dst_rank)
python/triton_dist/test/nvidia/test_pp.py:82
↓ 2 callersFunctionshow_generated_code
(kernel_name="basic")
python/little_kernel/design/flashcomm_dispatch.py:373
↓ 2 callersFunctionshow_generated_code
Show generated CUDA code.
python/little_kernel/design/flashcomm_barrier.py:100
↓ 2 callersFunctionshow_generated_code
()
python/little_kernel/design/flashcomm_postprocess.py:172
↓ 2 callersFunctionshow_generated_code
(kernel_name="offset")
python/little_kernel/design/flashcomm_compute.py:329
↓ 2 callersFunctionshow_generated_code
(kernel_name="preprocess")
python/little_kernel/design/flashcomm_combine.py:531
↓ 2 callersFunctionshow_generated_code
()
python/little_kernel/design/flashcomm_dispatch_chunk.py:349
↓ 2 callersFunctionsilu
(x)
python/triton_dist/kernels/nvidia/swiglu.py:57
↓ 2 callersFunctionsp_intra_node_barrier
(ctx)
python/triton_dist/kernels/nvidia/ulysses_sp_infer_gemm_a2a.py:389
↓ 2 callersFunctionsplits_to_cumsum
(splits: torch.Tensor)
python/triton_dist/test/nvidia/test_ep_moe_inference.py:224
↓ 2 callersFunctionsplits_to_cumsum
(splits: torch.Tensor)
tutorials/04-deepseek-infer-all2all.py:279
↓ 2 callersFunctionst
( ptr, val, scope: core.constexpr = core.constexpr("gpu"), semantic: core.cons
python/triton_dist/language/extra/cuda/language_extra.py:687
↓ 2 callersFunctionswiglu_forward
(AB, scale=None, sm_margin=0, use_aot=False)
python/triton_dist/kernels/nvidia/swiglu.py:239
↓ 2 callersFunctionswizzle_ag_gemm_imperfect
(original_pid_m, M, rank, world_size, CHUNK_SIZE_M: tl.constexpr, BLOCK_SIZE_M:
python/triton_dist/kernels/amd/allgather_gemm.py:487
↓ 2 callersMethodsync
Not in use; reserved for future path (symmetric buffers are allocated eagerly in ``__init__``).
python/triton_dist/layers/amd/ep_a2a_fused_layer.py:140
↓ 2 callersFunctiontanh
(x)
python/triton_dist/mega_triton_kernel/kernels/utils.py:30
↓ 2 callersFunctiontanh
(x)
python/triton_dist/kernels/nvidia/flash_decode.py:65
↓ 2 callersFunctiontest
(kernel, M, N, K, label="")
python/little_kernel/benchmark/gemm_sm90/gemm_v2.py:141
↓ 2 callersFunctiontest
Test correctness.
python/little_kernel/benchmark/gemm_sm90/gemm_v1.py:183
↓ 2 callersFunctiontest_ag_gemm_tma_intra_node
(args, autotune=False, use_tma=True)
python/triton_dist/test/metax/test_ag_gemm_intra_node.py:48
↓ 2 callersFunctiontest_mori_shmem_bandwidth
Test MoRI SHMEM P2P bandwidth using putmem_nbi Args: src_rank: Source PE rank dst_rank: Destination PE rank num_
python/triton_dist/test/amd/test_mori_shmem_bw.py:201
↓ 2 callersFunctiontest_perf_ag_gemm_tma_intra_node
(args, autotune=False, use_tma=True)
python/triton_dist/test/metax/test_ag_gemm_intra_node.py:136
↓ 2 callersFunctiontest_reduce_topk
(ntokens, topk, N, dtype: torch.dtype, n_split)
python/triton_dist/test/nvidia/test_moe_utils.py:87
↓ 2 callersFunctiontest_simt_add_vec
(n, dtype)
python/triton_dist/test/common/test_simt_vec_add.py:60
↓ 2 callersFunctionthreadblock_swizzle_allgather_gemm_kernel
( tiled_m, M, rank, WORLD_SIZE: tl.constexpr, NNODES: tl.constexpr, BLOCK_SIZE_M: tl.c
python/triton_dist/kernels/nvidia/ag_gemm_threadblock_swizzle.py:52
↓ 2 callersFunctiontile_kernel_gather_combine_token_intra_node
( pid, num_pid, counter_ptr, # symm buffer, [max_tokens, hidden_size // gemm_block_size_n] ba
python/triton_dist/kernels/nvidia/ep_all2all_fused.py:308
↓ 2 callersFunctiontile_kernel_moe_grouped_gemm_nk_const
( pid, num_pid, barriers_ptr, # symm uint64 [experts_per_rank * world_size] a_ptr, # this ra
python/triton_dist/kernels/amd/ep_all2all_fused.py:219
↓ 2 callersFunctiontile_kernel_scatter_token_intra_node
( pid, num_pid, barriers_ptr, # symm buffer, per token barrier [max_tokens * topk * local_world_s
python/triton_dist/kernels/nvidia/ep_all2all_fused.py:406
↓ 2 callersFunctiontile_kernel_transposed_moe_grouped_gemm_nk_const
( pid, num_pid, grad_output_ptr, orig_input_ptr, grad_weight_ptr, split_size_ptr,
python/triton_dist/kernels/nvidia/ep_all2all_fused.py:750
↓ 2 callersMethodto_torch
Copy contents to a PyTorch tensor for inspection.
python/little_kernel/design/test_flashcomm_multi_gpu.py:122
↓ 2 callersMethodto_torch
Copy to a PyTorch tensor for inspection.
python/little_kernel/design/test_flashcomm_torchrun.py:120
↓ 2 callersFunctiontorch_ag_gemm
return C = all_gather(A) @ B.T
python/triton_dist/test/amd/test_ag_gemm_intra_node.py:53
↓ 2 callersMethodtorch_ag_gemm
Reference PyTorch forward pass using AllGather-GEMM.
python/triton_dist/layers/amd/tp_mlp.py:152
↓ 2 callersFunctiontorch_ag_group_gemm
( pg: torch.distributed.ProcessGroup, local_input: torch.Tensor, local_weight: torch.Tensor, f
python/triton_dist/test/nvidia/test_ag_moe.py:61
↓ 2 callersMethodtorch_gemm_rs
Reference PyTorch forward pass using GEMM-ReduceScatter.
python/triton_dist/layers/amd/tp_mlp.py:182
↓ 2 callersFunctiontorch_post_attn_all_to_all_transpose
(sp_group, input, a2a_only, is_dp, seq_lens_cpu=None)
python/triton_dist/test/nvidia/test_llm_ulysess_post_attn_all2all_intra_node.py:119
↓ 2 callersFunctiontorch_pre_attn_qkv_pack_a2a
(sp_group, input, bs, seq_len, nh, head_dim, gqa, seq_lens_cpu=None)
python/triton_dist/test/nvidia/test_llm_ulysess_gemm_all2all_intra_node.py:145
↓ 2 callersFunctiontorch_pre_attn_qkv_pack_a2a
(sp_group, input, bs, seq_len, nh, head_dim, gqa, seq_lens_cpu=None)
python/triton_dist/test/nvidia/test_llm_ulysess_pre_attn_all2all_intra_node.py:175
↓ 2 callersFunctiontranspose_2d
(mat: list[list[int]])
python/triton_dist/tools/tune/tune_gemm.py:214
↓ 2 callersFunctiontranspose_kernel
(ptr_x, ptr_y, M, N, BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr)
python/triton_dist/kernels/nvidia/all_to_all_vdev_2d_offset_inter_node.py:46
↓ 2 callersFunctiontransposed_moe_grouped_gemm
( grad_output, original_input, split_size, split_size_cum_per_expert, grad_weight=None,
python/triton_dist/kernels/nvidia/group_gemm.py:988
↓ 2 callersFunctiontriton_all_gather
(ag_buffer)
python/triton_dist/test/nvidia/test_all_gather.py:62
↓ 2 callersFunctiontriton_dist_post_attn_a2a
( input, seq_lens_cpu=None, num_comm_sm=-1, )
python/triton_dist/test/nvidia/test_llm_ulysess_post_attn_all2all_intra_node.py:72
↓ 2 callersFunctiontriton_dist_post_attn_a2a_gemm
( attention_outputs, weight, seq_lens_cpu=None, bias=None, outputs=None, a2a_outputs=N
python/triton_dist/test/nvidia/test_llm_ulysess_all2all_gemm_intra_node.py:75
↓ 2 callersFunctiontriton_dist_post_attn_a2a_no_cpy
( input, seq_lens_cpu=None, num_comm_sm=-1, comm_buf_idx=0, )
python/triton_dist/test/nvidia/test_llm_ulysess_post_attn_all2all_intra_node.py:85
↓ 2 callersFunctiontriton_dist_pre_attn_gemm_a2a
( attention_inputs, weight, seq_lens_cpu=None, bias=None, outputs=None, num_comm_sms=1
python/triton_dist/test/nvidia/test_llm_ulysess_gemm_all2all_intra_node.py:75
↓ 2 callersFunctiontriton_dist_pre_attn_qkv_pack_a2a
( inputs, seq_lens_cpu=None, num_comm_sm=-1, comm_buf_idx=0, )
python/triton_dist/test/nvidia/test_llm_ulysess_pre_attn_all2all_intra_node.py:105
↓ 2 callersFunctiontriton_dist_pre_attn_qkv_pack_a2a_no_cpy
( inputs, seq_lens_cpu=None, num_comm_sm=-1, comm_buf_idx=0, )
python/triton_dist/test/nvidia/test_llm_ulysess_pre_attn_all2all_intra_node.py:121
↓ 2 callersFunctiontriton_dist_sp_barrier_all
()
python/triton_dist/test/nvidia/test_llm_ulysess_post_attn_all2all_intra_node.py:67
↓ 2 callersFunctionunpack_bf16x2_f32
(v1, v2, v3, v4, _semantic=None)
python/triton_dist/kernels/nvidia/memory_ops.py:202
↓ 2 callersMethodunwrap_modifier
Unwrap a loop modifier call to get the inner iterator. Args: node: AST Call node representing the modifier call
python/little_kernel/core/passes/utils/registries/loop_modifier_registry.py:153
↓ 2 callersFunctionval_cast
Cast a value to a specific type. Usage: val_cast(value, ll.uint32) # Cast value to uint32 ll.uint32(value) # Same as above
python/little_kernel/language/intrin/dtype.py:80
↓ 2 callersMethodvisit_Constant
Process constants (map to C++ syntax).
python/little_kernel/codegen/visitors/expression_codegen.py:89
↓ 2 callersMethodvisit_EnumDef
Handle EnumDef IR node.
python/little_kernel/core/passes/utils/enhanced_unparse.py:250
↓ 2 callersMethodvisit_SpecialStructDef
Handle SpecialStructDef IR node.
python/little_kernel/core/passes/utils/enhanced_unparse.py:279
↓ 2 callersFunctionwait_eq
(barrier_ptr, thread_idx, value, scope: core.constexpr)
python/triton_dist/language/extra/cuda/language_extra.py:811
↓ 2 callersFunctionwarp_prefix_sum_kernel
(value, lane_id, len)
python/triton_dist/kernels/nvidia/ag_gemm_threadblock_swizzle.py:40
↓ 2 callersFunctionworkspace_bytes_per_in_byte
(world_size, method: AllReduceMethod)
python/triton_dist/kernels/nvidia/allreduce.py:52
↓ 2 callersMethodwriteln_builtin
Write a line to builtin buffer (for mixin compatibility).
python/little_kernel/codegen/special_struct/translator.py:100
↓ 2 callersMethodzero_
(self)
python/little_kernel/design/test_flashcomm_multi_gpu.py:118
↓ 2 callersFunctionzero_vec_f32
(vec_size: tl.constexpr, _semantic=None)
python/triton_dist/kernels/nvidia/memory_ops.py:155
↓ 1 callersFunctionCUDALaunchKernel
python/triton_dist/tools/runtime/triton_aot_runtime.cc:190
↓ 1 callersFunctionCUDAModuleGetFunction
python/triton_dist/tools/runtime/triton_aot_runtime.cc:153
↓ 1 callersFunctionCUDAModuleLoadData
python/triton_dist/tools/runtime/triton_aot_runtime.cc:143
↓ 1 callersFunctionCUDAModuleUnload
python/triton_dist/tools/runtime/triton_aot_runtime.cc:164
↓ 1 callersMethodUnload
python/triton_dist/tools/runtime/triton_aot_runtime.cc:65
↓ 1 callersMethod__enter__
(self)
python/triton_dist/profiler_utils.py:232
↓ 1 callersMethod__enter__
(self)
python/triton_dist/language/simt_ops.py:253
↓ 1 callersMethod__eq__
(self, other: object)
python/little_kernel/core/type_system.py:79
↓ 1 callersMethod__exit__
(self, exc_type, exc_val, exc_tb)
python/triton_dist/profiler_utils.py:236
↓ 1 callersMethod__exit__
(self, exc_type, exc_value, traceback)
python/triton_dist/language/simt_ops.py:256
↓ 1 callersMethod__init__
(self, spec: LazyTensorSpec, allocator: 'LazyAllocator')
python/triton_dist/utils.py:1105
↓ 1 callersMethod__init__
( self, pg: torch.distributed.ProcessGroup, batch_size: int, q_head: int,
python/triton_dist/test/nvidia/test_sp_ag_attention_intra_node.py:43
↓ 1 callersMethod__init__
( self, pg: torch.distributed.ProcessGroup, local_world_size: int, nnodes: int
python/triton_dist/test/nvidia/test_sp_ag_attention_inter_node.py:41
↓ 1 callersMethod__post_init__
(self)
python/triton_dist/kernels/nvidia/allreduce.py:97
↓ 1 callersMethod__repr__
(self)
python/triton_dist/language/simt_ops.py:165
↓ 1 callersFunction__shfl_sync_i32
(mask, value, laneid)
python/triton_dist/language/extra/cuda/language_extra.py:842
↓ 1 callersFunction__shfl_xor_sync_i32
Shuffle xor (butterfly): each lane reads from (laneid ^ mask).
python/triton_dist/language/extra/hip/language_extra.py:535
↓ 1 callersMethod__str__
(self)
python/triton_dist/mega_triton_kernel/core/task_base.py:174
↓ 1 callersMethod_add_basic_headers
Add C++ headers for std types.
python/little_kernel/codegen/codegen_base.py:93
↓ 1 callersFunction_all_reduce_with_output
(x)
python/triton_dist/test/nvidia/test_allreduce.py:108
↓ 1 callersMethod_analyze_return_from_annotation
Analyze return type from annotation (e.g., -> Tuple[bool, int, int]).
python/little_kernel/codegen/special_struct/struct_converter.py:736
↓ 1 callersMethod_analyze_return_value
Analyze return statements to determine return value structure.
python/little_kernel/codegen/special_struct/struct_converter.py:593
↓ 1 callersFunction_attn_fwd
( tile_id, q_ptr, k_ptr, v_ptr, out_ptr, # N_CTX, # H_Q: tl.constexpr, H_KV:
python/triton_dist/mega_triton_kernel/kernels/flash_attn.py:190
↓ 1 callersFunction_autotune_always_tune
()
python/triton_dist/tune.py:61
↓ 1 callersFunction_bench_fn
(self: Autotuner, *args, config, **meta)
python/triton_dist/autotuner.py:134
↓ 1 callersFunction_bench_threadblock_swizzle
()
python/triton_dist/kernels/nvidia/ag_gemm_threadblock_swizzle.py:346
↓ 1 callersMethod_build_tasks_impl
(cls, device_prop, layer_id: int, dependency: TaskDependency, io_tensors, extra_params,
python/triton_dist/mega_triton_kernel/tasks/norm.py:166
↓ 1 callersMethod_build_tasks_impl
(cls, device_prop, layer_id: int, dependency: TaskDependency, io_tensors, extra_params,
python/triton_dist/mega_triton_kernel/tasks/activation.py:74
← previousnext →901–1,000 of 5,103, ranked by callers