Backport CK develop onto release/rocm-rel-7.0.2.1 for the ROCm 7.0.2.1 delivery - #3773
Open
jdcampbe wants to merge 10 commits into
Open
Backport CK develop onto release/rocm-rel-7.0.2.1 for the ROCm 7.0.2.1 delivery#3773jdcampbe wants to merge 10 commits into
jdcampbe wants to merge 10 commits into
Conversation
… delivery branch This delivery branch carries recent CK to consumers through the existing ROCm 7.0.2.1 delivery path, because CI cannot yet consume CK from its new rocm-libraries monorepo location. The tree is replaced wholesale with projects/composablekernel as it stands in rocm-libraries at dce1b137e55, which is a byte-identical copy of develop@223bff09f: the CK subtree object is e7fd2ba at both commits, so this import introduces no edits of its own. Relative to release/rocm-rel-7.0.2.1 this adds 3656 files, removes 747, and modifies 3634 of the 3685 shared paths. The removals are old instance, test and example sources that CK develop has since deleted or renamed; retaining them would break the build against the new headers. One file is deliberately preserved from release/rocm-rel-7.0.2.1 rather than taken from the monorepo: .azuredevops/rocm-ci.yml. The monorepo subtree has no .azuredevops directory, since the monorepo drives CI its own way, and that pipeline definition is what delivers this branch. Every other CI and packaging input -- Jenkinsfile, .github, groovy/, CMakeLists.txt, requirements.txt -- exists in the monorepo subtree and is taken from it. The four convolution fixes authored on top of this import are applied as separate commits that follow, so they stay individually reviewable.
The group_count=1 fast path in DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffle_v1 builds a flat AK0-M-AK1 descriptor whenever num_group_ == 1 and K >= AK1, regardless of whether N has been split across workgroups (num_workgroups_per_Conv_N_ > 1). When N is split, the flat descriptor still describes the full, unsplit N extent, so a workgroup that only owns one N-slice ends up consuming/producing data for the whole N dimension through it, silently ignoring its own split. Guard population of flat_a_container_/flat_b_container_/flat_c_container_ on a_g_n_k_wos_lengths[1] / conv_N_per_block_ == 1, i.e. only build the flat descriptor when the whole N extent fits in a single workgroup's block. This is the same condition IsSupportedArgument's dispatch guard already checks via num_workgroups_per_Conv_N_ == 1; it is recomputed inline here because num_workgroups_per_Conv_N_ itself is not assigned until later in this function. Once the container is never populated for split-N cases, the dispatch guard's existing "!arg.flat_a_container_.empty()" clause already excludes them, so no other guard needs to change. On the 14-shape backward-data corpus with the size gate bypassed: before, 10 of 14 shapes fail with error 0.49-0.72 against a 0.082 tolerance; after, 14 of 14 shapes pass at 1.1e-05 to 2.1e-05. The same ten shape indices fail before the fix on both gfx950 (MI355X) and gfx942 (MI300X).
Four grouped backward-weight device ops -- DeviceGroupedConvBwdWeight_Xdl_CShuffleV3, DeviceGroupedConvBwdWeight_Xdl_CShuffle, DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle, and DeviceGroupedConvBwdWeight_Xdl_WaveletModel_CShuffleV3 -- each carry a real stride_overflow check in IsSupportedArgument that rejects arguments whose tensors overflow a 32-bit stride. That check is unreachable on the path MIOpen actually calls: the index_t-typed MakeArgumentPointer overload (and, for the two_stage class, specifically its index_t MakeArgumentPointer's non-LargeTensors branch) never computes stride_overflow and leaves it defaulted to false, so the guard silently never fires no matter how large the tensors are. Compute tensor_exceeds_2gb<T>() over the B/C/A tensors (the weight tensor uses tensor_exceeds_2gb<float> in the two_stage class, matching that file's own long_index_t overload) and thread the result into Argument's trailing stride_overflow parameter at each of the four call sites, mirroring the long_index_t overloads in the same files that already do this correctly. Before the fix, the guard's own unsupported-stride_overflow diagnostic fires 0 times while a shape returns a wrong answer; after, it fires 1,254 times. With the size gate relaxed, six previously-miscomputed backward-weight shapes move to CK solver 156 and verify correctly, 124x-459x faster on gfx950 (MI355X) and 144x-343x faster on gfx942 (MI300X).
DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffleV3 and its sibling DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffle_v1 each have a long_index_t-typed MakeArgument/MakeArgumentPointer overload that correctly computes stride_overflow from the A/B/E/Ds tensors and threads it into Argument. The index_t-typed overloads never do this: they leave stride_overflow defaulted to false, so IsSupportedArgument's overflow guard can never reject an oversized argument built through them. The DirectLoad instance is the one path affected in practice: it bypasses the normal global-to-register-to-LDS path and materializes a raw 32-bit flat coordinate offset for its global-to-LDS load, so it is the path that actually overflows once a tensor exceeds 2GB. Device-side instrumentation on the DirectLoad global-load coordinate captured the overflow directly: a negative src_offset of -1090087936 alongside values clustered near +1.83e9, i.e. int32 wraparound. On the failing shape, output batches roughly 100-239 of 240 come back exactly zero (never written, not miscomputed) while the lower batches are largely correct. Size-bisected on the same instance and template parameters: batchsize 32 (~931 MiB, under 2^31 bytes) verifies at 1.35962e-05; batchsize 240 (~6.82 GiB) fails at 0.565379. Compute stride_overflow from the A/B/E/Ds tensors and thread it into Argument at both index_t call sites in each file, mirroring each file's own long_index_t overload. Both files are required: fixing v3 alone leaves the shape failing through a v1 instance carrying the identical bug. After the fix, the oversized shape verifies at 1.34e-05 under competitive MIOPEN_FIND_ENFORCE=SEARCH, 3 of 3 trials, with no instance filtering or solver pinning. Force-pinning the DirectLoad instance on the oversized shape now makes it self-reject at IsSupportedArgument and fall back correctly. A sub-2GB regression check at batchsize 32 still selects the DirectLoad instance and verifies, confirming no regression below the threshold. Verified on gfx950 (MI355X) and gfx942 (MI300X).
…criptor
The packed descriptor fast path computes K0 = K / AK1 with truncating
integer division and performs no remainder handling and no padding of the
dropped tail, so when out_channels is not a multiple of AK1 the trailing
K % AK1 channels are omitted from the backward-data reduction entirely.
They are not computed with wrong values, they are never touched.
The existing guard only requires K >= AK1, which admits every shape whose
channel count is above the tile size but not aligned to it. For the
AK1=8/BK1=8/MPerBlock=64/NPerBlock=64 instance the heuristic commonly
selects, the error scales with the fraction of channels dropped: 4 of 12
gives 0.184, 4 of 20 gives 0.050, 4 of 260 gives 0.0079.
Verified on hardware with an instrumented build. Three affected shapes go
from 0.18357, 0.00788791 and 0.05016 to 2.93617e-06, 2.5624e-06 and
5.90636e-07, landing on stock ROCm 7.0.2's values to three significant
figures. Six shapes that already satisfied the new clause are unchanged at
1e-06, confirming the guard costs nothing where the fast path is safe.
This is a pre-existing defect in CK rather than a consequence of this
branch: it is absent from 7.0.2 and present in develop, introduced with the
flat descriptor mechanism in ROCm/composable_kernel#6208. It remains
unfixed on the develop tip.
The guard declines the fast path rather than repairing the remainder, which
is the conservative interim shape. The proper fix is remainder handling
inside MakeADescriptor_AK0_M_AK1_Packed and MakeBDescriptor_BK0_N_BK1_Packed
and belongs with the CK team. Static impact on the target corpus is 11 of
337 backward-data shapes, 3.3 percent; the 14 headline forward shapes all
have out_channels in {512, 1024, 2048, 4096} and cannot be affected.
The ROCm 7.0.2.1 release CI builds rhel-8.8 with an alternative interpreter at /opt/Python-3.8.13. Two generator scripts imported by CMake at configure time use PEP 585 builtin generics in annotations, which Python only accepts from 3.9 onward. Both raise "TypeError: 'type' object is not subscriptable" while the module is imported, so configure never completes. example/ck_tile/01_fmha/codegen/ops/fmha_bwd.py annotated _api_inners with tuple[str, str]; the module already imports typing.Tuple, so this uses that instead. This is the failure reported by build #41, which aborted at library/src/tensor_operation_instance/gpu/mha/CMakeLists.txt:46 with "CK Tile MHA FAILED to generate a list of kernels via Python". test/ck_tile/gemm_streamk/generate_test_files.py has the same defect and is invoked from test/ck_tile/gemm_streamk/CMakeLists.txt:90. CI never reached it because configure had already failed on the FMHA generator. It gets "from __future__ import annotations", which keeps its several list[...] and dict[...] annotations lazy rather than rewriting each one. script/patch_prefetch_offset.py already carries the __future__ import and needs no change. Verified against a real Python 3.8.20 interpreter, not by inspection. A full CK configure with the CI architecture list (gfx908;gfx90a;gfx942;gfx950;gfx1030;gfx1100;gfx1101;gfx1102;gfx1150;gfx1151;gfx1200;gfx1201) fails identically to CI before these changes, fails on the gemm_streamk generator after only the first is applied, and completes with exit 0 with both. The FMHA generator alone was also run under python:3.8-slim with the exact CMake command line: exit 1 before, exit 0 after.
…_loads instances The 2D grouped convolution backward-data f32 and f32_tf32 "optimized_loads" instances segfault the ROCm 7.0.2.x backend while generating gfx908 code. The crash is in the "SI Form memory clauses" pass and is what failed CI build #41 on rhel-9.5 and ubuntu-22.04: clang++: error: unable to execute command: Segmentation fault (core dumped) clang++: error: clang frontend command failed due to signal Running pass 'SI Form memory clauses' ... -target-cpu gfx908 These instances arrived with the CK develop import; release/rocm-rel-7.0.2.1 has no optimized_loads sources at all, so nothing that shipped previously is affected. Reproduced and measured with the toolchain CI actually uses, in the rocm/dev-ubuntu-22.04:7.0.2 container (AMD clang 20, roc-7.0.2), not with a newer compiler. Per-architecture results at -O3 for both translation units: gfx908 CRASH gfx90a ok gfx942 ok gfx950 ok gfx1030 ok gfx1100 ok gfx1200 ok The defect is specific to gfx908, so gfx908 is the only thing dropped and every other target keeps -O3. -O2 crashes as well; -O1 compiles but would have cost code quality on gfx942 and gfx950, and -mllvm -amdgpu-max-memory-clause=1 and =2 do not avoid the crash. Removing gfx908 from the per-source target list is not sufficient on its own. The global CXX flags already carry a --offload-arch for every GPU target and clang unions all of them, so the architecture has to be taken away again with --no-offload-arch=gfx908, which is emitted after the global flags and wins. A gfx908-only build has no other target to fall back on. There the two translation units keep gfx908 and are compiled at -O1 instead, which is the one setting verified to survive the bug. Builds that also target gfx942 or gfx950 never take that path. Both instance registration functions now return early on gfx908. Without that, the instances would still be registered on a device whose code object is absent, and a heuristic that picked one would fail the launch with hipErrorNoBinaryForGpu. Verification with roc-7.0.2: - a full build of every device instance for gfx908 produced 1895 objects and exactly these two failures, so no other translation unit is affected - full single-architecture builds for gfx950 (2273 objects), gfx942 (2264) and gfx90a (1897) all completed with zero crashes - after this change a gfx908-only build reconfigures and builds the grouped_conv2d_bwd_data_nhwgc instance library with zero crashes - compile_commands.json for a CI-shaped 12-architecture configure carries --no-offload-arch=gfx908 on exactly these two translation units Revert once the backend bug is fixed.
…ances The 3D grouped conv bwd-data f32 and f32_tf32 "optimized_loads" translation units hit the same "SI Form memory clauses" backend segfault as the 2D pair, but the guard added in the previous commit matches conv2d only, so they were still generating gfx908 code at -O3. CI build #41 stopped at 2116/2788 before reaching them, which is why the first pass looked like only two affected translation units. Compiling each of the eight optimized_loads sources for gfx908 individually at -O3 with roc-7.0.2 gives: conv2d bf16 ok conv3d bf16 ok conv2d f16 ok conv3d f16 ok conv2d f32 CRASH conv3d f32 CRASH conv2d f32_tf32 CRASH conv3d f32_tf32 CRASH This widens the existing source_name match to cover the 3D names and adds the same gfx908 registration guard to the two 3D sources, so a heuristic cannot select an instance that has no device image. Verified against a full 11-target configure: all four translation units now carry --no-offload-arch=gfx908, and both 3D units compile with the flags cmake emits for them.
CK now compiles cleanly in CI, and build #43 failed on the last of 2788 steps: linking bin/ckProfiler. FAILED: bin/ckProfiler ld.lld: error: relocation R_X86_64_PC32 out of range: -2172146889 is not in [-2147483648, 2147483647] 20 relocations overflow, all a few percent past the limit. ckProfiler links every device instance library and each object carries embedded device images for all 12 release GPU targets, so after the CK develop import the distance from .text to .rodata exceeds the 2 GB reach of x86-64 PC-relative relocations. It is not a code defect; the binary simply outgrew the small code model. This never appeared before because nothing had reached the link. Build #41 stopped at 2116/2788 and #42 at 2785/2788, both on the gfx908 backend crash; #43 is the first run to compile all of CK, and the link is step 2788. Default BUILD_CK_PROFILER to OFF here rather than passing a flag, because the release pipeline does not set the option at all: its CK configure line passes 22 -D arguments and none of them mention BUILD_CK_PROFILER, so the value comes entirely from this default. Turning it off needs no CI change. Nothing depends on the binary. profiler/ produces a single add_executable and no library, no other CMakeLists references the ckProfiler target, and MIOpen has zero references to it anywhere in its sources, headers, CMake or scripts. MIOpen tunes through its own driver and perf db. This was also exercised end to end: CK was installed with BUILD_CK_PROFILER=OFF and MIOpen then built against it with zero unresolved symbols. Packaging stays consistent because rocm_package_setup_component(profiler ...) and add_subdirectory(profiler) both sit inside if(BUILD_CK_PROFILER), so with the option off the component is never registered and CPack sees no empty component. The only visible effect is that the separate ckprofiler package is no longer produced. Alternative considered: -mcmodel=medium on the profiler target, which would keep the binary. It was not taken because it only buys headroom -- the overflow is already only a few percent over -- and it would have to cover the instance libraries too, not just the final link. Restore with -DBUILD_CK_PROFILER=ON once the binary fits or the code model is addressed.
…lified CK is green in build #47 on all four OS targets; hipTensor now fails instead: hipTensor/library/src/contraction/contraction_cpu_reference_impl.hpp:428:21: error: use of undeclared identifier 'make_ParallelTensorFunctor' make_ParallelTensorFunctor is a CK symbol and it still exists at the same path it always did, include/ck/library/utility/host_tensor.hpp. What changed is its namespace. On release/rocm-rel-7.0.2.1 that header contains no namespace at all and everything in it is global; the CK develop import wraps the whole file in namespace ck. hipTensor is pinned at ff1e8d59, the 7.0.2.1 revision, and calls these names unqualified, so they no longer resolve. Two of the CK headers hipTensor includes made that move. Comparing all 22 CK headers it includes, between the release branch and this one: ck/host_utility/kernel_launch.hpp global -> ck ck/library/utility/host_tensor.hpp global -> ck Rather than fix only the symbol CI happened to stop on, every top-level name declared in those two headers was cross-referenced against every hipTensor source for unqualified use. Three names are affected, not one: make_ParallelTensorFunctor library/src/contraction/contraction_cpu_reference_impl.hpp launch_and_time_kernel library/src/reduction/device/hiptensor_device_reduce_multiblock.hpp joinable_thread library/src/reduction/device/hiptensor_reference_reduce.hpp CI would have surfaced these one build at a time, as happened with the 2D and 3D halves of the gfx908 workaround. All three are re-exported at global scope with using-declarations placed after the namespace closes. In kernel_launch.hpp the shim sits inside the existing #ifndef __HIPCC_RTC__ guard, since the declarations only exist there. Verified against the toolchain CI uses (rocm/dev-ubuntu-22.04:7.0.2, AMD clang 20, roc-7.0.2) with a translation unit that includes both headers and calls the names unqualified exactly as hipTensor does. Without the shims it reproduces hipTensor's failure verbatim -- "unknown type name 'joinable_thread'" and "use of undeclared identifier 'make_ParallelTensorFunctor'". With them it compiles, exit 0. The using-declarations live in the headers, so a successful include also proves each name resolves inside namespace ck. This is deliberately a shim, not the fix. The correct change is for hipTensor to qualify these with ck::, after which this commit should be reverted. It is kept minimal, covering only names hipTensor actually uses, so that reverting it is mechanical. It is also possible that further API breaks lie behind this one; this unblocks the build so the next ones can be found.
jdcampbe
requested review from
a team,
afagaj,
andriy-ca,
aosewski,
bartekxk,
carlushuang,
coderfeli,
ddembeckAMD,
geyyer,
illsilin,
poyenc,
qianfengz,
shumway,
tenpercent and
vidyasagar-amd
as code owners
September 11, 2026 20:14
Author
|
CI is green: hipTensor remained pinned at Description updated with the full build history. |
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
2 participants
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.
Delivers recent CK to ROCm 7.0.2.1 through the existing release path, because CI cannot yet consume CK from its new
rocm-librariesmonorepo location.What this is
The tree is replaced wholesale with
projects/composablekernelas it stands inrocm-librariesatdce1b137e55, which is a byte-identical copy ofdevelop@223bff09f: the CK subtree object ise7fd2bab5f58352dc28f4e1409219b219952e0d7at both commits, so the import introduces no edits of its own. Four grouped-convolution correctness fixes and the changes needed to build under the 7.0.2.1 toolchain sit on top.Relative to
release/rocm-rel-7.0.2.1this adds 3656 files, removes 747, and modifies 3634 of the 3685 shared paths. The removals are instance, test and example sources that CK develop has since deleted or renamed; keeping them would break the build against the new headers.One file is deliberately preserved from the release branch rather than taken from the monorepo:
.azuredevops/rocm-ci.yml. The monorepo subtree has no.azuredevops/directory, and that pipeline is what delivers this branch. Every other CI and packaging input —Jenkinsfile,.github/,groovy/,CMakeLists.txt,requirements.txt— exists in the monorepo subtree and is taken from it.Commits
54f25023erocm-librariesdevelop@223bff09f)4e6e78185a4202aa8cstride_overflowguard in grouped bwd-weight opsdb2cb698bstride_overflowguard in grouped bwd-data ops345a24fe4ad65e2f7fddc242d51optimized_loads(2D)4fee0a6ad4f72e96c9ckProfileron this branch476aad69aThe bulk import is isolated in
54f25023e, so reverting the merge backs the whole thing out in one step, and the four correctness fixes stay individually reviewable.Build fixes, and why each was needed
Python 3.8 (
ad65e2f7f). The rhel-8.8 leg builds with/opt/Python-3.8.13. Two generator scripts imported by CMake at configure time used PEP 585 builtin generics in annotations, which Python only accepts from 3.9, so configure never completed. Verified against a real 3.8.20 interpreter, not by inspection.gfx908 backend crash (
ddc242d51,4fee0a6ad). Fouroptimized_loadstranslation units segfault clang in the AMDGPU "SI Form memory clauses" pass when generating gfx908 code. Measured on the toolchain CI uses (rocm/dev-ubuntu-22.04:7.0.2, AMD clang 20roc-7.0.2):gfx908 is dropped from those four TUs only; every other architecture keeps
-O3. A full gfx908 build of every device instance produced 1895 objects and exactly these failures. The matching registration functions self-reject on gfx908 at runtime, so no kernel without a device image can be launched.ckProfilerlink (4f72e96c9).ckProfilerlinks every device instance library, each carrying embedded device images for all 12 release targets, and after the import the binary exceeded the 2 GB reach of x86-64 PC-relative relocations — 20 ×relocation R_X86_64_PC32 out of range, about 1% over.--offload-compresswas already enabled on all 74 profiler objects and every instance library, so there was no size headroom left. Nothing linksckProfilerand MIOpen has zero references to it, so not building it costs nothing at runtime; it only drops the separateckprofilerpackage.hipTensor compatibility (
476aad69a). Two CK headers moved from the global namespace intonamespace ckin develop. hipTensor is pinned at the 7.0.2.1 revision and calls three of their symbols unqualified. All three are re-exported at global scope.Status
compute-rocm-rel-7.0.2.1build #48 is SUCCESS on all four OS targets, withcomposable_kernelpinned at476aad69a— the head of this branch.Notably hipTensor is still pinned at
ff1e8d59, the unmodified 7.0.2.1 revision, so the compatibility shim in476aad69ais what allows it to build against the imported CK.For the record, the path there:
optimized_loadsoptimized_loadsckProfilerlink, relocation overflowTwo failures seen along the way were diagnosed as not caused by this branch, and both were fixed elsewhere:
lightningfailed on ubuntu-24.04 with an identicalllvm-projectrevision in the passing and failing builds, androcshmemfailed on an orphaned upstream OpenMPI ref.Known caveats
476aad69ais a shim, not a fix. The correct change is for hipTensor to qualify withck::, after which that commit should be reverted. It is minimal and mechanical to revert.ckprofilerpackage is no longer produced. Restore with-DBUILD_CK_PROFILER=ONonce the binary fits or the code model is addressed. With compression already maxed out, it will overflow again as CK grows; the durable fixes are-mcmodel=mediumon the profiler, splitting it, or moving the release to generic GPU targets.