JonathanLichtnerAMD · Oct 6, 2026 · Oct 2, 2026 · Oct 2, 2026 · Oct 2, 2026 · Oct 2, 2026
@@ -8,12 +8,20 @@ Full documentation for MIOpen is available [here](https://rocm.docs.amd.com/proj
 ### Added
 * [Conv] Added gfx950 depthwise backward-weights (fp16/bf16) and gfx1250 depthwise
   kernels to the `ConvHipConv` solver (hipconv v0.3.1).
+* [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).
 
 ### Removed
 * [Conv] Removed gfx803 convolution solver `ConvBinWinogradRxSFused` and its kernel sources.
 * [Conv] Removed the gfx803 code paths and kernel sources from `ConvBinWinograd3x3U` and `ConvBinWinogradRxS`.
 * [Conv] Removed the gfx803 find-database files.
 
+### Resolved Issues
+* [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).
+
 ## MIOpen 3.6.1 for ROCm 10.1.0
 
 ### Resolved Issues

@@ -9,8 +9,8 @@ if(MIOPEN_ENABLE_SQLITE)
 endif()
 
 if(MIOPEN_USE_HIPCONV)
-    # MIOpen provides its own 1x1/pointwise path, so drop hipconv's hipBLASLt-backed one.
-    set(HIPCONV_ENABLE_POINTWISE OFF CACHE BOOL "" FORCE)
+    # MIOpen provides its own 1x1/pointwise path, so drop hipconv's hipBLASLt-backed ones.
+    set(HIPCONV_ENABLE_EXPLICIT_GEMM OFF CACHE BOOL "" FORCE)
     add_subdirectory(hipconv)
 endif()
 

@@ -1,5 +1,65 @@
 # Changelog for hipconv
 
+## v0.3.2 - 2026-10-02
+
+### Added
+
+| Kernel                   | Architecture     | Layer type             | Direction           | Data type           | Filter, stride                        |
+|--------------------------|------------------|------------------------|---------------------|---------------------|---------------------------------------|
+| `patch_embed`            | CDNA 4 (gfx950)  | patch-embedding conv2d | fprop, dgrad, wgrad | fp16, bf16; dW fp32 | stride = filter, no padding           |
+| `depthwise_wgrad_hankel` | CDNA 5 (gfx1250) | depthwise conv2d       | wgrad               | fp16, bf16; dW fp32 | 3×3 … 11×11 odd square, stride 1 or 2 |
+
+* tf32 on three CDNA 4 kernels, stored as fp32 and computed as three bf16 MFMAs on a
+  (big, small) bf16 split: `depthwise_1d_toeplitz` fprop and dgrad (#223), `direct_l1`
+  fprop and dgrad at 2×2 … 5×5 (#226), and `direct_wgrad` at 2×2, 3×1, 3×3, 4×4 and 5×5
+  (#234).
+* In builds with `HIPCONV_ENABLE_EXPLICIT_GEMM=ON` (default OFF; MIOpen leaves it off),
+  tf32 on pointwise layers, stored as fp32 and computed by hipBLASLt's
+  `HIPBLAS_COMPUTE_32F_FAST_TF32`. That compute type asks for "at least tf32 precision"
+  and leaves the method to the backend, so the accuracy it delivers is
+  architecture-dependent: an architecture without xf32 MFMA decomposes into bf16 and
+  lands inside the error bound hipconv derives for tf32, while one with xf32 computes a
+  single lower-precision pass that the bound does not cover (#229).
+* `find_arch()`, which resolves a GFX name to an `ArchHandle` without the device-code
+  check `resolve_arch()` makes, and `arch_names()`, which lists the architectures the
+  registry holds. Neither needs a device or a HIP context, so a host with no gfx1250 can
+  enumerate which layers the gfx1250 kernels claim. Call `resolve_arch()` before a launch
+  (#245).
+
+### Changed
+
+* Renamed the `pointwise` algorithm to `explicit_gemm`, which now serves non-overlapping
+  patch-embedding convolutions as well as 1×1 through hipBLASLt. Its CMake option
+  `HIPCONV_ENABLE_POINTWISE` is now `HIPCONV_ENABLE_EXPLICIT_GEMM`, still default OFF and
+  left off by MIOpen. CMake ignores the old name without a warning, so a consumer that sets
+  it must rename it (#237).
+* CDNA 4 `depthwise_wgrad_hankel` stores each partition's dW with plain stores into a
+  channel-major workspace, and a fold kernel sums the partitions, replacing per-tap
+  atomics: geometric mean +88% over fifteen shapes (#227). With `weight_grad_type` equal
+  to the input type, it writes dW as fp16 or bf16 in that final store, so a caller needs
+  no separate cast kernel (#235).
+* CDNA 5 `direct` launches a one-dimensional grid and unfolds the block index into batch,
+  group, spatial and K tiles. It rejects depthwise layers (`k == c`, one channel per
+  group) and configs whose thread count exceeds 2^32 (#274).
+
+### Fixed
+
+* CDNA 5 `direct` put groups × batch tiles in the grid's z dimension, so a layer with more
+  than 65535 of them crashed. Reported as ALMIOPEN-2624 (#265, #274).
+* CDNA 5 `direct`: a TDM load could overwrite an LDS buffer another wave was still
+  reading, because the compiler placed that wave's `s_wait_dscnt` in the following stage.
+  The kernel now drains `dscnt` after each tile load, and the odd-K dgrad restriction
+  from #224 is lifted (#248).
+* CDNA 5 `direct` sized its output staging LDS for one image of a batch-folded tile
+  (`tile_size_n > 1`).
+* CDNA 4 `direct_wgrad` ordered its staging reuse only between partitions, so on
+  single-partition configs clang 24 let a lane's next-position stores reach LDS words
+  other lanes were still draining, and dW came back holding the next position's gradient.
+  Already in MIOpen develop through rocm-libraries #12406; this release is its first
+  hipconv version (#231).
+* Fixed builds with compilers that reject implicit conversions between vector types
+  (`-flax-vector-conversions=none`): builtin operands are now converted explicitly (#264).
+
 ## v0.3.1 - 2026-09-17
 
 ### Added

@@ -91,6 +91,17 @@ endif()
 option(HIPCONV_ENABLE_CDNA4_DW_2D_TOEPLITZ
     "Build the experimental cdna4 depthwise 2D patch-Toeplitz kernel (slower than 1D; not auto-selected)" OFF)
 
+# The explicit-GEMM families (1x1 and patch embedding) are backed by hipBLASLt.
+# Off by default so a build takes on that dependency only when it asks for it,
+# and so a framework providing its own GEMM-based convolutions can drop the code.
+#
+# An option() for this was lost in the arch-registry rewrite (350e6ef), leaving the
+# `if(HIPCONV_ENABLE_EXPLICIT_GEMM)` guards below reading an undeclared variable.
+# Declaring it here restores `cmake -LH` discoverability and the misspelling check;
+# the default matches what the build has actually been doing since then.
+option(HIPCONV_ENABLE_EXPLICIT_GEMM
+    "Build the hipBLASLt-backed explicit-GEMM convolution paths (1x1 and patch embedding)" 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")
@@ -145,12 +156,13 @@ set(_hipconv_srcs
     src/tolerance.cpp
     "${_arch_registry}"
 )
-if(HIPCONV_ENABLE_POINTWISE)
+if(HIPCONV_ENABLE_EXPLICIT_GEMM)
     find_package(hipblaslt REQUIRED)
     list(APPEND _hipconv_srcs
-        src/pointwise/pointwise_conv.cpp
-        src/pointwise/pointwise_fp16bf16.cpp
-        src/pointwise/hipblaslt_matmul.cpp
+        src/explicit_gemm/explicit_gemm_algorithm.cpp
+        src/explicit_gemm/pointwise/pointwise_kernel.cpp
+        src/explicit_gemm/hipblaslt_matmul.cpp
+        src/explicit_gemm/patch_embed/patch_embed.cpp
     )
 endif()
 
@@ -169,7 +181,7 @@ 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
                                        ${HIPCONV_DEVICE_COMPILE_OPTIONS})
-if(HIPCONV_ENABLE_POINTWISE)
+if(HIPCONV_ENABLE_EXPLICIT_GEMM)
     target_link_libraries(hipconv PUBLIC roc::hipblaslt)
-    target_compile_definitions(hipconv PUBLIC HIPCONV_ENABLE_POINTWISE)
+    target_compile_definitions(hipconv PUBLIC HIPCONV_ENABLE_EXPLICIT_GEMM)
 endif()
@@ -1,11 +1,7 @@
-EMU version: v0.3.1-0-g93b98ccf
-EMU commit: 93b98ccf7819f92558118d546717bebad7fcca1d
-Snapshot date: 2026-09-17
+EMU version: v0.3.2-0-ga6c85d2a
+EMU commit: a6c85d2a8592b918d9c79f94b4d3bf6473021fb3
+Snapshot date: 2026-10-02
 Arches: cdna3,cdna4,cdna5
-Out-of-snapshot patch: src/arch/cdna4/direct/direct_wgrad/epilogue.h carries hipconv
-  e1777ac9 (AMD-ROCm-Core-Perf-Team/hipconv#231), a cross-lane ordering fix taken ahead of
-  the next sync. Any snapshot from hipconv main at or after that commit supersedes it;
-  do not publish an earlier revision over this tree.
 Files:
   ./CHANGELOG.md
   ./.clang-format
@@ -76,6 +72,11 @@ Files:
   ./src/arch/cdna4/direct/direct_wgrad/row_loader.h
   ./src/arch/cdna4/direct/direct_wgrad/row_schedule.h
   ./src/arch/cdna4/direct/direct_wgrad/workgroup_tiles.h
+  ./src/arch/cdna4/direct/patch_embed/config_desc.cpp
+  ./src/arch/cdna4/direct/patch_embed/config_desc.h
+  ./src/arch/cdna4/direct/patch_embed/config.h
+  ./src/arch/cdna4/direct/patch_embed/config_table.h
+  ./src/arch/cdna4/direct/patch_embed/kernel.h
   ./src/arch/cdna4/grouped_backend.cpp
   ./src/arch/cdna4/grouped/grouped_16c.cpp
   ./src/arch/cdna4/grouped/grouped_16c_wgrad_fp16bf16.cpp
@@ -99,6 +100,16 @@ Files:
   ./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/depthwise/depthwise_wgrad_hankel/config_matcher.cpp
+  ./src/arch/cdna5/depthwise/depthwise_wgrad_hankel/config_matcher.hpp
+  ./src/arch/cdna5/depthwise/depthwise_wgrad_hankel/config_table.h
+  ./src/arch/cdna5/depthwise/depthwise_wgrad_hankel/delta_issuer.h
+  ./src/arch/cdna5/depthwise/depthwise_wgrad_hankel/input_issuer.h
+  ./src/arch/cdna5/depthwise/depthwise_wgrad_hankel/kernel.h
+  ./src/arch/cdna5/depthwise/depthwise_wgrad_hankel/lds.h
+  ./src/arch/cdna5/depthwise/depthwise_wgrad_hankel/reduction.cpp
+  ./src/arch/cdna5/depthwise/depthwise_wgrad_hankel/reduction.hpp
+  ./src/arch/cdna5/depthwise/depthwise_wgrad_hankel/wave.h
   ./src/arch/cdna5/direct_backend.cpp
   ./src/arch/cdna5/direct/config.hpp
   ./src/arch/cdna5/direct/config_matcher.cpp
@@ -127,6 +138,12 @@ Files:
   ./src/detail.h
   ./src/direct_algorithm_stub.h
   ./src/direct_conv_kernel.h
+  ./src/explicit_gemm/explicit_gemm_algorithm.cpp
+  ./src/explicit_gemm/hipblaslt_matmul.cpp
+  ./src/explicit_gemm/hipblaslt_matmul.hpp
+  ./src/explicit_gemm/patch_embed/patch_embed.cpp
+  ./src/explicit_gemm/patch_embed/patch_embed.hpp
+  ./src/explicit_gemm/pointwise/pointwise_kernel.cpp
   ./src/grouped_algorithm_stub.h
   ./src/grouped_conv_kernel.h
   ./src/hip_util.h
@@ -140,13 +157,11 @@ Files:
   ./src/matrix_layout_print.h
   ./src/memory.h
   ./src/packed_ops.h
+  ./src/patch_embed_conv_kernel.h
   ./src/pointwise_conv_kernel.h
-  ./src/pointwise/hipblaslt_matmul.cpp
-  ./src/pointwise/hipblaslt_matmul.hpp
-  ./src/pointwise/pointwise_conv.cpp
-  ./src/pointwise/pointwise_fp16bf16.cpp
   ./src/registry.cpp
   ./src/swizzle.h
+  ./src/tensor_view.hpp
   ./src/tolerance.cpp
   ./src/tolerance.h
   ./src/transpose_lds_layout.h

@@ -67,6 +67,7 @@ function(hipconv_add_arch_lib name)
         "direct/direct_l1/*.cpp"
         "direct/direct/*.cpp"
         "direct/direct_wgrad/*.cpp"
+        "direct/patch_embed/*.cpp"
     )
     list(APPEND sources "${variant_sources}")
     add_library(hipconv_arch_${name} OBJECT ${sources})

@@ -39,7 +39,9 @@ enum class Algorithm
     // Depthwise (1 channel per group, groups == C == K).
     Depthwise,
     Direct,
-    Pointwise
+    // Layers run as an explicit GEMM through hipBLASLt, as opposed to the
+    // implicit GEMM the direct family convolves with.
+    ExplicitGemm
 };
 
 // Every Direction value, in declaration order.
@@ -49,7 +51,7 @@ inline constexpr std::array all_directions{Direction::Fprop, Direction::Dgrad, D
 inline constexpr std::array all_algorithms{Algorithm::Grouped,
                                            Algorithm::Depthwise,
                                            Algorithm::Direct,
-                                           Algorithm::Pointwise};
+                                           Algorithm::ExplicitGemm};
 
 // DataType values the driver accepts as input tags, in declaration order.
 inline constexpr std::array input_data_types{DataType::fp16, DataType::bf16, DataType::tf32};
@@ -65,7 +67,8 @@ auto to_string(Algorithm algo) -> char const*;
 
 // Parse an algorithm name (the inverse of to_string(Algorithm)).
 //
-// Returns nullopt for any string that is not one of grouped|depthwise|direct|pointwise.
+// Returns nullopt for any string that is not one of
+// grouped|depthwise|direct|explicit_gemm.
 std::optional<Algorithm> parse_algorithm(std::string_view name);
 
 // Parse a supported input data-type name (fp16|bf16|tf32).

@@ -44,6 +44,20 @@ using ConvKernelHandle = ConvKernel*;
 // Returns nullopt if this build has no support for that architecture.
 HIPCONV_API std::optional<ArchHandle> resolve_arch(std::string_view name);
 
+// Look up a GFX arch name in the registry without asking a device.
+//
+// Reports what the registry serves rather than what this binary can run: it
+// makes no HIP call and needs no context, so a caller can enumerate coverage for
+// a named architecture off-device. A caller that intends to launch wants
+// resolve_arch, which adds the device-code check.
+HIPCONV_API std::optional<ArchHandle> find_arch(std::string_view name);
+
+// Every architecture in this build's registry, in registry order.
+//
+// A name here is one find_arch accepts. Whether this binary carries device code
+// for it is a separate question; see resolve_arch.
+HIPCONV_API std::vector<std::string_view> arch_names();
+
 // Error code returned by hipconv's launch API. Currently a hipError_t and
 // carries the standard hipError_t values; may be extended with hipconv-specific
 // codes in the future.

@@ -4,16 +4,16 @@
 #include "arch_registry.h"
 
 // clang-format off
-#ifdef HIPCONV_ENABLE_POINTWISE
-extern const hipconv::ConvAlgorithm pointwise_algo;
+#ifdef HIPCONV_ENABLE_EXPLICIT_GEMM
+extern const hipconv::ConvAlgorithm explicit_gemm_algo;
 #endif
 extern const hipconv::ConvAlgorithm grouped_@HIPCONV_ARCH_NAME@;
 extern const hipconv::ConvAlgorithm depthwise_@HIPCONV_ARCH_NAME@;
 extern const hipconv::ConvAlgorithm direct_@HIPCONV_ARCH_NAME@;
 
 const hipconv::ArchAlgorithm algos_@HIPCONV_ARCH_NAME@[] = {
-#ifdef HIPCONV_ENABLE_POINTWISE
-    { hipconv::Algorithm::Pointwise, &pointwise_algo },
+#ifdef HIPCONV_ENABLE_EXPLICIT_GEMM
+    { hipconv::Algorithm::ExplicitGemm, &explicit_gemm_algo },
 #endif
     { hipconv::Algorithm::Grouped, &grouped_@HIPCONV_ARCH_NAME@ },
     { hipconv::Algorithm::Depthwise, &depthwise_@HIPCONV_ARCH_NAME@ },

@@ -13,6 +13,12 @@ hipconv_autoshard(ARCH         cdna4
                   CONFIG_TABLE direct/direct_l1/config_table.h
                   KERNEL       direct/direct_l1/kernel.h
                   NUM_SHARDS   ${HIPCONV_NUM_SHARDS})
+hipconv_autoshard(ARCH         cdna4
+                  NAMESPACE    patch_embed
+                  KERNEL_CLASS PatchEmbed_ConvKernel
+                  CONFIG_TABLE direct/patch_embed/config_table.h
+                  KERNEL       direct/patch_embed/kernel.h
+                  NUM_SHARDS   ${HIPCONV_NUM_SHARDS})
 hipconv_autoshard(ARCH         cdna4
                   NAMESPACE    direct_wgrad
                   KERNEL_CLASS DirectWgrad_ConvKernel