Skip to content

fix(tensilelite): fail the benchmark client on GPU errors during timing - #13117

Merged
caio96 merged 1 commit into
developfrom
users/csalvado/tensilelite-exit-on-gpu-fault
Oct 8, 2026
Merged

caio96 merged 1 commit into
developfrom
users/csalvado/tensilelite-exit-on-gpu-fault

Conversation

@caio96

@caio96 caio96 commented Oct 5, 2026 •

Copy link
Copy Markdown
Contributor

Motivation

If a kernel crashed on the GPU during benchmarking, the TensileLite client did not notice. It reported meaningless timings and exited successfully, so Tensile showed the run as passing.

Technical Details

  • BenchmarkTimer: errors from the HIP event calls used for timing were ignored. They are now checked, so a GPU error fails the solution being timed.
  • main.cpp: after a solution fails, the client checks whether the GPU can still be used. If it can't, the client stops and exits with a nonzero code instead of running the remaining solutions, which would all fail or crash. Other failures (for example, running out of memory) still continue as before.

Test Plan

Ran Tensile with a kernel known to crash on the GPU for one problem size:

  • that size alone
  • that size followed by a size that works
  • a size that works, as a control

Test Result

  • Crashing size: before, the client exited 0 and Tensile reported PASS. Now the client exits 1 and Tensile reports a failure.
  • Crashing size followed by a good one: the client stops after the crash with a clear error and exits 1. Without the early stop, it crashed (SIGABRT) on the next problem.
  • Good size: unchanged, passes with normal timings.

Submission Checklist

Risk level

Low. This only affects the benchmark client, and only when a GPU error happens. Runs without GPU errors behave the same.

JIRA ID : AIHPBLAS-5062

The event-based timer ignored HIP return codes, so a kernel fault went
unnoticed and the client exited 0 with bogus timings. Check them, and
stop the run once the device is unusable instead of continuing on a
dead context.
@caio96
caio96 requested a review from a team as a code owner October 5, 2026 21:22
@therock-pr-bot

therock-pr-bot Bot commented Oct 5, 2026 •

Copy link
Copy Markdown

⏳ Waiting on Required Checks

⏳ Required checks have not all concluded yet. This bot run is exiting without failing so those checks can report independently.

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ⚠️ Warning Error: Source/code files changed without an accompanying unit test.
Expected: add at least one test file named like test_<name>.py / test_<name>.cpp (or <name>_test.*).
Current: code file(s) changed: projects/hipblaslt/tensilelite/client/main.cpp, projects/hipblaslt/tensilelite/client/src/BenchmarkTimer.cpp; no test file found
🔎 pre-commit ⏳ Pending ⏳ Still running…
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —

⏳ Required checks are still running or have not started yet.
This PR is not ready for review until they complete.

📖 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 Oct 5, 2026 •

Copy link
Copy Markdown

⏳ Waiting for required checks to start or finish:

  • ⏳ pre-commit

@codecov

codecov Bot commented Oct 5, 2026 •

Copy link
Copy Markdown

Codecov Report

❌ Patch coverage is 0% with 8 lines in your changes missing coverage. Please review.

Files with missing lines Patch % Lines
...ipblaslt/tensilelite/client/src/BenchmarkTimer.cpp 0.00% 8 Missing ⚠️

❌ Your project check has failed because the head coverage (49.02%) is below the target coverage (80.00%). You can increase the head coverage or adjust the target coverage.

Additional details and impacted files
@@           Coverage Diff            @@
##           develop   #13117   +/-   ##
========================================
  Coverage    72.12%   72.12%           
========================================
  Files         2861     2861           
  Lines       472095   472095           
  Branches     69805    69813    +8     
========================================
+ Hits        340452   340453    +1     
  Misses      107911   107911           
+ Partials     23732    23731    -1     
Flag Coverage Δ *Carryforward flag
TensileLite-CPP 49.02% <0.00%> (+0.01%) ⬆️
TensileLite-Unit 83.28% <ø> (ø)
hipBLAS 90.62% <ø> (ø) Carriedforward from efe215b
hipBLASLt 38.65% <ø> (ø)
hipCUB 82.68% <ø> (ø) Carriedforward from efe215b
hipDNN 86.64% <ø> (ø) Carriedforward from efe215b
hipFFT 54.32% <ø> (ø) Carriedforward from efe215b
hipRAND 76.12% <ø> (ø) Carriedforward from efe215b
hipSOLVER 68.96% <ø> (ø) Carriedforward from efe215b
hipSPARSE 86.99% <ø> (ø) Carriedforward from efe215b
rocBLAS 48.25% <ø> (ø) Carriedforward from efe215b
rocFFT 48.78% <ø> (ø) Carriedforward from efe215b
rocRAND 57.45% <ø> (ø) Carriedforward from efe215b
rocSOLVER 76.83% <ø> (ø) Carriedforward from efe215b
rocSPARSE 74.75% <ø> (ø) Carriedforward from efe215b
rocThrust 93.15% <ø> (ø) Carriedforward from efe215b

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

Files with missing lines Coverage Δ
...ipblaslt/tensilelite/client/src/BenchmarkTimer.cpp 0.00% <0.00%> (ø)

... and 1 file with indirect coverage changes

🚀 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.

@jaopaulolc jaopaulolc added the rocm:gemm algos Label to identify PRs of the GEMM Algorithms team. label Oct 6, 2026
@caio96
caio96 merged commit 895474b into develop Oct 8, 2026
800 of 846 checks passed
@caio96
caio96 deleted the users/csalvado/tensilelite-exit-on-gpu-fault branch October 8, 2026 14:57
@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.

3 participants