[bartekxk]

Motivation

Add bf16 support for direct conv.

Technical Details

Extend direct conv implementation with bf16 instructions

Test Plan

test_*_conv_corectness.py

Test Result

passed

Submission Checklist

• Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.

JIRA ID: AICK-1781

[therock-pr-bot]

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🔎 pre-commit	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—
🤖 therock-pr-bot	✅ Pass	—

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

[therock-pr-bot]

🎉 All checks passed! This PR is ready for review.

[Copilot]

Copilot review overview

🟡 Changes recommended

The Python direct-conv kernel builder contains a mismatched helper call that will raise a TypeError during kernel construction (runtime-k-loop path).

Get a fresh assessment by requesting another Copilot review.

Review effort: Lite
Findings: 1 · 1

Open (2)

• Remove extra io_type argument from _buf_load_vN call · New
• Add direct-conv bf16 correctness coverage · New

What changed in this PR

Adds bf16 I/O support for the grouped direct-convolution kernel family in rocke, updating both the Python and C++ engines to select bf16 buffer ops, mfma atoms, kernel naming flags, and spec validation. This fits into the rocke “byte-identity” dual-engine design by extending the same kernel family across both emitters.

Changes:

• Add dtype to the direct-conv problem/specs and gate kernels/validation on "fp16" vs "bf16".
• Teach Python + C++ direct-conv builders to emit bf16 buffer load/store + bf16 MFMA ops and bf16-aware kernel names.
• Update the direct-conv benchmark CLI/parsing to run bf16 and reject unsupported fp32/int8 modes.

File	Description
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​instances/​common/​conv_direct_grouped_spec_and_validate.cpp	Adds dtype defaulting, validation, bf16 kernel-name flags, and bf16 MFMA capability checks; introduces dtype-aware signature helper.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​instances/​common/​conv_direct_grouped_build_8c.cpp	Switches 8c builder I/O types and ops to bf16 when requested (loads/stores/MFMA).
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​instances/​common/​conv_direct_grouped_build_32c.cpp	Switches 32c builder I/O types and ops to bf16 when requested (loads/stores/MFMA).
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​instances/​common/​conv_direct_grouped_build_16c_phases.cpp	Switches 16c builder prologue/weights/LDS helpers to dtype-selected IR types.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​instances/​common/​conv_direct_grouped_build_16c_loop_closures.cpp	Switches 16c loop closures to dtype-dispatched buffer loads + generic LDS ops and bf16 MFMA/store paths.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​include/​rocke/​instance_conv_direct_grouped.h	Documents dtype support and adds rocke_direct_conv_signature_for_dtype() declaration.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​include/​rocke/​instance_conv_direct_grouped_internal.h	Extends direct-conv builder contexts with io_type for dtype-selected IR typing.
dnn-providers/​hip-kernel-provider/​rocke/​library/​kernels/​common/​conv_direct_grouped.py	Adds bf16 dtype plumbing, dtype-dispatched buffer ops/MFMA selection, and bf16 naming/validation across direct-conv variants.
dnn-providers/​hip-kernel-provider/​rocke/​library/​benchmarks/​common/​benchmark_direct_conv.py	Adds bf16 benchmark execution via --dtype and stricter miopen driver dtype handling.

---

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

[cgmillette]

Review: direct conv bf16 support - Claude code 4.7

NO blocking issues. The notes below are mostly cleanup, plus one behavior change worth a second look.

Legend: 🟡 = worth fixing before/with merge · ⚪ = minor / optional

#	Where	What's the problem (plain English)	Why it matters	Suggestion
1 🟡	benchmarks/common/benchmark_direct_conv.py (~L120–127)	fp32 and convint8 now raise an error. Previously fp32 was accepted and convint8 fell back to fp16 with a warning.	Any existing script/command passing conv (fp32) or convint8 will now hard-fail instead of running.	Confirm nothing relies on the old behavior and note the change in the PR — or warn instead of erroring for one release.
2 🟡	kernels/common/conv_direct_grouped.py — _buf_store_vN (L112)	A new store helper was added but is never called anywhere (the matching load helper _buf_load_vN is used 15×).	Dead code is confusing and looks like a half-finished refactor.	Use it at the store sites, or delete it.
3 🟡	conv_direct_grouped.py (~9 spots: L814, L1627, L2069, L2949, …)	The "if bf16 use the _bf16 op, else the _f16 op" choice is hand-written at every MFMA/store site.	Easy to miss or flip one branch and silently emit the wrong precision.	Add one small dispatch helper (like _buf_load_vN) and call it everywhere.
4 🟡	platform/cpp/instances/common/conv_direct_grouped_build_*.cpp (~15+ spots)	Same copy-paste, in C: if (dtype && strcmp(dtype,"bf16")==0) … else … at every load/store/MFMA. io_type is already resolved once in each prologue, but each site re-checks the string.	Fragile — one mistyped/missing check silently emits wrong-precision math.	Compute a single is_bf16 (or branch on the already-resolved io_type) once and reuse.
5 🟡	include/rocke/instance_conv_direct_grouped.h — rocke_direct_conv_problem_t.dtype	dtype was appended as the last struct field. A caller that builds the struct without rocke_direct_conv_problem_default() (e.g. zero-init) gets dtype == NULL, which silently means fp16.	A bf16 request could be silently dropped to fp16; it's also a struct-layout/ABI change.	Make sure all construction sites set dtype or go through _default(); call out the layout change.
6 ⚪	conv_direct_grouped.py (L86–112)	Three new helper functions are defined in the middle of the import block.	Style/readability.	Move them below all imports.
7 ⚪	conv_direct_grouped_build_8c.cpp (f16ptr, lds_total_fp16, fp16x4_zero)	Names still say "fp16" even though they now hold bf16 too. (build_16c_phases.cpp already renamed to io_ptr.)	Misleads future readers about element width.	Rename to dtype-neutral, like 16c_phases did.
8 ⚪	conv_direct_grouped.py (L986); benchmark (L83, L86)	A new error message and comments use a non-ASCII em-dash.	rocKE convention is ASCII.	Use a plain -.
9 ⚪	benchmark_direct_conv.py (L83/86, L207)	The map values "conv":"fp32" / "convint8":"int8" are never used (they raise first), and the else "fp16" at L207 is unreachable.	Dead/misleading code.	Remove the unused entries/branch.
10 ⚪	tests/test_direct_conv_correctness.py	There's no test asserting cpg=4 + bf16 is rejected (only a docstring note); and the _TOL_BF16 comment "~3× less mantissa" is imprecise (bf16 has 3 fewer mantissa bits than fp16, i.e. ~8× coarser).	Small coverage/accuracy gaps.	Add a tiny negative test; fix the comment.

-Claude Code 4.7

[aledudek]

one comment otherwise LGTM

[cgmillette]

LGTM

[assistant-librarian]

TheRock Submodule Bump Activity

Newest first

• 2026-09-25 — Included in TheRock via ROCm/TheRock#8472.