Repository navigation
Conversation
Building CK for gfx1250 together with any gfx11 target fails at link: ld.lld: error: local memory (131552) exceeds limit (65536) in function 'kernel_grouped_conv_fwd_grouped_gemm_wmma_cshuffle_v3' 152 such errors across 5 objects, LDS 65648..197104 B against a 65536 B limit. Root cause: the gfx1250 grouped-conv instances added in #10620 are gated on CK_USE_GFX1250, a host-side macro that CMakeLists.txt defines globally whenever the target list contains gfx1250. HIP compiles each TU once per --offload-arch, so the preprocessor state is identical for every device pass; the kernel bodies are guarded only by `__gfx11__ || __gfx12__`, which is true on gfx1101/gfx1151. A host macro cannot express "this offload architecture", so tiles sized for the 320 KB LDS of gfx1250 are emitted for 64 KB parts. Move those instances into dedicated *_gfx1250*.cpp translation units and restrict them to --offload-arch=gfx1250* with the existing per-source filter in library/src/tensor_operation_instance/gpu/CMakeLists.txt (same mechanism already used to confine gemm_wmma_universal f8 to gfx12). CK_USE_GFX1250 is retained only on the host-side declarations and calls, where "this build includes gfx1250" is what it actually means. Also restores two bf16 bwd_weight instances that `#ifndef CK_USE_GFX1250` silently removed from gfx11 builds, and adds gfx1250 to the multi-arch Jenkins build that would have caught this. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
…e, real CI wiring Review of the first commit found three blockers; all fixed here. 1. The Jenkinsfile edit was a no-op. `execute_args` is never read by any groovy code; `ck.runBuildInstancesOnly` supplies its own hardcoded arch list at ck.groovy:1492. Reverted the Jenkinsfile and added gfx1250 where the list is actually consumed. 2. Removing `#ifndef CK_USE_GFX1250` from the bwd_weight WMMA list was wrong. That guard was a K-granularity exclusion, not an LDS one: the two instances it covers are the only ones with KPerBlock=48, and gridwise_gemm_wmma_cshuffle_v3.hpp:617 rejects KPerBlock % (KPerWmmaBlk * 2) != 0 on gfx12 (48 % 32 = 16). Restored the guard with a comment explaining why it exists. 3. The XDL large_k list was left behind at 196608 B of LDS -- the largest offender in the original incident, and the fifth failing object. Its two pipev1 TUs are now gfx1250-only, which also removes the gfx9 wave64 offload passes where the CShuffle static_assert fires. Also from review: - Export CK_USE_GFX1250 via config.h.in. It was only an add_definitions(), so every new guard in an installed public header compiled to off for downstream consumers, silently dropping all gfx1250 instances for MIOpen. - Gate the factory calls on ck::is_gfx1250_supported() as well as the macro. A gfx1200;gfx1250 build defines the macro but carries no gfx1200 code object for these TUs, which would fail at launch with hipErrorNoBinaryForGpu. - Match the CMake filter against the full source path with a (^|[_/]) boundary. - Use SUPPORTED_GPU_TARGETS, not GPU_TARGETS, for the -mllvm scheduler flags; GPU_TARGETS is empty in a GPU_ARCHS build. - Give all 13 gfx1250 entry points a _gfx1250 token so one grep audits them. - clang-format 18.1.3 clean. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Code review pass — 15 findings, 11 addressed in f22a696Ran an adversarial review over the first revision. It found three blockers; the PR Blockers (fixed)
Also fixed
Deliberate, called out
Open, deferred to follow-ups
Checked cleanPreprocessor nesting in all touched files; CMake Still not compiled locally — no multi-arch toolchain available. CI remains the real check. |
…ally binds Per-source --offload-arch flags are additive. When the caller sets GPU_TARGETS (rather than GPU_ARCHS) CMakeLists.txt:361-365 does not unset it before find_package(hip), so hip::device injects --offload-arch for every target globally and the per-source filtering in add_instance_library can only widen, never narrow. Emit --no-offload-arch for the excluded targets instead. Verified: a 131552-byte __shared__ built for gfx1101+gfx1250 fails, and succeeds with --no-offload-arch=gfx1101. Also from review: revert the filter to source_name (generate_sharded_instantiations appends absolute paths, so a build dir containing gfx1250 matched every shard); drop the config.h.in export of CK_USE_GFX1250 (it desynced CK_ENABLE_CLUSTER_LAUNCH, ck.hpp:205); use the pre-existing is_gfx125_supported() instead of a duplicate predicate; skip a TU whose target list is empty rather than compiling it for the default arch. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The instance-factory headers (grouped_convolution_*.hpp and their .inc files) are installed, and this PR guarded their gfx1250 registrations with CK_USE_GFX1250. That macro is only an add_definitions() flag on CK's own build; it is not in config.h, so a downstream consumer compiling against the installed headers never sees it and every gfx1250 registration compiles out. Preprocessing grouped_convolution_forward.hpp as a consumer: 0 gfx1250 registration calls survive with the old guard, 6 with the new one. CK_USE_GFX1250 itself cannot be exported: ck/ck.hpp:205 keys the host half of CK_ENABLE_CLUSTER_LAUNCH on it and the device half on __gfx125__, so exporting it would desynchronise the two on every non-gfx1250 offload pass. Add CK_HAS_GFX1250_INSTANCES, set alongside CK_USE_GFX1250 and exported through config.h.in, meaning only "this CK build contains the gfx1250 instance TUs". Swap all 26 registration guards this PR added to it. The runtime is_gfx125_supported() check is unchanged. CK_USE_GFX1250 and CK_ENABLE_CLUSTER_LAUNCH are untouched. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
The stage runs only from the nightly cron, uses the standard CK image (gfx1250 stages use a separate image and compiler), and configures with GPU_ARCHS rather than the GPU_TARGETS path MIOpen uses. CI coverage for gfx1250 alongside other targets belongs in its own change. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
The empty-target skip and the --no-offload-arch subtraction in add_instance_library applied
to every instance source. That changed the arch set of hundreds of unrelated TUs and broke
configurations that worked before: configure failures with no GPU_TARGETS or with
CK_USE_FP8_ON_UNSUPPORTED_ARCH=ON, dropped WMMA f8 TUs for gfx942;gfx1100, FP8 gemm TUs
losing gfx90a, WMMA/XDL TUs losing archs that host code still queries ("invalid device
function"), and MIOpen per-arch unbundling dropping objects.
Only _gfx1250 sources are now restricted, by --no-offload-arch=all followed by their gfx1250
targets. =all also clears archs hip::device injected from auto-detection or a stale cache,
which an explicit subtraction list could not see. Every other source's flags are identical to
develop's.
Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
The extra compile options are set on _gfx1250 sources only, and add_instance_library already drops those sources when gfx1250 is not targeted, so the surrounding if(SUPPORTED_GPU_TARGETS MATCHES "gfx1250") re-derived a decision made elsewhere. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Each gfx1250 registration in the five installed grouped-conv factory headers was a 7-line block pairing #ifdef CK_HAS_GFX1250_INSTANCES with a runtime is_gfx125_supported() check, pasted 13 times, and evaluated the device query once per call site on every GetInstances(). Move the runtime check into the 13 add_..._gfx1250_instances() functions, as ..._data_cache_prefetch_instance.cpp already does. The headers keep only the compile-time guard and no longer include device_prop.hpp or make HIP runtime calls. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
gemm_universal.hpp and gemm_universal_xdl.inc are installed, but gated the gfx1250 v3 prefetch instances on CK_USE_GFX1250, which is only an add_definitions() flag on CK's own build. A consumer compiling against the installed headers never saw those registrations. The backing TU is built for every XDL target and checks is_gfx125_supported() internally. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…o choice The comment on the retained #ifndef CK_USE_GFX1250 in the bf16 bwd-weight WMMA list said the two KPerBlock=48 instances stay available on gfx1100-gfx1201 and cited a line that is a closing brace. CK_USE_GFX1250 is build-wide, so in any build that includes gfx1250 the rows are removed for every architecture (preprocessing the gfx1101 device pass: 0 rows with the macro, 2 without). State what the guard actually does; behaviour is unchanged. Also record why CK_HAS_GFX1250_INSTANCES is not CK_GFX1250_SUPPORT: the ctypes dispatcher undefines and redefines the latter per target architecture. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
CK now builds some instance objects for gfx1250 only. With MIOPEN_CK_UNBUNDLE_PER_ARCH=ON, unbundle_archive.cmake mishandled them in two ways: - An object with no device code for ARCH was dropped entirely, host code included, so the per-arch CK plugin failed to dlopen with undefined add_device_*gfx1250* symbols. - For ARCH=gfx1250, the thin fatbin of a gfx1250;gfx1250-strict object came out larger than the original and the script stopped with FATAL_ERROR. Keep such objects unchanged in both cases. The device code that rides along is never launched on another architecture; CK registers these instances only when is_gfx125_supported(). Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Whitespace and line breaks only; the token stream is unchanged. CK's static-checks stage runs clang-format-18 over every .h/.hpp/.cpp and fails on any diff. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Add a second barrier against the AICK-2468 failure mode, independent of the _gfx1250 file split and CMake flags: a WMMA conv tile too large for an architecture's LDS now compiles to an empty kernel there and is rejected at run time, instead of failing the device link with "ld.lld: local memory (N) exceeds limit (65536)". - Device: the four WMMA conv kernels (fwd, fwd large-tensor, bwd-data, bwd-weight) place the __shared__ array and body inside if constexpr(LDS_size <= get_lds_size(get_device_arch())). - Host: GridwiseGemm_wmma_cshuffle_v3_base::CheckValidity rejects an instance whose GetSharedMemoryNumberOfByteOnHost() exceeds the running device's LDS, so an emptied kernel is never selected. The helpers it needs become __host__ __device__. Tiles that fit are unaffected: device code for them is byte-identical on gfx1151 and gfx1250. Borrowed from Enrico Degregori's users/enricodeg/ck/multi-arch-conv (ef12db1), without its instance-guard removals, XDL changes, KPerBlock checks or LdsTranspose rejection. Adds the device_prop.hpp include that gridwise_gemm_xdl_cshuffle_v3.hpp already has; without it the b_scale, ab_scale and blockscale WMMA instances fail to compile. Co-Authored-By: Enrico Degregori <enrico@streamhpc.com> Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
|
This pull request has been inactive for 25 days and will be marked as stale. If you would like to keep this PR open, please:
This PR will be automatically closed in 5 days if no further activity occurs. |
|
This pull request has been automatically closed due to inactivity (30 days with no updates). If you'd like to continue working on this, feel free to reopen the PR or create a new one. |
Fixes AICK-2468. Unblocks ALMIOPEN-2820.
Problem
Any build combining gfx1250 with another target fails at link:
Reproduced locally at
bc47ae72ea8: 108local memoryerrors and 13static assertion failedacross 5 objects, LDS 65648..197104 B against a 65536 B limit. All ten distinct LDS sizes
match MIOpen develop nightly #1589. Nightlies #1589-#1592 are red.
Root cause
The gfx1250 grouped-conv instances added in #10620 are gated on
CK_USE_GFX1250, whichCMakeLists.txt:511defines globally when the target list contains gfx1250. HIP compileseach TU once per
--offload-arch, so every device pass sees the same preprocessor state, andthe kernel bodies are guarded only by
__gfx11__ || __gfx12__— true on gfx1101/gfx1151. A hostmacro cannot express "this offload architecture", so tiles sized for the 320 KB LDS of gfx1250
are emitted for 64 KB parts.
Not caught by CI because no job anywhere builds gfx1250 next to another family:
ck.groovy:1492(the only multi-arch CK configure) omitted gfx1250, and TheRock GHA builds one family per job.
Change
Move every gfx1250-exclusive instance list into dedicated
*_gfx1250*.cppTUs, and compile thoseTUs only for gfx1250. Every other instance TU keeps exactly the architectures it has on develop.
Restricting a TU to gfx1250 needs
--no-offload-archCK's per-source filter in
library/src/tensor_operation_instance/gpu/CMakeLists.txtsets--offload-arch=flags on each TU viaset_source_files_properties. Those flags are additive.CMakeLists.txt:361-365unsetsGPU_TARGETSbeforefind_package(hip)only when the callerused
GPU_ARCHS. When the caller setsGPU_TARGETSinstead,hip::deviceinjects--offload-archfor every target globally, and the per-source flags are appended afterwards —so they can only widen the arch set, never narrow it.
GPU_ARCHS(ck.groovy:1492) → the filter works.GPU_TARGETS(Dockerfile:425) → the filter has no effect.Same TU, clang 24 and clang 22:
What the PR does
_gfx1250sources with--no-offload-arch=all. For those sources only,add_instance_libraryemits--no-offload-arch=allfollowed by their gfx1250 targets.=allalsoclears archs
hip::devicetook from GPU auto-detection or a stale CMake cache, which no explicitexclusion list can see. Every other source's flags are unchanged from develop.
_gfx1250sources when gfx1250 is not targeted. The filter matches the file name, notthe full path:
generate_sharded_instantiationsproduces absolute paths, so matching the pathwould restrict every shard whenever the build directory contains
gfx1250.(
grouped_convolution_*.hpp,gemm_universal.hppand their.incfiles) are installed, so theirguards are compiled by downstream consumers. Each gfx1250 registration sits under
CK_HAS_GFX1250_INSTANCES, exported throughconfig.h.inand set only when the gfx1250 TUs arebuilt.
CK_USE_GFX1250cannot serve: it is build-local, and exporting it would make the host anddevice halves of
CK_ENABLE_CLUSTER_LAUNCH(ck/ck.hpp:205) disagree.add_function withck::is_gfx125_supported(), as…_data_cache_prefetch_instance.cppalready does. Agfx1200;gfx1250build contains the gfx1250 TUs but no gfx1200 code object for them, so they mustnot register on a gfx1200 card.
MIOPEN_CK_UNBUNDLE_PER_ARCH=ON,projects/miopen/cmake/unbundle_archive.cmake): an object with no code for the arch is keptrather than dropped, and an object that thinning would not shrink is kept rather than failing.
pipev1TUs to gfx1250 also removes the gfx9 wave64 offload passes, whichis where the
threadwise_tensor_slice_transfer_v6r1.hpp:60CShufflestatic_assertfires — sothat object builds too.
WMMA conv kernels (fwd, fwd large-tensor, bwd-data, bwd-weight) place their
__shared__array andbody inside
if constexpr(LDS_size <= get_lds_size(get_device_arch())), so such a tile compiles toan empty kernel on that architecture instead of failing the device link. The host side matches:
GridwiseGemm_wmma_cshuffle_v3_base::CheckValidityrejects an instance whose LDS footprint exceedsthe running device's, so an emptied kernel is never selected. Device code for every tile that fits
is unchanged. This is a second barrier, independent of the
_gfx1250file split and CMake flags.Design by Enrico Degregori (co-author of that commit).
All 13 gfx1250 entry points carry a
_gfx1250token, sogit grep 'add_device_.*gfx1250'auditsthe set.
Test Plan / Test Results
Built locally against TheRock 10.2.0a20260929 (AMD clang 24.0.0git),
32-core host, gfx1151 GPU. This PR (
7d4e581998e) against develop —bc47ae72ea8(where the failure was reproduced) forTest A; the merge base
df07597821bfor Test E — with an identical configure:GPU_TARGETS="gfx908;gfx90a;gfx942;gfx950;gfx1101;gfx1151;gfx1250;gfx1250-strict"—GPU_TARGETSrather than
GPU_ARCHSis required to reproduce, for the reason given under Change.bc47ae72ea8local memoryerrorsstatic assertion failedAll ten distinct develop LDS sizes match MIOpen develop nightly #1589.
ninja -k 0: exit 0, counts above.gfx1250;gfx1250-strict: 13 gfx1250 TUs, each built for exactly those two.gfx1101;gfx1151: 0 gfx1250 TUs, base TUs kept.gfx90a+ FP8,gfx942;gfx1100,gfx90a;gfx1030, …): every non-gfx1250 instance TU gets exactly develop's arch flags; all configure cleanly.add_function registers 0 ops.408ed62b668/ after7d4e581998e):IsSupportedArgument→ rejected ("Instance tile too large for LDS size of target device!"); 10 regular instances still accepted.Logs, scripts, repro commands and per-configuration detail:
https://gist.github.com/afagaj/fe34335e1f1ea4baf85d11dedddab531
Limits. No gfx9 or gfx1250 GPU here, so numerical correctness on gfx1250 still needs CI or real
hardware. Test A rebuilt 7 of the ~190 objects in these three libraries; Test E checks every instance
TU's flags but compiles nothing; Test K compares device code, not runtime behaviour — please run the full multi-arch
build in CI before merging.
🤖 Generated with Claude Code