warp_shuffle_op: avoid sub-int narrow↔wide roundtrip (IGC fp16 miscompile)#3
Merged
Conversation
3 tasks
pvelesko
added a commit
that referenced
this pull request
Apr 15, 2026
Adds a minimal regression test that exercises the rocprim::warp_shuffle_down chain on __half with a divergent consumer. This is the exact control-flow shape emitted by warp_reduce_shuffle for sizeof(T)<4 types: a pair of chained OpSubgroupShuffleINTEL ops on a value that round-trips through OpUConvert ushort<->uint, with the narrowed result consumed inside a block entered via OpBranchConditional. IGC on Intel Arc miscompiles this pattern: the second shuffle returns 0 instead of the neighbor-lane value, so reductions are off by the lost terms. With the current warp_shuffle_op implementation the test fails with 16/16 mismatches (one per logical warp). The test pairs with the rocprim.warp_reduce Floating/10 and /11 (ReduceSum on __half, LogicalWarpSize 4 and 8) failures.
…compile IGC on Intel Arc A770 miscompiles chained OpSubgroupShuffleINTEL when the shuffled value undergoes narrow->wide->shuffle->narrow (OpUConvert ushort<->uint around the shuffle) and the narrowed consumer sits in a block entered via OpBranchConditional. The second chained shuffle returns 0 instead of the neighbor lane, so rocPRIM warp_reduce_shuffle on __half with LogicalWarpSize < 32 produces systematically wrong sums. The miscompile is deterministic and reproducible with a pure Level Zero + SPIR-V assembler harness (see upstream IGC issue, to be filed). This patch adds a chipStar/SPIR-V-only warp_shuffle_op specialization for trivially-copyable T with sizeof(T) < sizeof(int) (covering __half and hip_bfloat16) that writes the post-shuffle int into a volatile stack slot and loads the low sizeof(T) bytes out via byte-wise memcpy from a volatile pointer. This emits OpBitcast + OpLoad of the narrow type rather than a scalar-register i32->i16 trunc, bypassing IGC's bad codegen path. Upstream rocPRIM / AMD targets are unaffected: the specialization is guarded on __HIP_PLATFORM_SPIRV__ and does not change the AMD code path. Resolves failures: - test_warp_reduce Floating/10,11 (__half ReduceSum) 2 -> 0 - test_warp_scan Floating/11,12,20,21,22 (__half/bf16) 7 -> 0 - test_block_reduce Floating/9 (block_params<__half,__half,64>) 1 -> 0 - test_warp_shuffle_fp16 (new regression test) FAIL -> PASS Upstream IGC bug to be filed at intel/intel-graphics-compiler.
pvelesko
force-pushed
the
fix-warp-shuffle-fp16
branch
from
April 16, 2026 09:53
bc73c2f to
b4d8bf0
Compare
pvelesko
marked this pull request as ready for review
April 17, 2026 05:38
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
IGC on Intel Arc A770 miscompiles
OpSubgroupShuffleINTELwhen all three conditions hold simultaneously:OpUConvert %uint <- %ushort(widen before shuffle)OpUConvert %ushort <- %uintOpBranchConditionalResult: the narrowed value reads as 0 instead of the neighbor lane's value.
Upstream IGC bug: intel/intel-graphics-compiler#402 (filed with 206-line pure Level Zero + spvasm reproducer).
rocPRIM's generic
warp_shuffle_op<T>forsizeof(T) < 4hits this pattern for__half/bf16warp operations. Fix: for sub-int types, bounce the post-shuffle int through avolatilestack slot and load the lowsizeof(T)bytes via byte-wise memcpy — emitsOpBitcast+OpLoadinstead of the offendingOpUConvert ushort <- uint, dodging the IGC bug.Impact (Arc A770, chipStar HEAD, Level Zero)
test_warp_reduce__half ReduceSum)test_warp_scan__half/bf16 (Inclusive)ScanReduce)test_block_reduceblock_params<__half,__half,64>)test_warp_shuffle_fp16(new regression)Likely also resolves a further ~20 rocThrust
__half-bearing tests downstream (not individually minimized).Commits
a8ad010test: reproduce IGC fp16 narrow-wide shuffle miscompile in warp_shufflebc73c2ffix: specialize warp_shuffle_op for sizeof(T)<4 to avoid IGC fp16 miscompileScope / compatibility
#ifdef __HIP_PLATFORM_SPIRV__— AMD path is unchangedTest plan
test_warp_shuffle_fp16PASS on Arc A770 / Level Zerotest_warp_reduce/test_warp_scan/test_block_reducefully green on Arc A770