Do not stamp reqd_sub_group_size on kernels with indirect calls (#1380) - #1386
Open
pvelesko wants to merge 1 commit into
Open
Do not stamp reqd_sub_group_size on kernels with indirect calls (#1380)#1386pvelesko wants to merge 1 commit into
pvelesko wants to merge 1 commit into
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
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