MCPcopy Create free account

hub / github.com/Tencent/hpc-ops / functions

Functions10,542 in github.com/Tencent/hpc-ops

⨍ Functions 10,542◇ Types & classes 8,025↳ Endpoints 13
↓ 1 callersFunctionnaive_attn_with_kvcache_qkpv_sparse_fixed_pscale
Sparse qkpertoken/vperhead reference for the fixed P-scale=256 FP8 path.
tests/test_attention_blocksparse_qkpertoken_perhead_vperhead_fp8.py:31
↓ 1 callersFunctionnaive_attn_with_kvcache_sparse_fixed_pscale
Sparse reference for the kernel's fixed P-scale=256 FP8 path.
tests/test_attention_blocksparse_qpertoken_perhead_kvpertensor_fp8.py:43
↓ 1 callersFunctionnaive_deepep_input_format_group_gemm_blockwise_fp8
(x, w, seqlens, cu_seqlens, xscale, wscale)
tests/test_group_gemm_blockwise.py:151
↓ 1 callersFunctionnaive_fuse_moe_blockwise_fp8
( x, x_scale, gate_up_weight, gate_up_weight_scale, down_weight, down_weight_scale,
tests/test_fuse_moe_blockwise.py:212
↓ 1 callersFunctionnaive_fuse_moe_pertensor_fp8
( x, gate_up_weight, down_weight, gate_up_scale, down_scale, act_and_mul_scale, to
tests/test_fuse_moe_cp_async.py:115
↓ 1 callersFunctionnaive_fuse_moe_pertensor_fp8
( x, gate_up_weight, down_weight, gate_up_scale, down_scale, act_and_mul_scale, to
tests/test_fuse_moe_pertensor.py:119
↓ 1 callersFunctionnaive_gather_expert_inputs
(x, topk_ids, num_expert, rank_ep)
tests/test_fuse_moe_cp_async.py:15
↓ 1 callersFunctionnaive_gather_expert_inputs
(x, x_scale, topk_ids, num_expert, rank_ep)
tests/test_fuse_moe_blockwise.py:23
↓ 1 callersFunctionnaive_gather_expert_inputs
(x, topk_ids, num_expert, rank_ep)
tests/test_fuse_moe_pertensor.py:23
↓ 1 callersFunctionnaive_group_gemm
(x, w, seqlens, cu_seqlens, xscale, wscale)
tests/test_group_gemm_blockwise.py:20
↓ 1 callersFunctionnaive_group_gemm_pertensor_fp8
(x, w, seqlens, cu_seqlens, scale)
tests/test_group_gemm_pertensor.py:20
↓ 1 callersFunctionnaive_reduce
(x_bf16, topk_pos, topk_scale, shared_output=None)
tests/test_fuse_moe_cp_async.py:103
↓ 1 callersFunctionnaive_reduce
(x_bf16, topk_pos, topk_scale, shared_output=None)
tests/test_fuse_moe_blockwise.py:194
↓ 1 callersFunctionnaive_reduce
(x_bf16, topk_pos, topk_scale, shared_output=None)
tests/test_fuse_moe_pertensor.py:106
↓ 1 callersMethodnan_with_sign
3rd/cutlass/include/cutlass/exmy_base.h:597
↓ 1 callersFunctionnormalized_m
src/gemm/sm90/entry.cc:21
↓ 1 callersFunctionnratio
3rd/cutlass/include/cute/numeric/integral_ratio.hpp:105
↓ 1 callersFunctionnum_digits
3rd/cutlass/include/cute/util/print.hpp:44
↓ 1 callersFunctionnum_gemm_k_filter_r
//////////////////////////////////////////////////////////// Methods used for strided dgrad implementation ///////////////////////////////////////////
3rd/cutlass/include/cutlass/conv/conv2d_problem_size.h:304
↓ 1 callersFunctionnum_gemm_k_filter_s
Number of filter s positions to accumulate in gemm-k dim
3rd/cutlass/include/cutlass/conv/conv2d_problem_size.h:310
↓ 1 callersMethodoperand_A_ref
Returns a TensorRef to the A operand
3rd/cutlass/include/cutlass/gemm/threadblock/mma_planar_complex_base.h:158
↓ 1 callersMethodoperand_A_ref
Returns a TensorRef to the A operand
3rd/cutlass/include/cutlass/gemm/threadblock/mma_base.h:186
↓ 1 callersMethodoperand_A_ref
Returns a TensorRef to the A operand
3rd/cutlass/include/cutlass/gemm/threadblock/mma_softmax_mainloop_fusion_multistage.h:159
↓ 1 callersMethodoperand_A_ref
Returns a TensorRef to the A operand
3rd/cutlass/include/cutlass/gemm/threadblock/mma_sparse_base.h:212
↓ 1 callersMethodoperand_A_ref
Returns a TensorRef to the A operand
3rd/cutlass/include/cutlass/gemm/threadblock/mma_layernorm_mainloop_fusion_multistage.h:194
↓ 1 callersMethodoperand_A_ref
Returns a TensorRef to the A operand
3rd/cutlass/include/cutlass/conv/threadblock/implicit_gemm_fprop_fusion_multistage.h:215
↓ 1 callersMethodoperand_A_ref
Returns a TensorRef to the A operand
3rd/cutlass/include/cutlass/conv/threadblock/implicit_gemm_wgrad_fusion_multistage.h:199
↓ 1 callersMethodoperand_A_ref
Returns a TensorRef to the A operand
3rd/cutlass/include/cutlass/conv/threadblock/depthwise_mma_base.h:188
↓ 1 callersMethodoperator()
3rd/cutlass/include/cutlass/transform/collective/sm90_wgmma_transpose.hpp:163
↓ 1 callersMethodoperator()
3rd/cutlass/include/cutlass/transform/collective/sm90_wgmma_transpose.hpp:328
↓ 1 callersMethodoperator()
3rd/cutlass/include/cutlass/transform/collective/sm90_wgmma_transpose.hpp:546
↓ 1 callersMethodoperator++
3rd/cutlass/include/cutlass/transform/threadblock/regular_tile_access_iterator_pitch_linear.h:394
↓ 1 callersFunctionoperator<<
3rd/cutlass/include/cute/swizzle.hpp:317
↓ 1 callersFunctionoutput_tile_index
Returns the linearized index of the output tile corresponding to the tile with offset [L, M, K]
3rd/cutlass/include/cutlass/gemm/kernel/sm90_tile_scheduler_stream_k.hpp:596
↓ 1 callersFunctionpad_decode_inputs_to_align8
Pad decode batch/rows to a multiple of 8 (simulates CUDA-graph padding).
tests/test_rope.py:118
↓ 1 callersFunctionparse_args
()
benchmark/route_gemm/benchmark_gemm_bf16xfp32.py:368
↓ 1 callersFunctionparse_args
()
benchmark/sampler/benchmark_sampler.py:440
↓ 1 callersFunctionparse_args
()
benchmark/attention_decode/bench_attention_decode_fp8.py:476
↓ 1 callersFunctionparse_int_digits
3rd/cutlass/include/cute/numeric/integral_constant.hpp:500
↓ 1 callersFunctionparse_int_list
(text)
benchmark/route_gemm/benchmark_gemm_bf16xfp32.py:25
↓ 1 callersFunctionparse_nvtx_csv
(path: Path)
benchmark/sampler/benchmark_sampler.py:326
↓ 1 callersMethodparticipates_in_reduction
Returns true if the threadblock computes the reduction
3rd/cutlass/include/cutlass/epilogue/thread/linear_combination_drelu.h:154
↓ 1 callersFunctionpartition_accumulator_shape
Construct A Single Stage's Accumulator Shape
3rd/cutlass/include/cutlass/gemm/collective/sm100_blockscaled_mma_array_warpspecialized.hpp:627
↓ 1 callersFunctionpartition_accumulator_shape
Construct A Single Stage's Accumulator Shape
3rd/cutlass/include/cutlass/gemm/collective/sm100_mma_array_warpspecialized_planar_complex.hpp:460
↓ 1 callersFunctionpartition_accumulator_shape
Construct A Single Stage's Accumulator Shape
3rd/cutlass/include/cutlass/gemm/collective/sm100_mma_mixed_tma_cpasync_warpspecialized.hpp:355
↓ 1 callersFunctionpartition_accumulator_shape
Construct A Single Stage's Accumulator Shape
3rd/cutlass/include/cutlass/gemm/collective/sm100_blockscaled_mma_mixed_tma_cpasync_warpspecialized.hpp:473
↓ 1 callersFunctionpartition_accumulator_shape
Construct A Single Stage's Accumulator Shape
3rd/cutlass/include/cutlass/gemm/collective/sm100_mma_array_warpspecialized_rcggemm.hpp:448
↓ 1 callersFunctionpartition_accumulator_shape
Construct A Single Stage's Accumulator Shape
3rd/cutlass/include/cutlass/gemm/collective/sm100_blockscaled_mma_array_warpspecialized_rcggemm.hpp:598
↓ 1 callersFunctionpipeline_is_consumer
3rd/cutlass/include/cutlass/pipeline/sm90_pipeline.hpp:74
↓ 1 callersFunctionpipeline_is_producer
3rd/cutlass/include/cutlass/pipeline/sm90_pipeline.hpp:58
↓ 1 callersMethodpossibly_transpose_problem
3rd/cutlass/include/cutlass/gemm/kernel/grouped_problem_visitor.h:154
↓ 1 callersFunctionprint_csv
(rows)
benchmark/route_gemm/benchmark_gemm_bf16xfp32.py:294
↓ 1 callersFunctionprint_latex_copy
3rd/cutlass/include/cute/util/print_latex.hpp:338
↓ 1 callersFunctionprint_latex_mma
3rd/cutlass/include/cute/util/print_latex.hpp:196
↓ 1 callersFunctionprint_svg_mma
3rd/cutlass/include/cute/util/print_svg.hpp:81
↓ 1 callersFunctionprint_table
(rows: list[dict])
benchmark/fused_moe/benchmark_fuse_moe.py:414
↓ 1 callersFunctionprint_table
(rows: list[dict])
benchmark/sampler/benchmark_sampler.py:401
↓ 1 callersFunctionprint_tflops_table
(rows, providers)
benchmark/route_gemm/benchmark_gemm_bf16xfp32.py:281
↓ 1 callersFunctionprint_tuple
3rd/cutlass/include/cute/container/tuple.hpp:641
↓ 1 callersFunctionproducer_commit_2x1SM
3rd/cutlass/include/cutlass/pipeline/sm100_pipeline.hpp:879
↓ 1 callersFunctionprofile_one
( backend: str, roots: Roots, gpu_id: int, model: str, bs: int, tp: int, ep: int, output_dir: Path
benchmark/fused_moe/benchmark_fuse_moe.py:320
↓ 1 callersFunctionprologue
GEMM prologue. Bootstrap the global->shared memory pipeline by fetching the global fragments needed by the first kStages-1 threadblock mainloop itera
3rd/cutlass/include/cutlass/gemm/threadblock/mma_multistage.h:362
↓ 1 callersFunctionquant_paged_cache_perhead
(cache: torch.Tensor, block_size: int)
benchmark/attention_decode/bench_attention_decode_fp8.py:101
↓ 1 callersFunctionquant_paged_cache_perhead
(cache, block_size)
tests/test_attention_decode_qkpertoken_perhead_vperhead_fp8.py:39
↓ 1 callersFunctionquant_paged_cache_pertoken
(cache: torch.Tensor, block_size: int)
benchmark/attention_decode/bench_attention_decode_fp8.py:80
↓ 1 callersFunctionquant_paged_cache_pertoken
(cache, block_size)
tests/test_attention_decode_qkpertoken_perhead_vperhead_fp8.py:14
↓ 1 callersFunctionraked_product
3rd/cutlass/include/cute/layout.hpp:1750
↓ 1 callersFunctionrank
3rd/cutlass/include/cute/int_tuple.hpp:74
↓ 1 callersMethodraw
Obtains raw bits
3rd/cutlass/include/cutlass/bfloat16.h:200
↓ 1 callersMethodred
3rd/cutlass/include/cutlass/complex.h:207
↓ 1 callersMethodreduce
Reduce
3rd/cutlass/include/cutlass/block_striped.h:220
↓ 1 callersMethodref
Returns a TensorRef pointing to the first element of the tensor.
3rd/cutlass/include/cutlass/tensor_view.h:202
↓ 1 callersMethodref
Returns a TensorRef pointing to the first element of the tensor.
3rd/cutlass/include/cutlass/tensor_view_planar_complex.h:194
↓ 1 callersFunctionref_allreduce_rmsnorm
(input_list, residual, weight, eps)
benchmark/fuse_allreduce_rmsorm/benchmark_fuse_allreduce_rmsnorm.py:59
↓ 1 callersFunctionref_allreduce_rmsnorm
(input_list, residual, weight, rms_norm_eps)
tests/test_fuse_allreduce_rmsnorm_high_throughput.py:22
↓ 1 callersFunctionref_allreduce_rmsnorm
(input_list, residual, weight, rms_norm_eps)
tests/test_fuse_allreduce_rmsnorm_low_latency.py:23
↓ 1 callersFunctionref_attn_with_paged_kvcache_func
( q, k, v, kvcache, block_ids, nblocks, seqlenq, cu_seqlenq, num_seq_kvcache )
tests/test_attention_decode_bf16.py:22
↓ 1 callersFunctionref_attn_with_paged_kvcache_func
( q, k, v, kvcache, block_ids, nblocks, seqlenq, cu_seqlenq, num_seq_kvcac
tests/test_attention_decode_qpertoken_perhead_kvpertensor_fp8.py:14
↓ 1 callersFunctionref_attn_with_paged_kvcache_func
( q, k, v, kvcache, block_ids, nblocks, seqlenq, cu_seqlenq, num_seq_kvcac
tests/test_attention_decode_qkpertoken_perhead_vperhead_fp8.py:55
↓ 1 callersFunctionref_masked_act_mul_and_blockwise_quant
(gate_up, num_per_expert)
tests/test_act.py:105
↓ 1 callersFunctionref_masked_act_mul_and_quant
(gate_up, scale, num_per_expert)
tests/test_act.py:31
↓ 1 callersFunctionreference_torch_rmsnorm
(x, weight, eps)
tests/test_normalization.py:23
↓ 1 callersFunctionreport_prefix
(output_dir: Path, tag: str)
benchmark/sampler/benchmark_sampler.py:255
↓ 1 callersFunctionrequired_roots
Which source roots a backend needs.
benchmark/fused_moe/benchmark_fuse_moe.py:119
↓ 1 callersFunctionrequires_fixup
Returns whether fixup is needed for `work_tile_info`.
3rd/cutlass/include/cutlass/gemm/kernel/sm100_tile_scheduler_stream_k.hpp:592
↓ 1 callersFunctionrequires_fixup
Returns whether fixup is needed for `work_tile_info`.
3rd/cutlass/include/cutlass/gemm/kernel/sm90_tile_scheduler_stream_k.hpp:383
↓ 1 callersMethodresize
Changes the size of the view without affecting pointer or layout
3rd/cutlass/include/cutlass/tensor_view_planar_complex.h:168
↓ 1 callersFunctionresolve_roots
(args)
benchmark/fused_moe/benchmark_fuse_moe.py:104
↓ 1 callersFunctionrmsnorm
(x, w, eps)
benchmark/fuse_allreduce_rmsorm/benchmark_fuse_allreduce_rmsnorm.py:54
↓ 1 callersFunctionrmsnorm
(x, w, rms_norm_eps)
tests/test_fuse_allreduce_rmsnorm_high_throughput.py:15
↓ 1 callersFunctionrmsnorm
(x, w, rms_norm_eps)
tests/test_fuse_allreduce_rmsnorm_low_latency.py:16
↓ 1 callersFunctionrotl
3rd/cutlass/include/cute/numeric/math.hpp:205
↓ 1 callersFunctionrows_from_events
(args, cells: list[dict])
benchmark/sampler/benchmark_sampler.py:381
↓ 1 callersFunctionrows_from_nsys
(args, cells: list[dict], output_dir: Path, tag: str)
benchmark/sampler/benchmark_sampler.py:361
↓ 1 callersMethodrun
Runs the kernel using initialized state.
3rd/cutlass/include/cutlass/gemm/device/rank_2k.h:313
↓ 1 callersMethodrun
Runs the kernel using initialized state.
3rd/cutlass/include/cutlass/gemm/device/ell_gemm.h:499
↓ 1 callersMethodrun
Runs the kernel using initialized state.
3rd/cutlass/include/cutlass/gemm/device/rank_k.h:290
↓ 1 callersMethodrun
Runs the kernel using initialized state.
3rd/cutlass/include/cutlass/gemm/device/trmm.h:479
↓ 1 callersMethodrun
Runs the kernel using initialized state.
3rd/cutlass/include/cutlass/gemm/device/gemm_splitk_parallel.h:334
↓ 1 callersMethodrun
Runs the kernel using initialized state.
3rd/cutlass/include/cutlass/gemm/device/gemm_batched.h:407
← previousnext →1,301–1,400 of 10,542, ranked by callers