MCPcopy Create free account

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

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

↓ 15 callersMethod_lltype_to_cpp_type
Convert LLType to C++ type string.
python/little_kernel/codegen/special_struct/struct_converter.py:792
↓ 15 callersMethodcreate
(max_seq, q_nheads, kv_nheads, k_head_dim, v_head_dim, rank, world_size, local_world_size, data_dtype,
python/triton_dist/kernels/nvidia/ulysses_sp_infer_gemm_a2a.py:75
↓ 15 callersFunctionget_base_dir
()
python/build_helpers.py:32
↓ 15 callersFunctionget_special_struct_codegen
Get the code generators for a special struct type.
python/little_kernel/codegen/registries/special_struct_registry.py:142
↓ 15 callersFunctionis_hip
()
python/triton_dist/utils.py:75
↓ 15 callersMethodunparse
Unparse a node with custom attributes.
python/little_kernel/core/passes/utils/enhanced_unparse.py:307
↓ 14 callersFunctionaddNamedAttrs
pass named attrs (e.g., tt.contiguity) from Triton to Triton
lib/Conversion/TritonDistributedToTritonGPU/METAX/TritonDistributedToTritonGPU.cpp:64
↓ 14 callersMethodis_tensor
Check if the type is a tensor (ND Array)
python/little_kernel/core/type_system.py:55
↓ 13 callersMethod__new__
(cls, inner_type: LLType)
python/little_kernel/core/type_system.py:263
↓ 13 callersMethod_convert_op
(self, op_type: str, layer_id: int, io_tensors: List[List[torch.Tensor]], extra_params: Di
python/triton_dist/mega_triton_kernel/models/model_builder.py:199
↓ 13 callersMethod_create_task
(cls, layer_id: int, task_id: int, tile_id_or_start: int, num_tiles: int, config: ConfigBase,
python/triton_dist/mega_triton_kernel/core/builder.py:48
↓ 13 callersMethod_writeln
Write a line with proper indentation.
python/little_kernel/codegen/special_struct/struct_converter.py:70
↓ 13 callersFunctionaddNamedAttrs
pass named attrs (e.g., tt.contiguity) from Triton to Triton
lib/Conversion/TritonDistributedToTritonGPU/TritonDistributedToTritonGPU.cpp:63
↓ 13 callersFunctionbarrier_all_intra_node_atomic_cas_block
NOTE: this function should only be called with atomic support. memory over PCI-e does not support atomic r/w. DON'T use this function on such pla
python/triton_dist/kernels/nvidia/common_ops.py:154
↓ 13 callersFunctioncdiv
(x: int, y: int)
python/triton_dist/mega_triton_kernel/tasks/utils.py:53
↓ 13 callersMethodcreate_config
(cls, **kwargs)
python/triton_dist/mega_triton_kernel/core/builder.py:73
↓ 13 callersMethodget_task_id
(cls, layer_id: int)
python/triton_dist/mega_triton_kernel/core/task_base.py:92
↓ 13 callersFunctionpromote_type
Promote types following C++ implicit conversion rules. Returns: (promoted_left, promoted_right, was_promoted): The promoted type
python/little_kernel/core/passes/utils/type_inference/type_inference_visitors.py:210
↓ 13 callersMethodrelease_tile
(self, task_base_info: TaskBaseInfo, tile_id)
python/triton_dist/mega_triton_kernel/kernels/task_context.py:141
↓ 13 callersMethodresolve
Resolve a name to its value. Priority: 1. If name is local (defined in kernel), return None (don't resolve)
python/little_kernel/core/passes/utils/scope_manager.py:266
↓ 13 callersFunctionthreads_per_warp
(_semantic=None)
python/triton_dist/language/extra/language_extra.py:99
↓ 12 callersFunction__shfl_sync_i32
Shuffle idx: read *value* from *lane* (broadcast if lane is constant).
python/triton_dist/language/extra/hip/language_extra.py:505
↓ 12 callersMethod_check_error
Check CUDA error and raise exception if needed.
python/little_kernel/runtime/cuda_runtime.py:333
↓ 12 callersMethod_init_ctx
Initializes contexts for triton_dist AllGather-GEMM and GEMM-ReduceScatter operations.
python/triton_dist/layers/amd/tp_mlp.py:89
↓ 12 callersFunctioninit_model_cpu
(model_name: str, dtype: torch.dtype)
python/triton_dist/models/utils.py:108
↓ 12 callersFunctionnum_warps
(_semantic=None)
python/triton_dist/language/extra/language_extra.py:114
↓ 11 callersMethod_ensure_materialized
Ensure the tensor is materialized before access.
python/triton_dist/utils.py:1167
↓ 11 callersFunction_set_signal_cuda
(signal_tensor: torch.Tensor, signal: int, stream: Optional[torch.cuda.Stream] = None)
python/triton_dist/kernels/nvidia/common_ops.py:387
↓ 11 callersMethoddefine_from_ast
Helper to define variables from AST nodes (targets of assignments, args, etc.) Handles: - ast.Name: Simple variable
python/little_kernel/core/passes/utils/scope_manager.py:206
↓ 11 callersFunctionexport_to_perfetto_trace
(profiler_buffer: torch.Tensor, task_names: List[str], file_name: str, verbose: b
python/triton_dist/tools/profiler/viewer.py:115
↓ 11 callersFunctionhas_amdsmi
()
python/triton_dist/amd_utils.py:47
↓ 11 callersFunctionlaneid
(_semantic=None)
python/triton_dist/language/extra/hip/language_extra.py:474
↓ 11 callersFunctionld_b32
(ptr, _semantic=None)
python/triton_dist/language/extra/cuda/language_extra.py:672
↓ 11 callersMethodrun
(self)
python/setup.py:619
↓ 11 callersMethodswitch_off_use
(self)
python/little_kernel/core/passes/constfold.py:93
↓ 11 callersFunctiontry_build
(name, build_fn)
python/little_kernel/design/test_flashcomm_compile_all.py:42
↓ 10 callersFunctionHIP_CHECK
(call_result)
python/triton_dist/utils.py:482
↓ 10 callersFunctionMACA_CHECK
(err)
python/triton_dist/utils.py:493
↓ 10 callersFunction_correctness_check
Build a kernel from the given module and verify correctness.
python/little_kernel/tests/integration/test_sm90_gemm.py:41
↓ 10 callersFunction_take_a_meta
(x)
python/triton_dist/tools/compile_aot.py:446
↓ 10 callersFunction_wait_eq_cuda
(signal_tensor: torch.Tensor, signal: int, stream: Optional[torch.cuda.Stream] = None, requi
python/triton_dist/kernels/nvidia/common_ops.py:364
↓ 10 callersMethodadd_variable_to_scope
Add a variable to the active scope
python/little_kernel/core/passes/utils/type_inference/infer_type.py:87
↓ 10 callersFunctionatomic_cas
(ptr, cmp_val, target_val, scope, semantic)
python/triton_dist/language/extra/language_extra.py:53
↓ 10 callersMethodgetThreadsPerWarp
lib/Conversion/TritonDistributedToTritonGPU/TritonDistributedToTritonGPU.cpp:974
↓ 10 callersFunctionget_shmem_backend
()
python/triton_dist/utils.py:97
↓ 10 callersMethodis_last_stage
Check if this is the last PP stage
python/triton_dist/layers/nvidia/pp_block.py:213
↓ 10 callersFunctionlayer_norm
Applies RMS Normalization
python/triton_dist/layers/amd/tp_attn.py:52
↓ 10 callersFunctionmembar
(scope: core.constexpr = core.constexpr("cta"), _semantic=None)
python/triton_dist/language/extra/cuda/language_extra.py:958
↓ 10 callersMethodrun
(self)
python/triton_dist/mega_triton_kernel/models/model_builder.py:557
↓ 10 callersMethodwriteln_struct
Write line to struct buffer with newline.
python/little_kernel/codegen/codegen_base.py:136
↓ 9 callersFunction_correctness_check
Build a kernel from the given module and verify correctness.
python/little_kernel/tests/integration/test_sm100_gemm.py:41
↓ 9 callersFunction_int_constraint
(bitwidth: core.constexpr, _semantic=None)
python/triton_dist/language/extra/cuda/language_extra.py:33
↓ 9 callersFunctionalloc_profiler_buffer
(max_num_profile_slots)
python/triton_dist/tools/profiler/context.py:50
↓ 9 callersFunctionatomic_add_per_warp
Warp-level atomic add: lane 0 performs atomic_add, result is broadcast to all lanes. AMD equivalent of the NVIDIA atomic_add_per_warp. Uses wavef
python/triton_dist/language/extra/hip/language_extra.py:552
↓ 9 callersFunctionbarrier_all_on_stream
Call shmem barrier on stream: mori_shmem when backend is mori_shmem, else rocshmem.
python/triton_dist/kernels/amd/common_ops.py:182
↓ 9 callersFunctionbin_op_kind_str
(kind)
python/little_kernel/core/expr.py:343
↓ 9 callersFunctionbitwise_equal
(x: torch.Tensor, y: torch.Tensor)
python/triton_dist/test/utils.py:96
↓ 9 callersFunctionget_special_attrs
Get all special attributes from a node as a dictionary. Args: node: AST node to extract attributes from Returns:
python/little_kernel/core/passes/utils/preserve_attributes.py:95
↓ 9 callersMethodget_variable_type
Get type of a variable from scope/context.
python/little_kernel/core/passes/utils/type_inference/infer_type.py:93
↓ 9 callersFunctionhas_fullmesh_nvlink
()
python/triton_dist/nv_utils.py:295
↓ 9 callersMethodis_local
Check if name is defined in any local scope (shadows global). This is the key method for determining symbol resolution prior
python/little_kernel/core/passes/utils/scope_manager.py:249
↓ 9 callersFunctionload_v4_u32
(ptr, _semantic=None)
python/triton_dist/language/extra/cuda/language_extra.py:181
↓ 9 callersFunctionmultimem_st_b64
(ptr, val0, _semantic=None)
python/triton_dist/language/extra/cuda/language_extra.py:278
↓ 9 callersFunctionntid
(axis: core.constexpr, _semantic=None)
python/triton_dist/language/extra/hip/language_extra.py:93
↓ 9 callersMethodreset
(self)
python/triton_dist/mega_triton_kernel/models/model_builder.py:511
↓ 9 callersMethodswitch_on_use
(self)
python/little_kernel/core/passes/constfold.py:89
↓ 9 callersFunctionwith_pynvml
()
python/triton_dist/nv_utils.py:63
↓ 8 callersMethod__init__
(self, name: str, expr_type: LLType, defining_stmt=None, is_shared: bool = False, is_global:
python/little_kernel/core/expr.py:249
↓ 8 callersFunction_format_value
Format a value for display in dump.
python/little_kernel/core/passes/utils/enhanced_dump.py:117
↓ 8 callersFunction_ptx_suffix_to_constraint
(suffix: core.constexpr, _semantic=None)
python/triton_dist/language/extra/cuda/language_extra.py:51
↓ 8 callersFunctionbarrier_all_kernel
(rank, num_ranks, comm_buf_ptr)
python/triton_dist/kernels/amd/common_ops.py:168
↓ 8 callersMethodclear
(self)
python/triton_dist/models/kv_cache.py:62
↓ 8 callersMethoddata_ptr
(self)
python/little_kernel/design/test_flashcomm_torchrun.py:114
↓ 8 callersMethoddefine
Define a name in the current scope. This marks the name as locally defined, which means it shadows any global defini
python/little_kernel/core/passes/utils/scope_manager.py:182
↓ 8 callersFunctiondownload_and_copy
(name, src_func, dst_path, variable, version, url_func)
python/setup.py:424
↓ 8 callersFunctionelement_at
(x: tl.tensor, idx: int)
python/triton_dist/kernels/nvidia/group_gemm.py:35
↓ 8 callersMethodfinalize
(self)
tutorials/06-inter-node-reduce-scatter.py:124
↓ 8 callersFunctionfn
()
python/triton_dist/test/nvidia/test_moe_utils.py:103
↓ 8 callersMethodinference
(self, input_ids: torch.LongTensor, position_ids: torch.LongTensor, kv_cache: KV_Cache, wo_l
python/triton_dist/models/dense.py:215
↓ 8 callersFunctionis_ir_node
Check if a node is an IR node (EnumDef or SpecialStructDef).
python/little_kernel/core/passes/utils/ir_nodes.py:91
↓ 8 callersFunctionlayer_norm
Applies RMS Normalization using flashinfer.
python/triton_dist/layers/nvidia/tp_attn.py:61
↓ 8 callersFunctionnext_power_of_2
(n: tl.constexpr)
python/triton_dist/kernels/nvidia/common_ops.py:271
↓ 8 callersFunctionnvsmi
(attrs, device_id=0, dtype: type = int)
python/triton_dist/nv_utils.py:70
↓ 8 callersFunctionpack_b32_v2
Pack two 32-bit values into one 64-bit value (val0: low 32, val1: high 32).
python/triton_dist/language/extra/hip/language_extra.py:544
↓ 8 callersFunctionraise_pass_error
Raise a pass error with formatted message. Args: message: Error message node: AST node where the error occurred
python/little_kernel/core/passes/utils/error_report.py:267
↓ 8 callersFunctionregister_pass
Convenience function to register a pass.
python/little_kernel/core/passes/pass_registry.py:96
↓ 8 callersFunctionrmsnorm_ref
(x, w, eps=1e-6)
python/triton_dist/mega_triton_kernel/test/torch_impl_utils.py:41
↓ 8 callersMethodsp_group_barrier_all_intra_node
(self, stream=None)
python/triton_dist/kernels/nvidia/sp_ulysess_o_all2all_gemm.py:627
↓ 8 callersFunctiontriton_dist_get_input_comm_buf
( tensor, comm_buf_idx=0, )
python/triton_dist/test/nvidia/test_llm_ulysess_pre_attn_all2all_intra_node.py:136
↓ 8 callersFunctiontriton_dist_sp_group_barrier_all
(stream=None)
python/triton_dist/test/nvidia/test_llm_ulysess_pre_attn_all2all_intra_node.py:144
↓ 8 callersFunctionupdate_context
(ctx, name, obj)
python/little_kernel/core/passes/utils/update_context.py:27
↓ 7 callersFunction__shfl_up_sync_i32
Shuffle up: each lane reads from (laneid - delta), clamped to 0.
python/triton_dist/language/extra/hip/language_extra.py:512
↓ 7 callersFunction_check_rocm_smi_json
(cmd)
python/triton_dist/amd_utils.py:62
↓ 7 callersFunction_get_pynvml_device_id
(device_id: int | None = None)
python/triton_dist/nv_utils.py:157
↓ 7 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/layers/amd/tp_mlp.py:64
↓ 7 callersFunction_is_ascend_platform
Checks if 'npu-smi' is available on the system's PATH.
python/setup.py:87
↓ 7 callersFunction_optional_or
(arg, default)
python/triton_dist/test/nvidia/perf_ring_reduce_tma.py:180
↓ 7 callersFunction_run_straggler
(ctx, straggler_option)
python/triton_dist/kernels/nvidia/allreduce.py:138
↓ 7 callersMethod_scan_for_assigned_vars
Recursively scan statements for variables assigned anywhere (including nested For/If/While). Returns dict of var_name -> (var_cpp_type, var_ll
python/little_kernel/codegen/visitors/control_flow_codegen.py:49
↓ 7 callersFunctionag_gemm
allgather gemm C = all_gather(A) * B Args: A (torch.Tensor<float>): local matmul A matrix. shape: [M_per_rank, K] B (torch.Te
python/triton_dist/kernels/nvidia/allgather_gemm.py:570
← previousnext →101–200 of 5,103, ranked by callers