[MIOpen] Sync hipconv v0.3.1

projects/miopen/src/hipconv/CHANGELOG.md

@@ -0,0 +1,160 @@
+# Changelog for hipconv
+
+## v0.3.1 - 2026-09-17
+
+### Added
+
+* `gfx1250-strict` among the CDNA 5 targets, so a build serving MI455 A0 gets its own
+  offload bundle and registry row, and the `GPU_TARGETS` filter keeps the `-strict` suffix
+  it had dropped as a non-offload target. rocm-libraries #12085 made this change to the
+  vendored copy; carrying it here is what keeps it across snapshots.
+* `CHANGELOG.md` ships with the snapshot (#225).
+
+### Changed
+
+* `-fno-offload-lto` on hipconv's own device targets, since device LTO costs up to 8x on
+  these kernels. `-DHIPCONV_OFFLOAD_LTO=ON` opts back in. Reported as
+  ROCm/llvm-project#4434 (#221).
+
+### Fixed
+
+* CDNA 5 `direct` dgrad rejects shapes whose input or weight extent ends mid-dword, where
+  a TDM load corrupts a concurrent `ds_load_tr16_b128` (#224).
+* Moved `direct_wgrad`'s ladder identity out of a `static_assert` and into a test. The
+  compile-time search cost four seconds of each shard's front end and sat close enough to
+  clang's constexpr step limit that GCC 14's libstdc++ pushed it over. Reported as
+  rocm-libraries #11990, and patched there in the vendored copy by rocm-libraries #11989
+  (#218).
+
+## v0.3.0 - 2026-09-11
+
+### Added
+
+| Kernel                   | Architecture     | Layer type       | Direction    | Data type           | Filter, stride                        |
+|--------------------------|------------------|------------------|--------------|---------------------|---------------------------------------|
+| `depthwise_wgrad_hankel` | CDNA 4 (gfx950)  | depthwise conv2d | wgrad        | fp16, bf16; dW fp32 | 3×3 … 11×11 odd square, stride 1 or 2 |
+| `depthwise_1d_toeplitz`  | CDNA 5 (gfx1250) | depthwise conv2d | fprop, dgrad | fp16, bf16, tf32    | 3×3 … 11×11 odd square, stride 1 or 2 |
+
+* tf32 on every CDNA 5 kernel, stored as fp32 and computed by a three-way bf16
+  decomposition (#201).
+* conv3d descriptors, and the unfolding that reduces a degenerate 3D convolution to an
+  equivalent conv2d so those layers reach the existing 2D kernels (#189).
+* An option to emit DWARF line tables for device code (#196).
+
+### Changed
+
+* Take an unserved architecture's offload target from a declared GPU (#211).
+* Removed 1×1 from the CDNA 4 `direct` filter set (#213, #214).
+
+### Fixed
+
+* Worked around the gfx1250 WMMA C-operand hazard (#215).
+
+## v0.2.8 - 2026-09-08
+
+### Changed
+
+* Take an unserved architecture's offload target from a declared GPU, rather than the
+  build machine's.
+
+## v0.2.7 - 2026-08-24
+
+* clang-format over the CDNA 5 grouped kernels.
+
+## v0.2.6 - 2026-08-24
+
+### Fixed
+
+* CDNA 5 grouped convolution: `s_wait_tensorcnt` does not order a `ds_store` against a TDM
+  into the same slot, which corrupted the staged tile.
+* CDNA 5 grouped wgrad failure on gfx1250.
+
+## v0.2.5 - 2026-08-20
+
+### Changed
+
+* Build the kernels for `GPU_TARGETS` rather than the build machine's devices (#190).
+* Derive `direct_wgrad`'s error tolerance from its blocked accumulation (#169).
+
+## v0.2.4 - 2026-08-16
+
+### Changed
+
+* Compile the device code against Microsoft's standard library (#188).
+
+## v0.2.3 - 2026-08-16
+
+### Changed
+
+* Build `libautoshard` as a static library (#187).
+
+## v0.2.2 - 2026-08-14
+
+### Fixed
+
+* Guard `direct_wgrad`'s kernel entry on the gfx950 builtins it needs, so it no longer
+  fails to compile for another target (#186).
+
+## v0.2.1 - 2026-08-14
+
+* Formatting (#185).
+
+## v0.2.0 - 2026-08-14
+
+### Added
+
+| Kernel                  | Architecture    | Layer type       | Direction    | Data type           | Filter, stride                        |
+|-------------------------|-----------------|------------------|--------------|---------------------|---------------------------------------|
+| `direct_wgrad`          | CDNA 4 (gfx950) | dense conv2d     | wgrad        | fp16, bf16; dW fp32 | 2×2, 3×3, 4×4, 5×5, 3×1; stride 1     |
+| `depthwise_1d_toeplitz` | CDNA 4 (gfx950) | depthwise conv2d | fprop, dgrad | fp16, bf16          | 3×3 … 11×11 odd square, stride 1 or 2 |
+| `depthwise_2d_toeplitz` | CDNA 4 (gfx950) | depthwise conv2d | fprop        | fp16, bf16          | 3×3, 5×5, 7×7; stride 1 or 2          |
+
+* The `Depthwise` algorithm and its backend gate, `groups == C == K` (#128, #157).
+  `depthwise_2d_toeplitz` is opt-in behind `HIPCONV_ENABLE_CDNA4_DW_2D_TOEPLITZ` and never
+  wins the auto-pick even when built.
+* CDNA 5 `direct` filters to 5×5, up from 4×4, on a row ring buffer.
+
+### Changed
+
+* Replaced the preprocessor architecture guard with feature guards at the call sites, and
+  built every architecture in every configuration rather than only the one the build
+  targets.
+* Reworked the CDNA 5 `direct` inner loop: NHW folding, a Gray-code MFMA walk order (also
+  enabled on CDNA 4), modulo LDS indexing in place of the prologue and the TDM zero
+  writes, and bank-conflict removal on fprop and dgrad.
+* Let spec layers declare their own directions and data type (#154).
+
+## v0.1.2 - 2026-07-24
+
+### Changed
+
+* Ship `.clang-format` with the snapshot so the consumer's format check applies hipconv's
+  style, and adopt the pre-commit framework (#147).
+
+## v0.1.1 - 2026-07-22
+
+### Changed
+
+* Dropped references to an internal document from shipped comments (#144).
+
+## v0.1.0 - 2026-07-22
+
+### Added
+
+| Kernel          | Architecture     | Layer type       | Direction           | Data type                 | Filter, stride                                   |
+|-----------------|------------------|------------------|---------------------|---------------------------|--------------------------------------------------|
+| `pointwise`     | CDNA 3, 4, 5     | pointwise conv2d | fprop, dgrad, wgrad | fp16, bf16; dW fp32       | 1×1, stride 1                                    |
+| `direct_l1`     | CDNA 4 (gfx950)  | dense conv2d     | fprop, dgrad        | fp16, bf16                | 2×2 … 5×5, stride 1                              |
+| `direct`        | CDNA 4 (gfx950)  | dense conv2d     | fprop, dgrad        | fp16, bf16                | 1×1, 3×3 … 7×7; stride 1                         |
+| `grouped`       | CDNA 4 (gfx950)  | grouped conv2d   | fprop, dgrad        | fp16, bf16, tf32          | 3×3, stride 1 or 2; 4/8/16/32 channels per group |
+| `grouped_wgrad` | CDNA 4 (gfx950)  | grouped conv2d   | wgrad               | fp16, bf16, tf32; dW fp32 | 3×3, stride 1; 4/8/16/32 channels per group      |
+| `direct`        | CDNA 5 (gfx1250) | dense conv2d     | fprop, dgrad        | fp16, bf16                | 1×1 … 4×4, stride 1                              |
+| `grouped`       | CDNA 5 (gfx1250) | grouped conv2d   | fprop, dgrad        | fp16, bf16, tf32          | 3×3, stride 1 or 2; 4/8/16/32 channels per group |
+| `grouped_wgrad` | CDNA 5 (gfx1250) | grouped conv2d   | wgrad               | fp16, bf16, tf32; dW fp32 | 3×3, stride 1; 4/8/16/32 channels per group      |
+
+* `pointwise` (#115) is a hipBLASLt GEMM behind the `ConvKernel` interface, off by default.
+  A library vendoring hipconv serves 1×1 through its own GEMM.
+* The multi-architecture registry mapping a GFX name to an `ArchHandle` and a
+  per-architecture algorithm list, the `ConvKernel` interface with its weighted throughput
+  index, per-kernel configuration descriptors, the autoshard build that splits each kernel
+  family across translation units, and CDNA 3 stub backends.

projects/miopen/src/hipconv/CMakeLists.txt

@@ -7,17 +7,21 @@ cmake_minimum_required(VERSION 3.21)
 # why the stand-in must be a GPU hipconv declares.
 set(HIPCONV_BASELINE_GPU gfx942)
 
-# Take the GPU list from GPU_TARGETS, which is what a ROCm consumer sets.
+# Take the GPU list from GPU_TARGETS when nothing has already named the GPUs.
 #
-# Enabling the HIP language is what fills CMAKE_HIP_ARCHITECTURES, from the build
-# machine's own devices, so this has to precede project(). Entries that are not
-# offload targets are dropped: GPU_TARGETS also carries TheRock family names such as
-# gfx94X-dcgpu, and hipconv serves no generic target.
+# GPU_TARGETS is ROCm's GPU list, set by the consumer or filled from amdgpu-arch by
+# find_package(hip) when unset, and MIOpen enables no HIP language, so it names its
+# GPUs this way alone. CMake's HIP language reads CMAKE_HIP_ARCHITECTURES and fills
+# it from the build machine's own devices, so this has to precede project(). A
+# consumer that sets CMAKE_HIP_ARCHITECTURES itself, as TheRock's toolchain does,
+# skips this. Entries naming no offload target are dropped, hipconv serving no
+# generic target. A '-strict' suffix names a specific silicon revision rather than a
+# generic family, so it is kept.
 if(NOT CMAKE_HIP_ARCHITECTURES AND GPU_TARGETS)
     set(_hipconv_gpus "")
     foreach(_hipconv_gpu IN LISTS GPU_TARGETS)
         string(REGEX REPLACE ":.*" "" _hipconv_logical "${_hipconv_gpu}")
-        if(_hipconv_logical MATCHES "^gfx[0-9a-f]+$")
+        if(_hipconv_logical MATCHES "^gfx[0-9a-f]+(-strict)?$")
             list(APPEND _hipconv_gpus "${_hipconv_gpu}")
         else()
             message(STATUS
@@ -87,6 +91,25 @@ endif()
 option(HIPCONV_ENABLE_CDNA4_DW_2D_TOEPLITZ
     "Build the experimental cdna4 depthwise 2D patch-Toeplitz kernel (slower than 1D; not auto-selected)" OFF)
 
+option(HIPCONV_LINE_TABLES "Emit DWARF line tables for device code (-gline-tables-only)" OFF)
+if(HIPCONV_LINE_TABLES)
+    string(APPEND CMAKE_HIP_FLAGS " -gline-tables-only")
+endif()
+
+# Device LTO costs up to 8x on the kernels this library exists to provide.
+#
+# Some ROCm clangs default the AMDGCN LTO mode to full, and under that pipeline the
+# register allocator spills into direct_wgrad's software-pipelined row loop. The
+# computation is unchanged and no test can see it. docs/rocm-device-lto-regression.md
+# has the measurements; ROCm/llvm-project#4434 is the upstream report.
+#
+# Carried on the targets rather than in CMAKE_HIP_FLAGS because a consumer may compile
+# these sources as CXX -- MIOpen does -- and then the HIP flags never reach them.
+option(HIPCONV_OFFLOAD_LTO "Leave the compiler's device LTO default alone" OFF)
+if(NOT HIPCONV_OFFLOAD_LTO)
+    set(HIPCONV_DEVICE_COMPILE_OPTIONS -fno-offload-lto)
+endif()
+
 # Every arch in the tree is built, whatever this build targets: the set of arch
 # directories is a property of the snapshot, the GPU list is not.
 file(GLOB _hipconv_arch_directories RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} "src/arch/*")
@@ -115,7 +138,7 @@ set(_arch_registry "src/arch_registry.cpp")
 hipconv_make_arch_registry("${_arch_targets}" "${_arch_registry}")
 
 set(_hipconv_srcs
-    src/conv2d_params.cpp
+    src/conv_params.cpp
     src/conv_kernel.cpp
     src/kv_descriptor.cpp
     src/registry.cpp
@@ -144,7 +167,8 @@ target_include_directories(hipconv
 # Position-independent so the static lib can link into libMIOpen.so
 set_target_properties(hipconv PROPERTIES POSITION_INDEPENDENT_CODE ON)
 target_link_libraries(hipconv PUBLIC hip::device)
-target_compile_options(hipconv PRIVATE -fvisibility=hidden -fvisibility-inlines-hidden)
+target_compile_options(hipconv PRIVATE -fvisibility=hidden -fvisibility-inlines-hidden
+                                       ${HIPCONV_DEVICE_COMPILE_OPTIONS})
 if(HIPCONV_ENABLE_POINTWISE)
     target_link_libraries(hipconv PUBLIC roc::hipblaslt)
     target_compile_definitions(hipconv PUBLIC HIPCONV_ENABLE_POINTWISE)

projects/miopen/src/hipconv/HIPCONV_VERSION

@@ -1,16 +1,17 @@
-EMU version: v0.2.8-0-g4bec0aad
-EMU commit: 4bec0aad76eb76d0650c80be8a4502466e4c1f21
-Snapshot date: 2026-09-08
+EMU version: v0.3.1-0-g93b98ccf
+EMU commit: 93b98ccf7819f92558118d546717bebad7fcca1d
+Snapshot date: 2026-09-17
 Arches: cdna3,cdna4,cdna5
 Files:
+  ./CHANGELOG.md
   ./.clang-format
   ./cmake/autoshard.cpp.in
   ./cmake/GpuArch.cmake
   ./cmake/libautoshard.cpp
   ./cmake/libautoshard.hpp
   ./CMakeLists.txt
   ./HIPCONV_VERSION
-  ./include/hipconv/conv2d_params.hpp
+  ./include/hipconv/conv_params.hpp
   ./include/hipconv/export.hpp
   ./include/hipconv/hipconv.hpp
   ./include/hipconv/tolerance.hpp
@@ -28,6 +29,10 @@ Files:
   ./src/arch/cdna4/depthwise/depthwise_1d_toeplitz/kernel.h
   ./src/arch/cdna4/depthwise/depthwise_2d_toeplitz/config_table.h
   ./src/arch/cdna4/depthwise/depthwise_2d_toeplitz/kernel.h
+  ./src/arch/cdna4/depthwise/depthwise_wgrad_hankel/config_desc.cpp
+  ./src/arch/cdna4/depthwise/depthwise_wgrad_hankel/config_desc.h
+  ./src/arch/cdna4/depthwise/depthwise_wgrad_hankel/config_table.h
+  ./src/arch/cdna4/depthwise/depthwise_wgrad_hankel/kernel.h
   ./src/arch/cdna4/direct_backend.cpp
   ./src/arch/cdna4/direct/direct/config.hpp
   ./src/arch/cdna4/direct/direct/config_matcher.cpp
@@ -86,6 +91,10 @@ Files:
   ./src/arch/cdna5/bunnies_mi400.hpp
   ./src/arch/cdna5/CMakeLists.txt
   ./src/arch/cdna5/depthwise_backend.cpp
+  ./src/arch/cdna5/depthwise/depthwise_1d_toeplitz/config_matcher.cpp
+  ./src/arch/cdna5/depthwise/depthwise_1d_toeplitz/config_matcher.hpp
+  ./src/arch/cdna5/depthwise/depthwise_1d_toeplitz/config_table.h
+  ./src/arch/cdna5/depthwise/depthwise_1d_toeplitz/kernel.h
   ./src/arch/cdna5/direct_backend.cpp
   ./src/arch/cdna5/direct/config.hpp
   ./src/arch/cdna5/direct/config_matcher.cpp
@@ -105,10 +114,10 @@ Files:
   ./src/arch_registry.cpp.in
   ./src/arch_registry.h
   ./src/bunnies.hpp
-  ./src/conv2d_params.cpp
   ./src/conv_kernel.cpp
   ./src/conv_kernel.h
   ./src/conv_kernel_table.h
+  ./src/conv_params.cpp
   ./src/depthwise_algorithm_stub.h
   ./src/depthwise_conv_kernel.h
   ./src/detail.h

projects/miopen/src/hipconv/cmake/GpuArch.cmake

@@ -56,11 +56,13 @@ function(hipconv_add_arch_lib name)
     # One .cpp per kernel header, globbed here.
     #
     # hipconv_autoshard generates the per-config shard .cpp files for direct,
-    # direct_l1 and depthwise_1d/2d_toeplitz into the build tree and adds them
-    # with target_sources, so the glob sees only their host helper .cpp.
+    # direct_l1, depthwise_1d/2d_toeplitz and depthwise_wgrad_hankel into the
+    # build tree and adds them with target_sources, so the glob sees only their
+    # host helper .cpp.
     file(GLOB variant_sources CONFIGURE_DEPENDS
         "grouped/*.cpp"
         "depthwise/*.cpp"
+        "depthwise/*/*.cpp"
         "direct/*.cpp"
         "direct/direct_l1/*.cpp"
         "direct/direct/*.cpp"
@@ -74,6 +76,7 @@ function(hipconv_add_arch_lib name)
         "${HIPCONV_ROOT}/include"
         "${CMAKE_CURRENT_SOURCE_DIR}")
     target_link_libraries(hipconv_arch_${name} PRIVATE hip::device)
+    target_compile_options(hipconv_arch_${name} PRIVATE ${HIPCONV_DEVICE_COMPILE_OPTIONS})
     # Position-independent so the objects can link into libMIOpen.so.
     #
     # hipconv_set_stand_in_offload sets HIP_ARCHITECTURES, once every architecture has

projects/miopen/src/hipconv/cmake/libautoshard.cpp

@@ -36,7 +36,7 @@ int make_autoshard(std::string const& target_dir,
         for(int cfg = shard; cfg < num_configs; cfg += num_shards)
         {
             file << "template void " << the_namespace << "::launch_impl<configs[" << cfg
-                 << "]>(const LaunchParams&, const Conv2dParams&, const void*, const void*, void*, "
+                 << "]>(const LaunchParams&, const ConvParams&, const void*, const void*, void*, "
                     "void*, hipStream_t);";
         }
         file << "}\n";
@@ -55,7 +55,7 @@ int make_autoshard(std::string const& target_dir,
         for(int cfg = 0; cfg < num_configs; ++cfg)
         {
             file << "extern template void " << the_namespace << "::launch_impl<configs[" << cfg
-                 << "]>(const LaunchParams&, const Conv2dParams&, const void*, const void*, void*, "
+                 << "]>(const LaunchParams&, const ConvParams&, const void*, const void*, void*, "
                     "void*, hipStream_t);";
         }
         file << "inline auto kernels=std::array<" << kernel_class << "," << num_configs << ">{";