conv_wino_fury_RxS for gfx1151 crashes

[assistant-librarian]

[Migrated from original issue] ROCm/MIOpen#3735

Original issue author: @stellaraccident

This is related to the draft PR to disable the kernel: ROCm/MIOpen#3685

This convolution is an important component in a number of GenAI models like SDXL and various ComfyUI workflows.

Iiuc, this is rooted in a low level issue specific to this chip.

Can we get this tracked and fixed?

[assistant-librarian]

Comment by @jammm on 2025-05-16T17:17:42Z

Some more context:

Currently, diffusion models like SDXL, SD2/3/3.5, flux etc. rely on certain convolution operations during the VAE decode phase that MIOpen automatically selects to be executed via. conv_wino_fury_RxS. However, due to the error as described in ROCm/MIOpen#3685, we're unable to run this particular kernel. By disabling this kernel, MIOpen ends up selecting the naive conv kernels like naive_conv_ab_nonpacked_fwd_nchw_float_double_float which is much slower and ends up giving TDR's on windows with gfx1151.

A quick reproducer is to try the following command with MIOpenDriver.exe:
MIOpenDriver.exe conv -n 1 -c 256 -H 1024 -W 1024 -k 256 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1

This consistently gives me TDR's.

I believe that using the optimized winograd kernel can help solve this TDR as it'll ideally run faster, thus avoiding the long kernel execution times that causes TDR's.

[assistant-librarian]

Comment by @Cfouria on 2025-05-16T23:26:41Z

export MIOPEN_DEBUG_AMD_WINOGRAD_FURY_RXS_F2X3=0

my 780m (gfx1103) only needs f2x3 disabled and not f3x2, but both might help gfx1151
export MIOPEN_DEBUG_AMD_WINOGRAD_FURY_RXS_F3X2=0

[assistant-librarian]

Comment by @jammm on 2025-05-17T17:44:02Z

export MIOPEN_DEBUG_AMD_WINOGRAD_FURY_RXS_F2X3=0

my 780m (gfx1103) only needs f2x3 disabled and not f3x2, but both might help gfx1151 export MIOPEN_DEBUG_AMD_WINOGRAD_FURY_RXS_F3X2=0

This helps avoid those kernels. The workaround is to disable this for gfx1151 in the code itself, so users don't have to manually set these environment variables.

[assistant-librarian]

Comment by @jammm on 2025-05-17T19:29:00Z

Some more context:

Currently, diffusion models like SDXL, SD2/3/3.5, flux etc. rely on certain convolution operations during the VAE decode phase that MIOpen automatically selects to be executed via. conv_wino_fury_RxS. However, due to the error as described in #3685, we're unable to run this particular kernel. By disabling this kernel, MIOpen ends up selecting the naive conv kernels like naive_conv_ab_nonpacked_fwd_nchw_float_double_float which is much slower and ends up giving TDR's on windows with gfx1151.

A quick reproducer is to try the following command with MIOpenDriver.exe: MIOpenDriver.exe conv -n 1 -c 256 -H 1024 -W 1024 -k 256 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1

This consistently gives me TDR's.

I believe that using the optimized winograd kernel can help solve this TDR as it'll ideally run faster, thus avoiding the long kernel execution times that causes TDR's.

When running the above command with the following env vars set, you can also notice some PAL related issues:

$env:MIOPEN_DEBUG_CONV_DIRECT_NAIVE_CONV_FWD=0
$env:MIOPEN_DEBUG_GCN_ASM_KERNELS=1

MIOpen Forward Conv. Algorithm: 3, Solution: 84/ConvBinWinogradRxSf2x3g1
GPU Kernel Time Forward Conv. Elapsed: 58.212578 ms (average)
stats: name, n, c, ho, wo, x, y, k, flopCnt, bytesRead, bytesWritten, GFLOPs, GB/s, timeMs
stats: fwd-conv3x3u1, 1, 256, 1024, 1024, 3, 3, 256,  1236950581248, 1076101120, 1073741824, 21249, 37, 58.212578
:3:hip_device_runtime.cpp   :641 : 0303374421 us:   hipDeviceSynchronize (  )
:3:hip_device_runtime.cpp   :645 : 0303374446 us:  hipDeviceSynchronize: Returned hipSuccess :
:3:hip_memory.cpp           :792 : 0303374492 us:   hipMemcpy ( 000001AC776AC060, 0000000347900000, 1073741824, hipMemcpyDeviceToHost )
:3:hip_memory.cpp           :793 : 0303476371 us:  hipMemcpy: Returned hipSuccess : : duration: 101879 us
:3:hip_device_runtime.cpp   :688 : 0303476631 us:   hipSetDevice ( 0 )
:3:hip_device_runtime.cpp   :692 : 0303476672 us:  hipSetDevice: Returned hipSuccess :
:3:hip_device_runtime.cpp   :161 : 0303476717 us:   hipDeviceGetAttribute ( 00000014FB38D8F0, 63, 0 )
:3:hip_device_runtime.cpp   :452 : 0303476738 us:  hipDeviceGetAttribute: Returned hipSuccess :
:3:hip_module.cpp           :60  : 0303483684 us:   hipModuleLoadData ( 00000014FB38D7C0, 000001AC33DE3080 )
:3:devprogram.cpp           :2621: 0303486088 us:  Using Code Object V5.
:3:hip_module.cpp           :61  : 0303491645 us:  hipModuleLoadData: Returned hipSuccess :
:3:hip_module.cpp           :76  : 0303491722 us:   hipModuleGetFunction ( 00000014FB38DD90, 000001AC3423DD90, naive_conv_ab_nonpacked_fwd_nchw_float_double_float )
:3:hip_module.cpp           :90  : 0303491754 us:  hipModuleGetFunction: Returned hipSuccess :
:3:hip_device_runtime.cpp   :688 : 0303491801 us:   hipSetDevice ( 0 )
:3:hip_device_runtime.cpp   :692 : 0303491820 us:  hipSetDevice: Returned hipSuccess :
:3:hip_event.cpp            :333 : 0303491847 us:   hipEventCreate ( 00000014FB38DD90 )
:3:hip_event.cpp            :339 : 0303491866 us:  hipEventCreate: Returned hipSuccess : event:000001AC2856F660
:3:hip_event.cpp            :333 : 0303491884 us:   hipEventCreate ( 00000014FB38DD90 )
:3:hip_event.cpp            :339 : 0303491902 us:  hipEventCreate: Returned hipSuccess : event:000001AC2856EC40
:3:hip_module.cpp           :537 : 0303491932 us:   hipExtModuleLaunchKernel ( 0x000001AC26770DC0, 65536, 1, 1, 256, 1, 1, 0, stream:000001AC285B8010, char array:<null>, 00000014FB38DEE0, event:000001AC2856F660, event:000001AC2856EC40, 0 )
:3:hip_module.cpp           :549 : 0303491971 us:  hipExtModuleLaunchKernel: Returned hipSuccess :
:3:hip_event.cpp            :454 : 0303492005 us:   hipEventSynchronize ( event:000001AC2856EC40 )
:3:palvirtual.cpp           :2661: 0303492073 us:  !    kernel : naive_conv_ab_nonpacked_fwd_nchw_float_double_float

:3:device/pal/palvirtual.hpp:739 : 0303492102 us:   = ptr:[0000000307900000-0000000347900000] size:[1073741824] heap[1]
:3:device/pal/palvirtual.hpp:739 : 0303492125 us:   = ptr:[0000000300400000-0000000300640000] size:[67108864] heap[1]
:3:device/pal/palvirtual.hpp:739 : 0303492142 us:   = ptr:[0000000347900000-0000000387900000] size:[1073741824] heap[1]
:2:device/pal/palvirtual.hpp:164 : 0309493855 us:  PAL fence isn't ready!
:3:hip_event.cpp            :474 : 0310120824 us:  hipEventSynchronize: Returned hipSuccess :
:3:hip_event.cpp            :369 : 0310121007 us:   hipEventElapsedTime ( 000001AC34240C94, event:000001AC2856F660, event:000001AC2856EC40 )
:3:hip_event.cpp            :386 : 0310121104 us:  hipEventElapsedTime: Returned hipSuccess : Elapsed Time = , 6628.64
:3:hip_event.cpp            :344 : 0310121175 us:   hipEventDestroy ( event:000001AC2856EC40 )
:3:hip_event.cpp            :364 : 0310121247 us:  hipEventDestroy: Returned hipSuccess :
:3:hip_event.cpp            :344 : 0310121287 us:   hipEventDestroy ( event:000001AC2856F660 )
:3:hip_event.cpp            :364 : 0310121314 us:  hipEventDestroy: Returned hipSuccess :
:3:hip_device_runtime.cpp   :641 : 0310121365 us:   hipDeviceSynchronize (  )
:3:hip_device_runtime.cpp   :645 : 0310121406 us:  hipDeviceSynchronize: Returned hipSuccess :
:3:hip_memory.cpp           :792 : 0310121442 us:   hipMemcpy ( 000001ACB76B2060, 0000000347900000, 1073741824, hipMemcpyDeviceToHost )
:1:palvirtual.cpp           :430 : 0310121529 us:  PAL failed to submit CMD!
:1:palvirtual.cpp           :383 : 0310121565 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310122760 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310142970 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310143422 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310144163 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310144655 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310145160 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310145592 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310146071 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310146480 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310146654 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310147069 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310147381 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310147859 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310147935 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310148252 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310148333 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310150290 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310150680 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310151157 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310151294 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310151663 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310151881 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310152212 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310152290 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310152773 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310152908 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310153584 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310153714 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310154188 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310154456 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310154799 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310154893 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310155249 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310155334 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310155687 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310155794 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310156157 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310156249 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310156565 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310156633 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310156959 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310157050 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310157971 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310158349 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310158789 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310158913 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310159273 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310159378 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310159813 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310159948 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310160396 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310160572 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310160998 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310161143 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310161540 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310161644 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310161680 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310161710 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310161733 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310161763 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310161833 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310161945 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310162046 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310162343 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310162442 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310162528 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310162598 us:  PAL failed to finalize a command buffer!
:1:palvirtual.cpp           :383 : 0310162649 us:  PAL failed to finalize a command buffer!
:3:hip_memory.cpp           :793 : 0310162680 us:  hipMemcpy: Returned hipSuccess : : duration: 41238 us
Forward Convolution FAILED: 0.203646 > 1.5e-06
:3:hip_memory.cpp           :818 : 0311302395 us:   hipMemPtrGetInfo ( 0000000347900000, 00000014FB38F7E8 )
:3:hip_memory.cpp           :829 : 0311302547 us:  hipMemPtrGetInfo: Returned hipSuccess :
:3:hip_memory.cpp           :769 : 0311302616 us:   hipFree ( 0000000347900000 )
:1:palvirtual.cpp           :383 : 0311302701 us:  PAL failed to finalize a command buffer!
:3:hip_memory.cpp           :771 : 0311302773 us:  hipFree: Returned hipSuccess :
:3:hip_memory.cpp           :818 : 0311302917 us:   hipMemPtrGetInfo ( 0000000300400000, 00000014FB38F7E8 )
:3:hip_memory.cpp           :829 : 0311302973 us:  hipMemPtrGetInfo: Returned hipSuccess :
:3:hip_memory.cpp           :769 : 0311302999 us:   hipFree ( 0000000300400000 )
:3:hip_memory.cpp           :771 : 0311303043 us:  hipFree: Returned hipSuccess :
:3:hip_memory.cpp           :818 : 0311327134 us:   hipMemPtrGetInfo ( 0000000307900000, 00000014FB38F7E8 )
:3:hip_memory.cpp           :829 : 0311327287 us:  hipMemPtrGetInfo: Returned hipSuccess :
:3:hip_memory.cpp           :769 : 0311327333 us:   hipFree ( 0000000307900000 )
:3:hip_memory.cpp           :771 : 0311327384 us:  hipFree: Returned hipSuccess :
:3:hip_module.cpp           :46  : 0311327538 us:   hipModuleUnload ( 000001AC3423E880 )
:3:hip_module.cpp           :50  : 0311327678 us:  hipModuleUnload: Returned hipSuccess :
:3:hip_module.cpp           :46  : 0311327699 us:   hipModuleUnload ( 000001AC3423DD90 )
:3:hip_module.cpp           :50  : 0311330343 us:  hipModuleUnload: Returned hipSuccess :
:3:hip_memory.cpp           :769 : 0311330415 us:   hipFree ( 0000000306000000 )
:3:hip_memory.cpp           :771 : 0311330453 us:  hipFree: Returned hipSuccess :
:3:hip_memory.cpp           :769 : 0311330475 us:   hipFree ( 0000000304000000 )
:3:hip_memory.cpp           :771 : 0311330495 us:  hipFree: Returned hipSuccess :

Not sure if this is a byproduct of the TDR process itself, but thought i'd share just in case. The MIOpenDriver command still runs the naive conv kernels here.

[assistant-librarian]

Comment by @jammm on 2025-05-18T09:18:10Z

Setting just $env:MIOPEN_DEBUG_CONV_DIRECT_NAIVE_CONV_FWD=0 seems to have helped avoid TDR's in stable-diffusion webui/forge

[assistant-librarian]

Comment by @jammm on 2025-05-22T14:17:03Z

I can confirm this issue doesn't happen on Windows. Tested on stable diffusion webui forge:

:3:hip_device_runtime.cpp   :671 : 38582772997 us:  ^[[32m hipGetDevice ( 00000027433EA864 ) ^[[0m
:3:hip_device_runtime.cpp   :683 : 38582772999 us:  hipGetDevice: Returned hipSuccess : 0
:3:hip_device_runtime.cpp   :671 : 38582773001 us:  ^[[32m hipGetDevice ( 00000027433EA674 ) ^[[0m
:3:hip_module.cpp           :550 : 38582773001 us:  hipExtModuleLaunchKernel: Returned hipSuccess : 
:3:hip_device_runtime.cpp   :683 : 38582773003 us:  hipGetDevice: Returned hipSuccess : 0
:3:palvirtual.cpp           :2662: 38582773002 us:  !    kernel : 
miopenSp3AsmConvFury_v2_4_1_gfx11_1536vgprs_fp16_fp16acc_f2x3_c32_stride1

:3:hip_device_runtime.cpp   :671 : 38582773006 us:  ^[[32m hipGetDevice ( 00000027433EA31C ) ^[[0m
:3:hip_device_runtime.cpp   :683 : 38582773717 us:  hipGetDevice: Returned hipSuccess : 0
:3:device/pal/palvirtual.hpp:744 : 38582773008 us:  data_addr = ptr:[0000000C84970000-0000000C89570000] 

So it indeed runs the winograd furty kernel without error.

Note that I still had to set MIOPEN_DEBUG_CONV_DIRECT_NAIVE_CONV_FWD=0 or it'd cause TDR's - might make a separate issue for that.

[assistant-librarian]

Comment by @stellaraccident on 2025-05-22T14:26:41Z

I can confirm this issue doesn't happen on Windows. Tested on stable diffusion webui forge:

:3:hip_device_runtime.cpp   :671 : 38582772997 us:  �[32m hipGetDevice ( 00000027433EA864 ) �[0m
:3:hip_device_runtime.cpp   :683 : 38582772999 us:  hipGetDevice: Returned hipSuccess : 0
:3:hip_device_runtime.cpp   :671 : 38582773001 us:  �[32m hipGetDevice ( 00000027433EA674 ) �[0m
:3:hip_module.cpp           :550 : 38582773001 us:  hipExtModuleLaunchKernel: Returned hipSuccess : 
:3:hip_device_runtime.cpp   :683 : 38582773003 us:  hipGetDevice: Returned hipSuccess : 0
:3:palvirtual.cpp           :2662: 38582773002 us:  !    kernel : 
miopenSp3AsmConvFury_v2_4_1_gfx11_1536vgprs_fp16_fp16acc_f2x3_c32_stride1

:3:hip_device_runtime.cpp   :671 : 38582773006 us:  �[32m hipGetDevice ( 00000027433EA31C ) �[0m
:3:hip_device_runtime.cpp   :683 : 38582773717 us:  hipGetDevice: Returned hipSuccess : 0
:3:device/pal/palvirtual.hpp:744 : 38582773008 us:  data_addr = ptr:[0000000C84970000-0000000C89570000] 

So it indeed runs the winograd furty kernel without error.

Note that I still had to set MIOPEN_DEBUG_CONV_DIRECT_NAIVE_CONV_FWD=0 or it'd cause TDR's - might make a separate issue for that.

Thanks. We isolated this issue down to a driver issue that was specific to Linux (a misconfigured GPU setting specific to how it works on Linux). There will be a fix coming for that but still getting details on how and when.

We were scratching our heads to find the issue also on Windows as we couldn't see how that would be the case. Thanks for confirming.