[MIOpen] Sync hipconv v0.4.0

projects/miopen/CHANGELOG.md

@@ -11,6 +11,8 @@ Full documentation for MIOpen is available [here](https://rocm.docs.amd.com/proj
 * [Conv] Added gfx950 patch-embedding kernels (fp16/bf16), gfx950 TF32 kernels for dense
   and depthwise forward/backward-data and dense backward-weights, and gfx1250 depthwise
   backward-weights kernels (fp16/bf16) to the `ConvHipConv` solver (hipconv v0.3.2).
+* [Conv] Added a gfx1250 dense backward-weights kernel (fp16/bf16) to the `ConvHipConv`
+  solver (hipconv v0.4.0).
 
 ### Removed
 * [Conv] Removed gfx803 convolution solver `ConvBinWinogradRxSFused` and its kernel sources.
@@ -21,6 +23,9 @@ Full documentation for MIOpen is available [here](https://rocm.docs.amd.com/proj
 * [Conv] Fixed a `ConvHipConv` crash on gfx1250 for grouped convolutions with more than
   65535 group-batch tiles, and an LDS race in its gfx1250 direct kernel
   (hipconv v0.3.2).
+* [Conv] Fixed `ConvHipConv` losing its tuned config on a perf-db reload, which could fail
+  execution with "not enough workspace for direct kernel" after an exhaustive Find. Records
+  now carry the hipconv minor version, and records from another version are discarded.
 
 ## MIOpen 3.6.1 for ROCm 10.1.0
 

projects/miopen/src/hipconv/CHANGELOG.md

@@ -1,5 +1,40 @@
 # Changelog for hipconv
 
+## v0.4.0 - 2026-10-06
+
+### Added
+
+| Kernel         | Architecture     | Layer type            | Direction | Data type           | Filter, stride                          | PR         |
+|----------------|------------------|-----------------------|-----------|---------------------|-----------------------------------------|------------|
+| `direct_wgrad` | CDNA 5 (gfx1250) | dense, grouped conv2d | wgrad     | fp16, bf16; dW fp32 | 2×2, 3×3, 4×4, 5×5, 3×1, 1×3; stride 1  | #212, #297 |
+
+* `describe_config(kernel, full)`: `full` includes default values, so the description selects
+  exactly that config through `matches_descriptor()` (#291).
+* `version()`, this build's version as `git describe --tags` names it (#291).
+* `config_version()`, the major and minor version (e.g. `v0.4`), for keying a stored config
+  description. Within a minor version, a full description always selects the same config
+  (#291).
+
+### Fixed
+
+* Add the missing config descriptors to `grouped`, `grouped_wgrad`, and `depthwise_2d_toeplitz`.
+  Change the `depthwise_1d_toeplitz` descriptors to use the standard form. Add the missing
+  `row_pad` field to the `patch_embed` descriptor (#291).
+* Every config descriptor registers every field of its kernel's `Config`, so a `--config` spec
+  can name any of them. Fields added to `direct` (CDNA 4 and CDNA 5), `direct_l1`,
+  `direct_wgrad`, and `depthwise_wgrad_hankel` (CDNA 4 and CDNA 5). The short description of
+  every `direct` config now includes `kh` and `kw`, and on CDNA 4 also `tile_size_c` (#291).
+* `ConvParams::is_valid()` rejects a layer whose dilated filter is wider than its padded
+  input on some axis (`output_is_empty()`). Such a layer reached kernel matching and scored
+  a NaN weighted throughput index (#167, #288).
+* `ConvLaunch::make` returns `nullopt`, and `launch()` returns `hipErrorNotSupported`, for
+  an explicit-GEMM pointwise or patch-embedding layer whose GEMM hipBLASLt has no algorithm
+  for on the current device. The launch failed with `HIPBLAS_STATUS_INTERNAL_ERROR` instead,
+  and ended the driver's run. Selection is unchanged and still needs no GPU; the driver
+  binds the next ranked kernel or reports `no_match` (#244, #281).
+* `launch()` and `ConvLaunch::launch` return `hipErrorUnknown` when any other hipBLASLt call
+  in an explicit-GEMM launch fails. The hipBLASLt exception escaped both instead (#281).
+
 ## v0.3.2 - 2026-10-02
 
 ### Added

projects/miopen/src/hipconv/CMakeLists.txt

@@ -151,6 +151,8 @@ hipconv_make_arch_registry("${_arch_targets}" "${_arch_registry}")
 set(_hipconv_srcs
     src/conv_params.cpp
     src/conv_kernel.cpp
+    src/data_type_field.cpp
+    src/direction_field.cpp
     src/kv_descriptor.cpp
     src/registry.cpp
     src/tolerance.cpp
@@ -166,8 +168,23 @@ if(HIPCONV_ENABLE_EXPLICIT_GEMM)
     )
 endif()
 
+# hipconv::version(), regenerated at build time so a new commit needs no reconfigure;
+# cmake/write_version.cmake says where the version comes from.
+set(_version_cpp "${CMAKE_CURRENT_BINARY_DIR}/version.cpp")
+add_custom_target(hipconv_version
+    COMMAND "${CMAKE_COMMAND}"
+            "-DSOURCE_DIR=${CMAKE_CURRENT_SOURCE_DIR}"
+            "-DTEMPLATE=${CMAKE_CURRENT_SOURCE_DIR}/cmake/version.cpp.in"
+            "-DOUTPUT=${_version_cpp}"
+            -P "${CMAKE_CURRENT_SOURCE_DIR}/cmake/write_version.cmake"
+    BYPRODUCTS "${_version_cpp}"
+    VERBATIM)
+list(APPEND _hipconv_srcs "${_version_cpp}")
+
 add_library(hipconv STATIC ${_hipconv_srcs} "${_reg_gen}")
 set_source_files_properties(${_hipconv_srcs} "${_reg_gen}" PROPERTIES LANGUAGE HIP)
+set_source_files_properties("${_version_cpp}" PROPERTIES GENERATED TRUE)
+add_dependencies(hipconv hipconv_version)
 foreach(_target IN LISTS _arch_targets)
     target_sources(hipconv PRIVATE $<TARGET_OBJECTS:${_target}>)
 endforeach()

projects/miopen/src/hipconv/HIPCONV_VERSION

@@ -1,6 +1,6 @@
-EMU version: v0.3.2-0-ga6c85d2a
-EMU commit: a6c85d2a8592b918d9c79f94b4d3bf6473021fb3
-Snapshot date: 2026-10-02
+EMU version: v0.4.0
+EMU commit: d6e403450af5fabd2f38c6bebb58cab2b8274bf8
+Snapshot date: 2026-10-06
 Arches: cdna3,cdna4,cdna5
 Files:
   ./CHANGELOG.md
@@ -10,6 +10,8 @@ Files:
   ./cmake/libautoshard.cpp
   ./cmake/libautoshard.hpp
   ./CMakeLists.txt
+  ./cmake/version.cpp.in
+  ./cmake/write_version.cmake
   ./HIPCONV_VERSION
   ./include/hipconv/conv_params.hpp
   ./include/hipconv/export.hpp
@@ -25,8 +27,12 @@ Files:
   ./src/arch/cdna4/bunnies_cdna4.hpp
   ./src/arch/cdna4/CMakeLists.txt
   ./src/arch/cdna4/depthwise_backend.cpp
+  ./src/arch/cdna4/depthwise/depthwise_1d_toeplitz/config_desc.cpp
+  ./src/arch/cdna4/depthwise/depthwise_1d_toeplitz/config_desc.h
   ./src/arch/cdna4/depthwise/depthwise_1d_toeplitz/config_table.h
   ./src/arch/cdna4/depthwise/depthwise_1d_toeplitz/kernel.h
+  ./src/arch/cdna4/depthwise/depthwise_2d_toeplitz/config_desc.cpp
+  ./src/arch/cdna4/depthwise/depthwise_2d_toeplitz/config_desc.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
@@ -116,10 +122,19 @@ Files:
   ./src/arch/cdna5/direct/config_matcher.hpp
   ./src/arch/cdna5/direct/config_table.hpp
   ./src/arch/cdna5/direct/kernel.hpp
+  ./src/arch/cdna5/direct_wgrad/config.hpp
+  ./src/arch/cdna5/direct_wgrad/config_matcher.cpp
+  ./src/arch/cdna5/direct_wgrad/config_matcher.hpp
+  ./src/arch/cdna5/direct_wgrad/config_table.hpp
+  ./src/arch/cdna5/direct_wgrad/kernel.hpp
   ./src/arch/cdna5/grouped_backend.cpp
+  ./src/arch/cdna5/grouped/grouped_multi_g/config_desc.cpp
+  ./src/arch/cdna5/grouped/grouped_multi_g/config_desc.h
   ./src/arch/cdna5/grouped/grouped_multi_g/config.hpp
   ./src/arch/cdna5/grouped/grouped_multi_g/config_table.hpp
   ./src/arch/cdna5/grouped/grouped_multi_g/kernel.hpp
+  ./src/arch/cdna5/grouped/grouped_multi_g_wgrad/config_desc.cpp
+  ./src/arch/cdna5/grouped/grouped_multi_g_wgrad/config_desc.h
   ./src/arch/cdna5/grouped/grouped_multi_g_wgrad/config.hpp
   ./src/arch/cdna5/grouped/grouped_multi_g_wgrad/config_table.hpp
   ./src/arch/cdna5/grouped/grouped_multi_g_wgrad/kernel.hpp
@@ -133,11 +148,15 @@ Files:
   ./src/conv_kernel.h
   ./src/conv_kernel_table.h
   ./src/conv_params.cpp
+  ./src/data_type_field.cpp
+  ./src/data_type_field.h
   ./src/depthwise_algorithm_stub.h
   ./src/depthwise_conv_kernel.h
   ./src/detail.h
   ./src/direct_algorithm_stub.h
   ./src/direct_conv_kernel.h
+  ./src/direction_field.cpp
+  ./src/direction_field.h
   ./src/explicit_gemm/explicit_gemm_algorithm.cpp
   ./src/explicit_gemm/hipblaslt_matmul.cpp
   ./src/explicit_gemm/hipblaslt_matmul.hpp
@@ -167,3 +186,4 @@ Files:
   ./src/transpose_lds_layout.h
   ./src/types.h
   ./src/unreachable.h
+  ./VERSION

projects/miopen/src/hipconv/VERSION

@@ -0,0 +1 @@
+v0.4.0

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

@@ -61,9 +61,11 @@ function(hipconv_add_arch_lib name)
     # host helper .cpp.
     file(GLOB variant_sources CONFIGURE_DEPENDS
         "grouped/*.cpp"
+        "grouped/*/*.cpp"
         "depthwise/*.cpp"
         "depthwise/*/*.cpp"
         "direct/*.cpp"
+        "direct_wgrad/*.cpp"
         "direct/direct_l1/*.cpp"
         "direct/direct/*.cpp"
         "direct/direct_wgrad/*.cpp"

projects/miopen/src/hipconv/cmake/version.cpp.in

@@ -0,0 +1,10 @@
+#include "hipconv/hipconv.hpp"
+
+namespace hipconv
+{
+
+std::string_view version() { return "@HIPCONV_VERSION@"; }
+
+std::string_view config_version() { return "@HIPCONV_CONFIG_VERSION@"; }
+
+} // namespace hipconv

projects/miopen/src/hipconv/cmake/write_version.cmake

@@ -0,0 +1,50 @@
+# Writes version.cpp for hipconv::version() and config_version(), run at build time by the
+# hipconv_version target.
+#
+# Inputs: SOURCE_DIR (the library root), TEMPLATE, OUTPUT.
+#
+# A published snapshot carries a VERSION file and is read first: it sits inside the consumer's
+# repository, where git describe would name the consumer's tags. A checkout asks git the way
+# publish/publish_to_miopen.sh does, so the two agree.
+
+set(HIPCONV_VERSION "unknown")
+if(EXISTS "${SOURCE_DIR}/VERSION")
+    file(STRINGS "${SOURCE_DIR}/VERSION" HIPCONV_VERSION LIMIT_COUNT 1)
+else()
+    find_package(Git QUIET)
+    if(GIT_FOUND)
+        execute_process(
+            COMMAND "${GIT_EXECUTABLE}" describe --tags --match "v[0-9]*"
+            WORKING_DIRECTORY "${SOURCE_DIR}"
+            OUTPUT_VARIABLE described
+            OUTPUT_STRIP_TRAILING_WHITESPACE
+            RESULT_VARIABLE described_status
+            ERROR_QUIET)
+        if(described_status EQUAL 0)
+            set(HIPCONV_VERSION "${described}")
+        else()
+            execute_process(
+                COMMAND "${GIT_EXECUTABLE}" rev-parse --short=7 HEAD
+                WORKING_DIRECTORY "${SOURCE_DIR}"
+                OUTPUT_VARIABLE sha
+                OUTPUT_STRIP_TRAILING_WHITESPACE
+                RESULT_VARIABLE sha_status
+                ERROR_QUIET)
+            if(sha_status EQUAL 0)
+                set(HIPCONV_VERSION "v0.0.0-0-g${sha}")
+            endif()
+        endif()
+    endif()
+endif()
+
+# v0.0.0 is the pseudo-version of a build no release tag precedes, which promises nothing
+# about its configs.
+set(HIPCONV_CONFIG_VERSION "unknown")
+if(NOT HIPCONV_VERSION MATCHES "^v0\\.0\\.0" AND HIPCONV_VERSION MATCHES "^(v[0-9]+\\.[0-9]+)\\.")
+    # The last MATCHES evaluated sets CMAKE_MATCH_1.
+    set(HIPCONV_CONFIG_VERSION "${CMAKE_MATCH_1}")
+endif()
+
+# configure_file rewrites OUTPUT only when its content changes, so an unchanged version
+# rebuilds nothing.
+configure_file("${TEMPLATE}" "${OUTPUT}" @ONLY)

projects/miopen/src/hipconv/include/hipconv/conv_params.hpp

@@ -214,7 +214,22 @@ struct ConvParams
         {
             return false;
         }
-        return true;
+        return !output_is_empty();
+    }
+
+    // True when, on some used axis, the dilated filter is wider than the padded input.
+    //
+    // Computed without division, so it is defined for any stride. compute_output_size
+    // cannot detect this case: its division truncates toward zero, so it reports an
+    // extent of 1 for an overhang shorter than the stride.
+    bool output_is_empty() const
+    {
+        auto overhangs = [](long long x, long long kx, long long pad, long long dilation) {
+            return dilation * (kx - 1) + 1 > x + 2 * pad;
+        };
+        return overhangs(w, kw, pad_w, dilation_w) ||
+               (dims >= 2 && overhangs(h, kh, pad_h, dilation_h)) ||
+               (dims >= 3 && overhangs(d, kd, pad_d, dilation_d));
     }
 
     // The conv2d layer this one is equivalent to, or a copy when there is none.

projects/miopen/src/hipconv/include/hipconv/hipconv.hpp

@@ -40,6 +40,18 @@ using ArchHandle = const ArchEntry*;
 // lifetime of the program.
 using ConvKernelHandle = ConvKernel*;
 
+// This build's hipconv version, as `git describe --tags` names it.
+//
+// "v0.4.0" at the tag, "v0.4.0-2-g1a2b3c4d" past it; "unknown" when neither a VERSION file
+// nor git provides one.
+HIPCONV_API std::string_view version();
+
+// The major and minor version of version(), e.g. "v0.4".
+//
+// "unknown" when no release tag precedes this build. For keying a stored config
+// description: within a minor version, a full description always selects the same config.
+HIPCONV_API std::string_view config_version();
+
 // Resolve a GFX arch name (e.g. "gfx950", "gfx950:sramecc+:xnack-") to a handle.
 // Returns nullopt if this build has no support for that architecture.
 HIPCONV_API std::optional<ArchHandle> resolve_arch(std::string_view name);
@@ -71,6 +83,10 @@ inline constexpr std::size_t ALL_RANKED_CONFIGS = std::numeric_limits<std::size_
 
 // All valid kernels for the given params, best first. Empty if unsupported.
 //
+// Selection reads `par` and the architecture alone, so it needs neither a GPU nor the HIP
+// runtime. A kernel that launches through a runtime library can still be declined on the
+// current device by ConvLaunch::make or launch(); a caller then tries the next kernel.
+//
 // `par` is matched as written; no folding happens here. A caller with a conv1d
 // or conv3d layer passes ConvParams::unfolded() instead, which is what gets a
 // degenerate conv3d served by a conv2d kernel; without it such a layer matches
@@ -112,7 +128,9 @@ HIPCONV_API Algorithm algorithm(ConvKernelHandle kernel);
 //
 // E.g. "waves_k=2,wave_k16=4,kh=3,kw=3,direction=fprop", exactly a spec that
 // matches_descriptor() accepts. A family with no descriptor fields returns "".
-HIPCONV_API std::string describe_config(ConvKernelHandle kernel);
+// `full` renders every field; otherwise default values are omitted. Only the full
+// form is guaranteed to select this config alone through matches_descriptor().
+HIPCONV_API std::string describe_config(ConvKernelHandle kernel, bool full = false);
 
 // Does this kernel satisfy a descriptor constraint string?
 //
@@ -137,7 +155,9 @@ HIPCONV_API float get_weighted_throughput_index(ConvKernelHandle kernel, const C
 // Enqueue the kernel on the stream, returning the launch-time HIP status.
 //
 // hipSuccess means the launch was submitted; otherwise it is the launch error
-// (hipErrorInvalidValue if kernel is null). The kernel runs asynchronously, so
+// (hipErrorInvalidValue if kernel is null, hipErrorNotSupported if the runtime library the
+// kernel launches through has no algorithm for `par` on the current device, as for
+// ConvLaunch::make). The kernel runs asynchronously, so
 // an execution fault (e.g. an out-of-bounds access) surfaces at a later
 // synchronization, not here.
 HIPCONV_API hipconvError_t launch(ConvKernelHandle kernel,
@@ -176,7 +196,9 @@ class HIPCONV_API ConvLaunch
     // Build a ConvLaunch for `kernel` bound to `par`.
     //
     // Returns nullopt if the kernel does not support `par` (same predicate as
-    // is_applicable). Throws std::invalid_argument if `kernel` is null.
+    // is_applicable), or if the runtime library it launches through, such as hipBLASLt
+    // for the explicit-GEMM kernels, has no algorithm for `par` on the current device.
+    // Throws std::invalid_argument if `kernel` is null.
     static std::optional<ConvLaunch> make(ConvKernelHandle kernel, ConvParams par);
 
     ConvLaunch(ConvLaunch&&) noexcept;

projects/miopen/src/hipconv/src/arch/cdna4/depthwise/depthwise_1d_toeplitz/config_desc.cpp

@@ -0,0 +1,26 @@
+#include "direction_field.h"
+#include "config_desc.h"
+
+namespace hipconv::cdna4::depthwise_1d_toeplitz
+{
+
+// A stride-2 dgrad entry is stride=1,dilation=2 here, the Config's encoding, not the layer's.
+hipconv::KVDescriptor config_fields(const Config& cfg)
+{
+    hipconv::KVDescriptor d;
+    d.int_field("kh", cfg.kh);
+    d.int_field("kw", cfg.kw);
+    d.int_field("stride", cfg.stride);
+    d.int_field("dilation", cfg.dilation, /*default=*/1);
+    direction_field(d, cfg.direction);
+    d.int_field("waves_per_wg", cfg.waves_per_wg);
+    d.int_field("group_size", cfg.group_size, /*default=*/8);
+    d.bool_field("narrow_c", cfg.narrow_c, /*default=*/false);
+    d.int_field("w_fold", cfg.w_fold, /*default=*/1);
+    d.int_field("lds_buffers", cfg.lds_buffers, /*default=*/3);
+    d.int_field("n_fold", cfg.n_fold, /*default=*/8);
+    d.int_field("elem_bytes", cfg.elem_bytes, /*default=*/2);
+    return d;
+}
+
+} // namespace hipconv::cdna4::depthwise_1d_toeplitz

projects/miopen/src/hipconv/src/arch/cdna4/depthwise/depthwise_1d_toeplitz/config_desc.h

@@ -0,0 +1,9 @@
+#pragma once
+
+#include "kv_descriptor.h"
+#include "config_table.h"
+
+namespace hipconv::cdna4::depthwise_1d_toeplitz
+{
+hipconv::KVDescriptor config_fields(const Config& cfg);
+} // namespace hipconv::cdna4::depthwise_1d_toeplitz