Skip to content

Relicense to Apache-2.0 (dual-license; external-contributor files remain MIT) - #408

Merged
Anerudhan merged 1 commit into
NVIDIA:developfrom
Anerudhan:relicense-apache-2.0
Jul 30, 2026
Merged

Relicense to Apache-2.0 (dual-license; external-contributor files remain MIT)#408
Anerudhan merged 1 commit into
NVIDIA:developfrom
Anerudhan:relicense-apache-2.0

Conversation

@Anerudhan

@Anerudhan Anerudhan commented Jul 17, 2026

Copy link
Copy Markdown
Collaborator

Summary

Relicenses NVIDIA-authored code in cudnn-frontend from MIT to the Apache License 2.0, using the exact NVIDIA OSS SPDX header. Files that still carry contributions from external contributors whose permission has not been established remain under MIT. Rebased onto current develop (b950af1).

License assignment

Every source file carries an SPDX SPDX-License-Identifier: tag.

License Files Which
Apache-2.0 541 NVIDIA-authored, plus files cleared by consent or employment
MIT (consent pending) 50 Surviving lines from external contributors who have not yet consented
MIT (third-party-derived) 29 FlashAttention / QuACK-derived files carrying external authors' copyright

Per-file mapping — including the commit that introduced each surviving external line — is in LICENSING.md.

How the split is computed

Blame-based, not touch-based. A file stays MIT only if a not-yet-cleared external contributor's lines survive in current develop. Files whose external lines were fully overwritten by NVIDIA, or whose contributors are cleared, go to Apache-2.0.

Licensing files

  • LICENSE.txt (Apache-2.0), LICENSE-MIT.txt (MIT subset)
  • LICENSING.md — manifest: every MIT file, its pending contributor, and the introducing commit; plus the consent/employment roster
  • THIRD_PARTY_LICENSES.txt — nlohmann/json, FlashAttention, CUTLASS, Megatron-LM, QuACK, labml.ai, dlpack, pybind11, Catch2
  • NOTICE, pyproject.toml (license = "Apache-2.0 AND MIT"), README.md

Verification

  • Every source file has exactly one SPDX tag; none unlabeled
  • All Python compiles (py_compile); git diff -U0 confirms every source change is comment/header-only — compile-neutral by construction

🤖 Generated with Claude Code

@coderabbitai

coderabbitai Bot commented Jul 17, 2026

Copy link
Copy Markdown
Contributor

Important

Review skipped

Too many files!

This PR contains 621 files, which is 321 over the limit of 300.

To get a review, narrow the scope:
• coderabbit review --committed # exclude uncommitted changes
• coderabbit review --dir # limit to a subdirectory
• coderabbit review --base # compare against a closer base

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: f0e9742c-6cd9-409c-ad38-0626dc4dda39

📥 Commits

Reviewing files that changed from the base of the PR and between b950af1 and 0c08141.

⛔ Files ignored due to path filters (6)
  • include/cudnn_frontend/generated/rms_norm_silu/sm100/ln_fwd_silu_kernel.h is excluded by !**/generated/**
  • include/cudnn_frontend/generated/rms_norm_silu/sm100/ln_headers.h is excluded by !**/generated/**
  • include/cudnn_frontend/generated/sdpa/sm100/prefill/full_seqlens/d128_fprop_kernel.h is excluded by !**/generated/**
  • include/cudnn_frontend/generated/sdpa/sm100/prefill/full_seqlens/d64_fprop_kernel.h is excluded by !**/generated/**
  • include/cudnn_frontend/generated/sdpa/sm90/prefill/full_seqlens/d128_fprop_kernel.h is excluded by !**/generated/**
  • include/cudnn_frontend/generated/sdpa/sm90/prefill/full_seqlens/d64_fprop_kernel.h is excluded by !**/generated/**
📒 Files selected for processing (621)
  • CMakeLists.txt
  • LICENSE-MIT.txt
  • LICENSE.txt
  • LICENSING.md
  • NOTICE
  • README.md
  • THIRD_PARTY_LICENSES.txt
  • benchmark/bench_moe.py
  • benchmark/csa/bench_csa_compressor.py
  • benchmark/csa/gate_csa_compressor_r128.py
  • benchmark/csa/reg_probe_csa_compressor_r128.py
  • benchmark/cutedsl_fusion_kernels/cutedsl_fusion_benchmarks.py
  • benchmark/dsa/benchmark_dsa_sparse_attention_backward.py
  • benchmark/norms/Dockerfile
  • benchmark/norms/__init__.py
  • benchmark/norms/benchmark_single_norm.py
  • benchmark/norms/charts.py
  • benchmark/norms/config_types.py
  • benchmark/norms/configs/__init__.py
  • benchmark/norms/configs/all_models.py
  • benchmark/norms/configs/deepseek_v3.py
  • benchmark/norms/configs/gpt3_175b.py
  • benchmark/norms/configs/llama31_405b.py
  • benchmark/norms/configs/llama3_70b.py
  • benchmark/norms/configs/llama3_8b.py
  • benchmark/norms/configs/llama4_e16.py
  • benchmark/norms/configs/mixtral_8x22b.py
  • benchmark/norms/configs/mixtral_8x7b.py
  • benchmark/norms/configs/nemotronh_56b.py
  • benchmark/norms/configs/qwen3_235b.py
  • benchmark/norms/configs/qwen3_30b.py
  • benchmark/norms/runner.py
  • benchmark/sdpa_benchmark_training/Dockerfile
  • benchmark/sdpa_benchmark_training/__init__.py
  • benchmark/sdpa_benchmark_training/bench_ar_dit_peak.py
  • benchmark/sdpa_benchmark_training/benchmark_single_sdpa.py
  • benchmark/sdpa_benchmark_training/charts.py
  • benchmark/sdpa_benchmark_training/config_types.py
  • benchmark/sdpa_benchmark_training/configs/__init__.py
  • benchmark/sdpa_benchmark_training/configs/auto_regressive_dit.py
  • benchmark/sdpa_benchmark_training/configs/dsv3.py
  • benchmark/sdpa_benchmark_training/configs/gpt_oss.py
  • benchmark/sdpa_benchmark_training/configs/kimiK26.py
  • benchmark/sdpa_benchmark_training/configs/llama.py
  • benchmark/sdpa_benchmark_training/configs/ltx2.py
  • benchmark/sdpa_benchmark_training/configs/qwen35.py
  • benchmark/sdpa_benchmark_training/configs/wan22.py
  • benchmark/sdpa_benchmark_training/runner.py
  • cmake/cuDNN.cmake
  • cudnn_frontend-config.cmake.in
  • include/cudnn_backend_base.h
  • include/cudnn_frontend.h
  • include/cudnn_frontend/backend/backend_descriptor.h
  • include/cudnn_frontend/backend/device_properties.h
  • include/cudnn_frontend/backend/execution_helpers.h
  • include/cudnn_frontend/backend/kernel_cache.h
  • include/cudnn_frontend/backend/plan_helpers.h
  • include/cudnn_frontend/context.h
  • include/cudnn_frontend/cudnn_interface.h
  • include/cudnn_frontend/experimental/attention_utils.h
  • include/cudnn_frontend/experimental/nvrtc_shim.h
  • include/cudnn_frontend/experimental/oss_engine_interface.h
  • include/cudnn_frontend/experimental/sm100_rms_norm_silu_engine.h
  • include/cudnn_frontend/experimental/sm100_rms_norm_silu_knobs.h
  • include/cudnn_frontend/experimental/sm100_sdpa_prefill_engine.h
  • include/cudnn_frontend/experimental/sm90_sdpa_prefill_engine.h
  • include/cudnn_frontend/graph_helpers.h
  • include/cudnn_frontend/graph_interface.h
  • include/cudnn_frontend/graph_properties.h
  • include/cudnn_frontend/knobs.h
  • include/cudnn_frontend/node/adaptive_layernorm.h
  • include/cudnn_frontend/node/batchnorm.h
  • include/cudnn_frontend/node/batchnorm_inference.h
  • include/cudnn_frontend/node/block_scale_dequantize.h
  • include/cudnn_frontend/node/block_scale_quantize.h
  • include/cudnn_frontend/node/bn_finalize.h
  • include/cudnn_frontend/node/concatenate.h
  • include/cudnn_frontend/node/conv_dgrad.h
  • include/cudnn_frontend/node/conv_fprop.h
  • include/cudnn_frontend/node/conv_wgrad.h
  • include/cudnn_frontend/node/dbn.h
  • include/cudnn_frontend/node/dbn_weight.h
  • include/cudnn_frontend/node/diagonal_band_mask.h
  • include/cudnn_frontend/node/dln.h
  • include/cudnn_frontend/node/genstats.h
  • include/cudnn_frontend/node/instancenorm.h
  • include/cudnn_frontend/node/layernorm.h
  • include/cudnn_frontend/node/matmul.h
  • include/cudnn_frontend/node/matmul_fp8.h
  • include/cudnn_frontend/node/moe_grouped_matmul.h
  • include/cudnn_frontend/node/moe_grouped_matmul_bwd.h
  • include/cudnn_frontend/node/paged_cache_load.h
  • include/cudnn_frontend/node/pointwise.h
  • include/cudnn_frontend/node/reduction.h
  • include/cudnn_frontend/node/resample.h
  • include/cudnn_frontend/node/reshape.h
  • include/cudnn_frontend/node/rmsnorm.h
  • include/cudnn_frontend/node/rng.h
  • include/cudnn_frontend/node/rope.h
  • include/cudnn_frontend/node/rope_backward.h
  • include/cudnn_frontend/node/scaled_dot_product_flash_attention.h
  • include/cudnn_frontend/node/sdpa_fp8_bwd.h
  • include/cudnn_frontend/node/sdpa_support_surface.h
  • include/cudnn_frontend/node/slice.h
  • include/cudnn_frontend/node/softmax.h
  • include/cudnn_frontend/node/transpose.h
  • include/cudnn_frontend/node_interface.h
  • include/cudnn_frontend/plans.h
  • include/cudnn_frontend/utils/attn_score_modifiers.h
  • include/cudnn_frontend/utils/serialize.h
  • include/cudnn_frontend_ConvDesc.h
  • include/cudnn_frontend_Engine.h
  • include/cudnn_frontend_EngineConfig.h
  • include/cudnn_frontend_EngineConfigGenerator.h
  • include/cudnn_frontend_EngineFallbackList.h
  • include/cudnn_frontend_Errata.h
  • include/cudnn_frontend_ExecutionPlan.h
  • include/cudnn_frontend_ExecutionPlanCache.h
  • include/cudnn_frontend_Filters.h
  • include/cudnn_frontend_Heuristics.h
  • include/cudnn_frontend_Logging.h
  • include/cudnn_frontend_MatMulDesc.h
  • include/cudnn_frontend_Operation.h
  • include/cudnn_frontend_OperationGraph.h
  • include/cudnn_frontend_PointWiseDesc.h
  • include/cudnn_frontend_ReductionDesc.h
  • include/cudnn_frontend_Reorder_Tensor.h
  • include/cudnn_frontend_Resample.h
  • include/cudnn_frontend_Rng.h
  • include/cudnn_frontend_Tensor.h
  • include/cudnn_frontend_VariantPack.h
  • include/cudnn_frontend_find_plan.h
  • include/cudnn_frontend_get_plan.h
  • include/cudnn_frontend_shim.h
  • include/cudnn_frontend_utils.h
  • include/cudnn_frontend_version.h
  • pyproject.toml
  • python/CMakeLists.txt
  • python/cudnn/__init__.py
  • python/cudnn/_experimental_warnings.py
  • python/cudnn/_pygraph.py
  • python/cudnn/api_base.py
  • python/cudnn/block_sparse_attention/__init__.py
  • python/cudnn/block_sparse_attention/_interface.py
  • python/cudnn/block_sparse_attention/api.py
  • python/cudnn/block_sparse_attention/csrc/__init__.py
  • python/cudnn/block_sparse_attention/csrc/bwd/__init__.py
  • python/cudnn/block_sparse_attention/csrc/bwd/bsa_bwd_postprocess.py
  • python/cudnn/block_sparse_attention/csrc/bwd/bsa_bwd_prepost.py
  • python/cudnn/block_sparse_attention/csrc/bwd/bsa_bwd_preprocess.py
  • python/cudnn/block_sparse_attention/csrc/bwd/bucketed_k2q_csr.py
  • python/cudnn/block_sparse_attention/csrc/bwd/sm100_blk128/__init__.py
  • python/cudnn/block_sparse_attention/csrc/bwd/sm100_blk128/bsa_bwd_sm100.py
  • python/cudnn/block_sparse_attention/csrc/bwd/sm100_blk64/__init__.py
  • python/cudnn/block_sparse_attention/csrc/bwd/sm100_blk64/bsa_bwd_sm100.py
  • python/cudnn/block_sparse_attention/csrc/bwd/sm90_blk64/bsa_bwd_sm90.py
  • python/cudnn/block_sparse_attention/csrc/fwd/__init__.py
  • python/cudnn/block_sparse_attention/csrc/fwd/sm100_blk128/__init__.py
  • python/cudnn/block_sparse_attention/csrc/fwd/sm100_blk128/bsa_fwd_sm100.py
  • python/cudnn/block_sparse_attention/csrc/fwd/sm100_blk64/bsa_fwd_combine.py
  • python/cudnn/block_sparse_attention/csrc/fwd/sm100_blk64/bsa_fwd_helpers.py
  • python/cudnn/block_sparse_attention/csrc/fwd/sm100_blk64/bsa_fwd_sm100.py
  • python/cudnn/block_sparse_attention/csrc/fwd/sm120_blk64/bsa_fwd_sm120.py
  • python/cudnn/block_sparse_attention/csrc/fwd/sm90_blk64/bsa_fwd_sm90.py
  • python/cudnn/block_sparse_attention/csrc/utils/__init__.py
  • python/cudnn/block_sparse_attention/csrc/utils/batched_static_scheduler.py
  • python/cudnn/block_sparse_attention/csrc/utils/block_info.py
  • python/cudnn/block_sparse_attention/csrc/utils/block_sparse_tile_scheduler.py
  • python/cudnn/block_sparse_attention/csrc/utils/copy_utils.py
  • python/cudnn/block_sparse_attention/csrc/utils/cute_dsl_utils.py
  • python/cudnn/block_sparse_attention/csrc/utils/kernel_utils.py
  • python/cudnn/block_sparse_attention/csrc/utils/layout_utils.py
  • python/cudnn/block_sparse_attention/csrc/utils/mma_sm100_desc.py
  • python/cudnn/block_sparse_attention/csrc/utils/named_barrier.py
  • python/cudnn/block_sparse_attention/csrc/utils/pack_gqa.py
  • python/cudnn/block_sparse_attention/csrc/utils/pipeline.py
  • python/cudnn/block_sparse_attention/csrc/utils/seqlen_info.py
  • python/cudnn/block_sparse_attention/csrc/utils/sm90_utils.py
  • python/cudnn/block_sparse_attention/csrc/utils/softmax.py
  • python/cudnn/block_sparse_attention/csrc/utils/tcgen05_mma_helpers.py
  • python/cudnn/block_sparse_attention/csrc/utils/tile_scheduler.py
  • python/cudnn/collect_env.py
  • python/cudnn/csa/__init__.py
  • python/cudnn/csa/compressor/__init__.py
  • python/cudnn/csa/compressor/api.py
  • python/cudnn/csa/compressor/compressor_sm100.py
  • python/cudnn/csa/compressor/compressor_sm100_r128.py
  • python/cudnn/datatypes.py
  • python/cudnn/deepseek_sparse_attention/__init__.py
  • python/cudnn/deepseek_sparse_attention/indexer_backward/__init__.py
  • python/cudnn/deepseek_sparse_attention/indexer_backward/api.py
  • python/cudnn/deepseek_sparse_attention/indexer_backward/dense_indexer_backward_sm100.py
  • python/cudnn/deepseek_sparse_attention/indexer_backward/dense_indexer_backward_sm90.py
  • python/cudnn/deepseek_sparse_attention/indexer_backward/indexer_backward_sm100.py
  • python/cudnn/deepseek_sparse_attention/indexer_backward/indexer_backward_sm90.py
  • python/cudnn/deepseek_sparse_attention/indexer_forward/__init__.py
  • python/cudnn/deepseek_sparse_attention/indexer_forward/_interface.py
  • python/cudnn/deepseek_sparse_attention/indexer_forward/_interface_sm90.py
  • python/cudnn/deepseek_sparse_attention/indexer_forward/api.py
  • python/cudnn/deepseek_sparse_attention/indexer_forward/indexer_fwd_sm100.py
  • python/cudnn/deepseek_sparse_attention/indexer_forward/indexer_fwd_sm90.py
  • python/cudnn/deepseek_sparse_attention/indexer_top_k/__init__.py
  • python/cudnn/deepseek_sparse_attention/indexer_top_k/api.py
  • python/cudnn/deepseek_sparse_attention/indexer_top_k/block_scan.py
  • python/cudnn/deepseek_sparse_attention/indexer_top_k/compactify.py
  • python/cudnn/deepseek_sparse_attention/indexer_top_k/indexer_top_k_decode_varlen.py
  • python/cudnn/deepseek_sparse_attention/indexer_top_k/indexer_top_k_varlen_util.py
  • python/cudnn/deepseek_sparse_attention/indexer_top_k/local_to_global_dsl.py
  • python/cudnn/deepseek_sparse_attention/score_recompute/__init__.py
  • python/cudnn/deepseek_sparse_attention/score_recompute/_interface_sm100.py
  • python/cudnn/deepseek_sparse_attention/score_recompute/_interface_sm90.py
  • python/cudnn/deepseek_sparse_attention/score_recompute/api.py
  • python/cudnn/deepseek_sparse_attention/score_recompute/dense_score_recompute_sm100.py
  • python/cudnn/deepseek_sparse_attention/score_recompute/dense_score_recompute_sm90.py
  • python/cudnn/deepseek_sparse_attention/score_recompute/pack_gqa.py
  • python/cudnn/deepseek_sparse_attention/score_recompute/sparse_score_recompute_sm100.py
  • python/cudnn/deepseek_sparse_attention/score_recompute/sparse_score_recompute_sm90.py
  • python/cudnn/deepseek_sparse_attention/sparse_attention_backward/__init__.py
  • python/cudnn/deepseek_sparse_attention/sparse_attention_backward/_interface_sm100.py
  • python/cudnn/deepseek_sparse_attention/sparse_attention_backward/_interface_sm90.py
  • python/cudnn/deepseek_sparse_attention/sparse_attention_backward/api.py
  • python/cudnn/deepseek_sparse_attention/sparse_attention_backward/dsa_bwd_sm100.py
  • python/cudnn/deepseek_sparse_attention/sparse_attention_backward/dsa_bwd_sm90.py
  • python/cudnn/deepseek_sparse_attention/utils/__init__.py
  • python/cudnn/deepseek_sparse_attention/utils/compiler.py
  • python/cudnn/deepseek_sparse_attention/utils/copy.py
  • python/cudnn/deepseek_sparse_attention/utils/runtime.py
  • python/cudnn/deepseek_sparse_attention/utils/seqlen.py
  • python/cudnn/deepseek_sparse_attention/utils/sm100/__init__.py
  • python/cudnn/deepseek_sparse_attention/utils/sm100/gemm.py
  • python/cudnn/deepseek_sparse_attention/utils/sm100/mma_desc.py
  • python/cudnn/deepseek_sparse_attention/utils/sm90/__init__.py
  • python/cudnn/deepseek_sparse_attention/utils/sm90/bwd_barriers.py
  • python/cudnn/deepseek_sparse_attention/utils/sm90/bwd_tile_scheduler.py
  • python/cudnn/deepseek_sparse_attention/utils/sm90/mma.py
  • python/cudnn/deepseek_sparse_attention/utils/sm90/primitives.py
  • python/cudnn/deepseek_sparse_attention/utils/tensor_conversion.py
  • python/cudnn/discrete_grouped_gemm/__init__.py
  • python/cudnn/discrete_grouped_gemm/discrete_grouped_gemm_dswiglu/__init__.py
  • python/cudnn/discrete_grouped_gemm/discrete_grouped_gemm_dswiglu/api.py
  • python/cudnn/discrete_grouped_gemm/discrete_grouped_gemm_dswiglu/discrete_B_blockscaled_grouped_gemm_dglu_dbias.py
  • python/cudnn/discrete_grouped_gemm/discrete_grouped_gemm_swiglu/__init__.py
  • python/cudnn/discrete_grouped_gemm/discrete_grouped_gemm_swiglu/api.py
  • python/cudnn/discrete_grouped_gemm/discrete_grouped_gemm_swiglu/discrete_B_blockscaled_grouped_gemm_glu_bias.py
  • python/cudnn/discrete_grouped_gemm/discrete_kernel_utils.py
  • python/cudnn/discrete_grouped_gemm/moe_persistent_scheduler.py
  • python/cudnn/discrete_grouped_gemm/moe_sched_extension.py
  • python/cudnn/discrete_grouped_gemm/moe_utils.py
  • python/cudnn/engines/__init__.py
  • python/cudnn/engines/base.py
  • python/cudnn/engines/engine_ids.py
  • python/cudnn/engines/reference_matmul_engine.py
  • python/cudnn/engines/router.py
  • python/cudnn/experimental/__init__.py
  • python/cudnn/experimental/ops/__init__.py
  • python/cudnn/experimental/ops/moe_grouped_matmul.py
  • python/cudnn/experimental/ops/sdpa.py
  • python/cudnn/gemm_amax/__init__.py
  • python/cudnn/gemm_amax/api.py
  • python/cudnn/gemm_amax/dense_blockscaled_gemm_persistent_amax.py
  • python/cudnn/gemm_dsrelu/__init__.py
  • python/cudnn/gemm_dsrelu/api.py
  • python/cudnn/gemm_dsrelu/dense_blockscaled_gemm_persistent_dsrelu_quant.py
  • python/cudnn/gemm_proj_rope_mxfp8/__init__.py
  • python/cudnn/gemm_proj_rope_mxfp8/api.py
  • python/cudnn/gemm_proj_rope_mxfp8/gemm_proj_rope_mxfp8_bf16in.py
  • python/cudnn/gemm_proj_rope_mxfp8/gemm_proj_rope_mxfp8_mxfp8in.py
  • python/cudnn/gemm_srelu/__init__.py
  • python/cudnn/gemm_srelu/api.py
  • python/cudnn/gemm_srelu/dense_blockscaled_gemm_persistent_srelu_quant.py
  • python/cudnn/gemm_swiglu/__init__.py
  • python/cudnn/gemm_swiglu/api.py
  • python/cudnn/gemm_swiglu/dense_blockscaled_gemm_persistent_swiglu_interleaved_quant.py
  • python/cudnn/gemm_swiglu/dense_gemm_persistent_swiglu.py
  • python/cudnn/graph.py
  • python/cudnn/graph_types.py
  • python/cudnn/grouped_gemm/__init__.py
  • python/cudnn/grouped_gemm/grouped_gemm_dglu/__init__.py
  • python/cudnn/grouped_gemm/grouped_gemm_dglu/_bf16_api.py
  • python/cudnn/grouped_gemm/grouped_gemm_dglu/_blockscaled_api.py
  • python/cudnn/grouped_gemm/grouped_gemm_dglu/api.py
  • python/cudnn/grouped_gemm/grouped_gemm_dglu/moe_blockscaled_grouped_gemm_dglu_dbias.py
  • python/cudnn/grouped_gemm/grouped_gemm_dglu/moe_blockscaled_grouped_gemm_dglu_rubin.py
  • python/cudnn/grouped_gemm/grouped_gemm_dglu/moe_grouped_gemm_dglu_dbias.py
  • python/cudnn/grouped_gemm/grouped_gemm_dsrelu/__init__.py
  • python/cudnn/grouped_gemm/grouped_gemm_dsrelu/api.py
  • python/cudnn/grouped_gemm/grouped_gemm_dsrelu/moe_blockscaled_grouped_gemm_dsrelu_quant.py
  • python/cudnn/grouped_gemm/grouped_gemm_dswiglu/__init__.py
  • python/cudnn/grouped_gemm/grouped_gemm_dswiglu/api.py
  • python/cudnn/grouped_gemm/grouped_gemm_dswiglu/grouped_gemm_dswiglu_quant.py
  • python/cudnn/grouped_gemm/grouped_gemm_glu/__init__.py
  • python/cudnn/grouped_gemm/grouped_gemm_glu/_bf16_api.py
  • python/cudnn/grouped_gemm/grouped_gemm_glu/_blockscaled_api.py
  • python/cudnn/grouped_gemm/grouped_gemm_glu/api.py
  • python/cudnn/grouped_gemm/grouped_gemm_glu/moe_blockscaled_grouped_gemm_glu_bias.py
  • python/cudnn/grouped_gemm/grouped_gemm_glu/moe_blockscaled_grouped_gemm_glu_rubin.py
  • python/cudnn/grouped_gemm/grouped_gemm_glu/moe_grouped_gemm_glu_bias.py
  • python/cudnn/grouped_gemm/grouped_gemm_glu_hadamard/__init__.py
  • python/cudnn/grouped_gemm/grouped_gemm_glu_hadamard/api.py
  • python/cudnn/grouped_gemm/grouped_gemm_glu_hadamard/hadamard_utils.py
  • python/cudnn/grouped_gemm/grouped_gemm_glu_hadamard/moe_blockscaled_grouped_gemm_glu_hadamard.py
  • python/cudnn/grouped_gemm/grouped_gemm_quant/__init__.py
  • python/cudnn/grouped_gemm/grouped_gemm_quant/api.py
  • python/cudnn/grouped_gemm/grouped_gemm_quant/grouped_gemm_quant.py
  • python/cudnn/grouped_gemm/grouped_gemm_quant/moe_blockscaled_grouped_gemm_quant_rubin.py
  • python/cudnn/grouped_gemm/grouped_gemm_srelu/__init__.py
  • python/cudnn/grouped_gemm/grouped_gemm_srelu/api.py
  • python/cudnn/grouped_gemm/grouped_gemm_srelu/moe_blockscaled_grouped_gemm_srelu_quant.py
  • python/cudnn/grouped_gemm/grouped_gemm_swiglu/__init__.py
  • python/cudnn/grouped_gemm/grouped_gemm_swiglu/api.py
  • python/cudnn/grouped_gemm/grouped_gemm_swiglu/grouped_gemm_swiglu_quant.py
  • python/cudnn/grouped_gemm/grouped_gemm_unfused/__init__.py
  • python/cudnn/grouped_gemm/grouped_gemm_unfused/_bf16_api.py
  • python/cudnn/grouped_gemm/grouped_gemm_unfused/api.py
  • python/cudnn/grouped_gemm/grouped_gemm_unfused/moe_grouped_gemm.py
  • python/cudnn/grouped_gemm/grouped_gemm_utils.py
  • python/cudnn/grouped_gemm/grouped_gemm_wgrad/__init__.py
  • python/cudnn/grouped_gemm/grouped_gemm_wgrad/_bf16_api.py
  • python/cudnn/grouped_gemm/grouped_gemm_wgrad/_blockscaled_api.py
  • python/cudnn/grouped_gemm/grouped_gemm_wgrad/api.py
  • python/cudnn/grouped_gemm/grouped_gemm_wgrad/moe_blockscaled_grouped_gemm_wgrad.py
  • python/cudnn/grouped_gemm/grouped_gemm_wgrad/moe_grouped_gemm_wgrad.py
  • python/cudnn/grouped_gemm/moe_kernel_helpers.py
  • python/cudnn/grouped_gemm/moe_persistent_scheduler.py
  • python/cudnn/grouped_gemm/moe_sched_extension.py
  • python/cudnn/grouped_gemm/moe_utils.py
  • python/cudnn/grouped_gemm/utils.py
  • python/cudnn/native_sparse_attention/__init__.py
  • python/cudnn/native_sparse_attention/compression/__init__.py
  • python/cudnn/native_sparse_attention/compression/api.py
  • python/cudnn/native_sparse_attention/compression/fmha.py
  • python/cudnn/native_sparse_attention/compression/fmha_helpers.py
  • python/cudnn/native_sparse_attention/selection/NSA_select_attn_fwd_hmma.py
  • python/cudnn/native_sparse_attention/selection/__init__.py
  • python/cudnn/native_sparse_attention/selection/api.py
  • python/cudnn/native_sparse_attention/sliding_window_attention/__init__.py
  • python/cudnn/native_sparse_attention/sliding_window_attention/api.py
  • python/cudnn/native_sparse_attention/top_k/__init__.py
  • python/cudnn/native_sparse_attention/top_k/api.py
  • python/cudnn/native_sparse_attention/top_k/nsa_top_k_reduction_fwd.py
  • python/cudnn/native_sparse_attention/utils.py
  • python/cudnn/nodes.py
  • python/cudnn/ops/__init__.py
  • python/cudnn/ops/causal_conv1d.py
  • python/cudnn/rmsnorm_rht_amax/__init__.py
  • python/cudnn/rmsnorm_rht_amax/api.py
  • python/cudnn/rmsnorm_rht_amax/kernel.py
  • python/cudnn/sdpa/__init__.py
  • python/cudnn/sdpa/bwd/__init__.py
  • python/cudnn/sdpa/bwd/api.py
  • python/cudnn/sdpa/bwd/fmha_backward_sm100_2kernel.py
  • python/cudnn/sdpa/bwd/fmha_dkdv_d256_sm100.py
  • python/cudnn/sdpa/bwd/fmha_dq_d256_sm100.py
  • python/cudnn/sdpa/fmha_utils.py
  • python/cudnn/sdpa/fwd/__init__.py
  • python/cudnn/sdpa/fwd/api.py
  • python/cudnn/sdpa/fwd/fmha_forward_sm100_d256.py
  • python/cudnn/sdpa/utils.py
  • python/cudnn/wrapper.py
  • python/cudnn/yarn.py
  • python/properties.cpp
  • python/pycudnn.cpp
  • python/pygraph/norm.cpp
  • python/pygraph/pointwise.cpp
  • python/pygraph/pygraph.cpp
  • python/pygraph/pygraph.h
  • python/pygraph/sdpa.cpp
  • samples/CMakeLists.txt
  • samples/cpp/CMakeLists.txt
  • samples/cpp/causal_conv1d/b2b_causal_conv1d.cpp
  • samples/cpp/causal_conv1d/causal_conv1d.cpp
  • samples/cpp/causal_conv1d/causal_conv1d_nwh.cpp
  • samples/cpp/convolution/conv_dynamic_shape_benchmark.cpp
  • samples/cpp/convolution/dgrads.cpp
  • samples/cpp/convolution/fp8_fprop.cpp
  • samples/cpp/convolution/fprop.cpp
  • samples/cpp/convolution/int8_fprop.cpp
  • samples/cpp/convolution/wgrads.cpp
  • samples/cpp/matmul/blackwell_nvfp4_mxfp8_block_scale_matmul.cpp
  • samples/cpp/matmul/complex_fp32_matmul.cpp
  • samples/cpp/matmul/fp8_matmul.cpp
  • samples/cpp/matmul/general_block_scale_matmul.cpp
  • samples/cpp/matmul/int8_matmul.cpp
  • samples/cpp/matmul/matmuls.cpp
  • samples/cpp/matmul/mixed_matmul.cpp
  • samples/cpp/membound/boolean_fusion.cpp
  • samples/cpp/membound/concat.cpp
  • samples/cpp/membound/membound_fusion.cpp
  • samples/cpp/membound/reshape.cpp
  • samples/cpp/membound/slice.cpp
  • samples/cpp/membound/transpose.cpp
  • samples/cpp/misc/autotuning.cpp
  • samples/cpp/misc/compile_time_constant_example.cpp
  • samples/cpp/misc/cudagraphs.cpp
  • samples/cpp/misc/custom_plan.cpp
  • samples/cpp/misc/deviceless_aot_compilation.cpp
  • samples/cpp/misc/parallel_compilation.cpp
  • samples/cpp/misc/pointwise.cpp
  • samples/cpp/misc/resample.cpp
  • samples/cpp/misc/serialization.cpp
  • samples/cpp/misc/slice.cpp
  • samples/cpp/misc/sm_carveout.cpp
  • samples/cpp/moe_grouped_matmul/moe_grouped_matmul.cpp
  • samples/cpp/norm/adaptive_layernorm.cpp
  • samples/cpp/norm/batchnorm.cpp
  • samples/cpp/norm/layernorm.cpp
  • samples/cpp/norm/layernorm_bitmask_relu.cpp
  • samples/cpp/norm/norm_block_scale.cpp
  • samples/cpp/norm/norm_zero_centered_gamma.cpp
  • samples/cpp/norm/rmsnorm.cpp
  • samples/cpp/sdpa/fp16_benchmark.cpp
  • samples/cpp/sdpa/fp16_bwd.cpp
  • samples/cpp/sdpa/fp16_bwd_with_cudagraphs.cpp
  • samples/cpp/sdpa/fp16_bwd_with_flexible_graphs.cpp
  • samples/cpp/sdpa/fp16_bwd_with_sink_token.cpp
  • samples/cpp/sdpa/fp16_cached.cpp
  • samples/cpp/sdpa/fp16_dynamic_shapes.cpp
  • samples/cpp/sdpa/fp16_fwd.cpp
  • samples/cpp/sdpa/fp16_fwd_paged_decode_and_prefill.cpp
  • samples/cpp/sdpa/fp16_fwd_with_block_mask.cpp
  • samples/cpp/sdpa/fp16_fwd_with_cu_seq_len.cpp
  • samples/cpp/sdpa/fp16_fwd_with_cudagraphs.cpp
  • samples/cpp/sdpa/fp16_fwd_with_custom_dropout.cpp
  • samples/cpp/sdpa/fp16_fwd_with_flexible_graphs.cpp
  • samples/cpp/sdpa/fp16_fwd_with_max_and_sum_exp.cpp
  • samples/cpp/sdpa/fp16_fwd_with_paged_caches.cpp
  • samples/cpp/sdpa/fp16_fwd_with_sink_token.cpp
  • samples/cpp/sdpa/fp8_bwd.cpp
  • samples/cpp/sdpa/fp8_bwd_bottom_right_causal_mask.cpp
  • samples/cpp/sdpa/fp8_bwd_with_current_scaling.cpp
  • samples/cpp/sdpa/fp8_fwd.cpp
  • samples/cpp/sdpa/fp8_fwd_bottom_right_causal_mask.cpp
  • samples/cpp/sdpa/fp8_fwd_current_scaling.cpp
  • samples/cpp/sdpa/mxfp8_bwd.cpp
  • samples/cpp/sdpa/mxfp8_fwd.cpp
  • samples/cpp/sdpa/prefill_oss_engine.cpp
  • samples/cpp/utils/helpers.h
  • samples/legacy_samples/CMakeLists.txt
  • samples/legacy_samples/conv_sample.cpp
  • samples/legacy_samples/conv_sample.h
  • samples/legacy_samples/cpu_references.h
  • samples/legacy_samples/fp16_dev.cu
  • samples/legacy_samples/fp16_emu.cpp
  • samples/legacy_samples/fp8_sample.cpp
  • samples/legacy_samples/fp8_sample.h
  • samples/legacy_samples/fusion_sample.cpp
  • samples/legacy_samples/fusion_sample.h
  • samples/legacy_samples/helpers.cpp
  • samples/legacy_samples/norm_samples.cpp
  • samples/legacy_samples/norm_samples.h
  • samples/legacy_samples/test_list.cpp
  • samples/legacy_samples/utils/error_util.h
  • samples/legacy_samples/utils/fp16_dev.h
  • samples/legacy_samples/utils/fp16_emu.h
  • samples/legacy_samples/utils/helpers.h
  • samples/llama/100_download_weight.py
  • samples/llama/101_hf_llama_tieout.py
  • samples/llama/102_torch_llama_tieout.py
  • samples/llama/103_cudnn_llama_tieout.py
  • samples/llama/104_torch_llama_nvtx.py
  • samples/llama/105_cudnn_llama_nvtx.py
  • samples/llama/decode_nvtx_profile.py
  • samples/llm_coverage/test_gqa_b+h+s+d.py
  • samples/llm_coverage/test_linear+swish_1+bs+d.py
  • samples/llm_coverage/test_linear_1+bs+d.py
  • samples/llm_coverage/test_linear_b+s+d.py
  • samples/llm_coverage/test_linear_b+s+d_fp32compute.py
  • samples/llm_coverage/test_rmsnorm_b+s+d.py
  • samples/llm_coverage/test_rmsnorm_bs+d.py
  • samples/llm_coverage/test_rope_b+s+h+d.py
  • samples/llm_coverage/test_swiglu_1+bs+d.py
  • samples/llm_coverage/test_swiglu_layer_1+bs+d.py
  • setup.py
  • test/CMakeLists.txt
  • test/cpp/CMakeLists.txt
  • test/cpp/get_engine_and_knobs.cpp
  • test/cpp/pointwise_tests.cpp
  • test/cpp/serialize.cpp
  • test/cpp/tensor.cpp
  • test/cpp/validate.cpp
  • test/cpp/version.cpp
  • test/python/conftest.py
  • test/python/fe_api/bsa/__init__.py
  • test/python/fe_api/bsa/bsa_reference.py
  • test/python/fe_api/bsa/bsa_utils.py
  • test/python/fe_api/bsa/test_BSA_attention_backward.py
  • test/python/fe_api/bsa/test_BSA_attention_forward.py
  • test/python/fe_api/csa/test_CSA_compressor.py
  • test/python/fe_api/dsa/dsa_reference.py
  • test/python/fe_api/dsa/dsa_utils.py
  • test/python/fe_api/dsa/test_DSA_dense_indexer_backward.py
  • test/python/fe_api/dsa/test_DSA_dense_score_recompute.py
  • test/python/fe_api/dsa/test_DSA_indexer_backward.py
  • test/python/fe_api/dsa/test_DSA_indexer_forward.py
  • test/python/fe_api/dsa/test_DSA_indexer_top_k.py
  • test/python/fe_api/dsa/test_DSA_runtime.py
  • test/python/fe_api/dsa/test_DSA_sparse_attention_backward.py
  • test/python/fe_api/dsa/test_DSA_sparse_score_recompute.py
  • test/python/fe_api/gemm/test_gemm_amax.py
  • test/python/fe_api/gemm/test_gemm_amax_utils.py
  • test/python/fe_api/gemm/test_gemm_dsrelu.py
  • test/python/fe_api/gemm/test_gemm_dsrelu_utils.py
  • test/python/fe_api/gemm/test_gemm_proj_rope_mxfp8.py
  • test/python/fe_api/gemm/test_gemm_proj_rope_mxfp8_utils.py
  • test/python/fe_api/gemm/test_gemm_srelu.py
  • test/python/fe_api/gemm/test_gemm_srelu_utils.py
  • test/python/fe_api/gemm/test_gemm_swiglu.py
  • test/python/fe_api/gemm/test_gemm_swiglu_utils.py
  • test/python/fe_api/grouped_gemm/test_discrete_grouped_gemm_dswiglu.py
  • test/python/fe_api/grouped_gemm/test_discrete_grouped_gemm_dswiglu_utils.py
  • test/python/fe_api/grouped_gemm/test_discrete_grouped_gemm_swiglu.py
  • test/python/fe_api/grouped_gemm/test_discrete_grouped_gemm_swiglu_utils.py
  • test/python/fe_api/grouped_gemm/test_grouped_gemm_dglu.py
  • test/python/fe_api/grouped_gemm/test_grouped_gemm_dsrelu.py
  • test/python/fe_api/grouped_gemm/test_grouped_gemm_dsrelu_utils.py
  • test/python/fe_api/grouped_gemm/test_grouped_gemm_dswiglu.py
  • test/python/fe_api/grouped_gemm/test_grouped_gemm_dswiglu_utils.py
  • test/python/fe_api/grouped_gemm/test_grouped_gemm_glu.py
  • test/python/fe_api/grouped_gemm/test_grouped_gemm_glu_hadamard.py
  • test/python/fe_api/grouped_gemm/test_grouped_gemm_quant.py
  • test/python/fe_api/grouped_gemm/test_grouped_gemm_quant_utils.py
  • test/python/fe_api/grouped_gemm/test_grouped_gemm_srelu.py
  • test/python/fe_api/grouped_gemm/test_grouped_gemm_srelu_utils.py
  • test/python/fe_api/grouped_gemm/test_grouped_gemm_swiglu.py
  • test/python/fe_api/grouped_gemm/test_grouped_gemm_swiglu_utils.py
  • test/python/fe_api/grouped_gemm/test_grouped_gemm_wgrad.py
  • test/python/fe_api/grouped_gemm/test_grouped_gemm_wgrad_utils.py
  • test/python/fe_api/norm/test_rmsnorm_rht_amax.py
  • test/python/fe_api/nsa/nsa_reference.py
  • test/python/fe_api/nsa/nsa_utils.py
  • test/python/fe_api/nsa/test_NSA_compression_attention.py
  • test/python/fe_api/nsa/test_NSA_selection_attention.py
  • test/python/fe_api/nsa/test_NSA_swa.py
  • test/python/fe_api/nsa/test_NSA_topk_reduction.py
  • test/python/fe_api/sdpa/test_sdpa_bwd.py
  • test/python/fe_api/sdpa/test_sdpa_bwd_utils.py
  • test/python/fe_api/sdpa/test_sdpa_fwd.py
  • test/python/fe_api/sdpa/test_sdpa_fwd_utils.py
  • test/python/fe_api/test_api_base_logging.py
  • test/python/fe_api/test_fe_api_utils.py
  • test/python/fe_api/test_grouped_gemm_bf16.py
  • test/python/fe_api/test_grouped_gemm_bf16_utils.py
  • test/python/fe_api/test_grouped_gemm_dglu_bf16_utils.py
  • test/python/fe_api/test_grouped_gemm_glu_bf16_utils.py
  • test/python/fe_api/test_grouped_gemm_wgrad_bf16_utils.py
  • test/python/fe_api/test_rubin_kernel_dispatch.py
  • test/python/sdpa/blocked.py
  • test/python/sdpa/fp16.py
  • test/python/sdpa/fp16_ref.py
  • test/python/sdpa/fp8.py
  • test/python/sdpa/fp8_ref.py
  • test/python/sdpa/helpers.py
  • test/python/sdpa/mxfp8.py
  • test/python/sdpa/mxfp8_ref.py
  • test/python/sdpa/random_config.py
  • test/python/test_api_signature_parity.py
  • test/python/test_apply_rope.py
  • test/python/test_batchnorm.py
  • test/python/test_block_scale_quantize.py
  • test/python/test_block_scale_quantize_dynamic_shape.py
  • test/python/test_collect_env.py
  • test/python/test_conv_bias.py
  • test/python/test_conv_fprop.py
  • test/python/test_conv_fuzzer.py
  • test/python/test_conv_genstats.py
  • test/python/test_conv_reduction.py
  • test/python/test_cudnn_sdpa_op.py
  • test/python/test_deviceless_aot_compilation.py
  • test/python/test_engine_router.py
  • test/python/test_flexible_sdpa.py
  • test/python/test_flexible_sdpa_bprop.py
  • test/python/test_graph_native.py
  • test/python/test_instancenorm.py
  • test/python/test_kernel_cache.py
  • test/python/test_layernorm.py
  • test/python/test_low_precision_matmul.py
  • test/python/test_matmul_bias_relu.py
  • test/python/test_matmul_fuzzer.py
  • test/python/test_mhas.py
  • test/python/test_mhas_v2.py
  • test/python/test_moe_grouped_matmul.py
  • test/python/test_moe_grouped_matmul_op.py
  • test/python/test_native_backend_lowering.py
  • test/python/test_norm_fuzzer.py
  • test/python/test_oss_rope.py
  • test/python/test_rmsnorm.py
  • test/python/test_sdpa_chunked_prefill.py
  • test/python/test_sdpa_custom_features.py
  • test/python/test_sdpa_edge_cases.py
  • test/python/test_sdpa_fp32_rejected.py
  • test/python/test_sdpa_thd.py
  • test/python/test_sdpa_with_caching.py
  • test/python/test_silu_and_mul.py
  • test/python/test_slice.py
  • test/python/test_sm100_prefill_oss_engine.py
  • test/python/test_sm100_rms_norm_silu_graph_api.py
  • test/python/test_sm90_prefill_oss_engine.py
  • test/python/test_utils.py
  • test/python/test_wgrads.py
  • test/python/test_yarn_rope.py
  • tools/cudnn_repro/cudnn_repro/__init__.py
  • tools/cudnn_repro/cudnn_repro/__main__.py
  • tools/cudnn_repro/cudnn_repro/log_parser.py
  • tools/cudnn_repro/cudnn_repro/operations.py
  • tools/cudnn_repro/cudnn_repro/repro_command.py
  • tools/cudnn_repro/cudnn_repro/sdpa_bwd.py
  • tools/cudnn_repro/cudnn_repro/sdpa_fp8_bwd.py
  • tools/cudnn_repro/cudnn_repro/sdpa_fp8_fwd.py
  • tools/cudnn_repro/cudnn_repro/sdpa_fwd.py
  • tools/cudnn_repro/cudnn_repro/utils.py
  • tools/cudnn_repro/tests/__init__.py
  • tools/cudnn_repro/tests/helpers.py
  • tools/cudnn_repro/tests/test_cudnn_repro_bwd.py
  • tools/cudnn_repro/tests/test_cudnn_repro_cli.py
  • tools/cudnn_repro/tests/test_cudnn_repro_closed_loop.py
  • tools/cudnn_repro/tests/test_cudnn_repro_fp8.py
  • tools/cudnn_repro/tests/test_cudnn_repro_fp8_closed_loop.py
  • tools/cudnn_repro/tests/test_cudnn_repro_log_parser.py
  • tools/cudnn_repro/tests/test_cudnn_repro_mxfp8_closed_loop.py
  • tools/cudnn_repro/tests/test_cudnn_repro_schema.py
  • tools/cudnn_repro/tests/test_cudnn_repro_utils.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
✨ Finishing Touches
🧪 Generate unit tests (beta)
  • Create PR with unit tests

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

@Anerudhan Anerudhan mentioned this pull request Jul 23, 2026
18 tasks
@Anerudhan
Anerudhan force-pushed the relicense-apache-2.0 branch from 158ed02 to 9f5a3cb Compare July 23, 2026 19:18
@Anerudhan Anerudhan changed the title Relicense from MIT to Apache 2.0 Relicense to Apache-2.0 (dual-license; external-contributor files remain MIT) Jul 23, 2026
@Anerudhan
Anerudhan force-pushed the relicense-apache-2.0 branch 3 times, most recently from c85e9df to 1fdbcc4 Compare July 30, 2026 06:08
…MIT)

Rebased onto current develop (b950af1). Relicenses NVIDIA-authored code from
MIT to Apache-2.0 using the exact NVIDIA OSS SPDX header, keeping under MIT
only files that still carry contributions from external contributors whose
permission has not been established.

License assignment (per-file SPDX tag on every source file):
- Apache-2.0: 541 files
- MIT (50): surviving lines from external contributors, consent pending
- MIT (29): FlashAttention/QuACK-derived files carrying external authors'
  copyright

Cleared, freeing 16 files to Apache-2.0:
- Written consent on issue NVIDIA#431 (10): take-cheeze, fallintoplace, zianglih,
  JackRao123, zkyue, Hyaloid, haowen-han, junaire, szluyu99, dimitar-asenov.
- NVIDIA employment (2), commits under personal email addresses:
  HollowMan6, and hxbai (Hongxiao Bai, PR NVIDIA#410) -- frees
  .../indexer_top_k/indexer_top_k_varlen_util.py and
  test/python/fe_api/dsa/test_DSA_indexer_top_k.py.
Files touched by both a cleared and a still-pending contributor remain MIT.

PR NVIDIA#427 (CSA fused Compressor kernels ported from Megatron-LM, author @zkyue
who consented): its 13 files had no license header at all; they now carry the
NVIDIA Apache-2.0 SPDX header. Megatron-LM added to THIRD_PARTY_LICENSES.txt
as an NVIDIA Apache-2.0 provenance note.

Still pending (7): ConnorBaker, DrDirk, EmilienM, jyknight, sbcd90, valgur,
and co-author Benjamin Leff.

Every change is comment/header-only; all Python compiles.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
@Anerudhan
Anerudhan force-pushed the relicense-apache-2.0 branch from 1fdbcc4 to 0c08141 Compare July 30, 2026 06:13
@Anerudhan Anerudhan self-assigned this Jul 30, 2026
@Anerudhan Anerudhan added cat-feature Requests for new functionality, APIs, examples, or behavior improvements. mod-frontend cuDNN frontend APIs, operation graph construction, plans, and user-facing wrappers. orig-nv-eng Reported or requested by NVIDIA engineering. labels Jul 30, 2026
@Anerudhan Anerudhan added this to the Frontend 1.27.0 milestone Jul 30, 2026
@Anerudhan
Anerudhan marked this pull request as ready for review July 30, 2026 17:42
@Anerudhan
Anerudhan requested a review from vedaanta July 30, 2026 17:47
@Anerudhan
Anerudhan merged commit 57dc579 into NVIDIA:develop Jul 30, 2026
1 check passed
YangXu1990uiuc added a commit to YangXu1990uiuc/cudnn-frontend that referenced this pull request Aug 9, 2026
sdpa/fwd/heuristics.py was created by copying the header from config_sm120.py,
which is MIT -- so the new file inherited a tag that does not apply to it.

Per LICENSING.md, the repo relicensed MIT -> Apache-2.0 in NVIDIA#408 and a file is
kept under MIT for exactly two reasons: surviving lines from an external
contributor who has not consented to relicensing, or derivation from
third-party source. A file written from scratch at NVIDIA has neither, so
Apache-2.0 is the correct tag -- as it already is on every other file this
change adds content to (engines/heuristics.py, engines/manifest.py,
sdpa/fwd/engine.py, and the tests). The MIT neighbours in sdpa/fwd are
pre-existing files this PR only edits, and editing does not move a file
between licenses.

Also switches to the SPDX-FileCopyrightText form the Apache-2.0 files use.
Anerudhan pushed a commit that referenced this pull request Aug 9, 2026
…_sort (#528)

* Take the backend's plans one heuristic mode at a time

Ranking the two sides against each other needs to know which backend entries
are mode-A recommendations and which are fallbacks -- "the backend's A ahead of
ours, its fallbacks behind" cannot be said about one opaque list. Until now the
whole thing arrived from a single create_execution_plans([A, FALLBACK]).

No C++ change is needed. C++ appends each query to the same plan list, and
get_execution_plan_count() already exists, so asking one mode at a time and
reading the count after each gives the boundaries. Measured on a 512^3 bf16
matmul (sm90, cuDNN 9.25): A -> plans[0:15], all knob-bearing; FALLBACK ->
plans[15:17], bare eng0/eng7 with no knobs; the two segments do not overlap.

A mode with no configs raises, which is not a decline while another mode still
has entries -- an OPENSOURCE-only query legitimately leaves the cuDNN modes
empty. Only every mode failing means the backend has nothing, and then the last
error is re-raised so the caller still reports why.

* Move plan ranking out of the engines and into one heuristics function

An engine cannot rank. It sees neither its siblings nor the backend's entries,
so propose_plans could only ever order its own knobs -- and then something
downstream had to merge the two sides anyway, which heuristics_sort did by
concatenating and calling it ranking. All four in-tree propose_plans were the
base class's default copied verbatim: the hook has never decided anything.

create_execution_plans() now gathers the inputs (parsed facts, the family's
offered ids, the backend's entries tagged by mode) and hands all of it to the
graph's family in ONE call. What comes back IS graph.plans, position for
position. An engine answers two questions: can I serve this graph
(check_support), and compile me this config (build_plan).

sdpa/fwd/heuristics.py is the first such hook, and it is deliberately a frame
with no tuning in it: one entry per eligible cell at the config its capability
row declares. Mode A and FALLBACK differ only in which backend entries they
carry; OPENSOURCE is mode A without the backend's recommendation, since these
cells ARE the open-source implementation. Real per-cell rules land on top.

Deleted, all superseded or never used:
  BaseEngine.propose_plans + its 4 implementations
  BaseEngine.default_knobs        only fed propose_plans
  heuristics_sort                 merging is part of ranking, not a step after
  engines/router.py entirely      Router / default_router / set_router /
                                  pygraph(router=) -- policy has one home now,
                                  and decline_types moved to base.py where the
                                  engine contract already lives
  engines.probe() (fwd + bwd)     superseded by check_support
  graph.engine                    pure alias of selected_engine, zero callers
  graph.from_serialized           zero callers; serialize/deserialize are the
                                  pybind-era API and stay

knobs=None no longer means "engine, pick for me" -- the heuristics name a
concrete config. A None field survives only on an axis whose capability row
declares no domain. That reading is what let one choice be made twice, once
when ranking and once inside the adapter.

* Update the dispatch tests to the ranking contract, and delete what it retired

Ranking has one home, so a test that wants a specific order replaces
heuristics.rank instead of subclassing Router. The _ranking() helper does that;
it is the same monkeypatch idiom the rest of the suite already uses.

Deleted rather than translated:
  test_set_router_frozen_after_planning   the API it tested is gone
  test_a_claiming_engine_is_tried_before_the_backend
                                          asserted that python plans always
                                          outrank the backend, which is a
                                          per-cell measurement, not a rule.
                                          FROST coverage rides on
                                          heur_mode.OPENSOURCE instead: ask for
                                          it and any graph still landing on a
                                          backend plan is one FROST cannot serve

Renamed for what they now test: test_mixed_ranking_dispatch,
test_empty_ranking_output_rejected, test_mixed_ranking_backend_slot_executes,
test_constructor_backends_validated_and_ranking_ids_checked.

One assertion changed meaning: the backend is queried once PER MODE now, so
_create_backend_plans records two create_execution_plans calls for [A, FALLBACK].

test_sdpa_graph_analyzer called engines.probe() twice; those two call
analyze_for directly, so no production API exists only for tests.

Six sdpa test files each carried a verbatim copy of _select_engine matching a
bare engine name. Plans now read <engine>[<knobs>] because the heuristics name
a concrete config for every entry, so they share frost_test_utils.select_engine,
which matches on the engine.

208 passed. test_a_replayed_plan_reports_its_own_notes still fails and also
fails on develop without this change -- C++ on 9.25 no longer raises for an
index one past the plan count.

* Update the design doc, and align the remaining sdpa test helpers

The doc still described a Router with three pluggability levels, engines that
propose their own plans, and heuristics_sort as the seam a cost model replaces.
Rewritten to what dispatch now does: one call per graph into the family's
heuristics hook, the backend's entries tagged by the mode that produced them,
and heur_mode.OPENSOURCE as the way FROST coverage is measured rather than
assumed.

Also states plainly what register_backend is and is not. It installs an engine
instance on one graph -- the hatch tests use to inject a fake. It does not make
an engine rankable: an out-of-tree engine declares no Capabilities, so nothing
can enumerate its configs or place it against the backend. The follow-up list
now names removing that concept, since an engine id is decodable from the
manifest alone.

Six sdpa test files each carried a verbatim copy of _select_engine matching a
bare engine name; the shared frost_test_utils.select_engine matches on the
engine, which is what plan names now carry a config suffix for.

208 passed locally. The one failure is test_a_replayed_plan_reports_its_own_notes,
which fails on develop without this change too.

* Decode an engine id from the manifest, with nothing registered first

An engine id is fully decodable from the manifest: the family owning the id
block, then the slot within it. _owners_for_id only ever looked inside the
graph's candidate set, so an id could be resolved only if something had already
put that engine there -- which made register_backend look like a prerequisite
for create_execution_plan() when it is really just one way to supply an
instance.

engine_for_id() closes that. _owners_for_id falls back to it, so replaying a
recorded (engine_id, knobs) works on a fresh graph, including for an engine
that is not a candidate for THAT graph -- there the replay is a deliberate pin,
not a routing decision. A gated-off slot still resolves to None rather than
being built.

Groundwork for removing the out-of-tree engine concept entirely.

* Remove the out-of-tree engine concept: the manifest is the only way in

Every python engine now exists exactly one way. register_backend,
pygraph(backends=), graph.backends and OUT_OF_TREE_ID_BASE are gone, and
_candidate_engines() is the graph's family and nothing else.

An out-of-tree engine could never be RANKED anyway: it declares no
Capabilities, so nothing could enumerate its configs or place it against the
backend. It was an entry point into the plan list, not into the decision. And
being a candidate had nothing to do with fitness -- an engine was in the list
because someone had registered it, so an engine that could not serve the graph
was still tried, and failed at build instead of at classification.

The linear_attention suites used register_backend to PIN an implementation --
cuTile rather than FROST. That is not what registration is for, and those
engines are in the manifest already, so the pin is now by name and applied
after planning through select_plan(): engine_utils.pin_engines() / apply_pin().
apply_pin raises when the pinned engine produced no plan, so a pin that stops
working fails the first op call. The cutile conftest used to check the pin by
inspecting the CANDIDATE list, which passes whether or not the pin took effect
-- which is how it ran for months against whichever engine the ranking picked
while the seam it pinned through was dead. That check is deleted; the pin
enforces itself.

heuristics: no engine sits outside a family now, so the "family-less engines go
last" branch is gone and _without_a_family is _unranked -- the case it covers
is a family that declares no heuristics hook, not an engine with no family.

test_engine_router.py -> test_dispatch.py. It never tested a Router; it tested
dispatch -- one plan list, the at-index APIs, select_plan's strict pin,
one-shot planning, how a decline advances the walk, note filters reaching
python plans, manifest classification, facts attachment. _offer(monkeypatch,
*engines) replaces register_backend by putting the fakes in a manifest family,
so the tests reach engines through the same path production does.

Six tests deleted with the concept they tested -- all checked registration-time
id validation, which has no subject now that engines never declare their own
ids: test_register_backend_validation,
test_engine_id_in_the_in_tree_region_is_rejected,
test_a_registered_in_tree_engine_is_not_offered_twice,
test_overlapping_declared_id_blocks_are_rejected,
test_a_lying_owns_id_cannot_capture_another_engines_plans,
test_constructor_backends_validated_and_ranking_ids_checked. What they
protected is covered by test_family_id_blocks_are_disjoint and
test_every_engine_spec_has_a_manifest_slot. BaseEngine.owns_id goes with them:
zero callers, and its docstring already called it a convenience.

Three tests needed real thought rather than a mechanical edit:

- The "no family, no facts payload" test built a bare relu graph. relu names no
  family, so there is no python candidate, and the backend declines a 2-D
  pass-by-value tensor -- planning raised before the assertion. The claim under
  test is about the payload, not about the graph being servable.
- The mutable-after-validate window was `not self._backends`: validate() lowers
  and freezes any graph the backend CAN lower, and registering an engine was
  the only way to skip that. With registration gone the window is exactly the
  ops with no backend lowering, which is what the property was always about.
- test_api_signature_parity asserted {"backends", "router"} were keyword-only.
  Both are gone, so the assertion had no subject; what it protected is that
  nothing pygraph-only is POSITIONAL, which is now asserted directly.

TorchMatmulEngine goes too. It reimplemented matmul, bias and relu in torch
inside a dispatch test: the numeric assertions proved torch, not dispatch, and
"torch_matmul" in plan names reads like something cuDNN ships. StubEngine
replaces it -- same claim on the graph, no arithmetic, and it RECORDS what
dispatch handed it, so the fusion test now asserts what was only implied
before: every node arrives in build order, each input port resolved to the
caller's storage, and the virtual intermediate carrying none.

* Give the SM120 SDPA-forward cell a real tile rule, as the worked example

The framework had no rule in it: every cell went to `_sole()` on each knob
axis, which answers None the moment a row declares more than one value. The
SM120 prefill row declares tile_ms={64,128}, tile_ns={64,128}, so its choice
fell through to api_dsl's `_SM120_Q_TILES[0]` default -- the choice being made
in the adapter is exactly what moving ranking out of the engines was meant to
stop, and it left the frame with nothing showing how a rule is added.

_sm120_tiles(facts) is that rule, and it is measured rather than invented:
regret 1.009 geomean / 1.054 worst against the best of the enumerated domain.
tile_n=128 always; tile_m=64 when the grid cannot fill the machine AND each CTA
has enough KV tiles to amortize the extra Q-tile loop, with a causal mask
counted as a halved effective grid because it halves the work per CTA. It reads
facts and nothing else -- device_sm_count is already on the record.

Shape a colleague can copy: write the function, list the cell in
_TILE_RULE_CELLS, put the measurement in the commit. A cell absent from that
set keeps the old behaviour (its row's sole point per axis), which is the
honest answer when nobody has timed it.

Mode A now emits the guess FIRST and the rest of the domain behind it, so a
caller who autotunes has the runners-up and a caller who does not gets the best
guess at index 0. FALLBACK takes the smallest tile the row admits -- the config
that asks least of the device; picking real fallback configs per cell is a
TODO left in the file.

* Ask one function whether an SM120 tile fits, not two

Naming tile_n=128 unconditionally broke D=208/224/240/256: the adapter's own
`if self.tile_n is None` branch was quietly shrinking the KV tile to whatever
fit SMEM, so leaving the knob None had been answering a CAPABILITY question,
not a tuning one. Requesting a value skips that branch, and the request then
fails the very check the branch existed to satisfy -- 106512 bytes wanted
against the part's 101376.

The fit arithmetic moves to config_sm120.smem_bytes(), beside the template it
describes, and both callers use it: the adapter's check and the ranking's
choice. The rule now reads "tile_n = the largest that fits, tile_m by
occupancy", and the runners-up it offers are filtered the same way -- a config
the kernel cannot fit is not a runner-up, it is an entry that sits in the list
to decline at build.

test_api_signature_parity asserted {"backends", "router"} were keyword-only.
Both are gone, so the assertion had no subject; what it protected is that
nothing pygraph-only is POSITIONAL, which is now asserted directly.

* Query the backend for the modes the ranking will actually place

The default mode list was written out twice -- once in _create_backend_plans,
once in heuristics.default_modes. They agree today; a change to one alone would
have the backend enumerate plans for a mode no family places, which reads as
the family losing entries rather than as the query asking for the wrong thing.

* Restore the #512 SDPA tests this branch had silently reverted

Four test files were carrying their PRE-#512 content while the production code
they exercise is post-#512. The branch is cherry-picked onto the github
develop, and the commit that consolidated the sdpa test helpers was authored
against a tree from before #512 landed -- so the cherry-pick took the whole
file, not the helper edit, and reverted #512's test additions with it.
api_dsl.py and engines.py were untouched by that, which is why nothing looked
wrong until an SM100 box ran the suite: 16 failures, all of them tests
asserting the old contract against the new kernels (a stats-less SM100 graph
now carves a dummy LSE, so get_workspace_size() is b*h*s*4, not 0).

Restored all four from gh/develop and re-applied only what this branch meant to
change:

- test_sdpa_fwd_dsl_sm100 / _sm120: the local verbatim copy of _select_engine
  -> frost_test_utils.select_engine.
- test_sdpa_frontend_integration: plan-name lookups made suffix-aware. The
  heuristics now name a concrete config for every entry, so a plan reads
  "<engine>[<knobs>]" and names.index(_FROST) raises ValueError.
- test_sdpa_graph_analyzer: engines.probe() is deleted, so _eligible asks
  analyze_for(...)[1] is None.

The ragged-Stats coverage #512 added (token-major and head-major layouts,
zero-length sequences, the analyzer acceptance test, the strict LSE presence
contract in both directions) is back verbatim.

* Bring the design doc to the architecture as it now stands

The dispatch tree, written out: what create_execution_plans does in order,
where the backend's per-mode entries come from, and where a family's rules sit.
That tree was the first thing anyone asked for and the doc did not have it.

Corrects three things the doc stated as settled that this PR changed:
the delegating entry leads the BACKEND's block and not the family's (it falls
through to native configs when the C++ OSS engine declines, so ahead of an
OPENSOURCE block it answers a coverage question with a native kernel); a plan's
identity is (engine_id, knobs) and never its cpp_index; whether a heuristic
mode succeeded is tracked per call, not inferred from plan spans.

Adds what each machine covers. The suites SKIP on the wrong arch rather than
fail, so a green sweep on one box says nothing about the others -- defaulting
to CUDA device 0 is how a whole SM100 run silently skips.

Follow-ups now name what is actually left: one tuning rule exists, FALLBACK is
a placeholder, _MEASURED_BEHIND is empty by design.

* Answer the backend's plan query once per distinct config

Two findings from the second review pass, both about APIs whose callers moved
under the branch.

Graph::create_execution_plans checks override_heuristics_query() FIRST and
returns before it reads the mode at all -- deterministic SDPA backward and FP8
backward both override. Asking one mode at a time therefore appends the SAME
engine-17 config once per mode, and backend_plan_entries() handed all of them
back. SDPA forward's recommend() would have deduped them; SDPA BACKWARD
declares no heuristics hook, so _unranked passed the duplicates straight into
graph.plans and build_plans(ALL) or an autotuner would compile and time one
config twice. Deduped at collection instead of in each family: a repeated
(engine, knobs) in the backend's own list is never two different things, and
the first index is the one whose mode span is real.

test_dsl_sm100_band_right_uncovered_tail_rejected called fwd_engines.probe().
That test arrived with #485, which this branch rebased onto after probe() was
already deleted here -- so it is a caller that did not exist when the deletion
was written, and it would have taken out the whole Blackwell L0 suite with an
AttributeError before reaching its assertion.

* Address the CodeRabbit pass: an id names one engine, and two tests could not fail

engine_for_id() matched an id exactly while _owners_for_id() matched a RANGE,
so a replay could resolve one way for a candidate engine and another way on a
fresh graph. Collapsed the other way from what was suggested: BaseEngine.id_end
and owned_id_range are deleted and _owners_for_id is an equality test. The range
existed so a REGISTERED engine could claim a block and registration could prove
two blocks disjoint; nothing registers now, no shipped engine ever set id_end,
and every range was [engine_id, engine_id + 1). Keeping it would have spread
dead machinery to fix an asymmetry that only that machinery created.
EngineFamily.id_end -- the family's block -- is a different thing and stays.

test_ranking_and_engine_read_the_same_record declared a probe_family by hand and
then called _offer(), whose own monkeypatch of MANIFEST won; the surviving
family had no analyzer. It passed anyway because both sides call
_facts_for(_probe_analyzer) directly, so the documented claim -- that the
ranking resolves the analyzer from EngineFamily.analyzer -- went unexercised.
Now declared through _offer, and it asserts the analyzer already ran BEFORE
ranking, which is the part only planning can do. Verified by mutation: drop the
analyzer declaration and the test fails.

select_engine(tiles=) matched the rendered plan name by substring, so a request
for tile_n=128 could select a tile_n=1280 plan and the test would pass having
run something else. Matches PlanConfig.knobs structurally now. No caller on this
branch -- the fp8 SM120 tile tests in #509 are the first, and they would have
been the ones to hit it.

Plus a cross-reference to a test renamed in this PR.

204 passed on the CPU suites; the one failure is the pre-existing
test_a_replayed_plan_reports_its_own_notes.

* License the one file this PR adds as Apache-2.0, not MIT

sdpa/fwd/heuristics.py was created by copying the header from config_sm120.py,
which is MIT -- so the new file inherited a tag that does not apply to it.

Per LICENSING.md, the repo relicensed MIT -> Apache-2.0 in #408 and a file is
kept under MIT for exactly two reasons: surviving lines from an external
contributor who has not consented to relicensing, or derivation from
third-party source. A file written from scratch at NVIDIA has neither, so
Apache-2.0 is the correct tag -- as it already is on every other file this
change adds content to (engines/heuristics.py, engines/manifest.py,
sdpa/fwd/engine.py, and the tests). The MIT neighbours in sdpa/fwd are
pre-existing files this PR only edits, and editing does not move a file
between licenses.

Also switches to the SPDX-FileCopyrightText form the Apache-2.0 files use.

* Compress the comments this PR adds

Six blocks broke the house rule that a call site explains only the non-obvious
load-bearing fact and rationale/measurements go in the MR description -- which
is where all of this already was, so it was duplicated, not lost.

Cut: measurement detail from _sm120_tiles (1.5x at 64 CTAs, 240-vs-320 CTAs,
2-4%, 106 KB vs 99) down to a pointer at PR #528, keeping the two thresholds a
reader needs and the warning that the rule is kernel-specific. _MEASURED_BEHIND
lost two antitheses ("deliberate but NOT a measurement", "an experiment, not an
edit") and a cross-reference the module docstring already makes.
_owners_for_id, _create_backend_plans and the backend-dedup comment lost
restated clauses. BaseEngine's note on the deleted id range stopped narrating
the deletion -- that belongs to the commit that made it, where it is verbatim.

31 fewer added lines, 8 fewer comment lines; no claim, number or caveat
dropped, only relocated to where it was already written.

199 passed; the one failure is the pre-existing
test_a_replayed_plan_reports_its_own_notes.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

cat-feature Requests for new functionality, APIs, examples, or behavior improvements. mod-frontend cuDNN frontend APIs, operation graph construction, plans, and user-facing wrappers. orig-nv-eng Reported or requested by NVIDIA engineering.

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants