feat(grouped-conv): port large-tensor / global load-store support (#9258)

[JonathanLichtnerAMD]

JIRA ID: ROCM-27526

Motivation

Grouped convolution kernels in CK previously indexed tensor memory with 32-bit index_t, which caps any single tensor descriptor at 2 GB. Workloads with large batch/spatial dimensions (common in modern vision and 3D models) exceed this limit and were either rejected outright or produced incorrect results due to integer overflow in address computation.

This PR is a backport/port of ROCm/rocm-libraries PR #9258 (ROCm/rocm-libraries#9258) onto release/rocm-rel-7.0.2.1. It brings large-tensor (>2 GB) / global load-store support to the grouped convolution forward, backward-data, and backward-weight paths, along with the supporting infrastructure (64-bit indexing in the descriptor/transform layer, a new XDL cshuffle v3 backward-data device op, GPU-side naive reference kernels, and a GPU verification harness) needed to build, dispatch, and validate these cases.

Technical Details

The core change is threading 64-bit long_index_t through the tensor-indexing machinery so address arithmetic no longer overflows past 2 GB, gated so that small tensors keep using 32-bit index_t and pay no register/perf penalty.

• Index-transform / descriptor layer (multi_index_transform.hpp, multi_index_transform_helper.hpp, tensor_descriptor.hpp, array_multi_index.hpp, statically_indexed_array_multi_index.hpp, dynamic_buffer.hpp, number.hpp, tuple.hpp): replace hard-coded index_t locals in the Merge/carry-check and offset paths with auto/decltype(...) so the accumulator type follows the (possibly 64-bit) upper index instead of truncating to 32 bits.
• Large-tensor dispatch: device grouped-conv ops now carry a LargeTensors template flag and select IndexType = std::conditional_t<LargeTensors, ck::long_index_t, ck::index_t>. When LargeTensors is set the ops build i64 length/stride arrays and skip the old "descriptor bigger than 2GB" rejection in AreDescriptorsSmallerThan2GB(); otherwise the pre-existing guard still fires. New helper include/ck/tensor_operation/gpu/device/tensor_size_check.hpp (tensor_exceeds_2gb<DataType>(lengths)) centralizes the 2 GB threshold check used to choose the path.
• New backward-data XDL cshuffle v3 op (device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle_v3.hpp, ~1.5k lines) and a large refactor/extraction of shared gridwise GEMM code into gridwise_gemm_xdl_cshuffle_common.hpp (~2.5k lines), with gridwise_gemm_xdl_cshuffle_v3_multi_d.hpp heavily reworked (~2.8k lines changed) to support the merged/large paths and NumGroupsToMerge.
• Architecture abstraction: new include/ck/utility/amd_arch.hpp provides compile-time arch tags (gfx9/gfx950/gfx10x/gfx11/gfx12x) and per-arch constants (LDS size, LDS banks, max VGPR count, VGPR-per-SIMD), used to size on-chip resources correctly across targets.
• GPU reference + verification: because a >2 GB host reference is impractical, new GPU-side naive conv kernels (naive_conv_fwd_gpu.hpp, naive_conv_bwd_data_gpu.hpp, naive_conv_bwd_weight_gpu.hpp, naive_conv_utils.hpp) and a gpu_verification.hpp harness (returning a GpuVerifyResult with error count / max error / all-zero detection, implicitly convertible to bool) were added. Correspondingly the profiler do_verification argument was widened from bool to int: 0=off, 1=host reference, 2=GPU reference, across profile_grouped_conv_{fwd,bwd_data,bwd_weight}.cpp and the profile impl headers.
• New instance/factory registrations for *_large_tensors_instance and *_comp_large_tensors_instance across conv2d/conv3d fwd, bwd-data, and bwd-weight (fp16/bf16/fp32), with matching .inc/CMakeLists.txt and add_device_operation_instance.hpp wiring; plus a new client example grouped_conv_bwd_data_xdl_v3_fp16.cpp and a ck4inductor op entry.
• ck_tile compiler-compat fixes (matching PR #9258): drop the illegal __host__ __device__ / CK_TILE_HOST_DEVICE_EXTERN qualifiers from deduction guides (ck_tile/core/numeric/math.hpp, core/tensor/sweep_tile.hpp, core/utility/unary_element_function.hpp, ops/reduce/block/block_reduce.hpp) and replace reinterpret_cast type-punning with __builtin_memcpy in core/utility/random.hpp. The orthogonal per-arch warp-size changes from #9258 (get_warp_size, arch.hpp, get_id.hpp, etc.) are intentionally excluded to keep this PR scoped to the grouped-conv large-tensor feature.

Trade-off: the 32-bit vs 64-bit path is a compile-time (LargeTensors) split rather than runtime, which multiplies instantiated instances and build time but keeps the small-tensor path register-lean. int8 was dropped from the large forward test type list (see Test Plan).

Risk Assessment

Risk Level: 🔴 High

This touches foundational, widely-included indexing/descriptor headers (multi_index_transform.hpp, tensor_descriptor.hpp, dynamic_buffer.hpp) that every CK kernel depends on, performs major refactoring/extraction of gridwise GEMM code, and spans 104 files with ~15k insertions / ~2k deletions.

Impacted Components

• Core tensor-description / index-transform layer (include/ck/tensor_description/*, include/ck/utility/*_multi_index.hpp, dynamic_buffer.hpp, number.hpp, tuple.hpp) — included by essentially all kernels.
• Grouped convolution device ops: fwd (dl / xdl cshuffle / xdl v3 / wmma), bwd-data (xdl v1 / xdl v3 / wmma), bwd-weight (dl / two-stage / explicit / xdl / xdl v3 / wmma).
• Gridwise GEMM: new gridwise_gemm_xdl_cshuffle_common.hpp, reworked gridwise_gemm_xdl_cshuffle_v3_multi_d.hpp and ..._conv_v3.hpp, plus block_to_ctile_map.hpp, gridwise_elementwise_2d.hpp, threadwise/thread-group slice-transfer helpers, and blockwise GEMM pipelines.
• conv-to-gemm transformers (fwd, bwd-data, bwd-weight v1/v2, ngchw↔nhwgc).
• New utilities: amd_arch.hpp, tensor_size_check.hpp.
• ck_tile core/ops headers (compat fixes): ck_tile/core/numeric/math.hpp, core/tensor/sweep_tile.hpp, core/utility/random.hpp, core/utility/unary_element_function.hpp, ops/reduce/block/block_reduce.hpp.
• Library reference + verification: GPU naive conv kernels, gpu_verification.hpp.
• Profilers, device-operation-instance registrations/factories, tests, one client example, and a ck4inductor op.

Potential Side Effects

• Changing index_t accumulators to auto/decltype in the hot Merge/carry-check paths could alter generated code (register pressure, VALU/SGPR selection) for existing small-tensor kernels; watch for perf regressions or numerical changes even when LargeTensors=false.
• The gridwise GEMM extraction/refactor risks behavioral drift in the shared code path used by non-large kernels.
• Widening profiler do_verification from bool to int is a semantic change to the CLI/API surface; callers passing 2 now select GPU reference.
• Larger instance set increases binary size and build time.

Mitigation Steps

• Small-tensor path is preserved behind the compile-time LargeTensors flag and the existing AreDescriptorsSmallerThan2GB() guard still runs for the default path.
• New dedicated large-case tests (fwd/bwd-data/bwd-weight) plus GPU-reference verification for cases the host reference cannot handle.
• This is a direct port of the already-reviewed upstream PR #9258, reducing novel-design risk.
• Recommended before merge: run the full grouped-conv correctness suite and a perf comparison of representative small-tensor instances against the base ref; SME review of the gridwise GEMM refactor.
• Revert plan: the branch is a single squashable feature port; revert the branch merge to restore prior behavior.

Test Plan

• New GoogleTest large-case suites:
  • test/grouped_convnd_fwd/test_grouped_convnd_fwd_large_cases_xdl.cpp (updated to use GPU reference, do_verification=2).
  • test/grouped_convnd_bwd_data/test_grouped_convnd_bwd_data_xdl_large_cases.cpp (updated).
  • test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight_large_cases.cpp (new, 164 lines) and CMake wiring.
• Profiler-driven manual verification via profile_grouped_conv_{fwd,bwd_data,bwd_weight} with do_verification=2 (GPU naive reference), which is the intended validation path for >2 GB tensors where a host reference is impractical.
• int8 was removed from the 2D large forward test type list (KernelTypes2d), matching upstream — the large-tensor path is validated for fp16/bf16/fp32.

Test Result

Correctness for large (>2 GB) grouped-conv fwd/bwd-data/bwd-weight cases is exercised through the new GPU-reference verification path in the added/updated tests and the profiler do_verification=2 mode. Functional parity is expected to match upstream ROCm/rocm-libraries PR #9258, from which this is ported. Known follow-ups: confirm no small-tensor perf regression from the auto/decltype indexing changes, and validate build time / binary size impact from the expanded large-tensor instance set on the 7.0.2.1 toolchain.

[Copilot]

Pull request overview

Ports upstream grouped-convolution “large tensor / global load-store” support into this release branch by extending indexing/descriptor infrastructure to support 64-bit addressing, adding large-tensor instance registrations, and introducing GPU-side naive reference + GPU verification paths for >2GB correctness validation.

Changes:

• Adds/threads 64-bit (long_index_t) indexing through tensor descriptors/transforms and selected kernel utilities to avoid >2GB address overflows.
• Expands grouped-conv instance/factory registrations with *_large_tensors_instance variants (2D/3D fwd/bwd-data/bwd-weight).
• Adds GPU naive reference conv kernels + GPU verification harness and widens profiler/test “verification” control from bool to int (0/1/2).

Reviewed changes

Copilot reviewed 107 out of 109 changed files in this pull request and generated 4 comments.

Show a summary per file

File	Description
test/grouped_convnd_fwd/test_grouped_convnd_fwd_large_cases_xdl.cpp	Switches large-case fwd tests to GPU reference verification; adjusts tested type list.
test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight.cpp	Updates split-k argument passing to match updated profiler API.
test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight_large_cases.cpp	Adds large-case bwd-weight gtests using GPU reference verification.
test/grouped_convnd_bwd_weight/CMakeLists.txt	Builds bwd-weight tests and adds a large-cases test executable.
test/grouped_convnd_bwd_data/test_grouped_convnd_bwd_data_xdl_large_cases.cpp	Updates bwd-data large-case tests (verification mode + reduced kernel type matrix).
test/conv_util/conv_util.cpp	Adds grouped-conv utility include for test helpers.
python/ck4inductor/grouped_conv_fwd/op.py	Adds large_tensor option to ck4inductor grouped-conv fwd op config.
profiler/src/profile_grouped_conv_fwd.cpp	Changes CLI parsing: do_verification widened from bool to int.
profiler/src/profile_grouped_conv_bwd_weight.cpp	Changes CLI parsing for do_verification (int) and split-k handling (string).
profiler/src/profile_grouped_conv_bwd_data.cpp	Changes CLI parsing: do_verification widened from bool to int.
profiler/include/profiler/profile_grouped_conv_fwd_impl.hpp	Adds GPU naive reference + GPU verification path and supports do_verification=2.
profiler/include/profiler/profile_grouped_conv_bwd_data_impl.hpp	Adds GPU naive reference + GPU verification path and switches descriptor arrays to long_index_t.
library/src/tensor_operation_instance/gpu/grouped_conv3d_fwd/xdl/comp/device_grouped_conv3d_fwd_xdl_ndhwgc_gkzyxc_ndhwgk_f32_comp_large_tensors_instance.cpp	Registers 3D fwd f32 “comp” large-tensor instances.
library/src/tensor_operation_instance/gpu/grouped_conv3d_fwd/xdl/comp/device_grouped_conv3d_fwd_xdl_ndhwgc_gkzyxc_ndhwgk_f16_comp_large_tensors_instance.cpp	Registers 3D fwd f16 “comp” large-tensor instances.
library/src/tensor_operation_instance/gpu/grouped_conv3d_fwd/xdl/comp/device_grouped_conv3d_fwd_xdl_ndhwgc_gkzyxc_ndhwgk_bf16_comp_large_tensors_instance.cpp	Registers 3D fwd bf16 “comp” large-tensor instances.
library/src/tensor_operation_instance/gpu/grouped_conv3d_fwd/CMakeLists.txt	Adds large-tensor fwd instance TU files to build.
library/src/tensor_operation_instance/gpu/grouped_conv3d_bwd_weight/xdl/ndhwgc_gkzyxc_ndhwgk/device_grouped_conv3d_bwd_weight_xdl_ndhwgc_gkzyxc_ndhwgk_f32_default_large_tensors_instance.cpp	Registers 3D bwd-weight f32 default large-tensor instances.
library/src/tensor_operation_instance/gpu/grouped_conv3d_bwd_weight/xdl/ndhwgc_gkzyxc_ndhwgk/device_grouped_conv3d_bwd_weight_xdl_ndhwgc_gkzyxc_ndhwgk_f16_default_large_tensors_instance.cpp	Registers 3D bwd-weight f16 default large-tensor instances.
library/src/tensor_operation_instance/gpu/grouped_conv3d_bwd_weight/xdl/ndhwgc_gkzyxc_ndhwgk/device_grouped_conv3d_bwd_weight_xdl_ndhwgc_gkzyxc_ndhwgk_bf16_default_large_tensors_instance.cpp	Registers 3D bwd-weight bf16 default large-tensor instances.
library/src/tensor_operation_instance/gpu/grouped_conv3d_bwd_weight/xdl/ndhwgc_gkzyxc_ndhwgk/device_grouped_conv3d_bwd_weight_two_stage_xdl_ndhwgc_gkzyxc_ndhwgk_f16_large_tensors_instance.cpp	Registers 3D bwd-weight f16 two-stage large-tensor instances.
library/src/tensor_operation_instance/gpu/grouped_conv3d_bwd_weight/xdl/ndhwgc_gkzyxc_ndhwgk/device_grouped_conv3d_bwd_weight_two_stage_xdl_ndhwgc_gkzyxc_ndhwgk_bf16_large_tensors_instance.cpp	Registers 3D bwd-weight bf16 two-stage large-tensor instances.
library/src/tensor_operation_instance/gpu/grouped_conv3d_bwd_weight/CMakeLists.txt	Adds large-tensor 3D bwd-weight instance TU files to build.
library/src/tensor_operation_instance/gpu/grouped_conv3d_bwd_data/xdl/device_grouped_conv3d_bwd_data_xdl_v3_ndhwgc_gkzyxc_ndhwgk_f32_large_tensors_instance.cpp	Registers 3D bwd-data v3 f32 large-tensor instances.
library/src/tensor_operation_instance/gpu/grouped_conv3d_bwd_data/xdl/device_grouped_conv3d_bwd_data_xdl_v3_ndhwgc_gkzyxc_ndhwgk_f16_large_tensors_instance.cpp	Registers 3D bwd-data v3 f16 large-tensor instances.
library/src/tensor_operation_instance/gpu/grouped_conv3d_bwd_data/xdl/device_grouped_conv3d_bwd_data_xdl_v3_ndhwgc_gkzyxc_ndhwgk_bf16_large_tensors_instance.cpp	Registers 3D bwd-data v3 bf16 large-tensor instances.
library/src/tensor_operation_instance/gpu/grouped_conv3d_bwd_data/CMakeLists.txt	Adds large-tensor 3D bwd-data v3 instance TU files to build.
library/src/tensor_operation_instance/gpu/grouped_conv2d_fwd/xdl/comp/device_grouped_conv2d_fwd_xdl_nhwgc_gkyxc_nhwgk_f32_comp_large_tensors_instance.cpp	Registers 2D fwd f32 “comp” large-tensor instances.
library/src/tensor_operation_instance/gpu/grouped_conv2d_fwd/xdl/comp/device_grouped_conv2d_fwd_xdl_nhwgc_gkyxc_nhwgk_f16_comp_large_tensors_instance.cpp	Registers 2D fwd f16 “comp” large-tensor instances.
library/src/tensor_operation_instance/gpu/grouped_conv2d_fwd/xdl/comp/device_grouped_conv2d_fwd_xdl_nhwgc_gkyxc_nhwgk_bf16_comp_large_tensors_instance.cpp	Registers 2D fwd bf16 “comp” large-tensor instances.
library/src/tensor_operation_instance/gpu/grouped_conv2d_fwd/CMakeLists.txt	Adds large-tensor 2D fwd instance TU files to build.
library/src/tensor_operation_instance/gpu/grouped_conv2d_bwd_weight/xdl/nhwgc_gkyxc_nhwgk/device_grouped_conv2d_bwd_weight_xdl_nhwgc_gkyxc_nhwgk_f32_default_large_tensors_instance.cpp	Registers 2D bwd-weight f32 default large-tensor instances.
library/src/tensor_operation_instance/gpu/grouped_conv2d_bwd_weight/xdl/nhwgc_gkyxc_nhwgk/device_grouped_conv2d_bwd_weight_xdl_nhwgc_gkyxc_nhwgk_f16_default_large_tensors_instance.cpp	Registers 2D bwd-weight f16 default large-tensor instances.
library/src/tensor_operation_instance/gpu/grouped_conv2d_bwd_weight/xdl/nhwgc_gkyxc_nhwgk/device_grouped_conv2d_bwd_weight_xdl_nhwgc_gkyxc_nhwgk_bf16_default_large_tensors_instance.cpp	Registers 2D bwd-weight bf16 default large-tensor instances.
library/src/tensor_operation_instance/gpu/grouped_conv2d_bwd_weight/xdl/nhwgc_gkyxc_nhwgk/device_grouped_conv2d_bwd_weight_two_stage_xdl_nhwgc_gkyxc_nhwgk_f16_large_tensors_instance.cpp	Registers 2D bwd-weight f16 two-stage large-tensor instances.
library/src/tensor_operation_instance/gpu/grouped_conv2d_bwd_weight/xdl/nhwgc_gkyxc_nhwgk/device_grouped_conv2d_bwd_weight_two_stage_xdl_nhwgc_gkyxc_nhwgk_bf16_large_tensors_instance.cpp	Registers 2D bwd-weight bf16 two-stage large-tensor instances.
library/src/tensor_operation_instance/gpu/grouped_conv2d_bwd_weight/CMakeLists.txt	Adds large-tensor 2D bwd-weight instance TU files to build.
library/src/tensor_operation_instance/gpu/grouped_conv2d_bwd_data/xdl/device_grouped_conv2d_bwd_data_xdl_v3_nhwgc_gkyxc_nhwgk_f32_large_tensors_instance.cpp	Registers 2D bwd-data v3 f32 large-tensor instances.
library/src/tensor_operation_instance/gpu/grouped_conv2d_bwd_data/xdl/device_grouped_conv2d_bwd_data_xdl_v3_nhwgc_gkyxc_nhwgk_f16_large_tensors_instance.cpp	Registers 2D bwd-data v3 f16 large-tensor instances.
library/src/tensor_operation_instance/gpu/grouped_conv2d_bwd_data/xdl/device_grouped_conv2d_bwd_data_xdl_v3_nhwgc_gkyxc_nhwgk_bf16_large_tensors_instance.cpp	Registers 2D bwd-data v3 bf16 large-tensor instances.
library/src/tensor_operation_instance/gpu/grouped_conv2d_bwd_data/CMakeLists.txt	Adds large-tensor 2D bwd-data v3 instance TU files to build.
library/include/ck/library/tensor_operation_instance/gpu/grouped_convolution_forward.hpp	Wires large-tensor forward instances into the forward factory.
library/include/ck/library/tensor_operation_instance/gpu/grouped_convolution_forward_comp_xdl.inc	Declares large-tensor forward “comp xdl” instance adders.
library/include/ck/library/tensor_operation_instance/gpu/grouped_convolution_backward_weight.hpp	Wires large-tensor bwd-weight instances into the bwd-weight factory.
library/include/ck/library/tensor_operation_instance/gpu/grouped_convolution_backward_weight_xdl.inc	Declares large-tensor bwd-weight xdl instance adders.
library/include/ck/library/tensor_operation_instance/gpu/grouped_convolution_backward_data.hpp	Wires large-tensor bwd-data v3 instances into the bwd-data factory.
library/include/ck/library/tensor_operation_instance/gpu/grouped_convolution_backward_data_xdl.inc	Declares large-tensor bwd-data v3 xdl instance adders.
library/include/ck/library/tensor_operation_instance/gpu/grouped_conv_bwd_weight/device_grouped_conv_bwd_weight_two_stage_xdl_instance.hpp	Adds two-stage large-tensor instance tuples for (b)f16.
library/include/ck/library/tensor_operation_instance/add_device_operation_instance.hpp	Refactors instance registration helper to default-construct tuple element types.
library/include/ck/library/reference_tensor_operation/gpu/naive_conv_utils.hpp	Adds shared GPU helper utilities for naive reference conv kernels.
include/ck/utility/tuple.hpp	Adds std::tuple compatibility to ck::tuple_element_t.
include/ck/utility/statically_indexed_array_multi_index.hpp	Generalizes MultiIndex to support alternative index types (e.g., long_index_t).
include/ck/utility/random_gen.hpp	Replaces type-punning casts with bit_cast for compiler compatibility.
include/ck/utility/number.hpp	Adds is_number/is_long_number type traits.
include/ck/utility/dynamic_buffer.hpp	Generalizes make_dynamic_buffer to support alternative index types; removes separate long helper.
include/ck/utility/array_multi_index.hpp	Generalizes MultiIndex to support alternative index types (e.g., long_index_t).
include/ck/utility/amd_arch.hpp	Adds compile-time AMD arch tags and per-arch constants.
include/ck/tensor_operation/operator_transform/transform_conv_fwd_to_gemm.hpp	Uses index-type-aware Number/LongNumber in transforms to preserve 64-bit indexing.
include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v6r1.hpp	Adds configurable coordinate index type for destination coordinate computation.
include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v3r2.hpp	Adds configurable coordinate index type and avoids repeated GetOffset calls.
include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v3r1.hpp	Adds long-index-safe coordinate handling for loads and scratch masking.
include/ck/tensor_operation/gpu/grid/gridwise_elementwise_2d.hpp	Adds IndexType support for batch counts/strides and global buffers.
include/ck/tensor_operation/gpu/grid/block_to_ctile_map.hpp	Adds IndexType support for block-to-tile mapping math.
include/ck/tensor_operation/gpu/device/tensor_size_check.hpp	Adds centralized 2GB tensor size threshold check utility.
include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_d_wmma_cshuffle.hpp	Adds 2GB overflow rejection plumbing for 32-bit path arguments.
include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp	Adds 2GB overflow rejection plumbing for 32-bit path arguments.
include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_dl_multiple_d_nhwc_kyxc_nhwk.hpp	Adds 2GB overflow rejection plumbing for 32-bit path arguments.
include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_xdl_cshuffle.hpp	Adds long-index MakeArgument overloads and 2GB overflow rejection plumbing.
include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_wmma_cshuffle.hpp	Adds long-index MakeArgument overloads and 2GB overflow rejection plumbing.
include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_explicit_xdl.hpp	Adds long-index MakeArgumentPointer overload that narrows to 32-bit path.
include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_dl.hpp	Adds long-index MakeArgument overloads and 2GB overflow rejection plumbing.
include/ck/tensor_operation/gpu/device/device_grouped_conv_bwd_weight.hpp	Extends base interface to accept long-index (64-bit) lengths/strides.
include/ck/tensor_operation/gpu/device/device_grouped_conv_bwd_data_multiple_d.hpp	Extends base interface to accept long-index (64-bit) lengths/strides.
include/ck/tensor_operation/gpu/block/thread_group_tensor_slice_transfer_v6r1.hpp	Propagates IndexType support through thread-group transfer wrapper.
include/ck/tensor_operation/gpu/block/thread_group_tensor_slice_transfer_v4r2.hpp	Propagates IndexType support through thread-group transfer wrapper.
include/ck/tensor_operation/gpu/block/thread_group_tensor_slice_transfer_v4r1.hpp	Propagates IndexType support through thread-group transfer wrapper.
include/ck/tensor_operation/gpu/block/blockwise_gemm_xdlops.hpp	Adjusts wave-size computation to avoid warp-size dependencies.
include/ck/tensor_operation/gpu/block/blockwise_gemm_smfmac_xdlops.hpp	Adjusts wave-size computation to avoid warp-size dependencies.
include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops.hpp	Adjusts wave-size computation to avoid warp-size dependencies.
include/ck/tensor_operation/gpu/block/blockwise_gemm_dpp.hpp	Adjusts wave-size computation to avoid warp-size dependencies.
include/ck/tensor_description/tensor_descriptor.hpp	Adds IndexType to TensorCoordinate and make_tensor_coordinate; returns offset as IndexType.
include/ck/tensor_description/multi_index_transform.hpp	Uses auto/decltype-based temporaries to avoid truncating 64-bit indices in merges/carry logic.
include/ck/tensor_description/multi_index_transform_helper.hpp	Selects merge transform variant compatible with 64-bit indices.
include/ck_tile/ops/reduce/block/block_reduce.hpp	Removes invalid host/device extern qualifiers from deduction guide.
include/ck_tile/core/utility/unary_element_function.hpp	Removes invalid host/device qualifiers from deduction guide.
include/ck_tile/core/utility/random.hpp	Replaces reinterpret_cast type punning with __builtin_memcpy.
include/ck_tile/core/tensor/sweep_tile.hpp	Removes invalid host/device extern qualifiers from deduction guide.
include/ck_tile/core/numeric/math.hpp	Removes invalid host/device qualifiers from deduction guides.
example/38_grouped_conv_bwd_data_multiple_d/grouped_conv_bwd_data_xdl_v3_fp16.cpp	Adds a reproducer for an ASAN issue using a v3 bwd-data instance with large-tensor indexing.

---

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.