Skip to content

fix(ck): convolution edge cases large tensors - #11754

Merged
EnricoDeg merged 12 commits into
developfrom
users/enricodeg/ck/fix-edge-cases
Sep 30, 2026
Merged

EnricoDeg merged 12 commits into
developfrom
users/enricodeg/ck/fix-edge-cases

Conversation

@EnricoDeg

@EnricoDeg EnricoDeg commented Sep 7, 2026 •

Copy link
Copy Markdown
Contributor

Motivation

Improve handling of large tensor edge cases in convolution

JIRA ID : ROCM-27829

Technical Details

Buffer resource is created with tensor element size as int but tensor byte size is stored as uint. In this way, for 16 bit types, a tensor byte size of exactly 2 GB works. So current limitations are:

  • tensor element size is stored in int32_t so max value is 2GB - 1
  • tensor element bytes is stored in uint32_t so max value is 4GB - 1.

Convolution fwd supports splitN so if strides are within int range, first splitN is applied. If after splitN, the element size is >= 2GB or the element byte size is > 2GB, the instance is not supported.

Convolution bwd data:

  • if the implementation supports splitN, the same workflow is used as convolution fwd
  • if the implementation doesn't support splitN, the tensor size is checked when the argument is created and stride_overflow is set. Then this is checked in IsSupportedArgument.

Convolution bwd weight doesn't have splitN support because N is in the GemmK dimension so we always check tensor size when we create the Argument and set stride_overflow. Then this is checked in IsSupportedArgument.

Bugs:

  • convolution bwd data xdl cshuffle v1 implementation is using a fast path but the conditions were not exhaustive:
    • num_workgroups_per_Conv_N_ must be 1
    • ConvK % AK1 == 0 && ConvK % BK1 == 0 because GemmK dimension is unmerged into K0 and K1 but K1 is not padded. I decided to not support this case instead of adding padding on the vector dimension which might create performance regressions on the fast path. This is the vector dimension of the descriptor but the actual vector dimension can be lower and this is where the problem comes from.
    • Modify flaky instance on gfx1250

Submission Checklist

@EnricoDeg
EnricoDeg requested review from a team as code owners September 7, 2026 10:42
@EnricoDeg EnricoDeg self-assigned this Sep 7, 2026
@github-actions github-actions Bot added project: composablekernel ck: convolution Used to tag composablekernel PRs that require approval from CK convolution review team. labels Sep 7, 2026
@EnricoDeg EnricoDeg added organization: streamhpc contributors from streamhpc and removed ck: convolution Used to tag composablekernel PRs that require approval from CK convolution review team. labels Sep 7, 2026
@EnricoDeg EnricoDeg changed the title fix(ck): conv edge cases large tensors fix(ck): convolution edge cases large tensors Sep 7, 2026
@therock-pr-bot

therock-pr-bot Bot commented Sep 7, 2026 •

Copy link
Copy Markdown

✅ All Checks Passed — Ready for Review

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ⚠️ Warning Error: Source/code files changed without an accompanying unit test.
Expected: add at least one test file named like test_<name>.py / test_<name>.cpp (or <name>_test.*).
Current: code file(s) changed: projects/composablekernel/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_data_multiple_d_wmma_cshuffle.hpp, projects/composablekernel/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_data_multiple_d_wmma_cshuffle_v3.hpp, projects/composablekernel/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle_v1.hpp, projects/composablekernel/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle_v3.hpp, projects/composablekernel/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_two_stage_wmma_cshuffle_v3.hpp (+10 more); no test file found
🔎 pre-commit ✅ Pass —
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —
🤖 therock-pr-bot ✅ Pass —

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

@therock-pr-bot

therock-pr-bot Bot commented Sep 7, 2026 •

Copy link
Copy Markdown

🎉 All checks passed! This PR is ready for review.

@chris-tsiaousis-hpc chris-tsiaousis-hpc left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM! Added some comments/questions

@EnricoDeg
EnricoDeg force-pushed the users/enricodeg/ck/fix-edge-cases branch from 766c1ed to 59677ec Compare September 9, 2026 07:31
@github-actions github-actions Bot added the ck: convolution Used to tag composablekernel PRs that require approval from CK convolution review team. label Sep 9, 2026

@JP-Fernando JP-Fernando left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This PR is correctly scoped and documented. The latest commits fix some earlier issues, already addressed by the other reviewer. Approved with a few comments/questions out of curiosity.

@EnricoDeg
EnricoDeg force-pushed the users/enricodeg/ck/fix-edge-cases branch from a3322c6 to 000b504 Compare September 14, 2026 06:58
@jakpiase
jakpiase requested a balanced review from Copilot September 16, 2026 13:38

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 Changes recommended

Several size checks can miss oversized tensors or use incorrect data types and descriptor-loop bounds.

Get a fresh assessment by requesting another Copilot review.

Pull request overview

Improves convolution handling for large tensors and tightens backward-data fast-path eligibility.

Changes:

  • Uses unsigned buffer byte sizes and adds 2 GiB tensor/descriptor checks.
  • Propagates overflow detection across backward convolution implementations.
  • Restricts an unsafe fast path and adjusts a BF16 forward instance.
File summaries
File Description
device_grouped_conv_fwd_xdl_instance.hpp Adjusts BF16 vector widths.
amd_buffer_addressing_builtins.hpp Uses unsigned resource sizes.
transform_conv_fwd_to_gemm.hpp Adds element-count limits.
tensor_size_check.hpp Expands large-tensor checks.
device_grouped_conv_bwd_weight_xdl_waveletmodel_cshuffle_v3.hpp Adds overflow rejection.
device_grouped_conv_bwd_weight_xdl_cshuffle.hpp Adds overflow rejection.
device_grouped_conv_bwd_weight_xdl_cshuffle_v3.hpp Adds overflow detection.
device_grouped_conv_bwd_weight_wmma_cshuffle.hpp Adds overflow detection.
device_grouped_conv_bwd_weight_wmma_cshuffle_v3.hpp Adds overflow detection.
device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp Checks accumulator workspace size.
device_grouped_conv_bwd_weight_two_stage_wmma_cshuffle_v3.hpp Checks accumulator workspace size.
device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle_v3.hpp Checks all tensor sizes.
device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle_v1.hpp Tightens fast path and descriptor checks.
device_grouped_conv_bwd_data_multiple_d_wmma_cshuffle.hpp Adds overflow detection.
device_grouped_conv_bwd_data_multiple_d_wmma_cshuffle_v3.hpp Adds descriptor-size validation.
Review details

Suppressed comments (2)

projects/composablekernel/include/ck/tensor_operation/gpu/device/tensor_size_check.hpp:19

  • This checks the logical element count, not the addressable element space implied by the supplied strides. These convolution APIs accept explicit stride arrays, so a padded tensor (or a long-index stride above INT32_MAX) can have a resource/descriptor span over the limit while stride_overflow remains false and the stride is later narrowed. Compute the maximum offset from both lengths and strides, or pass an already-built descriptor to this check, and update the callers accordingly.
    for(const auto& l : lengths)
        total *= l;
    long_index_t total_bytes = total * sizeof(DataType);

projects/composablekernel/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_wmma_cshuffle.hpp:912

  • The pointer API repeats the wrong GEMM-alias mapping for external tensors. For heterogeneous convolutions this can mark an oversized weight as safe or reject a valid output; use InDataType, WeiDataType, and OutDataType for the input, weight, and output lengths respectively, and make the corresponding correction in the long-index overload.
        const bool stride_ovf = tensor_exceeds_2gb<ADataType>(a_g_n_c_wis_lengths) ||
                                tensor_exceeds_2gb<BDataType>(b_g_k_c_xs_lengths) ||
                                tensor_exceeds_2gb<CDataType>(e_g_n_k_wos_lengths);
  • Files reviewed: 15/15 changed files
  • Comments generated: 4
  • Review effort level: Balanced

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

@EnricoDeg
EnricoDeg merged commit 820ba71 into develop Sep 30, 2026
213 of 215 checks passed
@EnricoDeg
EnricoDeg deleted the users/enricodeg/ck/fix-edge-cases branch September 30, 2026 06:31
assistant-librarian Bot pushed a commit to ROCm/composable_kernel that referenced this pull request Sep 30, 2026
fix(ck): convolution edge cases large tensors

## Motivation

Improve handling of large tensor edge cases in convolution

JIRA ID : ROCM-27829

## Technical Details

Buffer resource is created with tensor element size as int but tensor
byte size is stored as uint. In this way, for 16 bit types, a tensor
byte size of exactly 2 GB works. So current limitations are:
 - tensor element size is stored in int32_t so max value is 2GB - 1
 - tensor element bytes is stored in uint32_t so max value is 4GB - 1.

Convolution fwd supports splitN so if strides are within int range,
first splitN is applied. If after splitN, the element size is >= 2GB or
the element byte size is > 2GB, the instance is not supported.

Convolution bwd data:
- if the implementation supports splitN, the same workflow is used as
convolution fwd
- if the implementation doesn't support splitN, the tensor size is
checked when the argument is created and `stride_overflow` is set. Then
this is checked in `IsSupportedArgument`.

Convolution bwd weight doesn't have splitN support because N is in the
GemmK dimension so we always check tensor size when we create the
Argument and set `stride_overflow`. Then this is checked in
`IsSupportedArgument`.

Bugs:
- convolution bwd data xdl cshuffle v1 implementation is using a fast
path but the conditions were not exhaustive:
   - `num_workgroups_per_Conv_N_` must be 1
- `ConvK % AK1 == 0 && ConvK % BK1 == 0` because GemmK dimension is
unmerged into K0 and K1 but K1 is not padded. I decided to not support
this case instead of adding padding on the vector dimension which might
create performance regressions on the fast path. This is the vector
dimension of the descriptor but the actual vector dimension can be lower
and this is where the problem comes from.
   - Modify flaky instance on gfx1250

## Submission Checklist

- [x] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
@assistant-librarian

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

JonathanLichtnerAMD added a commit that referenced this pull request Oct 5, 2026
…es (#12396)

# revert(miopen): partial revert of CK byte span gate for large
instances

JIRA ID: ALMIOPEN-2645

## Motivation

PR #11585 widened `RequiresLargeTensorCKInstance()` to gate forward and
backward-data (not just backward-weights) on tensor byte span, to work
around a 32-bit byte-offset overflow in CK. CK PR #11754
(#11754) fixes that overflow
directly in CK, so the MIOpen-side forward/backward-data gate is no
longer needed and was unnecessarily excluding large-tensor problems from
faster int32 CK instances.

## Technical Details

Reverts `RequiresLargeTensorCKInstance()` in
`projects/miopen/src/ck_impl/implicitgemm_ck_util.hpp` to its pre-#11585
state: the large-tensor check is scoped back to backward-weights only,
and uses element-space (`GetElementSpace()`) rather than byte-span
(`GetNumBytes()`). Test expectations in
`projects/miopen/test/gtest/unit_implicitgemm_ck_util.cpp` are reverted
to match (forward/backward-data no longer expected to require
large-tensor instances at the >INT_MAX-byte-span shape), and the
byte-span-specific `ByteSpanOverflowGate` test and its
`MakeByteSpanConvProblem` helper are removed. The `.bz2` FDB/PDB
database files that PR #11585 also touched are intentionally left
untouched — this PR is source-only, with DB changes deferred to a later
database-tuning update.

## Risk Assessment

Source-only change to a solver-selection gate in a single header,
reverting an already-reviewed prior state; blast radius is confined to
CK instance selection for forward/backward-data conv problems with large
tensors, and is gated on a companion CK fix (#11754) landing first —
this PR is a **draft** and not intended to merge before that.

**Risk Level: 🟡 Medium**

### Impacted Components

- `projects/miopen/src/ck_impl/implicitgemm_ck_util.hpp`
(`RequiresLargeTensorCKInstance`)
- `projects/miopen/test/gtest/unit_implicitgemm_ck_util.cpp`

### Potential Side Effects

Until CK PR #11754 lands, large-tensor forward/backward-data problems
that trip the byte-span overflow (byte span > INT_MAX with element count
still < INT_MAX) could silently pick a narrowing int32 CK instance
again, reintroducing the correctness bug #11585 fixed for those two
directions.

## Test Plan

Built and ran `test_unit_implicitgemm_ck_util` in a docker dev container
(gfx942).

## Test Result

All 30 tests pass, including the two reverted large-tensor gate tests
(`WrwElementCountOverflowGate`, `Wrw3DElementCountOverflowGate`) back to
their pre-#11585 (element-space, backward-weights-only) expectations. A
broader multi-shape perf-eval sweep has been run as part of the MI455
performance benchmarking.
yassinsolim added a commit that referenced this pull request Oct 13, 2026
…abase (#13420)

Remove 20 ConvHipImplicitGemmGroupWrwXdlops FP16 grouped backward-weights entries from the
gfx1200 system database: the solver is dropped from 20 find-db lines and the 20 matching perf-db
rows are deleted. MIOPEN_DBSYNC_CLEAN=1 selects exactly these rows.

Their largest tensors hold 2^30 to 2^31 elements (up to 4,225,228,800 bytes). That is past CK's
tensor_exceeds_2gb byte guard from #11754 but below MIOpen's element-count large-tensor gate,
so CK rejects the stored non-large-tensor instances and the shapes fell back to untuned solvers.

Fixes Smoke/CPU_DBSync_NONE.StaticFDBSync/10 on gfx1200 (#13394). The source-side gap is
tracked in #13419.

JIRA ID: ROCM-29026
yassinsolim added a commit that referenced this pull request Oct 13, 2026
…rent gate no longer rejects (#13421)

Restore 228 of the 229 large-tensor CK entries #12088 removed from the gfx1201 system database
(165 backward-data, 63 forward find-db entries and 225 perf-db rows). Every restored line is
byte-identical to its pre-#12088 original; no lines are removed or newly written.

#12396 returned RequiresLargeTensorCKInstance to its backward-weight-only, element-count form once
#11754 gave CK its own per-instance overflow checks, so these forward and backward-data entries
resolve again. On gfx1201 the pre-#12088 database fails DBSync on 409 split-K entries only; the
split-K half is re-tuned separately in #13415. One entry stays out because it was slower.

Validated on an AMD Radeon AI PRO R9700 (gfx1201, 64 CU): StaticFDBSync and GPU-reference checks
pass for the base, this database and its union with #13415, and a paired MIOPEN_FIND_MODE=2 A/B
measured 1.43x aggregate kernel time across the 51 shapes these entries win (single run).

JIRA ID: ROCM-29026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

ck: convolution Used to tag composablekernel PRs that require approval from CK convolution review team. organization: streamhpc contributors from streamhpc project: composablekernel

Projects

None yet

Development

Successfully merging this pull request may close these issues.

5 participants