Skip to content

[hadamard-hip] Use type-dependant tolerance for validation#277

Open
maarquitos14 wants to merge 1 commit into
ORNL:masterfrom
maarquitos14:maronas/hadamard-hip-fix-validation
Open

[hadamard-hip] Use type-dependant tolerance for validation#277
maarquitos14 wants to merge 1 commit into
ORNL:masterfrom
maarquitos14:maronas/hadamard-hip-fix-validation

Conversation

@maarquitos14

Copy link
Copy Markdown
Contributor

The current tolerance is too restrictive for smaller types like half or __hip_bfloat16, causing validation errors for some optimization sets --e.g. gfx90a -O1. The current patch makes the tolerance type dependant, so the smaller types are more lenient, while the bigger types are still quite restrictive.

@zjin-lcf

zjin-lcf commented Jul 6, 2026

Copy link
Copy Markdown
Collaborator

Are there validation errors for "gfx90a -O3" in your test ? I ran the program on gfx90a and gfx942 devices and did not see errors. I assume this is a compiler issue instead of a benchmark issue. Am I right ? Thanks.

gfx90a -O3

./main 1024 100
batch size: 1024 | hidden dimension: 8 | Average kernel execution time : 3.842480 (us)
PASS
batch size: 1024 | hidden dimension: 8 | Average kernel execution time : 2.644990 (us)
PASS
batch size: 1024 | hidden dimension: 8 | Average kernel execution time : 3.080890 (us)
PASS
batch size: 1024 | hidden dimension: 64 | Average kernel execution time : 4.594580 (us)
PASS
batch size: 1024 | hidden dimension: 64 | Average kernel execution time : 3.545590 (us)
PASS
batch size: 1024 | hidden dimension: 64 | Average kernel execution time : 3.798980 (us)
PASS
batch size: 1024 | hidden dimension: 512 | Average kernel execution time : 7.335861 (us)
PASS
batch size: 1024 | hidden dimension: 512 | Average kernel execution time : 7.396870 (us)
PASS
batch size: 1024 | hidden dimension: 512 | Average kernel execution time : 7.832371 (us)
PASS
batch size: 1024 | hidden dimension: 4096 | Average kernel execution time : 36.296841 (us)
PASS
batch size: 1024 | hidden dimension: 4096 | Average kernel execution time : 30.766462 (us)
PASS
batch size: 1024 | hidden dimension: 4096 | Average kernel execution time : 31.833252 (us)
PASS
batch size: 1024 | hidden dimension: 32768 | Average kernel execution time : 964.972961 (us)
mPASS
akbatch size: 1024 | hidden dimension: 32768 | Average kernel execution time : 709.491638 (us)
e PASS
batch size: 1024 | hidden dimension: 32768 | Average kernel execution time : 1020.508423 (us)
PASS
./main 32768 100
batch size: 32768 | hidden dimension: 8 | Average kernel execution time : 40.878223 (us)
PASS
batch size: 32768 | hidden dimension: 8 | Average kernel execution time : 40.697422 (us)
PASS
batch size: 32768 | hidden dimension: 8 | Average kernel execution time : 41.196011 (us)
PASS
batch size: 32768 | hidden dimension: 64 | Average kernel execution time : 43.417702 (us)
PASS
batch size: 32768 | hidden dimension: 64 | Average kernel execution time : 42.325809 (us)
PASS
batch size: 32768 | hidden dimension: 64 | Average kernel execution time : 42.738903 (us)
PASS
batch size: 32768 | hidden dimension: 512 | Average kernel execution time : 107.993431 (us)
PASS
batch size: 32768 | hidden dimension: 512 | Average kernel execution time : 101.334152 (us)
PASS
batch size: 32768 | hidden dimension: 512 | Average kernel execution time : 102.482841 (us)
PASS
batch size: 32768 | hidden dimension: 4096 | Average kernel execution time : 894.906738 (us)
PASS
batch size: 32768 | hidden dimension: 4096 | Average kernel execution time : 772.407288 (us)
PASS
batch size: 32768 | hidden dimension: 4096 | Average kernel execution time : 799.625488 (us)
PASS
batch size: 32768 | hidden dimension: 32768 | Average kernel execution time : 30800.710938 (us)
PASS
batch size: 32768 | hidden dimension: 32768 | Average kernel execution time : 22580.835938 (us)
PASS
batch size: 32768 | hidden dimension: 32768 | Average kernel execution time : 32403.144531 (us)
PASS

@maarquitos14

Copy link
Copy Markdown
Contributor Author

I'm observing the issue on gfx90a using -O1, both for __half, at hidden dimension 512 and 32768. See output below.

batch size: 32768 | hidden dimension: 512 | Average kernel execution time : 167.719681 (us)
Mismatch at index 3507200 11.601562 11.593750
FAIL

batch size: 32768 | hidden dimension: 32768 | Average kernel execution time : 30491.154297 (us)
Mismatch at index 476938240 90.062500 90.125000
FAIL

If we look closer at the mismatch, for both it's exactly 1 ULP, which is the minimum difference we can find for a given format. However, the current tolerance is lower than 1 ULP for these two cases, so it fails validation.

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

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants