[Copilot]

Description

• Register BFloat16 for the GatedDeltaNet schema and CUDA kernel.
• Route BFloat16 execution through the recurrent engine while retaining FP16-only tensor-core paths.
• Add BFloat16 numerical and planner coverage.
• Guard mma.sync.m16n8k16 generation to SM80+, allowing pre-Ampere builds to compile.
• Update operator documentation.

Motivation and Context

GatedDeltaNet supported float and float16 CUDA inputs but not BFloat16. Its unguarded SM80 MMA instruction also caused ptxas failures when compiling for older targets such as compute_61.

[azure-pipelines]

Azure Pipelines:
There may be pipelines that require an authorized user to comment /azp run to run.

[Copilot]

Pull request overview

Adds CUDA BFloat16 support for GatedDeltaNet while retaining FP16-only tensor-core engines.

Changes:

• Registers and tests BFloat16 recurrent execution.
• Guards SM80 MMA assembly for older CUDA targets.
• Updates schemas and operator documentation.

Reviewed changes

Copilot reviewed 10 out of 10 changed files in this pull request and generated 1 comment.

Show a summary per file

File	Description
onnxruntime/test/contrib_ops/gated_delta_net_op_test.cc	Adds BFloat16 numerical and planner tests.
onnxruntime/core/graph/contrib_ops/bert_defs.cc	Adds BFloat16 to the schema.
onnxruntime/contrib_ops/cuda/cuda_contrib_kernels.cc	Registers the BFloat16 kernel.
onnxruntime/contrib_ops/cuda/bert/gated_delta_net.cc	Maps and dispatches BFloat16 inputs.
onnxruntime/contrib_ops/cuda/bert/gated_delta_net_split_impl.cu	Rejects unsupported BFloat16 split execution.
onnxruntime/contrib_ops/cuda/bert/gated_delta_net_mma.cuh	Guards SM80-only MMA assembly.
onnxruntime/contrib_ops/cuda/bert/gated_delta_net_impl.cu	Instantiates recurrent BFloat16 execution.
docs/OperatorKernels.md	Updates generated kernel documentation.
docs/ContribOperators.md	Updates generated schema documentation.
docs/contrib_ops/cuda/gdn.md	Documents BFloat16 constraints.

---

💡 Configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

[Tianlei Wu (tianleiwu)]

Nice, tightly-scoped enablement. Verified locally on SM90: the branch builds clean and all 36 GatedDeltaNet* / GatedDeltaNetPlanTest* tests pass.

A couple of things I checked and liked:

• The ToCudaType -> OrtToCudaType switch is required, not cosmetic: ToCudaType<BFloat16>::MappedType is ORT's own BFloat16, while OrtToCudaType<BFloat16>::type is __nv_bfloat16, which is what VariantPack<CudaT> and the (float) conversions in the kernels need. MLFloat16/float map identically in both traits, so existing types are unaffected.
• No planner change was needed because IoType::kBFloat16, type_ok = (io_type == kFloat16) and the "chunked engine requires float16 input" reject reason already existed. That means ORT_GDN_PLAN=chunked with bf16 fails with an actionable message instead of silently degrading. Good reuse.
• ToTensorType<T> is a net simplification of RunTypedCase, and the planner test is host-only so the routing contract is asserted even where the numeric tests skip.

Main concern

I agree with the existing thread on gated_delta_net_mma.cuh: the #else branch makes the tensor-core GEMM a silent no-op rather than a hard failure. Details and a suggested fix are in that thread.

Other notes

Inline comments below cover:

1. bf16 instantiating a GatedDeltaNetChunkedKernel that SelectPlan can never select (measured +14.6% object size).
2. The float64 reference being computed from unrounded fp32 inputs, so the 4e-2 tolerance mostly measures bf16 input quantization rather than kernel error, plus the bf16 coverage gaps.
3. A nitpick on the nested ternary for desc.io_type.

None of 1-3 are blocking.