Do not stamp reqd_sub_group_size on kernels with indirect calls (#1380) - #1386
Merged
Conversation
HipWarpsPass puts intel_reqd_sub_group_size on every kernel in the module as soon as one shuffle/ballot declaration appears anywhere in it, so that __shfl and friends see the warp width chipStar was built for. A stamped kernel silently loses the arguments of a device indirect call. The callee receives the first argument and every argument after it arrives as zero, with no error and no diagnostic. Kokkos declares the shuffle intrinsics from Kokkos_Core.hpp, so a single #include stamped an entire application and broke device-side indirect dispatch for every Kokkos program. Virtual calls lower to the same indirect call, and kynema-ugf / Nalu-Wind builds its whole Kernel, NodeKernel, EdgeKernel, MasterElement and CoeffApplier architecture on exactly that, so it linked and ran while computing garbage. Skip kernels that can reach an indirect call over direct edges, and leave every other kernel exactly as before. Narrowing the stamp the other way, to only those kernels that provably reach a sensitive function, was tried first and is wrong today: WarpSizeSensitiveFuncNames lists only some of the shuffle overloads, so Kokkos' own reductions lost the subgroup size they depend on and started accumulating each contribution four times. Reduced reproducer, plain HIP, no Kokkos: adding one unused __global__ that calls __shfl turns a correct indirect dispatch of f(7.0, 3) from 7036.0 into 7006.0. That is TestIndirectCallWithWarpPrimitive.hip. Fixes #1380
pvelesko
force-pushed
the
2026-07-28-warps-indirect-argdrop
branch
from
August 25, 2026 11:06
03bee9e to
99821fc
Compare
TestIndirectCall's callee takes a single argument, so it cannot observe #1380: a kernel stamped with intel_reqd_sub_group_size delivers the first argument of an indirect call and zero for every argument after it. Add the never-launched touchShfl kernel that puts the shuffle declaration in the module, and a two-argument callee reached through a kernel argument, so the same test covers both the plain indirect call and the argument drop. Drop the separate TestIndirectCallWithWarpPrimitive, whose only reason to exist was the missing second argument.
The blanket exclusion came in with #479, which disabled the test "for all platforms due to lack of runtime query for detecting indirect call support" while recording that it passes on Intel Level Zero and Intel OpenCL CPU & GPU. #822 carried it into known_failures.yaml as a bare entry and the reason was lost with the comment. known_failures.yaml has per-backend and per-host sections now, so the capability can be expressed directly: keep it excluded on PoCL, whose SPIR-V reader is not built with SPV_INTEL_function_pointers enabled, and on rusticl, whose spirv_to_nir implements no FunctionPointersINTEL capability. It runs everywhere else, which is where #1380's argument drop is observable in the first place.
pvelesko
force-pushed
the
2026-07-28-warps-indirect-argdrop
branch
from
August 25, 2026 11:08
99821fc to
9b55a4b
Compare
Collaborator
Author
|
/run-aurora-ci |
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.
HipWarpsPassstampsintel_reqd_sub_group_sizeon every kernel in a module as soon as one shuffle or ballot intrinsic is declared anywhere in it. A stamped kernel then silently drops every argument after the first of a device indirect call, with no error and no diagnostic. One#include <Kokkos_Core.hpp>is enough to trigger it, so every Kokkos application that dispatches on the device computed garbage.Skip kernels that can reach an indirect call over direct edges, and leave every other kernel stamped exactly as before.
Fixes #1380