Skip to content

Document the headers #3773 removed from 7.0.2.1, and restore the one that is a pure rename - #3786

Closed
yassinsolim wants to merge 1 commit into
release/rocm-rel-7.0.2.1from
users/ysoliman/ck-7021-removed-headers
Closed

yassinsolim wants to merge 1 commit into
release/rocm-rel-7.0.2.1from
users/ysoliman/ck-7021-removed-headers

Conversation

@yassinsolim

Copy link
Copy Markdown
Member

Closes the documentation half of #3784.

#3773 imported CK develop onto release/rocm-rel-7.0.2.1 and dropped 13 headers that composablekernel-devel shipped in ROCm 7.0.2.1 compute build 50. The StreamK GEMM device API is most of it. Nothing on this branch records that, so a consumer who compiles against the installed headers finds out at compile time after a patch-level update.

What is in here

Two things, and deliberately nothing else.

A 7.0.2.1 changelog section listing every removal with the replacement to move to. Measured on the published RPM file lists of rocm-osdb-rhel-9.x/compute-rocm-rel-7.0.2.1-{50,52}: all 13 are in the build-50 composablekernel-devel file list, none is in the build-52 one, and the package grew from 1,298 to 1,690 files over the same step, so this was targeted and not a packaging accident.

One restored header. device_grouped_conv_bwd_weight_explicit_xdl.hpp comes back, as a compatibility header rather than as the old file. That removal is a pure rename to device_grouped_conv_bwd_weight_explicit.hpp, with DeviceGroupedConvBwdWeight_Explicit_Xdl renamed to DeviceGroupedConvBwdWeight_Explicit and an identical template parameter list, so the old spelling is an alias template. Nothing is duplicated and the cost is one installed header.

Why the other twelve are not restored, measured

Every removed header was restored verbatim from 5b617e0a (build 50) into a copy of the abe93a40 tree, in place, so quoted sibling includes resolve exactly as they did before the import. Each was then compiled standalone with the release toolchain: AMD clang 20.0.0 roc-7.0.2 from the compute-rocm-rel-7.0 rhel-9.x CI image, -x hip --offload-arch=gfx942 -fsyntax-only -O3 -std=c++20. The five replacement headers that survived the import went through the same harness as a control and all five pass, so a failure below is the tree, not the harness.

header restored from build 50 result
ck/tensor_operation/gpu/device/device_gemm_streamk_v2.hpp fails, unknown type name 'StreamKReductionStrategy'
ck/tensor_operation/gpu/device/impl/device_gemm_xdl_streamk.hpp fails, use of undeclared identifier 'StreamKReductionStrategy'
ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_streamk_v3.hpp fails, unknown type name 'StreamKReductionStrategy'
ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_streamk.hpp fails, no member named 'device' in namespace 'ck::tensor_operation'
ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_streamk_v3.hpp fails, no member named 'cerr' in namespace 'std'
ck/library/tensor_operation_instance/gpu/gemm_universal_streamk.hpp fails, unknown type name 'StreamKReductionStrategy'
ck/library/reference_tensor_operation/gpu/naive_conv_fwd.hpp fails, unknown type name 'index_t'
ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_explicit_xdl.hpp fails, no member named 'EnvIsEnabled' in namespace 'ck'
ck/tensor_operation/gpu/device/device_gemm_streamk.hpp compiles
ck/library/tensor_operation_instance/gpu/gemm_streamk.hpp compiles
ck/library/tensor_operation_instance/gpu/grouped_conv_bwd_weight/device_grouped_conv_bwd_weight_wmma_instance.hpp compiles
ck_tile/ops/fmha/pipeline/block_fmha_bwd_pipeline_enum.hpp compiles
ck_tile/ops/reduce/block/block_reduce2d_default_policy.hpp compiles

Eight of thirteen cannot come back at all: StreamKReductionStrategy and the StreamK block-to-ctile helpers went with the import, and ck::EnvIsEnabled is gone, which is also why the old _explicit_xdl header is replaced by an alias instead of restored as it was.

The five that do compile would be worse than leaving them out:

  • device_gemm_streamk.hpp is the abstract interface for device classes that no longer exist.
  • gemm_streamk.hpp compiles because its 18 split-k declarations are still backed. Its one StreamK entry point, add_device_gemm_xdl_streamk_f16_f16_f16_mk_kn_mn_instances, has zero definitions left under library/src, so including it and using StreamK moves the failure from compile time to link time.
  • abe93a40 also deleted library/src/tensor_operation_instance/gpu/gemm_streamk/ and gemm_universal_streamk/, the example/01_gemm/gemm_xdl_*streamk* examples and the profiler op, so there is nothing to register.
  • block_reduce2d_default_policy.hpp compiles but builds a three-argument BlockReduce2dProblem and reads Vector_M/Vector_N from the block shape, where the replacement reads kOutputIndex and ThreadTile_M/ThreadTile_N. Restoring the name would hand callers a policy that silently does not match the new pipeline.

The same import adds a CK Tile StreamK GEMM under include/ck_tile/ops/gemm/kernel/streamk_gemm/, which is where this functionality went. None of the 13 exists on develop either, so the release branch is tracking upstream. The point of this PR is that a ROCm patch release should say so.

Checks

  • The compatibility header compiles against this branch at both -std=c++17 and -std=c++20, with a static_assert that the alias resolves to DeviceGroupedConvBwdWeight_Explicit, same toolchain and flags as above.
  • clang-format 18.1.3, the version pinned in .pre-commit-config.yaml, reports no change on it.
  • script/check_copyright_year.sh, script/check_ascii_only.sh and script/check_no_crlf.sh all exit 0 on it. The pre-commit local hooks cannot be run in this repository as configured: their entry paths point at projects/composablekernel/script/, which only exists in the monorepo layout.
  • No build input changes, so the shipped libraries are unaffected. The only artifact difference is one more header in composablekernel-devel.

One thing I did not touch

experimental/builder/test/test_instance_string_bwd_weight_grp_conv_wmma.cpp still includes device_grouped_conv_bwd_weight_wmma_instance.hpp, which abe93a40 deleted. It is the only remaining reference to any of the 13 anywhere in the tree. No CMakeLists.txt names that file, and CK_EXPERIMENTAL_BUILDER is OFF by default, so it is dead and cannot break a build. I left it alone rather than widen this diff. It should either be deleted or retargeted at device_grouped_conv_bwd_weight_v3_wmma_instance.hpp.

…ename

The CK develop import dropped 13 headers that composablekernel-devel shipped
in ROCm 7.0.2.1 compute build 50, including the whole
ck::tensor_operation::device StreamK GEMM API. Measured on the published RPM
file lists for builds 50 and 52: all 13 present at 50, none at 52, while the
package grew from 1,298 to 1,690 files.

Nothing on this branch records that, so a consumer compiling against the
installed headers finds out at compile time after a patch-level update.

Add a 7.0.2.1 changelog section listing every removal with its replacement,
and restore device_grouped_conv_bwd_weight_explicit_xdl.hpp as a thin
compatibility header. That one is a pure rename to
device_grouped_conv_bwd_weight_explicit.hpp with an identical template
parameter list, so the old spelling is an alias.

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.

🟢 Approval recommended

The alias preserves the template signature, existing installation rules include the header, and no blocking issues remain.

No open issues 🎉

What changed in this PR

Documents the header removals introduced by #3773 in ROCm 7.0.2.1 and restores compatibility for the grouped-convolution header rename.

Changes:

  • Lists all 13 removed headers with migration guidance where available.
  • Restores the old grouped-convolution header and type name through an alias.
File Description
include/​ck/​tensor_operation/​gpu/​device/​impl/​device_grouped_conv_bwd_weight_explicit_xdl.hpp Adds a compatibility alias to the renamed implementation.
CHANGELOG.md Records header removals, replacements, and compatibility support.

Review effort: Balanced

@yassinsolim

Copy link
Copy Markdown
Member Author

Closing: not needed for the 7.0.2.1 delivery.

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.

2 participants