Skip to content

Fix uint32 gridDim overflow in naive conv on large batch #3958

Open
chun-wan wants to merge 1 commit into
ROCm:developfrom
chun-wan:fix/naive-conv-uint32-grid-overflow
Open

Fix uint32 gridDim overflow in naive conv on large batch #3958
chun-wan wants to merge 1 commit into
ROCm:developfrom
chun-wan:fix/naive-conv-uint32-grid-overflow

Conversation

@chun-wan

Copy link
Copy Markdown

…figuration argument)

The naive forward convolution solver builds a 1-D launch grid with global work size grid_size * 256, where grid_size = n * k (default/NCHW layout) or n * ho (NHWC). This value is passed to hipExtModuleLaunchKernel() as a uint32_t. For large problems it reaches/exceeds 2^32 and is silently truncated, producing an illegal grid (e.g. gridDim.x == 0 when nk256 is a multiple of 2^32) and a HIP "invalid configuration argument" during Find (EvaluateInvokers). The asynchronous error then corrupts the HIP context and can make an unrelated later launch fail with miopenStatusUnknownError.

Reproduced on gfx950 (MI355X) with a ViT patch-embedding convolution:
Conv2d(3, 1024, kernel_size=14, stride=14, bias=False), input [N,3,14,14] bf16.
A stream of large N triggers it (e.g. N=32768: grid_size = n*k = 2^25, global work size = 2^25 * 256 = 2^33, truncated to 0 -> gridDim.x = 0).

Fix (two layers):

  • ConvDirectNaiveConvFwd::IsApplicable: declare the solver not applicable when the 1-D grid would overflow uint32 (grid_size * 256 >= 2^32), so a non-overflowing solver is selected instead.
  • HIPOCKernelInvoke::run: validate the launch geometry before the HIP call and throw (mirroring run_cooperative()'s WORKAROUND_SWDEV_448157) so an oversized launch can no longer silently truncate and corrupt the HIP context; Find drops the offending candidate cleanly.

Validated in a production-config build (MLIR + AI heuristics on): the repro's miopenStatusUnknownError is gone, every batch size in the repro sequence is served by a valid solver (ConvHipImplicitGemmGroupFwdXdlops) with correct output, and the conv-followed-by-another-kernel path no longer fails fatally.

Proposed changes

Please describe the motivation behind the pull request, whether it enables a new feature or fixes a bug. If there are associated pull requests or issues, please link them to the pull request.

Checklist

Please put an x into the boxes that apply. You can also fill these out after creating the PR. If you're not sure, please don't hesitate to ask.

  • I have added automated tests relevant to the introduced functionality
  • I have sufficient test coverage for the changes, and code coverage hasn't decreased as a result of my PR
  • I have ran the tests, and they are all passing locally
  • I have added relevant documentation for the changes
  • I have removed the stale documentation which is no longer relevant after this pull request
  • I have ran make format & make check_format to ensure the changes have been formatted

…figuration argument)

The naive forward convolution solver builds a 1-D launch grid with global work
size grid_size * 256, where grid_size = n * k (default/NCHW layout) or n * ho
(NHWC). This value is passed to hipExtModuleLaunchKernel() as a uint32_t. For
large problems it reaches/exceeds 2^32 and is silently truncated, producing an
illegal grid (e.g. gridDim.x == 0 when n*k*256 is a multiple of 2^32) and a HIP
"invalid configuration argument" during Find (EvaluateInvokers). The asynchronous
error then corrupts the HIP context and can make an unrelated later launch fail
with miopenStatusUnknownError.

Reproduced on gfx950 (MI355X) with a ViT patch-embedding convolution:
  Conv2d(3, 1024, kernel_size=14, stride=14, bias=False), input [N,3,14,14] bf16.
A stream of large N triggers it (e.g. N=32768: grid_size = n*k = 2^25,
global work size = 2^25 * 256 = 2^33, truncated to 0 -> gridDim.x = 0).

Fix (two layers):
- ConvDirectNaiveConvFwd::IsApplicable: declare the solver not applicable when the
  1-D grid would overflow uint32 (grid_size * 256 >= 2^32), so a non-overflowing
  solver is selected instead.
- HIPOCKernelInvoke::run: validate the launch geometry before the HIP call and
  throw (mirroring run_cooperative()'s WORKAROUND_SWDEV_448157) so an oversized
  launch can no longer silently truncate and corrupt the HIP context; Find drops
  the offending candidate cleanly.

Validated in a production-config build (MLIR + AI heuristics on): the repro's
miopenStatusUnknownError is gone, every batch size in the repro sequence is served
by a valid solver (ConvHipImplicitGemmGroupFwdXdlops) with correct output, and the
conv-followed-by-another-kernel path no longer fails fatally.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant