Description
MIOpen Gemm convolution solvers return workspace_size=0 on gfx1150 (Radeon 890M, RDNA 3.5) when called through PyTorch's ATen convolution path, causing solvers to be skipped silently.
Root Cause
PyTorch's ATen layer uses the legacy MIOpen Find1 API path (miopenConvolutionForwardGetWorkSpaceSize / miopenFindConvolutionForwardAlgorithm). This API queries workspace requirements, but for Gemm-based solvers on new GPU targets like gfx1150, the workspace query returns 0 bytes.
When ATen sees workspace_size=0, it allocates no workspace buffer. The Gemm solver then fails internally (needs workspace for im2col/GEMM scratch) and gets excluded from the algorithm list. The result is that only ConvDirectNaiveConvFwd remains available — which works but is extremely slow.
Expected Behavior
Gemm solvers should either:
- Report their actual workspace requirements through the Find1 API, or
- PyTorch should migrate to the Find2 / Immediate mode API which handles workspace allocation differently
Environment
- Hardware: AMD Ryzen AI 9 HX 370, Radeon 890M (gfx1150)
- OS: CachyOS, kernel 6.19.5
- ROCm: 7.2.0
- PyTorch: 2.12.0a0 (built from source for gfx1150)
- MIOpen: Built from amd-develop
Workaround
Setting MIOPEN_DEBUG_CONV_GEMM=0 disables the broken Gemm path, allowing other solvers (legacy CK, Winograd where applicable) to be selected.
Related
cc @jerryzh168 @jeffdaily @sunway513 @jithunnair-amd @pruthvistony @ROCmSupport @jataylo @hongxiayang @naromero77amd @pragupta @jerrymannil @xinyazhang
Description
MIOpen Gemm convolution solvers return
workspace_size=0on gfx1150 (Radeon 890M, RDNA 3.5) when called through PyTorch's ATen convolution path, causing solvers to be skipped silently.Root Cause
PyTorch's ATen layer uses the legacy MIOpen
Find1API path (miopenConvolutionForwardGetWorkSpaceSize/miopenFindConvolutionForwardAlgorithm). This API queries workspace requirements, but for Gemm-based solvers on new GPU targets like gfx1150, the workspace query returns 0 bytes.When ATen sees
workspace_size=0, it allocates no workspace buffer. The Gemm solver then fails internally (needs workspace for im2col/GEMM scratch) and gets excluded from the algorithm list. The result is that onlyConvDirectNaiveConvFwdremains available — which works but is extremely slow.Expected Behavior
Gemm solvers should either:
Environment
Workaround
Setting
MIOPEN_DEBUG_CONV_GEMM=0disables the broken Gemm path, allowing other solvers (legacy CK, Winograd where applicable) to be selected.Related
cc @jerryzh168 @jeffdaily @sunway513 @jithunnair-amd @pruthvistony @ROCmSupport @jataylo @hongxiayang @naromero77amd @pragupta @jerrymannil @xinyazhang