Skip to content

Megakernel backend for Nemotron-TTS - #68

Closed
adichatterje wants to merge 3 commits into
NVIDIA:mainfrom
adichatterje:adichatterje/megakernel/nemotron-tts
Closed

adichatterje wants to merge 3 commits into
NVIDIA:mainfrom
adichatterje:adichatterje/megakernel/nemotron-tts

Conversation

@adichatterje

Copy link
Copy Markdown

Adds Nemotron-TTS megakernel backend optimized for H100 for TTFA p95

  • Adds three reusable CUDA megakernels for prefill, stage 0 and stage 1.
  • Supports fixed-cohort and opt-in in-flight batching with text normalization and preprocessing.
  • The native serving path is independent of any framework like GGML / TensorRT.

This allows us to support up to batch size 512 within 200 ms TTFA p95.

Build and serving instructions: https://github.com/adichatterje/NeMo-Speech.cpp/blob/adichatterje/megakernel/nemotron-tts/docs/nemotron-megakernel.md#build

Measurements: https://docs.google.com/spreadsheets/d/1U0g30y94x6r5EaXxKM9SUn_YgL-NsZ2bcdHXGfzBC1k/edit?gid=326225413#gid=326225413

Output audios: https://drive.google.com/drive/folders/18dsr7XCHzZOfurTntVqwaLEC5Y5XbHVN

Jira: https://nvidia.atlassian.net/browse/RIVA-8275

@copy-pr-bot

copy-pr-bot Bot commented Oct 15, 2026

Copy link
Copy Markdown

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

@coderabbitai

coderabbitai Bot commented Oct 15, 2026

Copy link
Copy Markdown
Contributor

Important

Review skipped

Too many files!

This PR contains 639 files, which is 339 over the limit of 300.

To get a review, reduce the PR to 300 files or fewer by splitting it into smaller PRs or changing its base branch.

⚙️ Run configuration
  • Configuration used: Repository: NVIDIA/NeMo-Speech.cpp/.coderabbit.yaml
  • Review profile: CHILL
  • Plan: Enterprise
  • Run ID: ae4fd228-55e6-4858-bc34-f2eeb389c935
📥 Commits

Reviewing files that changed from the base of the PR and between 083861c and 721c67d.

⛔ Files ignored due to path filters: 1
  • native/text_preparation/Cargo.lock is excluded by !**/*.lock
📒 Files selected for processing: 639
  • .clang-format
  • .pre-commit-config.yaml
  • README.md
  • THIRD_PARTY_NOTICES.md
  • app/bench_tts.cpp
  • app/serve.cpp
  • app/synthesize.cpp
  • deploy/nemotron_native/LICENSE
  • deploy/nemotron_native/convert_model.py
  • deploy/nemotron_native/kernels/b1024/prefill.cu
  • deploy/nemotron_native/kernels/b1024/regions/prefill/native/prefill/dedup_native_prefill_impl.inc
  • deploy/nemotron_native/kernels/b1024/regions/prefill/native/prefill/dedup_predictor_mk2.cuh
  • deploy/nemotron_native/kernels/b1024/regions/prefill/native/prefill/mk2_abi.cuh
  • deploy/nemotron_native/kernels/b1024/regions/prefill/native/prefill/native_prefill_abi.cuh
  • deploy/nemotron_native/kernels/b1024/regions/prefill/native/prefill/prefill.cu
  • deploy/nemotron_native/kernels/b1024/regions/prefill/prefill.cu
  • deploy/nemotron_native/kernels/b1024/regions/stage0/native/stage0/mk2_abi.cuh
  • deploy/nemotron_native/kernels/b1024/regions/stage0/native/stage0/native_prefill_abi.cuh
  • deploy/nemotron_native/kernels/b1024/regions/stage0/native/stage0/native_prefill_impl.inc
  • deploy/nemotron_native/kernels/b1024/regions/stage0/native/stage0/predictor_mk2.cu
  • deploy/nemotron_native/kernels/b1024/regions/stage0/stage0.cu
  • deploy/nemotron_native/kernels/b1024/stage0.cu
  • deploy/nemotron_native/kernels/b128/prefill.cu
  • deploy/nemotron_native/kernels/b128/regions/prefill/native/prefill/dedup_native_prefill_impl.inc
  • deploy/nemotron_native/kernels/b128/regions/prefill/native/prefill/dedup_predictor_mk2.cuh
  • deploy/nemotron_native/kernels/b128/regions/prefill/native/prefill/mk2_abi.cuh
  • deploy/nemotron_native/kernels/b128/regions/prefill/native/prefill/native_prefill_abi.cuh
  • deploy/nemotron_native/kernels/b128/regions/prefill/native/prefill/prefill.cu
  • deploy/nemotron_native/kernels/b128/regions/prefill/prefill.cu
  • deploy/nemotron_native/kernels/b128/regions/stage0/native/stage0/mk2_abi.cuh
  • deploy/nemotron_native/kernels/b128/regions/stage0/native/stage0/native_prefill_abi.cuh
  • deploy/nemotron_native/kernels/b128/regions/stage0/native/stage0/native_prefill_impl.inc
  • deploy/nemotron_native/kernels/b128/regions/stage0/native/stage0/predictor_mk2.cu
  • deploy/nemotron_native/kernels/b128/regions/stage0/stage0.cu
  • deploy/nemotron_native/kernels/b128/stage0.cu
  • deploy/nemotron_native/kernels/b2048/prefill.cu
  • deploy/nemotron_native/kernels/b2048/regions/prefill/native/prefill/dedup_native_prefill_impl.inc
  • deploy/nemotron_native/kernels/b2048/regions/prefill/native/prefill/dedup_predictor_mk2.cuh
  • deploy/nemotron_native/kernels/b2048/regions/prefill/native/prefill/mk2_abi.cuh
  • deploy/nemotron_native/kernels/b2048/regions/prefill/native/prefill/native_prefill_abi.cuh
  • deploy/nemotron_native/kernels/b2048/regions/prefill/native/prefill/prefill.cu
  • deploy/nemotron_native/kernels/b2048/regions/prefill/prefill.cu
  • deploy/nemotron_native/kernels/b2048/regions/stage0/native/stage0/mk2_abi.cuh
  • deploy/nemotron_native/kernels/b2048/regions/stage0/native/stage0/native_prefill_abi.cuh
  • deploy/nemotron_native/kernels/b2048/regions/stage0/native/stage0/native_prefill_impl.inc
  • deploy/nemotron_native/kernels/b2048/regions/stage0/native/stage0/predictor_mk2.cu
  • deploy/nemotron_native/kernels/b2048/regions/stage0/stage0.cu
  • deploy/nemotron_native/kernels/b2048/stage0.cu
  • deploy/nemotron_native/kernels/b256/prefill.cu
  • deploy/nemotron_native/kernels/b256/regions/prefill/native/prefill/dedup_native_prefill_impl.inc
  • deploy/nemotron_native/kernels/b256/regions/prefill/native/prefill/dedup_predictor_mk2.cuh
  • deploy/nemotron_native/kernels/b256/regions/prefill/native/prefill/mk2_abi.cuh
  • deploy/nemotron_native/kernels/b256/regions/prefill/native/prefill/native_prefill_abi.cuh
  • deploy/nemotron_native/kernels/b256/regions/prefill/native/prefill/prefill.cu
  • deploy/nemotron_native/kernels/b256/regions/prefill/prefill.cu
  • deploy/nemotron_native/kernels/b256/regions/stage0/native/stage0/mk2_abi.cuh
  • deploy/nemotron_native/kernels/b256/regions/stage0/native/stage0/native_prefill_abi.cuh
  • deploy/nemotron_native/kernels/b256/regions/stage0/native/stage0/native_prefill_impl.inc
  • deploy/nemotron_native/kernels/b256/regions/stage0/native/stage0/predictor_mk2.cu
  • deploy/nemotron_native/kernels/b256/regions/stage0/stage0.cu
  • deploy/nemotron_native/kernels/b256/stage0.cu
  • deploy/nemotron_native/kernels/b32/codec.cu
  • deploy/nemotron_native/kernels/b32/native/codec/kernel_kf_codec.h
  • deploy/nemotron_native/kernels/b32/native/codec/legacy_codec.cuh
  • deploy/nemotron_native/kernels/b32/native/codec/native_stage1.cu
  • deploy/nemotron_native/kernels/b32/native/codec/native_stage1.h
  • deploy/nemotron_native/kernels/b32/native/codec/repack3.inc
  • deploy/nemotron_native/kernels/b32/native/codec/tail_kernel.inc
  • deploy/nemotron_native/kernels/b32/native/codec/wg2_kernel.inc
  • deploy/nemotron_native/kernels/b32/native/codec/wg3_kernel.inc
  • deploy/nemotron_native/kernels/b32/native/codec/wg4_step_widths.inc
  • deploy/nemotron_native/kernels/b32/native/codec/wg5_step_widths.inc
  • deploy/nemotron_native/kernels/b32/native/codec/wg_kernel.inc
  • deploy/nemotron_native/kernels/b32/native/codec/wgmma_macro.inc
  • deploy/nemotron_native/kernels/b32/native/prefill/dedup_native_prefill_impl.inc
  • deploy/nemotron_native/kernels/b32/native/prefill/dedup_predictor_mk2.cuh
  • deploy/nemotron_native/kernels/b32/native/prefill/mk2_abi.cuh
  • deploy/nemotron_native/kernels/b32/native/prefill/native_prefill_abi.cuh
  • deploy/nemotron_native/kernels/b32/native/prefill/prefill.cu
  • deploy/nemotron_native/kernels/b32/native/stage0/mk2_abi.cuh
  • deploy/nemotron_native/kernels/b32/native/stage0/native_prefill_abi.cuh
  • deploy/nemotron_native/kernels/b32/native/stage0/native_prefill_impl.inc
  • deploy/nemotron_native/kernels/b32/native/stage0/predictor_mk2.cu
  • deploy/nemotron_native/kernels/b32/native/stage0/sparse_moe_impl.cuh
  • deploy/nemotron_native/kernels/b32/prefill.cu
  • deploy/nemotron_native/kernels/b32/stage0.cu
  • deploy/nemotron_native/kernels/b4096/codec.cu
  • deploy/nemotron_native/kernels/b4096/prefill.cu
  • deploy/nemotron_native/kernels/b4096/regions/prefill/native/prefill/dedup_native_prefill_impl.inc
  • deploy/nemotron_native/kernels/b4096/regions/prefill/native/prefill/dedup_predictor_mk2.cuh
  • deploy/nemotron_native/kernels/b4096/regions/prefill/native/prefill/mk2_abi.cuh
  • deploy/nemotron_native/kernels/b4096/regions/prefill/native/prefill/native_prefill_abi.cuh
  • deploy/nemotron_native/kernels/b4096/regions/prefill/native/prefill/prefill.cu
  • deploy/nemotron_native/kernels/b4096/regions/prefill/prefill.cu
  • deploy/nemotron_native/kernels/b4096/regions/stage0/native/stage0/mk2_abi.cuh
  • deploy/nemotron_native/kernels/b4096/regions/stage0/native/stage0/native_prefill_abi.cuh
  • deploy/nemotron_native/kernels/b4096/regions/stage0/native/stage0/native_prefill_impl.inc
  • deploy/nemotron_native/kernels/b4096/regions/stage0/native/stage0/predictor_mk2.cu
  • deploy/nemotron_native/kernels/b4096/regions/stage0/stage0.cu
  • deploy/nemotron_native/kernels/b4096/regions/stage1/codec.cu
  • deploy/nemotron_native/kernels/b4096/regions/stage1/native/codec/kernel_kf_codec.h
  • deploy/nemotron_native/kernels/b4096/regions/stage1/native/codec/legacy_codec.cuh
  • deploy/nemotron_native/kernels/b4096/regions/stage1/native/codec/native_stage1.cu
  • deploy/nemotron_native/kernels/b4096/regions/stage1/native/codec/native_stage1.h
  • deploy/nemotron_native/kernels/b4096/regions/stage1/native/codec/repack3.inc
  • deploy/nemotron_native/kernels/b4096/regions/stage1/native/codec/tail_kernel.inc
  • deploy/nemotron_native/kernels/b4096/regions/stage1/native/codec/wg2_kernel.inc
  • deploy/nemotron_native/kernels/b4096/regions/stage1/native/codec/wg3_kernel.inc
  • deploy/nemotron_native/kernels/b4096/regions/stage1/native/codec/wg4_step_widths.inc
  • deploy/nemotron_native/kernels/b4096/regions/stage1/native/codec/wg5_step_widths.inc
  • deploy/nemotron_native/kernels/b4096/regions/stage1/native/codec/wg_kernel.inc
  • deploy/nemotron_native/kernels/b4096/regions/stage1/native/codec/wgmma_macro.inc
  • deploy/nemotron_native/kernels/b4096/stage0.cu
  • deploy/nemotron_native/kernels/b512/codec.cu
  • deploy/nemotron_native/kernels/b512/native/codec/kernel_kf_codec.h
  • deploy/nemotron_native/kernels/b512/native/codec/legacy_codec.cuh
  • deploy/nemotron_native/kernels/b512/native/codec/native_stage1.cu
  • deploy/nemotron_native/kernels/b512/native/codec/native_stage1.h
  • deploy/nemotron_native/kernels/b512/native/codec/repack3.inc
  • deploy/nemotron_native/kernels/b512/native/codec/tail_kernel.inc
  • deploy/nemotron_native/kernels/b512/native/codec/wg2_kernel.inc
  • deploy/nemotron_native/kernels/b512/native/codec/wg3_kernel.inc
  • deploy/nemotron_native/kernels/b512/native/codec/wg4_step_widths.inc
  • deploy/nemotron_native/kernels/b512/native/codec/wg5_step_widths.inc
  • deploy/nemotron_native/kernels/b512/native/codec/wg_kernel.inc
  • deploy/nemotron_native/kernels/b512/native/codec/wgmma_macro.inc
  • deploy/nemotron_native/kernels/b512/native/prefill/dedup_native_prefill_impl.inc
  • deploy/nemotron_native/kernels/b512/native/prefill/dedup_predictor_mk2.cuh
  • deploy/nemotron_native/kernels/b512/native/prefill/mk2_abi.cuh
  • deploy/nemotron_native/kernels/b512/native/prefill/native_prefill_abi.cuh
  • deploy/nemotron_native/kernels/b512/native/prefill/prefill.cu
  • deploy/nemotron_native/kernels/b512/native/stage0/mk2_abi.cuh
  • deploy/nemotron_native/kernels/b512/native/stage0/native_prefill_abi.cuh
  • deploy/nemotron_native/kernels/b512/native/stage0/native_prefill_impl.inc
  • deploy/nemotron_native/kernels/b512/native/stage0/predictor_mk2.cu
  • deploy/nemotron_native/kernels/b512/prefill.cu
  • deploy/nemotron_native/kernels/b512/stage0.cu
  • deploy/nemotron_native/kernels/codec/base/kernel_kf_codec.h
  • deploy/nemotron_native/kernels/codec/base/legacy_codec.cuh
  • deploy/nemotron_native/kernels/codec/base/native_stage1.cu
  • deploy/nemotron_native/kernels/codec/base/native_stage1.h
  • deploy/nemotron_native/kernels/codec/base/repack3.inc
  • deploy/nemotron_native/kernels/codec/base/tail_kernel.inc
  • deploy/nemotron_native/kernels/codec/base/wg2_kernel.inc
  • deploy/nemotron_native/kernels/codec/base/wg3_kernel.inc
  • deploy/nemotron_native/kernels/codec/base/wg4_step_widths.inc
  • deploy/nemotron_native/kernels/codec/base/wg5_step_widths.inc
  • deploy/nemotron_native/kernels/codec/base/wg_kernel.inc
  • deploy/nemotron_native/kernels/codec/base/wgmma_macro.inc
  • deploy/nemotron_native/kernels/codec/first_packet/codec.cu
  • deploy/nemotron_native/kernels/codec/first_packet/regions/stage1/codec.cu
  • deploy/nemotron_native/kernels/codec/first_packet/regions/stage1/native/codec/kernel_kf_codec.h
  • deploy/nemotron_native/kernels/codec/first_packet/regions/stage1/native/codec/legacy_codec.cuh
  • deploy/nemotron_native/kernels/codec/first_packet/regions/stage1/native/codec/native_stage1.cu
  • deploy/nemotron_native/kernels/codec/first_packet/regions/stage1/native/codec/native_stage1.h
  • deploy/nemotron_native/kernels/codec/first_packet/regions/stage1/native/codec/repack3.inc
  • deploy/nemotron_native/kernels/codec/first_packet/regions/stage1/native/codec/tail_kernel.inc
  • deploy/nemotron_native/kernels/codec/first_packet/regions/stage1/native/codec/wg2_kernel.inc
  • deploy/nemotron_native/kernels/codec/first_packet/regions/stage1/native/codec/wg3_kernel.inc
  • deploy/nemotron_native/kernels/codec/first_packet/regions/stage1/native/codec/wg4_step_widths.inc
  • deploy/nemotron_native/kernels/codec/first_packet/regions/stage1/native/codec/wg5_step_widths.inc
  • deploy/nemotron_native/kernels/codec/first_packet/regions/stage1/native/codec/wg_kernel.inc
  • deploy/nemotron_native/kernels/codec/first_packet/regions/stage1/native/codec/wgmma_macro.inc
  • deploy/nemotron_native/kernels/codec/small/kernel_kf_codec.h
  • deploy/nemotron_native/kernels/codec/small/legacy_codec.cuh
  • deploy/nemotron_native/kernels/codec/small/native_stage1.cu
  • deploy/nemotron_native/kernels/codec/small/native_stage1.h
  • deploy/nemotron_native/kernels/codec/small/repack3.inc
  • deploy/nemotron_native/kernels/codec/small/tail_kernel.inc
  • deploy/nemotron_native/kernels/codec/small/wg2_kernel.inc
  • deploy/nemotron_native/kernels/codec/small/wg3_kernel.inc
  • deploy/nemotron_native/kernels/codec/small/wg4_step_widths.inc
  • deploy/nemotron_native/kernels/codec/small/wg5_step_widths.inc
  • deploy/nemotron_native/kernels/codec/small/wg_kernel.inc
  • deploy/nemotron_native/kernels/codec/small/wgmma_macro.inc
  • deploy/nemotron_native/kernels/contraction/bf16/mk2_abi.cuh
  • deploy/nemotron_native/kernels/contraction/bf16/native_prefill_abi.cuh
  • deploy/nemotron_native/kernels/contraction/bf16/native_prefill_impl.inc
  • deploy/nemotron_native/kernels/contraction/bf16/predictor_mk2.cu
  • deploy/nemotron_native/kernels/contraction/bf16/sparse_moe_impl.cuh
  • deploy/nemotron_native/kernels/contraction/lane_transfer.cu
  • deploy/nemotron_native/kernels/contraction/mk2_abi.cuh
  • deploy/nemotron_native/kernels/inflight/progress.cuh
  • deploy/nemotron_native/kernels/inflight/runtime.cuh
  • deploy/nemotron_native/kernels/prefill/b1/dedup_native_prefill_impl.inc
  • deploy/nemotron_native/kernels/prefill/b1/dedup_predictor_mk2.cuh
  • deploy/nemotron_native/kernels/prefill/b1/mk2_abi.cuh
  • deploy/nemotron_native/kernels/prefill/b1/native_prefill_abi.cuh
  • deploy/nemotron_native/kernels/prefill/b1/prefill.cu
  • deploy/nemotron_native/kernels/prefill/b16/dedup_native_prefill_impl.inc
  • deploy/nemotron_native/kernels/prefill/b16/dedup_predictor_mk2.cuh
  • deploy/nemotron_native/kernels/prefill/b16/mk2_abi.cuh
  • deploy/nemotron_native/kernels/prefill/b16/native_prefill_abi.cuh
  • deploy/nemotron_native/kernels/prefill/b16/prefill.cu
  • deploy/nemotron_native/kernels/prefill/b2/dedup_native_prefill_impl.inc
  • deploy/nemotron_native/kernels/prefill/b2/dedup_predictor_mk2.cuh
  • deploy/nemotron_native/kernels/prefill/b2/mk2_abi.cuh
  • deploy/nemotron_native/kernels/prefill/b2/native_prefill_abi.cuh
  • deploy/nemotron_native/kernels/prefill/b2/prefill.cu
  • deploy/nemotron_native/kernels/prefill/b4/dedup_native_prefill_impl.inc
  • deploy/nemotron_native/kernels/prefill/b4/dedup_predictor_mk2.cuh
  • deploy/nemotron_native/kernels/prefill/b4/mk2_abi.cuh
  • deploy/nemotron_native/kernels/prefill/b4/native_prefill_abi.cuh
  • deploy/nemotron_native/kernels/prefill/b4/prefill.cu
  • deploy/nemotron_native/kernels/prefill/b64/dedup_native_prefill_impl.inc
  • deploy/nemotron_native/kernels/prefill/b64/dedup_predictor_mk2.cuh
  • deploy/nemotron_native/kernels/prefill/b64/mk2_abi.cuh
  • deploy/nemotron_native/kernels/prefill/b64/native_prefill_abi.cuh
  • deploy/nemotron_native/kernels/prefill/b64/prefill.cu
  • deploy/nemotron_native/kernels/prefill/b8/dedup_native_prefill_impl.inc
  • deploy/nemotron_native/kernels/prefill/b8/dedup_predictor_mk2.cuh
  • deploy/nemotron_native/kernels/prefill/b8/mk2_abi.cuh
  • deploy/nemotron_native/kernels/prefill/b8/native_prefill_abi.cuh
  • deploy/nemotron_native/kernels/prefill/b8/prefill.cu
  • deploy/nemotron_native/kernels/recipes.json
  • deploy/nemotron_native/kernels/stage0/b1/bb_router_selection.cuh
  • deploy/nemotron_native/kernels/stage0/b1/mk2_abi.cuh
  • deploy/nemotron_native/kernels/stage0/b1/native_prefill_abi.cuh
  • deploy/nemotron_native/kernels/stage0/b1/native_prefill_impl.inc
  • deploy/nemotron_native/kernels/stage0/b1/predictor_mk2.cu
  • deploy/nemotron_native/kernels/stage0/b16/mk2_abi.cuh
  • deploy/nemotron_native/kernels/stage0/b16/native_prefill_abi.cuh
  • deploy/nemotron_native/kernels/stage0/b16/native_prefill_impl.inc
  • deploy/nemotron_native/kernels/stage0/b16/predictor_mk2.cu
  • deploy/nemotron_native/kernels/stage0/b16/sparse_moe_impl.cuh
  • deploy/nemotron_native/kernels/stage0/b64/mk2_abi.cuh
  • deploy/nemotron_native/kernels/stage0/b64/native_prefill_abi.cuh
  • deploy/nemotron_native/kernels/stage0/b64/native_prefill_impl.inc
  • deploy/nemotron_native/kernels/stage0/b64/predictor_mk2.cu
  • deploy/nemotron_native/kernels/stage0/small/mk2_abi.cuh
  • deploy/nemotron_native/kernels/stage0/small/native_prefill_abi.cuh
  • deploy/nemotron_native/kernels/stage0/small/native_prefill_impl.inc
  • deploy/nemotron_native/kernels/stage0/small/predictor_mk2.cu
  • deploy/nemotron_native/kernels/stage0/small/sparse_moe_impl.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/LICENSE
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/common/base_ops.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/common/base_types.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/common/common.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/common/multimem.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/common/util.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/kittens.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/group.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/memory/memory.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/memory/tile/complex_global_to_register.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/memory/tile/complex_global_to_shared.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/memory/tile/complex_shared_to_register.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/memory/tile/global_to_register.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/memory/tile/global_to_shared.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/memory/tile/parallel_global_to_global.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/memory/tile/shared_to_register.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/memory/tile/tensor_to_register.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/memory/tile/tile.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/memory/tile/tma.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/memory/tile/tma_cluster.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/memory/vec/global_to_register.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/memory/vec/global_to_shared.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/memory/vec/shared_to_register.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/memory/vec/tma.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/memory/vec/tma_cluster.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/memory/vec/vec.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/mma/base/64x112.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/mma/base/64x128.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/mma/base/64x144.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/mma/base/64x16.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/mma/base/64x160.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/mma/base/64x176.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/mma/base/64x192.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/mma/base/64x208.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/mma/base/64x224.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/mma/base/64x240.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/mma/base/64x256.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/mma/base/64x32.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/mma/base/64x48.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/mma/base/64x64.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/mma/base/64x80.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/mma/base/64x96.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/mma/base/base.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/mma/mma.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/mma/tcgen05.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/mma/warp.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/mma/warpgroup.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/register/register.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/register/tile/complex_conversions.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/register/tile/complex_maps.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/register/tile/conversions.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/register/tile/maps.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/register/tile/reductions.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/register/tile/tile.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/register/vec/conversions.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/register/vec/maps.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/register/vec/reductions.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/register/vec/vec.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/shared/shared.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/shared/tile/conversions.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/shared/tile/maps.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/shared/tile/reductions.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/shared/tile/tile.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/shared/vec/conversions.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/shared/vec/maps.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/shared/vec/reductions.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/shared/vec/vec.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/util/sync.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/util/tma.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/util/tma_cluster.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/group/util/util.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/ops.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/thread/memory/memory.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/thread/memory/tile/shared_to_tensor.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/thread/memory/tile/tile.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/thread/memory/tile/tma.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/thread/memory/vec/tma.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/thread/memory/vec/vec.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/thread/mma/mma.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/thread/mma/tcgen05.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/thread/thread.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/thread/util/sync.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/thread/util/tma.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/ops/thread/util/util.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/pyutils/broker.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/pyutils/club.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/pyutils/parallel_tensor.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/pyutils/profiler.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/pyutils/pyutils.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/pyutils/torchutils.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/types/global/cgl.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/types/global/gl.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/types/global/global.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/types/global/tma.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/types/global/util.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/types/register/crt.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/types/register/crv.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/types/register/register.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/types/register/rt.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/types/register/rt_base.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/types/register/rt_layout.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/types/register/rv.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/types/register/rv_layout.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/types/shared/cst.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/types/shared/csv.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/types/shared/descriptor.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/types/shared/shared.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/types/shared/st.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/types/shared/sv.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/types/system/ipc.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/types/system/pgl.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/types/system/system.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/types/system/vmm.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/types/tensor/tensor.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/types/tensor/tt.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens-v4-prefill/include/types/types.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/LICENSE
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/common/base_ops.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/common/base_types.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/common/common.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/common/multimem.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/common/util.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/kittens.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/group.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/memory/memory.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/memory/tile/complex_global_to_register.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/memory/tile/complex_global_to_shared.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/memory/tile/complex_shared_to_register.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/memory/tile/global_to_register.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/memory/tile/global_to_shared.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/memory/tile/parallel_global_to_global.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/memory/tile/shared_to_register.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/memory/tile/tensor_to_register.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/memory/tile/tile.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/memory/tile/tma.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/memory/tile/tma_cluster.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/memory/vec/global_to_register.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/memory/vec/global_to_shared.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/memory/vec/shared_to_register.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/memory/vec/tma.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/memory/vec/tma_cluster.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/memory/vec/vec.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/mma/base/64x112.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/mma/base/64x128.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/mma/base/64x144.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/mma/base/64x16.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/mma/base/64x160.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/mma/base/64x176.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/mma/base/64x192.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/mma/base/64x208.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/mma/base/64x224.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/mma/base/64x240.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/mma/base/64x256.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/mma/base/64x32.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/mma/base/64x48.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/mma/base/64x64.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/mma/base/64x80.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/mma/base/64x96.impl
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/mma/base/base.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/mma/mma.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/mma/tcgen05.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/mma/warp.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/mma/warpgroup.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/register/register.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/register/tile/complex_conversions.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/register/tile/complex_maps.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/register/tile/conversions.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/register/tile/maps.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/register/tile/reductions.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/register/tile/tile.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/register/vec/conversions.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/register/vec/maps.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/register/vec/reductions.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/register/vec/vec.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/shared/shared.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/shared/tile/conversions.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/shared/tile/maps.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/shared/tile/reductions.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/shared/tile/tile.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/shared/vec/conversions.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/shared/vec/maps.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/shared/vec/reductions.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/shared/vec/vec.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/util/sync.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/util/tma.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/util/tma_cluster.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/group/util/util.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/ops.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/thread/memory/memory.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/thread/memory/tile/shared_to_tensor.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/thread/memory/tile/tile.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/thread/memory/tile/tma.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/thread/memory/vec/tma.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/thread/memory/vec/vec.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/thread/mma/mma.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/thread/mma/tcgen05.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/thread/thread.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/thread/util/sync.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/thread/util/tma.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/ops/thread/util/util.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/pyutils/broker.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/pyutils/club.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/pyutils/parallel_tensor.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/pyutils/profiler.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/pyutils/pyutils.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/pyutils/torchutils.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/types/global/cgl.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/types/global/gl.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/types/global/global.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/types/global/tma.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/types/global/util.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/types/register/crt.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/types/register/crv.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/types/register/register.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/types/register/rt.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/types/register/rt_base.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/types/register/rt_layout.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/types/register/rv.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/types/register/rv_layout.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/types/shared/cst.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/types/shared/csv.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/types/shared/descriptor.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/types/shared/shared.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/types/shared/st.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/types/shared/sv.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/types/system/ipc.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/types/system/pgl.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/types/system/system.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/types/system/vmm.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/types/tensor/tensor.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/types/tensor/tt.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/include/types/types.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/prototype/common/common.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/prototype/common/templates.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/prototype/common/util.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/prototype/interpreter/interpreter.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/prototype/interpreter/templates.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/prototype/lcf/lcf.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/prototype/lcf/templates.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/prototype/lcsc/lcsc.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/prototype/lcsf/lcsf.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/prototype/lcsf/templates.cuh
  • deploy/nemotron_native/kernels/third_party/thunderkittens/prototype/prototype.cuh
  • deploy/nemotron_native/model.lock.json
  • deploy/nemotron_native/profiles/playback-safe.json
  • deploy/nemotron_native/programs.json
  • deploy/nemotron_native/serve.py
  • docker/Dockerfile.nemotron
  • docs/nemotron-integration-results.md
  • docs/nemotron-megakernel.md
  • include/nemo_speech/tts.h
  • native/text_preparation/Cargo.toml
  • native/text_preparation/LICENSE.imported
  • native/text_preparation/NOTICE
  • native/text_preparation/crates/media-schema/Cargo.toml
  • native/text_preparation/crates/media-schema/src/lib.rs
  • native/text_preparation/crates/processor-abi/Cargo.toml
  • native/text_preparation/crates/processor-abi/src/lib.rs
  • native/text_preparation/crates/processor-abi/src/raw.rs
  • native/text_preparation/src/lib.rs
  • native/text_preparation/tests/qualification.py
  • native/text_preparation/text_preparation.h
  • native/tts-normalizer/.bazelrc
  • native/tts-normalizer/BUILD.bazel
  • native/tts-normalizer/LICENSE.legacy
  • native/tts-normalizer/LICENSE.riva
  • native/tts-normalizer/README.md
  • native/tts-normalizer/THIRD_PARTY_NOTICES.md
  • native/tts-normalizer/WORKSPACE.bazel
  • native/tts-normalizer/include/processor_abi.h
  • native/tts-normalizer/src/normalize.cc
  • native/tts-normalizer/src/normalize.h
  • native/tts-normalizer/src/speech_hints.cc
  • native/tts-normalizer/src/speech_hints.h
  • native/tts-normalizer/src/text_preprocessor.cc
  • native/tts-normalizer/src/text_preprocessor.h
  • native/tts-normalizer/src/text_preprocessor_c.cc
  • native/tts-normalizer/testdata/BUILD.bazel
  • native/tts-normalizer/testdata/en-US/abbr.txt
  • native/tts-normalizer/testdata/en-US/post_process.far
  • native/tts-normalizer/testdata/en-US/postprocessor.ascii_proto
  • native/tts-normalizer/testdata/en-US/sentence_boundary_exceptions.txt
  • native/tts-normalizer/testdata/en-US/sparrowhawk_configuration.ascii_proto
  • native/tts-normalizer/testdata/en-US/tokenize_and_classify.far
  • native/tts-normalizer/testdata/en-US/tokenizer.ascii_proto
  • native/tts-normalizer/testdata/en-US/verbalize.far
  • native/tts-normalizer/testdata/en-US/verbalizer.ascii_proto
  • native/tts-normalizer/tests/normalize_test.cc
  • native/tts-normalizer/tests/speech_hints_test.cc
  • native/tts-normalizer/tests/string_matcher_test.cc
  • native/tts-normalizer/tests/text_preprocessor_c_test.cc
  • native/tts-normalizer/tests/text_preprocessor_test.cc
  • native/tts-normalizer/third_party/BUILD.bazel
  • native/tts-normalizer/third_party/BUILD.openfst
  • native/tts-normalizer/third_party/BUILD.pugixml
  • native/tts-normalizer/third_party/BUILD.re2
  • native/tts-normalizer/third_party/BUILD.sparrowhawk_0_1
  • native/tts-normalizer/third_party/BUILD.thrax_1_2_7
  • native/tts-normalizer/third_party/sparrowhawk_normalizer.cc
  • native/tts-normalizer/third_party/sparrowhawk_rule_system.cc
  • native/tts-normalizer/third_party/sparrowhawk_string_matcher.h
  • python/nemotron_cpp_bridge/__init__.py
  • python/nemotron_cpp_bridge/__main__.py
  • python/nemotron_cpp_bridge/native_preparation.py
  • python/nemotron_cpp_bridge/normalization.py
  • python/nemotron_cpp_bridge/protocol.py
  • python/nemotron_cpp_bridge/server.py
  • python/nemotron_cpp_bridge/slabs.py
  • python/nemotron_native_backend/__init__.py
  • python/nemotron_native_backend/backend.py
  • python/nemotron_native_backend/batcher.py
  • python/nemotron_native_backend/bundle.py
  • python/nemotron_native_backend/execution/__init__.py
  • python/nemotron_native_backend/execution/chunk_boundary.py
  • python/nemotron_native_backend/execution/chunk_submission.py
  • python/nemotron_native_backend/execution/composition.py
  • python/nemotron_native_backend/execution/controller_rules.py
  • python/nemotron_native_backend/execution/engine.py
  • python/nemotron_native_backend/execution/helper_graphs.py
  • python/nemotron_native_backend/execution/inflight_engine.py
  • python/nemotron_native_backend/execution/inflight_policy.py
  • python/nemotron_native_backend/execution/inflight_reset.py
  • python/nemotron_native_backend/execution/lane_control.py
  • python/nemotron_native_backend/execution/lane_windows.py
  • python/nemotron_native_backend/execution/narrow_prefix_cache.py
  • python/nemotron_native_backend/execution/narrow_state_binding.py
  • python/nemotron_native_backend/execution/native_assets.py
  • python/nemotron_native_backend/execution/native_inputs.py
  • python/nemotron_native_backend/execution/native_stage1.py
  • python/nemotron_native_backend/execution/pcm_transfer.py
  • python/nemotron_native_backend/execution/prefix_cache.py
  • python/nemotron_native_backend/execution/prefix_snapshot.py
  • python/nemotron_native_backend/execution/program.py
  • python/nemotron_native_backend/execution/readiness.py
  • python/nemotron_native_backend/execution/request_noise.py
  • python/nemotron_native_backend/execution/research_program.py
  • python/nemotron_native_backend/execution/retained_narrow.py
  • python/nemotron_native_backend/execution/reusable_inputs.py
  • python/nemotron_native_backend/execution/reusable_runner.py
  • python/nemotron_native_backend/execution/small_width.py
  • python/nemotron_native_backend/execution/source_build.py
  • python/nemotron_native_backend/execution/split_native_runtime.py
  • python/nemotron_native_backend/execution/startup_binding.py
  • python/nemotron_native_backend/execution/state_binding.py
  • python/nemotron_native_backend/execution/token_staging.py
  • python/nemotron_native_backend/execution/wide_noise.py
  • python/nemotron_native_backend/hardware.py
  • python/nemotron_native_backend/host/__init__.py
  • python/nemotron_native_backend/host/cpu_pools.py
  • python/nemotron_native_backend/host/diagnostic.py
  • python/nemotron_native_backend/host/exact_admission.py
  • python/nemotron_native_backend/host/first_fence.py
  • python/nemotron_native_backend/host/ready_writer.py
  • python/nemotron_native_backend/host/serialization.py
  • python/nemotron_native_backend/host/slab_audit.py
  • python/nemotron_native_backend/host/slab_cancel_v2.py
  • python/nemotron_native_backend/host/slab_pool.py
  • python/nemotron_native_backend/host/transport.py
  • python/nemotron_native_backend/inflight_batcher.py
  • python/nemotron_native_backend/initial_publication.py
  • python/nemotron_native_backend/lifecycle.py
  • python/nemotron_native_backend/manifest.py
  • python/nemotron_native_backend/models.py
  • python/nemotron_native_backend/protocol.py
  • python/nemotron_native_backend/r31.py
  • python/nemotron_native_backend/receipts.py
  • python/nemotron_native_backend/server.py
  • python/nemotron_native_backend/streaming_batcher.py
  • python/nemotron_native_backend/streaming_contract.py
  • python/nemotron_native_backend/streaming_models.py
  • python/nemotron_native_backend/streaming_server.py
  • python/nemotron_native_backend/worker.py
  • scripts/check_license_headers.py
  • scripts/nemotron_benchmark.py
  • scripts/nemotron_native/build_kernels.py
  • scripts/nemotron_native/build_package.py
  • scripts/nemotron_native/launch.py
  • scripts/nemotron_native/requirements.txt
  • scripts/nemotron_native/runtime_dependencies.py
  • scripts/nemotron_native/tests/test_bridge.py
  • scripts/nemotron_native/tests/test_bridge_fairness.py
  • scripts/nemotron_native/tests/test_launch.py
  • scripts/nemotron_native/tests/test_launch_direct.py
  • scripts/nemotron_native/tests/test_launcher.py
  • scripts/nemotron_native/tests/test_native_preparation.py
  • server/http/CMakeLists.txt
  • server/http/http_server.cpp
  • src/server/riva_server.cc
  • src/tts/c_api.cpp
  • src/tts/magpietts/CMakeLists.txt
  • src/tts/magpietts/config.cpp
  • src/tts/magpietts/config.h
  • src/tts/nemotron/native_runtime.cpp
  • src/tts/nemotron/native_runtime.h
  • src/tts/nemotron/native_worker_protocol.cpp
  • src/tts/nemotron/native_worker_protocol.h
  • src/tts/nemotron/pcm_conversion.h
  • src/tts/synthesizer.cpp
  • src/tts/synthesizer.h
  • tests/cpp/CMakeLists.txt
  • tests/cpp/common/test_nemotron_pcm.cpp
  • tests/cpp/common/test_nemotron_worker_protocol.cpp
  • tests/nemotron_http_test.py
  • tests/nemotron_mock_worker.py
  • tests/nemotron_normalization_test.py

You can disable this status message by setting the reviews.review_status to false in the CodeRabbit configuration file.

Use the checkbox below for a quick retry:

  • 🔍 Trigger review
  • Autofix · Keep fixing CodeRabbit findings and required CI, and resolving merge conflicts

Comment @coderabbitai help to get the list of available commands.

@greptile-apps

greptile-apps Bot commented Oct 15, 2026 •

Copy link
Copy Markdown

RetriggerView in GreptileConfidence Score: 3/5

[High impact] Adds Nemotron native TTS backend with CUDA kernels and model conversion.

Fix C API creation and the dropped CLI runtime settings before merging.

Findings

  1. P1 Nemotron requires unused Magpie paths ▶
  2. P1 Configured seeds are ignored ▶

Summary

Adds an optional H100 Nemotron-TTS backend with CUDA kernels, native text preparation, Python batching, and C++ service integration.

  • Nemotron-TTS runs through CUDA kernels built for different batch widths.
  • The controller can batch Nemotron-TTS requests as they arrive.
  • C++ serving tools can use Nemotron-TTS and stream its audio.
  • Nemotron requests can be normalized and turned into model tokens before inference.
  • Operators can build and launch the Nemotron native backend from its deployment tools.

Diagram

%%{init: {'theme': 'neutral'}}%%
flowchart LR
    Client[CLI / C API / HTTP / gRPC] --> Synth[Synthesizer]
    Synth --> Magpie[Existing Magpie runtime]
    Synth --> Native[Native runtime and text preparation]
    Native --> Socket[Unix worker connection]
    Socket --> Batch[Fixed cohort or in-flight batcher]
    Batch --> GPU[CUDA prefill, stage 0 and codec]
    GPU --> PCM[Streamed PCM]
    PCM --> Client
Loading

Reviews (1) · Last reviewed commit: 721c67d · Reviewed by Greptile

Comment thread src/tts/c_api.cpp
Comment on lines +220 to +223
if (HAS_FIELD(model, nemo_speech_tts_model_config, model_backend) && model->model_backend)
cfg.backend = model->model_backend;
if (HAS_FIELD(model, nemo_speech_tts_model_config, native_socket))
cfg.native_socket = str_or_empty(model->native_socket);

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

P1 Nemotron requires unused Magpie paths

A C caller that sets model_backend="nemotron" and native_socket, but leaves magpie_model and codec_model empty, still fails in nemo_speech_tts_create before connecting to the worker. Creation unconditionally requires both Magpie paths even though Nemotron does not use them. Apply those checks only to the Magpie backend.

Knowledge Base Used: Text-to-speech synthesis

Comment thread app/synthesize.cpp
Comment on lines +65 to +72
if (parsed.backend == "nemotron") {
nemo_speech::tts::SynthesizerConfig config;
config.backend = parsed.backend;
config.native_socket = parsed.native_socket;
config.native_timeout_seconds = parsed.native_timeout_seconds;
config.default_language_code = parsed.default_language_code;
config.default_voice_name = parsed.default_voice_name;
return config;

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

P1 Configured seeds are ignored

The Nemotron branch drops parsed.runtime. For example, synthesize "Hello" --tts-backend nemotron --native-socket /tmp/nemotron.sock --tts.seed 42 accepts the setting but sends seed 0, because the new configuration keeps the default seed. Benchmarks reuse this helper and can report a seed that the worker did not use. Copy parsed.runtime into config.runtime before returning, as the serving path already does.

Suggested change
if (parsed.backend == "nemotron") {
nemo_speech::tts::SynthesizerConfig config;
config.backend = parsed.backend;
config.native_socket = parsed.native_socket;
config.native_timeout_seconds = parsed.native_timeout_seconds;
config.default_language_code = parsed.default_language_code;
config.default_voice_name = parsed.default_voice_name;
return config;
if (parsed.backend == "nemotron") {
nemo_speech::tts::SynthesizerConfig config;
config.backend = parsed.backend;
config.runtime = parsed.runtime;
config.native_socket = parsed.native_socket;
config.native_timeout_seconds = parsed.native_timeout_seconds;
config.default_language_code = parsed.default_language_code;
config.default_voice_name = parsed.default_voice_name;
return config;

@adichatterje
adichatterje deleted the adichatterje/megakernel/nemotron-tts branch October 15, 2026 14:58
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant