From 12e81dff234e9039cf86005cb025fb18ab919aef Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Tue, 14 Apr 2026 22:26:59 +0300 Subject: [PATCH 1/5] test: reproduce IGC fp16 narrow-wide shuffle miscompile in warp_shuffle 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. --- test/rocprim/CMakeLists.txt | 1 + test/rocprim/test_warp_shuffle_fp16.cpp | 119 ++++++++++++++++++++++++ 2 files changed, 120 insertions(+) create mode 100644 test/rocprim/test_warp_shuffle_fp16.cpp diff --git a/test/rocprim/CMakeLists.txt b/test/rocprim/CMakeLists.txt index 11bf1369d..1224097b2 100644 --- a/test/rocprim/CMakeLists.txt +++ b/test/rocprim/CMakeLists.txt @@ -383,6 +383,7 @@ add_rocprim_test("rocprim.warp_reduce" test_warp_reduce.cpp) add_rocprim_test("rocprim.warp_scan" test_warp_scan.cpp) add_rocprim_test("rocprim.warp_sort" test_warp_sort.cpp) add_rocprim_test("rocprim.warp_store" test_warp_store.cpp) +add_rocprim_test("rocprim.warp_shuffle_fp16" test_warp_shuffle_fp16.cpp) # GPU: 3/4 pass, TransformReduce fail. # CPU: likely stuck diff --git a/test/rocprim/test_warp_shuffle_fp16.cpp b/test/rocprim/test_warp_shuffle_fp16.cpp new file mode 100644 index 000000000..70cc1c684 --- /dev/null +++ b/test/rocprim/test_warp_shuffle_fp16.cpp @@ -0,0 +1,119 @@ +// MIT License +// +// Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +// Regression test for IGC fp16 narrow<->wide shuffle miscompile. +// +// Reproduces the pattern emitted by rocprim::warp_shuffle_op for +// sizeof(T) < sizeof(int): the SPIR-V translator inserts an +// OpUConvert %uint <- %ushort before OpSubgroupShuffleINTEL and a +// matching OpUConvert %ushort <- %uint afterwards. When the result +// is consumed inside a divergent block (OpBranchConditional), IGC +// on Intel Arc miscompiles the chain and the second shuffle yields 0. + +#include "../common_test_header.hpp" +#include "test_utils.hpp" + +#include +#include + +// Two chained warp_shuffles on __half followed by a divergent store. +// This is the minimum shape that triggers the IGC bug; a single +// unconditional shuffle passes. +template +__global__ __launch_bounds__(64) +void chained_half_shuffle_divergent_kernel(const __half* in, __half* out) +{ + const unsigned int tid = threadIdx.x; + __half v = in[tid]; + + // Mirror the rocprim::warp_reduce_shuffle inner loop: + // for(offset=1; offset h_in(BlockSize); + for(unsigned int i = 0; i < BlockSize; ++i) + { + // Values 2,3,4,5 per 4-lane group -> expected sum = 14. + h_in[i] = __half(static_cast((i % LogicalWarpSize) + 2)); + } + std::vector<__half> h_out(NumWarps, __half(0.0f)); + + __half* d_in = nullptr; + __half* d_out = nullptr; + HIP_CHECK(hipMalloc(&d_in, h_in.size() * sizeof(__half))); + HIP_CHECK(hipMalloc(&d_out, h_out.size() * sizeof(__half))); + HIP_CHECK(hipMemcpy(d_in, h_in.data(), h_in.size() * sizeof(__half), + hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(d_out, h_out.data(), h_out.size() * sizeof(__half), + hipMemcpyHostToDevice)); + + hipLaunchKernelGGL(HIP_KERNEL_NAME(chained_half_shuffle_divergent_kernel), + dim3(1), dim3(BlockSize), 0, 0, d_in, d_out); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipDeviceSynchronize()); + + HIP_CHECK(hipMemcpy(h_out.data(), d_out, h_out.size() * sizeof(__half), + hipMemcpyDeviceToHost)); + + constexpr float expected = 2.0f + 3.0f + 4.0f + 5.0f; // 14 + unsigned int mismatches = 0; + for(unsigned int w = 0; w < NumWarps; ++w) + { + const float got = static_cast(h_out[w]); + if(std::abs(got - expected) > 0.01f) + { + ++mismatches; + } + } + EXPECT_EQ(mismatches, 0u) + << "IGC fp16 narrow<->wide chained-shuffle miscompile triggered " + << "(got 0 for second shuffle under divergent consumer)."; + + HIP_CHECK(hipFree(d_in)); + HIP_CHECK(hipFree(d_out)); +} From d0d9385cd2e436f6b14f987af83fe6305ecb43d7 Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Tue, 14 Apr 2026 22:33:47 +0300 Subject: [PATCH 2/5] fix: specialize warp_shuffle_op for sizeof(T)<4 to avoid IGC fp16 miscompile 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. --- .../rocprim/intrinsics/warp_shuffle.hpp | 46 ++++++++++++++++++- 1 file changed, 45 insertions(+), 1 deletion(-) diff --git a/rocprim/include/rocprim/intrinsics/warp_shuffle.hpp b/rocprim/include/rocprim/intrinsics/warp_shuffle.hpp index 2055ddf0e..fd5f756f2 100644 --- a/rocprim/include/rocprim/intrinsics/warp_shuffle.hpp +++ b/rocprim/include/rocprim/intrinsics/warp_shuffle.hpp @@ -79,9 +79,53 @@ warp_shuffle_op(const T& input, ShuffleOp&& op) #endif } +#if defined(__HIP_PLATFORM_SPIRV__) && !defined(__HIP_CPU_RT__) +// Specialization for sub-int sized trivially-copyable types (e.g. __half). +// +// Works around an IGC miscompile on Intel Arc: the SPIR-V translator emits +// +// %w = OpUConvert %uint %ushort_input ; narrow -> wide +// %s = OpSubgroupShuffleINTEL %uint %w ... +// %n = OpUConvert %ushort %s ; wide -> narrow +// +// and IGC miscompiles the post-shuffle narrowing UConvert when the consumer +// lives in a block entered via OpBranchConditional (returns 0 instead of +// the shuffled value). See upstream IGC issue (to be filed). +// +// We defeat the trigger by bouncing the shuffle result through a +// memory-backed scratch slot and loading the low sizeof(T) bytes via a +// byte-wise memcpy from a volatile pointer. This emits OpBitcast + OpLoad +// of the narrow type instead of a scalar-register OpUConvert, which IGC +// handles correctly. template ROCPRIM_DEVICE ROCPRIM_INLINE -typename std::enable_if::value && (sizeof(T) % sizeof(int) == 0)), T>::type +typename std::enable_if< + std::is_trivially_copyable::value && (sizeof(T) < sizeof(int)), T>::type +warp_shuffle_op(const T& input, ShuffleOp&& op) +{ + int word = 0; // zero-init upper bytes that the memcpy below leaves untouched + __builtin_memcpy(&word, &input, sizeof(T)); + word = op(word); + + // Store the shuffle result into a volatile stack slot to force the + // narrowing to happen via a memory load rather than a register-level + // i32->i16 trunc (which the translator lowers to OpUConvert). + volatile int word_scratch = word; + T output; + __builtin_memcpy(&output, + const_cast(&word_scratch), + sizeof(T)); + return output; +} +#endif + +template +ROCPRIM_DEVICE ROCPRIM_INLINE +typename std::enable_if::value && (sizeof(T) % sizeof(int) == 0)) +#if defined(__HIP_PLATFORM_SPIRV__) && !defined(__HIP_CPU_RT__) + && !(std::is_trivially_copyable::value && (sizeof(T) < sizeof(int))) +#endif + , T>::type warp_shuffle_op(const T& input, ShuffleOp&& op) { constexpr int words_no = (sizeof(T) + sizeof(int) - 1) / sizeof(int); From b4d8bf020510a9c9f4fb9eddfd77b2582640aa63 Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Thu, 16 Apr 2026 12:53:41 +0300 Subject: [PATCH 3/5] ci: re-enable warp_reduce, warp_scan, block_shuffle (fixed by this PR) --- .ci/exclude.txt | 12 ++---------- 1 file changed, 2 insertions(+), 10 deletions(-) diff --git a/.ci/exclude.txt b/.ci/exclude.txt index abe9971b9..9847dd496 100644 --- a/.ci/exclude.txt +++ b/.ci/exclude.txt @@ -14,16 +14,8 @@ rocprim\.device_reduce_by_key # Mitigation: CHIP_L0_FORCE_SIMD16 or wait for IGC fix. rocprim\.device_run_length_encode -# IGC bug #402 (fp16 subgroup-shuffle narrowing miscompile). -# Upstream: https://github.com/intel/intel-graphics-compiler/issues/402 -# Fix pending in PR [CHIP-SPV/rocPRIM#3]; remove these two lines when #3 merges. -rocprim\.warp_reduce -rocprim\.warp_scan - -# fp16/bf16 block_shuffle failures on PoCL CPU (pastrami). Related to the same -# sub-int warp_shuffle roundtrip addressed by PR [CHIP-SPV/rocPRIM#3]; remove -# once the fp16 fix lands and is verified on CPU. -rocprim\.block_shuffle +# IGC bug #402 fp16 tests — REMOVED by PR #3 (warp_shuffle_op specialization). +# warp_reduce, warp_scan, block_shuffle should now pass on all platforms. # Mali GPU (salami) lacks fp64 + SubgroupDispatch SPIR-V capabilities. # These tests require hardware features the ARM Mali doesn't expose. From cbdc736cecb26f3e75f1f44ad850335a33ec014b Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Thu, 16 Apr 2026 15:55:42 +0300 Subject: [PATCH 4/5] ci: re-exclude block_shuffle (bf16 still fails on PoCL CPU) --- .ci/exclude.txt | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/.ci/exclude.txt b/.ci/exclude.txt index 9847dd496..8f775c8b2 100644 --- a/.ci/exclude.txt +++ b/.ci/exclude.txt @@ -14,8 +14,10 @@ rocprim\.device_reduce_by_key # Mitigation: CHIP_L0_FORCE_SIMD16 or wait for IGC fix. rocprim\.device_run_length_encode -# IGC bug #402 fp16 tests — REMOVED by PR #3 (warp_shuffle_op specialization). -# warp_reduce, warp_scan, block_shuffle should now pass on all platforms. +# IGC bug #402 fp16 tests — warp_reduce + warp_scan fixed by PR #3. +# block_shuffle bf16 still fails on PoCL CPU (pastrami) — separate code path +# not covered by warp_shuffle_op fix. +rocprim\.block_shuffle # Mali GPU (salami) lacks fp64 + SubgroupDispatch SPIR-V capabilities. # These tests require hardware features the ARM Mali doesn't expose. From 31c0c6ad6d5ae7a9a04e1b00226317505286bca5 Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Fri, 17 Apr 2026 08:38:50 +0300 Subject: [PATCH 5/5] Update exclude.txt --- .ci/exclude.txt | 1 + 1 file changed, 1 insertion(+) diff --git a/.ci/exclude.txt b/.ci/exclude.txt index 8f775c8b2..3b693c0d4 100644 --- a/.ci/exclude.txt +++ b/.ci/exclude.txt @@ -15,6 +15,7 @@ rocprim\.device_reduce_by_key rocprim\.device_run_length_encode # IGC bug #402 fp16 tests — warp_reduce + warp_scan fixed by PR #3. +# https://github.com/intel/intel-graphics-compiler/issues/402 # block_shuffle bf16 still fails on PoCL CPU (pastrami) — separate code path # not covered by warp_shuffle_op fix. rocprim\.block_shuffle