Skip to content

test(miopen): add CTC loss numel>INT_MAX scale coverage (ALMIOPEN-2148) - #10252

Open
peymanr wants to merge 4 commits into
ROCm:developfrom
peymanr:users/prazaghi/almiopen-2148-ctc-scale
Open

test(miopen): add CTC loss numel>INT_MAX scale coverage (ALMIOPEN-2148)#10252
peymanr wants to merge 4 commits into
ROCm:developfrom
peymanr:users/prazaghi/almiopen-2148-ctc-scale

Conversation

@peymanr

@peymanr peymanr commented Jul 31, 2026

Copy link
Copy Markdown

Summary

Add a new gtest file projects/miopen/test/gtest/ctc_large.cpp with a Scale-tier test that exercises MIOpen CTC loss at logit numel > INT_MAX.

Motivation

The existing gtest/ctc.cpp cases top out at approximately 64 M elements (batch=128, T=100, C=5001). MIOpen's CTC kernel accumulates the logit tensor dimensions through plain int locals in ctc.cpp:

int problog_offset = lb_prime_offset + batch_size * max_S_len;
int alpha_offset   = problog_offset + class_sz * batch_size * max_time_step;

When T * N * C exceeds INT_MAX (2 147 483 647), these multiplications wrap to negative values, producing out-of-bounds workspace reads/writes and silently wrong losses. Large-vocabulary or long-sequence ASR models (e.g. T=1000, vocab=50k, batch=128) are in this regime.

JIRA ID: ALMIOPEN-2148

What the test does

Parameter Value Notes
T (time steps) 1000 Matches ASR frame budget
N (batch) 2048 Large batch to cross the boundary
C (classes, excl. blank) 1048 class_sz = C+1 = 1049
numel = T × N × (C+1) 2 148 352 000 > INT_MAX = 2 147 483 647
Label length 2 per sample Kept small so alpha workspace is bounded
Input length 10 per sample Actual used time steps
  • Queries hipMemGetInfo and calls GTEST_SKIP() when free GPU memory < 19 GiB; safe to run on smaller devices.
  • Allocates logits/gradients/losses directly on GPU (no 8.59 GiB host buffer).
  • Probs initialised with hipMemset(0) — uniform logits give a valid softmax input.
  • Asserts every returned loss is finite. NaN or Inf signals integer overflow in the workspace offset computation.
  • No CPU reference comparison (host-side 8.59 GiB allocation is prohibitive).

MIOpenDriver equivalent for manual verification on TW325 (gfx942):

MIOpenDriver ctc -T 1000 -N 2048 -C 1048 --forw 1 -V 0 -t 1

Test tier

The test uses a TEST_F fixture with a Scale_GPU_CTC_numel_gt_INT_MAX case name. It is picked up automatically by file(GLOB TESTS *.cpp) in the gtest CMakeLists.txt — no CMakeLists.txt modification required. The Scale_ prefix and the memory guard prevent accidental execution in Smoke/Standard CI.

Checklist

  • New file only; no existing file modified
  • GTEST_SKIP() guard for insufficient GPU memory
  • static_assert confirms numel > INT_MAX at compile time
  • JIRA ID: ALMIOPEN-2148 included in commit message and PR body

🤖 Assisted by PR Pundit

Add projects/miopen/test/gtest/ctc_large.cpp with a single Scale-tier
test that runs MIOpen CTC loss forward pass on a shape where
T*N*C = 1000*2048*1049 = 2,148,352,000 > INT_MAX (2,147,483,647).

The existing gtest/ctc.cpp cases top out at ~64 M elements.  At numel >
INT_MAX the plain `int` dimension products inside ctc.cpp
(class_sz * batch_size * max_time_step) overflow, producing negative
workspace offsets and silent out-of-bounds access in the CTC kernel.

The new test:
- Skips gracefully via GTEST_SKIP() when free GPU memory < 19 GiB
  (logits alone require ~8.59 GiB float32; workspace adds ~1 GiB)
- Allocates logits / gradients / losses directly on GPU to avoid an
  equivalent host allocation
- Initialises probs with hipMemset(0) — uniform logits are valid CTC
  inputs (softmax(0,...,0) = uniform distribution)
- Asserts every loss value is finite; NaN/Inf flags index overflow
- No CPU reference comparison (host 8.59 GiB buffer is prohibitive)

Suggested MIOpenDriver equivalent for manual verification:
  MIOpenDriver ctc -T 1000 -N 2048 -C 1048 --forw 1 -V 0 -t 1

JIRA ID: ALMIOPEN-2148
@peymanr
peymanr requested a review from a team as a code owner July 31, 2026 23:18
@therock-pr-bot

therock-pr-bot Bot commented Jul 31, 2026

Copy link
Copy Markdown

✅ All Checks Passed — Ready for Review

Check Status Details
📝 PR Description ✅ Pass
Forbidden Files ✅ Pass
🧪 Unit Test ✅ Pass PR does not contain code files — Unit Test auto-passed
🔎 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

Copy link
Copy Markdown

Pre-commit check failed

pre-commit failed

Please run locally:

  • python -m pip install pre-commit
  • pre-commit install
  • pre-commit run --all-files --show-diff-on-failure

This repo uses .pre-commit-config.yaml.

Fix clang-format violations flagged by pre-commit CI.
Resolves: ALMIOPEN-2148
@therock-pr-bot

therock-pr-bot Bot commented Aug 1, 2026

Copy link
Copy Markdown

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

peymanr added 2 commits July 31, 2026 17:17
The three hipFree calls in the cleanup section of Run() were
silently discarding hipError_t (declared nodiscard), causing
-Werror,-Wunused-value build failures. Wrap with (void) cast
to explicitly acknowledge the return value in teardown context.

Fixes: MIOpen Build Hip Package CI failure (ALMIOPEN-2148).
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.

1 participant