↓ 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
↓ 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 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 callersFunctionmake_8x_ptrs(val0, val1, val2, val3, val4, val5, val6, val7)
python/triton_dist/kernels/amd/allgather.py:274
↓ 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 callersMethodwait_signal(self, pp_rank, buffer_id, value, stream=None, num_barriers=1)
python/triton_dist/layers/nvidia/p2p.py:159
↓ 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 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