Skip to content

feat(rocblas): Kernel Selection trace logging option - #8070

Merged
evedovelli merged 18 commits into
developfrom
users/evedovel/rocblas/kernel-selection-trace
Oct 29, 2026
Merged

evedovelli merged 18 commits into
developfrom
users/evedovel/rocblas/kernel-selection-trace

Conversation

@evedovelli

@evedovelli evedovelli commented Jun 4, 2026 •

Copy link
Copy Markdown
Contributor

Motivation

When a Level 3 call (trsm, trmm, syrk, ...) is decomposed into multiple GEMM sub-problems, there was no easy way to see which kernel ran for each one or to reproduce it in isolation. TENSILE_DB output doesn't cover hipBLASLt and doesn't map back to a benchmarkable command.

This PR adds ROCBLAS_LAYER=0x10 (rocblas_layer_mode_log_kernel_select), which writes one rocblas-bench-replayable line per internal GEMM sub-problem that launches a kernel.

ISSUE ID : AIROCBLAS-1288

Technical Details

  • New layer bit rocblas_layer_mode_log_kernel_select = 0x10 in rocblas-types.h. Output goes to the trace stream (ROCBLAS_LOG_TRACE_PATH).
  • Each line is a rocblas-bench command followed by a # comment:
    • The bench function is gemm_ex, gemm_batched_ex, or gemm_strided_batched_ex, matching the sub-problem layout. Strides are emitted only for strided-batched, --batch_count only for batched forms, and --api 1 when a dimension exceeds 32 bits.
    • source=: hipblaslt, tensile, source (rocBLAS source kernels), or gemv (m == 1 or n == 1 GEMM run as GEMV).
    • kernel=: the Tensile solution name or the source/GEMV launcher. Not reported for hipBLASLt.
    • fallback_from=: the backends that rejected the sub-problem first, e.g. hipblaslt, hipblaslt,tensile, or none.
    • parent_api=: the user-facing API, set by an RAII scope in each _impl.
  • Logged paths: hipBLASLt/Tensile (tensile_host.cpp), gemm and gemm_ex source fallback and scale path (gemm_templates.cpp, rocblas_gemm_ex_kernels_64.cpp), the 64-bit large-dimension source path (rocblas_gemm_kernels_64.cpp), and GEMV-in-GEMM (rocblas_gemm_ex.hpp).
  • Nothing is logged for fitness queries, device-memory-size queries, or rocblas_gemm_flags_check_solution_index, since no kernel launches.
  • All lines are formatted by one function, rocblas_internal_log_kernel_select() in logging.cpp, from a rocblas_kernel_select_problem description.
  • parent_api is set for gemm, hemm, herk, her2k, herkx, symm, syrk, syr2k, syrkx, trmm, trsm, trtri (all batched forms), gemm_ex variants, gemmt, syrk_ex, herk_ex, gemm_grouped_batched, and gemm_grouped_batched_ex.
  • The scope helper and per-call state are no-ops unless bit 0x10 is set.
  • Docs: docs/how-to/logging-in-rocblas.rst, docs/data/reference/logging-env.rst. CHANGELOG entry under 5.8.0.

How to use it

Turn it on:

  export ROCBLAS_LAYER=0x10
  export ROCBLAS_LOG_TRACE_PATH=./kernel_select.log   # without this, lines go to stderr
  ./your_app

You can combine it with other logging bits. For example, ROCBLAS_LAYER=0x18 adds internal backend logging (0x8) to the same file.

Each GEMM that runs produces one line:

./rocblas-bench -f gemm_ex --transposeA N --transposeB N -m 4096 -n 4096 -k 4096 --alpha 1 ... --flags 0 # source=tensile kernel=Cijk_Ailk_Bljk_SB_MT128x128x8 fallback_from=hipblaslt parent_api=rocblas_strsm

Reproduce one GEMM on its own by copying everything before the # and run it from the directory that contains rocblas-bench.

Useful one-liners:

grep -c "# source=" kernel_select.log                                  # how many GEMMs ran
grep -o "parent_api=[a-z_0-9]*" kernel_select.log | sort | uniq -c     # GEMMs per API call
grep -o "source=[a-z]*" kernel_select.log | sort | uniq -c             # which backends ran them
grep "fallback_from=[^n]" kernel_select.log                            # GEMMs where a backend fell back
grep "parent_api=rocblas_strsm" kernel_select.log | cut -d'#' -f1 > trsm_gemms.sh   # replayable commands for one API

Known limitations

  • hipBLASLt lines do not include a kernel identifier.
  • Non-GEMM kernel-selection steps do not emit lines yet.

Test Plan

  • logging_mode_gtest gains seven logging_kernel_select_* cases (category quick; harness in clients/include/testing_kernel_select.hpp):
    • trsm sub-problems all report parent_api=rocblas_strsm
    • gemm / gemm_batched / gemm_strided_batched replay as gemm_ex / gemm_batched_ex / gemm_strided_batched_ex
    • 0x8 and 0x10 combined on one stream
    • forced source GEMM (solution_index -2) reports source=source
    • forced GEMV (solution_index -9) reports source=gemv
    • Tensile with hipBLASLt disabled reports a Cijk_ kernel
    • check_solution_index logs nothing
  • The existing logging test includes a 0x10 smoke check for each real precision, pointer mode, and C/C_64 API.
  • Lines are checked for schema and path-specific fields, not golden-filed, since kernel names depend on arch and build.

Test Result

  • gfx90a: 27/27 pass (20 logging + 7 kernel-select), before and after moving the tests into logging_mode_gtest.
  • gfx908: 27/27 pass, before and after moving the tests into logging_mode_gtest.
  • gfx942 / gfx950 / gfx11xx / gfx12xx: covered by CI.

Submission Checklist

@codecov-commenter

codecov-commenter commented Jun 4, 2026 •

Copy link
Copy Markdown

Codecov Report

❌ Patch coverage is 54.44126% with 159 lines in your changes missing coverage. Please review.

Files with missing lines Patch % Lines
...ts/rocblas/library/src/blas_ex/rocblas_gemm_ex.hpp 0.00% 73 Missing ⚠️
projects/rocblas/library/src/tensile_host.cpp 0.00% 54 Missing ⚠️
...ibrary/src/src64/blas3/rocblas_gemm_kernels_64.cpp 0.00% 24 Missing ⚠️
projects/rocblas/library/src/logging.cpp 88.64% 1 Missing and 4 partials ⚠️
projects/rocblas/library/src/include/handle.hpp 88.89% 0 Missing and 2 partials ⚠️
...cblas/library/src/blas3/Tensile/gemm_templates.cpp 96.55% 0 Missing and 1 partial ⚠️
Additional details and impacted files
@@             Coverage Diff             @@
##           develop    #8070      +/-   ##
===========================================
- Coverage    77.26%   77.25%   -0.00%     
===========================================
  Files         2892     2892              
  Lines       485320   485624     +304     
  Branches     73315    73351      +36     
===========================================
+ Hits        374940   375153     +213     
- Misses       85629    85714      +85     
- Partials     24751    24757       +6     
Flag Coverage Δ *Carryforward flag
TensileLite-CPP 50.81% <ø> (ø) Carriedforward from 69f1cdb
TensileLite-Unit 84.02% <ø> (ø) Carriedforward from 69f1cdb
hipBLAS 90.62% <ø> (ø) Carriedforward from 69f1cdb
hipBLASLt 68.47% <ø> (ø) Carriedforward from 69f1cdb
hipCUB 82.68% <ø> (ø) Carriedforward from 69f1cdb
hipDNN 86.86% <ø> (ø) Carriedforward from 69f1cdb
hipFFT 38.02% <ø> (ø) Carriedforward from 69f1cdb
hipRAND 76.12% <ø> (ø) Carriedforward from 69f1cdb
hipSOLVER 69.67% <ø> (ø) Carriedforward from 69f1cdb
hipSPARSE 87.03% <ø> (ø) Carriedforward from 69f1cdb
rocBLAS 68.96% <54.44%> (+0.01%) ⬆️
rocFFT 47.34% <ø> (ø) Carriedforward from 69f1cdb
rocRAND 57.31% <ø> (ø) Carriedforward from 69f1cdb
rocSOLVER 76.74% <ø> (ø) Carriedforward from 69f1cdb
rocSPARSE 83.54% <ø> (ø) Carriedforward from 69f1cdb
rocThrust 93.15% <ø> (ø) Carriedforward from 69f1cdb

*This pull request uses carry forward flags. Click here to find out more.

Files with missing lines Coverage Δ
...s/rocblas/library/include/internal/rocblas-types.h 0.00% <ø> (ø)
...cblas/library/src/blas3/Tensile/gemm_templates.hpp 96.36% <100.00%> (+2.82%) ⬆️
...las/library/src/blas3/rocblas_gemm_batched_imp.hpp 71.74% <100.00%> (+0.15%) ⬆️
...ary/src/blas3/rocblas_gemm_grouped_batched_imp.hpp 46.04% <100.00%> (+0.27%) ⬆️
...cts/rocblas/library/src/blas3/rocblas_gemm_imp.hpp 94.48% <100.00%> (+0.03%) ⬆️
...ary/src/blas3/rocblas_gemm_strided_batched_imp.hpp 95.43% <100.00%> (+0.02%) ⬆️
...las/library/src/blas3/rocblas_hemm_batched_imp.hpp 73.48% <100.00%> (+0.15%) ⬆️
...cts/rocblas/library/src/blas3/rocblas_hemm_imp.hpp 74.58% <100.00%> (+0.14%) ⬆️
...ary/src/blas3/rocblas_hemm_strided_batched_imp.hpp 72.86% <100.00%> (+0.13%) ⬆️
...as/library/src/blas3/rocblas_her2k_batched_imp.hpp 73.66% <100.00%> (+0.14%) ⬆️
... and 47 more
🚀 New features to boost your workflow:
  • ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.
  • 📦 JS Bundle Analysis: Save yourself from yourself by tracking and limiting bundle sizes in JS merges.

@evedovelli evedovelli changed the title [rocblas] Kernel Selection trace logging option feat(rocblas): Kernel Selection trace logging option Jul 6, 2026
@therock-pr-bot

therock-pr-bot Bot commented Jul 6, 2026 •

Copy link
Copy Markdown

✅ 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

therock-pr-bot Bot commented Jul 6, 2026 •

Copy link
Copy Markdown

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

@evedovelli evedovelli self-assigned this Jul 6, 2026
@evedovelli
evedovelli marked this pull request as ready for review July 6, 2026 21:21
@evedovelli
evedovelli requested review from a team as code owners July 6, 2026 21:21
@TorreZuk

TorreZuk commented Jul 6, 2026

Copy link
Copy Markdown
Contributor

Lots to discuss to that we can go into during meeting tomorrow

@github-actions

github-actions Bot commented Aug 2, 2026

Copy link
Copy Markdown
Contributor

This pull request has been inactive for 25 days and will be marked as stale.

If you would like to keep this PR open, please:

  • Add new commits
  • Add a comment explaining why it should remain open

This PR will be automatically closed in 5 days if no further activity occurs.

@github-actions github-actions Bot added Stale PR has no activity for 25+ days and removed Stale PR has no activity for 25+ days labels Aug 2, 2026
@evedovelli

Copy link
Copy Markdown
Contributor Author

Lots to discuss to that we can go into during meeting tomorrow

@TorreZuk I've pushed the changes we had discussed a few weeks ago: the kernel ID is not printed as a comment to the reproducible command line log anymore, and I've minimized the amount of code executed in the "regular" path (no trace logging enabled).

@evedovelli
evedovelli requested a review from Ryker0627 August 19, 2026 14:34
@github-actions

Copy link
Copy Markdown
Contributor

This pull request has been inactive for 25 days and will be marked as stale.

If you would like to keep this PR open, please:

  • Add new commits
  • Add a comment explaining why it should remain open

This PR will be automatically closed in 5 days if no further activity occurs.

@github-actions github-actions Bot added Stale PR has no activity for 25+ days and removed Stale PR has no activity for 25+ days labels Sep 23, 2026

@TorreZuk TorreZuk left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

If tests pass looks good.
Added some optional things to consider.

@evedovelli
evedovelli merged commit 65d3c1d into develop Oct 29, 2026
234 of 238 checks passed
@evedovelli
evedovelli deleted the users/evedovel/rocblas/kernel-selection-trace branch October 29, 2026 21:07
@assistant-librarian

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

5 participants