MCPcopy Create free account

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

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

↓ 7 callersFunctioncodegen_func
Generated codegen function. args_cpp are the C++ code strings for function arguments.
python/little_kernel/language/simple_builtin.py:395
↓ 7 callersFunctioncontextual_autotune
(is_dist=False, n_repeat=5, n_warmup=3)
python/triton_dist/autotuner.py:97
↓ 7 callersFunctioncreate_ep_a2a_fused_context
( group: torch.distributed.ProcessGroup, max_tokens: int, hidden: int, topk: int, num_tot_
python/triton_dist/kernels/amd/ep_all2all_fused.py:468
↓ 7 callersMethoddispatch
(self, func: Callable)
python/triton_dist/language/extra/utils.py:49
↓ 7 callersMethodfinalize
(self)
python/triton_dist/mega_triton_kernel/models/model_builder.py:597
↓ 7 callersMethodforward
( self, input: torch.Tensor, # [M, K_per_rank] weight: torch.Tensor, # [K_per_rank,
python/triton_dist/test/nvidia/test_gemm_rs.py:117
↓ 7 callersMethodgeneric_visit
Default handler for unimplemented AST nodes.
python/little_kernel/codegen/codegen_base.py:244
↓ 7 callersFunctionget_bool_env
(env, default_value)
python/triton_dist/utils.py:899
↓ 7 callersMethodis_first_stage
Check if this is the first PP stage
python/triton_dist/layers/nvidia/pp_block.py:209
↓ 7 callersFunctionis_fp8_dtype
(dtype: torch.dtype)
python/triton_dist/utils.py:406
↓ 7 callersFunctionis_maca
()
python/triton_dist/utils.py:82
↓ 7 callersFunctionmori_shmem_barrier_all_on_stream
(stream: Optional[torch.cuda.Stream] = None)
python/triton_dist/utils.py:339
↓ 7 callersFunctionreset_profiler_buffer
(buf)
python/triton_dist/tools/profiler/context.py:56
↓ 7 callersMethodsend
(self, ctx, ts, rank, dst_rank)
python/triton_dist/test/nvidia/test_pp.py:193
↓ 7 callersFunctionset_rocshmem_ctx
(ctx, _semantic=None)
python/triton_dist/language/extra/hip/librocshmem_device.py:46
↓ 7 callersFunctionshard_local
(tensor: torch.Tensor, world_size: int, dim: int, local_rank: int)
python/triton_dist/layers/nvidia/tp_attn.py:47
↓ 7 callersMethodtorch_fwd
Reference PyTorch forward pass using Tensor Parallelism. Final output is AllReduced. x: input tensor, shape [batch_size * seq
python/triton_dist/layers/amd/tp_mlp.py:107
↓ 7 callersFunctionwarp_prefix_sum_kernel
(value, lane_id, len)
python/triton_dist/kernels/nvidia/gemm_rs_threadblock_swizzle.py:43
↓ 6 callersFunction_compute_pid
(tile_id, num_pid_in_group, num_pid_m, GROUP_SIZE_M, NUM_SMS)
python/triton_dist/kernels/nvidia/gemm.py:593
↓ 6 callersFunction_getmem_impl
(dest, source, nbytes, pe, SCOPE_SUFFIX: core.constexpr, NBI: core.constexpr = core.constexpr(""),
python/triton_dist/language/extra/hip/librocshmem_device.py:131
↓ 6 callersFunction_getmem_impl
(dest, source, nbytes, pe, SCOPE_SUFFIX: core.constexpr, NBI: core.constexpr = core.constexpr(""),
python/triton_dist/language/extra/cuda/libnvshmem_device.py:382
↓ 6 callersMethod_init_parameters
Initializes and shards MLP parameters for Tensor Parallelism. mlp: A standard nn.Module MLP (e.g., from HuggingFace Transformers).
python/triton_dist/mega_triton_kernel/models/layers/tp_mlp.py:61
↓ 6 callersFunction_is_maca_platform
Checks if 'mx-smi' is available on the system's PATH.
python/setup.py:77
↓ 6 callersFunction_make_tensor_desc
(desc_or_ptr, shape, strides, block_shape)
python/triton_dist/mega_triton_kernel/test/triton_impl_utils.py:149
↓ 6 callersFunction_make_tensor_desc
(desc_or_ptr, shape, strides, block_shape)
python/triton_dist/mega_triton_kernel/kernels/flash_attn.py:86
↓ 6 callersFunction_putmem_impl
(dest, source, nbytes, pe, SCOPE_SUFFIX: core.constexpr, NBI: core.constexpr = core.constexpr(""),
python/triton_dist/language/extra/hip/librocshmem_device.py:184
↓ 6 callersFunction_putmem_impl
(dest, source, nbytes, pe, SCOPE_SUFFIX: core.constexpr, NBI: core.constexpr = core.constexpr(""),
python/triton_dist/language/extra/cuda/libnvshmem_device.py:435
↓ 6 callersFunction_putmem_rma_impl
(dest, source, nbytes, pe, SCOPE_SUFFIX: core.constexpr, NBI: core.constexpr = core.constexpr(""),
python/triton_dist/language/extra/cuda/libnvshmem_device.py:850
↓ 6 callersFunction_putmem_signal_impl
(dest, source, nbytes, sig_addr, signal, sig_op, pe, SCOPE_SUFFIX: core.constexpr, NBI
python/triton_dist/language/extra/hip/librocshmem_device.py:237
↓ 6 callersFunction_putmem_signal_impl
(dest, source, nbytes, sig_addr, signal, sig_op, pe, SCOPE_SUFFIX: core.constexpr, NBI
python/triton_dist/language/extra/cuda/libnvshmem_device.py:488
↓ 6 callersFunction_putmem_signal_rma_impl
(dest, source, nbytes, sig_addr, signal, sig_op, pe, SCOPE_SUFFIX: core.constexpr,
python/triton_dist/language/extra/cuda/libnvshmem_device.py:905
↓ 6 callersFunction_recv_ll_block
split src/dest outside of _recv_ll. this function is designed for a threadblock num_ints: of the pre-LL-packed num_ints.
python/triton_dist/kernels/nvidia/low_latency_allgather.py:531
↓ 6 callersFunctionbarrier_all_intra_node_non_atomic
symm_flags is expected to: 1. of int32 dtype 2. has at least num_ranks * 2 + 1 elements 3. of symmetric pointer
python/triton_dist/kernels/nvidia/common_ops.py:200
↓ 6 callersFunctioncdiv
Ceiling division: (a + b - 1) // b Returns the same type as the first argument (or second if first is not LLType). If both arguments are const
python/little_kernel/language/intrin/arith.py:36
↓ 6 callersFunctionconstexpr
(s)
python/triton_dist/tools/compile/compile.py:59
↓ 6 callersFunctioncopy_1d_tilewise_kernel
(dst_ptr, src_ptr, # nelems, # BLOCK_SIZE: tl.conste
python/triton_dist/kernels/nvidia/ep_a2a.py:329
↓ 6 callersFunctioncopy_tensor
(dst_tensor: torch.Tensor, src_tensor: torch.Tensor, num_sms: int = -1, eager=False, persistent=True)
python/triton_dist/kernels/nvidia/memory_ops.py:452
↓ 6 callersMethoddevice
(self)
python/triton_dist/utils.py:1236
↓ 6 callersMethoddist_triton_fwd
triton_dist forward pass for TP. This version uses ag_gemm and gemm_rs. x: input tensor, shape [batch_size, seq_len, hidden_s
python/triton_dist/layers/amd/tp_mlp.py:122
↓ 6 callersMethodep_barrier_all
(self, stream: torch.cuda.Stream, intra_node_only: bool = False)
python/triton_dist/layers/nvidia/ep_a2a_layer.py:301
↓ 6 callersFunctionformat_error_message
Format an error message with location information and code snippet. Args: message: Error message node: AST node where th
python/little_kernel/core/passes/utils/error_report.py:170
↓ 6 callersFunctionfunc
()
python/triton_dist/test/metax/test_ag_gemm_intra_node.py:74
↓ 6 callersFunctionfunc
()
unittest/test_triton_tuple_input.py:114
↓ 6 callersFunctionfused_dispatch_token_moe_grouped_gemm
Run fused dispatch + grouped GEMM. Returns ``(gemm_output[M, N], meta)``. ``meta`` carries the layout (recv offsets, scatter idx, per-rank token
python/triton_dist/kernels/amd/ep_all2all_fused.py:851
↓ 6 callersFunctionfwd
()
python/triton_dist/test/amd/test_all_to_all.py:166
↓ 6 callersFunctionfwd
()
python/triton_dist/test/nvidia/test_all_to_all.py:164
↓ 6 callersMethodget_nvshmem_size_mb
Get the total nvshmem memory size in MB.
python/triton_dist/layers/nvidia/ep_a2a_fused_layer.py:253
↓ 6 callersFunctionget_pyrocshmem_device_ctx
()
python/triton_dist/kernels/amd/allgather.py:621
↓ 6 callersFunctionget_tensorcore_tflops
(dtype: torch.dtype)
python/triton_dist/kernels/nvidia/gemm_perf_model.py:187
↓ 6 callersFunctionget_triton_dist_world
()
python/triton_dist/utils.py:380
↓ 6 callersFunctionis_bitwise_match
(x: torch.Tensor, y: torch.Tensor)
python/triton_dist/test/nvidia/test_multimem_ld_reduce.py:104
↓ 6 callersMethodis_int
(self)
python/triton_dist/tools/tune/find_topk.py:109
↓ 6 callersFunctionis_nvshmem_multimem_supported
()
python/triton_dist/utils.py:766
↓ 6 callersFunctionmake_8x_ptrs
(val0, val1, val2, val3, val4, val5, val6, val7)
python/triton_dist/kernels/amd/allgather.py:274
↓ 6 callersFunctionmake_cuda_graph
(mempool, func)
python/triton_dist/test/amd/test_tp_mlp.py:65
↓ 6 callersFunctionmake_ptr_tensor
Create int64 tensor of data_ptr() values on target_device.
python/little_kernel/design/test_flashcomm_multi_gpu.py:192
↓ 6 callersMethodmatch
Checks if a given integer value matches the filter's rule. Args: val: The integer value to check. Returns:
python/triton_dist/tools/tune/tune_gemm.py:181
↓ 6 callersFunctionmultimem_ld_reduce_v4
Load data from global memory with PTX instructions multimem.ld_reduce Args: ptr: Pointer to the global memory acc_prec: Accum
python/triton_dist/language/extra/cuda/language_extra.py:436
↓ 6 callersFunctionsample_token
(logits: torch.Tensor, temperature=0.6, top_p=0.95, top_k=-1)
python/triton_dist/models/utils.py:86
↓ 6 callersFunctionsingle_block_prefix_sum_kernel_scan_scan
( split_ptr, partial_sum_ptr, res_ptr, M, N, num_warps: tl.constexpr, exclusive: t
python/triton_dist/kernels/nvidia/all_to_all_vdev_2d_offset.py:44
↓ 6 callersFunctionsort_by_vectors
(x)
python/triton_dist/test/nvidia/test_ep_a2a.py:78
↓ 6 callersMethodsp_group_barrier_all_intra_node
(self, stream=None)
python/triton_dist/kernels/nvidia/sp_ulysess_qkv_gemm_all2all.py:686
↓ 6 callersFunctionsupports_p2p_native_atomic
()
python/triton_dist/utils.py:544
↓ 6 callersFunctionto_allreduce_method
(method: str)
python/triton_dist/kernels/allreduce.py:64
↓ 6 callersFunctiontriton_dist_pre_attn_a2a
( inputs, seq_lens_cpu=None, num_comm_sm=-1, )
python/triton_dist/test/nvidia/test_llm_ulysess_pre_attn_all2all_intra_node.py:76
↓ 6 callersFunctiontriton_dist_pre_attn_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:90
↓ 6 callersFunctiontriton_func
()
python/triton_dist/test/metax/test_ag_gemm_inter_node.py:74
↓ 6 callersMethodupdate_kv_cache
(self, new_k_cache: torch.Tensor, new_v_cache: torch.Tensor, layer_idx: int)
python/triton_dist/models/kv_cache.py:49
↓ 6 callersMethodwait_signal
(self, pp_rank, buffer_id, value, stream=None, num_barriers=1)
python/triton_dist/layers/nvidia/p2p.py:159
↓ 6 callersMethodwriteln_header
Write a line to header buffer (for mixin compatibility).
python/little_kernel/codegen/special_struct/translator.py:106
↓ 5 callersMethodGetOrLoad
python/triton_dist/tools/runtime/triton_aot_runtime.cc:75
↓ 5 callersFunction_build_dispatch_metadata_device
Device-kernel version of ``_build_dispatch_metadata`` (same meta keys). Replaces the host ``bincount`` + ``all_gather_into_tensor`` + recv-offset
python/triton_dist/kernels/amd/ep_all2all_fused.py:784
↓ 5 callersMethod_build_tasks_impl
(cls, device_prop, layer_id: int, dependency: TaskDependency, io_tensors, extra_params,
python/triton_dist/mega_triton_kernel/tasks/linear.py:167
↓ 5 callersFunction_build_wgmma_64x64_fn_body
()
python/little_kernel/language/intrin/wgmma.py:270
↓ 5 callersFunction_get_amdsmi_device_index
(device_id: int | None)
python/triton_dist/amd_utils.py:272
↓ 5 callersFunction_get_sm_arch
Return SM architecture integer (e.g. 90, 100), or None.
python/little_kernel/tests/conftest.py:38
↓ 5 callersFunction_pack_ll_block
split src/dest outside of _recv_ll. this function is designed for a threadblock nbytes: of the pre-LL-packed bytes. BLOCK_SIZE: count by ints
python/triton_dist/kernels/nvidia/low_latency_allgather.py:549
↓ 5 callersFunction_ptx_suffix_to_tl_type
(suffix: core.constexpr, _semantic=None)
python/triton_dist/language/extra/cuda/language_extra.py:83
↓ 5 callersFunctionadd_links
(external_only, materialization=False)
python/setup.py:968
↓ 5 callersFunctionall_reduce
Performs an all-reduce operation on the input tensor `x`. This is the main entry point for the all-reduce collective. It selects the specifi
python/triton_dist/kernels/nvidia/allreduce.py:1130
↓ 5 callersFunctionall_to_all_v_offset_op
(ctx: AllToAllContext, rank_in_row: bool, input: torch.Tensor = None, output: torch
python/triton_dist/kernels/nvidia/all_to_all_vdev_2d_offset_inter_node.py:621
↓ 5 callersFunctionbincount
GPU bincount for AMD (no AOT). input_tensor: 1D int32 on device; output: length elements.
python/triton_dist/kernels/amd/ep_a2a.py:55
↓ 5 callersFunctioncheck_with_ntokens_per_rank_per_expert
(ntokens_per_rank_per_expert: np.ndarray, nexperts, TP_SIZE, LOCAL_TP_SIZE,
python/triton_dist/kernels/nvidia/threadblock_swizzle_ag_moe_triton.py:352
↓ 5 callersFunctioncheck_with_token_cnts
python/triton_dist/kernels/nvidia/threadblock_swizzle_ag_moe.cc:305
↓ 5 callersFunctionconsumer_gemm_persistent
(a, b, c, rank, num_ranks, barrier_tensor, needs_wait=True, barriers_per_block=1)
python/triton_dist/test/nvidia/test_distributed_wait.py:431
↓ 5 callersFunctionconsumer_gemm_persistent
(a, b, c, rank, num_ranks, barrier_tensor, needs_wait=True, barriers_per_block=1)
python/triton_dist/test/metax/test_distributed_wait.py:408
↓ 5 callersFunctioncreate_ag_gemm_context
create context for allgather gemm intra-node Args: rank (int): current rank num_ranks (int): total number of ranks max_M:
python/triton_dist/kernels/nvidia/allgather_gemm.py:511
↓ 5 callersMethoddump_trace
(self, trace_file_prefix="MEGA_KERNEL_TRACE")
python/triton_dist/mega_triton_kernel/models/model_builder.py:548
↓ 5 callersMethodenter_scope
Enter a new scope (e.g. function body, loop body). Args: initial_types: Optional initial type mappings for this
python/little_kernel/core/passes/utils/scope_manager.py:167
↓ 5 callersMethodep_barrier
(self)
python/triton_dist/kernels/amd/ep_all2all_fused.py:455
↓ 5 callersMethodep_barrier_all
(self, stream: torch.cuda.Stream, intra_node_only: bool = False)
python/triton_dist/layers/amd/ep_a2a_layer.py:274
↓ 5 callersFunctionf_run
(config)
python/triton_dist/autotuner.py:158
↓ 5 callersFunctionfast_all_to_all
low-latency all-to-all communication
python/triton_dist/kernels/nvidia/low_latency_all_to_all.py:198
↓ 5 callersMethodfinalize
(self)
python/triton_dist/layers/nvidia/sp_flash_decode_layer.py:79
↓ 5 callersFunctionfunc
()
python/triton_dist/test/nvidia/test_decode_attn.py:328
↓ 5 callersMethodfunc_dir
(self)
python/triton_dist/tune.py:306
↓ 5 callersFunctiongemm_swizzle2d_Nz
gemm swizzle Nz
python/triton_dist/language/extra/ascend/algorithm.py:40
↓ 5 callersFunctionget_all_registered_special_structs
Get all registered special structs.
python/little_kernel/codegen/registries/special_struct_registry.py:147
↓ 5 callersFunctionget_allreduce_methods
()
python/triton_dist/kernels/allreduce.py:70
↓ 5 callersMethodget_buffer
(self, buffer_id)
python/triton_dist/layers/nvidia/p2p.py:97
← previousnext →201–300 of 5,103, ranked by callers