[Issue]: ROCm 7.12.0 on gfx1150 gives HIP error launch failure partway through MIOpen optimization of SDXL VAE decode

[lostdisc]

Problem Description

While testing MIOpen-enabled SDXL VAE decode in ComfyUI with ROCm 7.12.0 on Windows and gfx1150 (collecting logs for Issue 2591), I encountered a "HIP error: unspecified launch failure" about 10-11 minutes into MIOpen's benchmarking runs. The detailed (level 7) log file is attached below. This did not occur with previous ROCm versions between 6.4 and 7.2.1.

ROCm7.12.0_sdxl_1600x1280_VAEdecodeonly_before-optimiz_withelapsedtime_miopen_output_logs_level7_ERROR.txt
ROCm7.12.0_sdxl_1600x1280_VAEdecodeonly_before-optimiz_hipblaslt_log.txt

I've pasted the error tracebacks below:

MIOpen(HIP)   0.273: Info [FindConvolution] miopenConvolutionFwdAlgoGEMM        458.512 4718592000
MIOpen(HIP)   0.168: Info [FillFindReturnParameters] FW Chosen Algorithm: GemmFwdRest , 4718592000, 458.512
C:\Users\<my username>\.conda\envs\rocmpreview\Lib\site-packages\torch\nn\modules\conv.py:548: UserWarning: HIP warning: unspecified launch failure (Triggered internally at B:/src/torch/aten/src\ATen/hip/impl/HIPGuardImplMasqueradingAsCUDA.h:83.)
  return F.conv2d(
!!! Exception during processing !!! HIP error: unspecified launch failure
Search for `hipErrorLaunchFailure' in https://rocm.docs.amd.com/projects/HIP/en/latest/index.html for more information.
HIP kernel errors might be asynchronously reported at some other API call, so the stacktrace below might be incorrect.
For debugging consider passing AMD_SERIALIZE_KERNEL=3
Compile with `TORCH_USE_HIP_DSA` to enable device-side assertions.

Traceback (most recent call last):
  File "C:\Users\<my username>\ComfyUI\execution.py", line 525, in execute
    output_data, output_ui, has_subgraph, has_pending_tasks = await get_output_data(prompt_id, unique_id, obj, input_data_all, execution_block_cb=execution_block_cb, pre_execute_cb=pre_execute_cb, v3_data=v3_data)
                                                              ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "C:\Users\<my username>\ComfyUI\execution.py", line 334, in get_output_data
    return_values = await _async_map_node_over_list(prompt_id, unique_id, obj, input_data_all, obj.FUNCTION, allow_interrupt=True, execution_block_cb=execution_block_cb, pre_execute_cb=pre_execute_cb, v3_data=v3_data)
                    ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "C:\Users\<my username>\ComfyUI\execution.py", line 308, in _async_map_node_over_list
    await process_inputs(input_dict, i)
  File "C:\Users\<my username>\ComfyUI\execution.py", line 296, in process_inputs
    result = f(**inputs)
  File "C:\Users\<my username>\ComfyUI\nodes.py", line 316, in decode
    images = vae.decode(latent)
  File "C:\Users\<my username>\ComfyUI\comfy\sd.py", line 972, in decode
    model_management.raise_non_oom(e)
    ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^^^
  File "C:\Users\<my username>\ComfyUI\comfy\model_management.py", line 289, in raise_non_oom
    raise e
  File "C:\Users\<my username>\ComfyUI\comfy\sd.py", line 965, in decode
    out = self.first_stage_model.decode(samples, **vae_options).to(device=self.output_device, dtype=self.vae_output_dtype(), copy=True)
          ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^^^^^^^^^^^^^^^^^^^^^^^^
  File "C:\Users\<my username>\ComfyUI\comfy\ldm\models\autoencoder.py", line 253, in decode
    dec = self.decoder(dec, **decoder_kwargs)
  File "C:\Users\<my username>\.conda\envs\rocmpreview\Lib\site-packages\torch\nn\modules\module.py", line 1776, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
           ~~~~~~~~~~~~~~~^^^^^^^^^^^^^^^^^
  File "C:\Users\<my username>\.conda\envs\rocmpreview\Lib\site-packages\torch\nn\modules\module.py", line 1787, in _call_impl
    return forward_call(*args, **kwargs)
  File "C:\Users\<my username>\ComfyUI\comfy\ldm\modules\diffusionmodules\model.py", line 813, in forward
    h1 = self.up[i_level].upsample(h1, conv_carry_in, conv_carry_out)
  File "C:\Users\<my username>\.conda\envs\rocmpreview\Lib\site-packages\torch\nn\modules\module.py", line 1776, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
           ~~~~~~~~~~~~~~~^^^^^^^^^^^^^^^^^
  File "C:\Users\<my username>\.conda\envs\rocmpreview\Lib\site-packages\torch\nn\modules\module.py", line 1787, in _call_impl
    return forward_call(*args, **kwargs)
  File "C:\Users\<my username>\ComfyUI\comfy\ldm\modules\diffusionmodules\model.py", line 137, in forward
    x = conv_carry_causal_3d([x], self.conv, conv_carry_in, conv_carry_out)
  File "C:\Users\<my username>\ComfyUI\comfy\ldm\modules\diffusionmodules\model.py", line 81, in conv_carry_causal_3d
    out = op(x)
  File "C:\Users\<my username>\.conda\envs\rocmpreview\Lib\site-packages\torch\nn\modules\module.py", line 1776, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
           ~~~~~~~~~~~~~~~^^^^^^^^^^^^^^^^^
  File "C:\Users\<my username>\.conda\envs\rocmpreview\Lib\site-packages\torch\nn\modules\module.py", line 1787, in _call_impl
    return forward_call(*args, **kwargs)
  File "C:\Users\<my username>\ComfyUI\comfy\ops.py", line 428, in forward
    return super().forward(*args, **kwargs)
           ~~~~~~~~~~~~~~~^^^^^^^^^^^^^^^^^
  File "C:\Users\<my username>\.conda\envs\rocmpreview\Lib\site-packages\torch\nn\modules\conv.py", line 553, in forward
    return self._conv_forward(input, self.weight, self.bias)
           ~~~~~~~~~~~~~~~~~~^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "C:\Users\<my username>\.conda\envs\rocmpreview\Lib\site-packages\torch\nn\modules\conv.py", line 548, in _conv_forward
    return F.conv2d(
           ~~~~~~~~^
        input, weight, bias, self.stride, self.padding, self.dilation, self.groups
        ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    )
    ^
torch.AcceleratorError: HIP error: unspecified launch failure
Search for `hipErrorLaunchFailure' in https://rocm.docs.amd.com/projects/HIP/en/latest/index.html for more information.
HIP kernel errors might be asynchronously reported at some other API call, so the stacktrace below might be incorrect.
For debugging consider passing AMD_SERIALIZE_KERNEL=3
Compile with `TORCH_USE_HIP_DSA` to enable device-side assertions.


Prompt executed in 00:10:46
Exception in thread Thread-9 (prompt_worker):
Traceback (most recent call last):
  File "C:\Users\<my username>\.conda\envs\rocmpreview\Lib\threading.py", line 1044, in _bootstrap_inner
    self.run()
    ~~~~~~~~^^
  File "C:\Users\<my username>\.conda\envs\rocmpreview\Lib\threading.py", line 995, in run
    self._target(*self._args, **self._kwargs)
    ~~~~~~~~~~~~^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "C:\Users\<my username>\ComfyUI\main.py", line 354, in prompt_worker
    comfy.model_management.soft_empty_cache()
    ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^^
  File "C:\Users\<my username>\ComfyUI\comfy\model_management.py", line 1779, in soft_empty_cache
    torch.cuda.synchronize()
    ~~~~~~~~~~~~~~~~~~~~~~^^
  File "C:\Users\<my username>\.conda\envs\rocmpreview\Lib\site-packages\torch\cuda\__init__.py", line 1107, in synchronize
    with torch.cuda.device(device):
         ~~~~~~~~~~~~~~~~~^^^^^^^^
  File "C:\Users\<my username>\.conda\envs\rocmpreview\Lib\site-packages\torch\cuda\__init__.py", line 546, in __init__
    self.idx = _get_device_index(device, optional=True)
               ~~~~~~~~~~~~~~~~~^^^^^^^^^^^^^^^^^^^^^^^
  File "C:\Users\<my username>\.conda\envs\rocmpreview\Lib\site-packages\torch\cuda\_utils.py", line 529, in _get_device_index
    return _torch_get_device_index(device, optional, allow_cpu)
  File "C:\Users\<my username>\.conda\envs\rocmpreview\Lib\site-packages\torch\_utils.py", line 876, in _get_device_index
    device_idx = _get_current_device_index()
  File "C:\Users\<my username>\.conda\envs\rocmpreview\Lib\site-packages\torch\_utils.py", line 813, in _get_current_device_index
    return _get_device_attr(lambda m: m.current_device())
  File "C:\Users\<my username>\.conda\envs\rocmpreview\Lib\site-packages\torch\_utils.py", line 798, in _get_device_attr
    return get_member(torch.cuda)
  File "C:\Users\<my username>\.conda\envs\rocmpreview\Lib\site-packages\torch\_utils.py", line 813, in <lambda>
    return _get_device_attr(lambda m: m.current_device())
                                      ~~~~~~~~~~~~~~~~^^
  File "C:\Users\<my username>\.conda\envs\rocmpreview\Lib\site-packages\torch\cuda\__init__.py", line 1095, in current_device
    return torch._C._cuda_getDevice()
           ~~~~~~~~~~~~~~~~~~~~~~~~^^
torch.AcceleratorError: HIP error: unspecified launch failure
Search for `hipErrorLaunchFailure' in https://rocm.docs.amd.com/projects/HIP/en/latest/index.html for more information.
HIP kernel errors might be asynchronously reported at some other API call, so the stacktrace below might be incorrect.
For debugging consider passing AMD_SERIALIZE_KERNEL=3
Compile with `TORCH_USE_HIP_DSA` to enable device-side assertions.

Operating System

Windows 11

CPU

AMD Ryzen AI 9 HX 370

GPU

AMD Radeon 890M

ROCm Version

7.12.0

ROCm Component

MIOpen

Steps to Reproduce

1. On Windows, install latest AMD Adrenalin (26.3.1). (No HIP SDK installed. Previous ROCm versions ran fine without it.)
2. Create Python 3.13 conda environment (not sure if it matters for the bug's purposes)
3. Install ROCm 7.12.0 following these instructions (except in conda env instead of venv)
4. Install PyTorch following these instructions (except in conda env instead of venv)
5. Install ComfyUI or use existing manual-install folder. Installation instructions here, except I also run pip install -r ComfyUI\manager_requirements.txt in addition to regular requirements.txt.
6. Add the following Windows system environment variables (mostly for logging). May need to close and reopen your terminal for them to take effect:

TORCH_ROCM_AOTRITON_ENABLE_EXPERIMENTAL 1
COMFYUI_ENABLE_MIOPEN 1
MIOPEN_ENABLE_LOGGING 1
MIOPEN_ENABLE_LOGGING_CMD 1
MIOPEN_ENABLE_LOGGING_ELAPSED_TIME 1
MIOPEN_LOG_LEVEL 7
HIPBLASLT_LOG_MASK 32
HIPBLASLT_LOG_FILE C:\hipblaslt_log.txt (use your preferred file path here)
TORCH_BLAS_PREFER_HIPBLASLT 1

7. Delete ".miopen" folder in your Windows account folder to clear existing MIOpen database and make it rerun benchmarks. Very important for this test.
8. Launch ComfyUI (I use the --enable-manager flag, but it shouldn't matter for the bug's purposes)
9. Drag-drop the workflow below into ComfyUI for decoding an empty latent:
   SDXL VAE Decode Empty Latent test.json
   Or recreate it like so:

Choose an SDXL model for the checkpoint.

10. Run the workflow. On my system it takes about 10-11 minutes to hit the error for a 1600x1280 image.

(Optional for Linux users) Output of /opt/rocm/bin/rocminfo --support

No response

Additional Information

No response

[astrelsky]

Does it also occur with a smaller image dimension? It might make it easier to narrow down the culprit since it'll be faster. I'll try as well tomorrow.

It would be great if error messages could at least tell you what kernel failed instead of leaving everyone guessing.

[lostdisc]

128x128 finished without problems in about 30 seconds.
256x256 finished in ~60s.
512x512 finished in ~230s.

So not every size... I'll see if I can find a small one that still errors out.

[astrelsky]

Oh wait sorry I'm dumb. The MIOPEN_ENABLE_LOGGING_CMD flag you used should let you find it pretty easily if you find the error in the log and backtrack upwards to the last command.

[lostdisc]

The error did not reappear when I retested by deleting the .miopen folder and decoding 1600x1280 again. Benchmarking finished in 37 minutes:
ROCm7.12.0_sdxl_1600x1280_VAEdecodeonly_before-optimiz_withelapsedtime_miopen_output_logs_level7.txt
ROCm7.12.0_sdxl_1600x1280_VAEdecodeonly_before-optimiz_hipblaslt_log.txt

So maybe the first time was a fluke. (I tested out the stock SDXL VAE too, and it also finished benching in 37 minutes.) I might close this bug unless/until it reoccurs.

That said, I encountered a different bug while testing post-optimization reruns.

The first rerun finished in ~20s without problems:
ROCm7.12.0_sdxl_1600x1280_VAEdecodeonly_after-optimiz_withelapsedtime_miopen_output_logs_level7_goodrun-only.txt
ROCm7.12.0_sdxl_1600x1280_VAEdecodeonly_after-optimiz_hipblaslt_log_goodrun-only.txt

But doing another rerun (after clearing ComfyUI's model/exec cache) maxed out my 32GB of RAM and froze Windows for several minutes before finishing:
ROCm7.12.0_sdxl_1600x1280_VAEdecodeonly_after-optimiz_withelapsedtime_miopen_output_logs_level7_RAMmaxedrun-only.txt
ROCm7.12.0_sdxl_1600x1280_VAEdecodeonly_after-optimiz_hipblaslt_log_RAMmaxedrun-only.txt

This occurred repeatedly when testing with Wai-illustrious-SDXL v16. If I was lucky, Windows would unfreeze after a while and finish the job (5-40 mins). If I was unlucky, Windows would crash and reboot itself. Again, these reruns are not doing MIOpen benchmarking. With past ROCm versions, they would take ~2 minutes. (Which is terrible, but I digress.)

The bug did not seem to occur with the stock SDXL VAE or the madebyollin fp16-compatible version. These finished reruns in ~6 seconds.

The "good" rerun log has a lot more database messages than the "bad" rerun, but somehow needed far less time and RAM to do the decode. Hope someone with more knowledge can figure out what it means. I will open a new issue for it.