↓ 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 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_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 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_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_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_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
↓ 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 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