[hadamard-hip] Use type-dependant tolerance for validation - #277
[hadamard-hip] Use type-dependant tolerance for validation#277maarquitos14 wants to merge 1 commit into
Conversation
|
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. |
|
I'm observing the issue on gfx90a using -O1, both for __half, at hidden dimension 512 and 32768. See output below. 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. |
|
With -O1, the gfx90a assembly similarly contains: v_fma_mixlo_f16 # incorrect direct rounding The selection rule is in llvm/lib/Target/AMDGPU/VOP3PInstructions.td: it maps fpround(fmul(f32, f32)) directly to V_FMA_MIXLO_F16, without checking floating-point permissions. The fix should remove/restrict the fpround(fmul) patterns; strict fmul + fptrunc must use v_mul_f32 + v_cvt_f16_f32. The mix instruction remains valid for an explicit fused operation. The tolerance improvement is still reasonable for a reduced-precision benchmark, but it may identify a real compiler miscompile. When the tolerance is type dependent, we probably would not look into assembly and find out the cause. I appreciate your explanation of the errors. |
2 participants
The current tolerance is too restrictive for smaller types like
halfor__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.