Draft (untested) install-n.bat

[sfinktah]

Let's just let this set here until we've tried in, and perhaps changed it a bit.

[patientx]

why are we updating to 2.8 again ? I don't usually update unless it is absolutely necessary. Also at the moment 6.2.4 is the latest official hip available for windows. I am not going to force people to install experimental / custom stuff. The rock build is an entirely optional thing and for that I am not even showing it in the main readme as you can see.

[sfinktah]

why are we updating to 2.8 again?

Because 2.7 is fundamentally incompatible with Triton >3.0. Also, it has much better memory management (it almost doesn't leak).

Also at the moment 6.2.4 is the latest official hip available for windows

Well, technically the hip DLL in c:\windows\system32 is HIP version 6.4 (courtesy of the latest adrenaline installer). Since I don't know where to get these special sources for 6.4, it was either 6.2 or 6.5.

The 6.5 requirement actually came from the original issue opened by alex (was that his name?) in your instructions.

[sfinktah]

Actually, I think the logic was that because the TheRock version being used was 6.5, we needed to use 6.5 HIP SDK. But that's not necessarily true, as only Triton actually uses the SDK, and it's not linked to a specific version of HIP as far as I know. Further testing will be required.

Postnote to my earlier comment: It loads the hipsdk 6.4 dll from windows\system32, and doesn't ever actually load the dll in the AMD\HIP directory.

Another postnote: Sorry, I guess that doesn't explain a great deal to a non-programmer. And in truth, I can't really find a way to explain exactly what it does mean, except that it's not like it's loading untested drivers into the kernel or anything nasty like that.

[patientx]

I am not going to update to latest torch unless it is absolutely necessary.
6.5 is from rock. Again I am not going to force people to install / try custom modified versions.
I was thinking about copying the libs folder and was about to add that.
Making 3.4 as default is good, and when it is used then , if your triton patch would work with 3.4 with any torches I could add that also.

So I can do only copying libs folder & patching triton (if that would work with the current setup) .

[sfinktah]

So the most important thing to know about HIP SDK 6.4 (and above) is that it adds support for gfx1201.

My triton patches for triton 3.4 will absolute work with torch 2.7, I just need to add the patches to fix torch 2.7. I'll tweak the install a bit and upload a new PR. But not right now.

[patientx]

So the most important thing to know about HIP SDK 6.4 (and above) is that it adds support for gfx1201.

My triton patches for triton 3.4 will absolute work with torch 2.7, I just need to add the patches to fix torch 2.7. I'll tweak the install a bit and upload a new PR. But not right now.

so you are saying just using that dll that comes with latest drivers in the windows folder , it would be hip 6.2.4 , with only one dll 1201 would be supported ? then it is only usable for the newer 9000 series , isn't it ? and can be added to readme / instructions if people using those gpu's want to try.

[sfinktah]

I'm saying that regardless of whether you install HIP SDK 6.5, the gfx1201 will work, because it will always be using HIP 6.4 DLL from c:\windows\systems.

But stuff that is compiled by Triton won't take advantage of any gfx1201 speed-ups or use fp8 latent space, because Triton won't know about them, unless there it has the 6.5 SDK installed.

Though how much of a difference that would actually make with pytorch 2.7 installed is another matter. Compilation is something that always improves with updates in pytorch.

[patientx]

added lib copying and triton install being only for 3.11 and 3.12 also added your patch, trying it now. at least installation works without problems. sage-attention works as usual. torch compile shows lots of errors. you could probably solve them though (it was kinda working with your previous "patch". added triton errors.)

triton.txt

[sfinktah]

Yeah, you didn't add the patch. And while the patch fixes those errors you pasted, that just opens the door for some wierd CUDA related errors that I have no idea have to solve, and they stop the job executing.

[sfinktah]

... though now I'm having the same problems with 2.8.0. So I'm just confused. Maybe you should test yourself with 2.7.0 by running this patch inside the venv.

pypatch-url apply https://raw.githubusercontent.com/sfinktah/amd-torch/refs/heads/main/patches/torch-2.7.0+cu118-cp311-cp311-win_amd64.patch -p 4 torch

[patientx]

Added your patches, and then added gfx code setting in zluda.py, deleted triton folder (to see it run again) and it works.

Here is the startup :

  ::  ------------------------ ZLUDA -----------------------  ::
  ::  Auto-detecting AMD GPU architecture for Triton...
  ::  Detected GPU via Windows registry: AMD Radeon RX 6800
  ::  Set TRITON_OVERRIDE_ARCH=gfx1030
  ::  Triton core imported successfully
  ::  Running Triton kernel test...
%%% [info] triton/runtime/build/platform_key: AMD64,Windows,64bit,WindowsPE
Using HIP SDK Clang.
%%% Triton is using the compiler 'C:\Program Files\AMD\ROCm\6.2\bin\clang.exe'
%%% [info] triton/runtime/build/_build: "C:\Program Files\AMD\ROCm\6.2\bin\clang.exe" D:\idm-temp\temp\tmpqe05mlq3\hip_utils.c -O3 -shared -Wno-psabi -march=native -mtune=native -o D:\idm-temp\temp\tmpqe05mlq3\hip_utils.cp311-win_amd64.pyd -LD:\sd\ComfyUI-Zluda\venv\libs "-LC:\Program Files\AMD\ROCm\6.2\lib" -ID:\sd\ComfyUI-Zluda\venv\Lib\site-packages\triton\backends\amd\include -ID:\idm-temp\temp\tmpqe05mlq3 -IC:\Users\aslan\AppData\Local\Programs\Python\Python311\Include
In file included from D:\idm-temp\temp\tmpqe05mlq3\hip_utils.c:8:
In file included from C:\Users\aslan\AppData\Local\Programs\Python\Python311\Include\Python.h:87:
C:\Users\aslan\AppData\Local\Programs\Python\Python311\Include\cpython/pytime.h:184:12: warning: declaration of
      'struct timeval' will not be visible outside of this function [-Wvisibility]
  184 |     struct timeval *tv,
      |            ^
C:\Users\aslan\AppData\Local\Programs\Python\Python311\Include\cpython/pytime.h:190:12: warning: declaration of
      'struct timeval' will not be visible outside of this function [-Wvisibility]
  190 |     struct timeval *tv,
      |            ^
2 warnings generated.
%%% [debug] make_llir: adding llvm option for amdgpu-waves-per-eu=1  -- triton/backends/amd/compiler.py
Using HIP SDK Clang.
%%% Triton is using the compiler 'C:\Program Files\AMD\ROCm\6.2\bin\clang.exe'
%%% [info] triton/runtime/build/_build: "C:\Program Files\AMD\ROCm\6.2\bin\clang.exe" D:\idm-temp\temp\tmpg3bcjclq\__triton_launcher.c -O3 -shared -Wno-psabi -march=native -mtune=native -o D:\idm-temp\temp\tmpg3bcjclq\__triton_launcher.cp311-win_amd64.pyd -LD:\sd\ComfyUI-Zluda\venv\libs "-LC:\Program Files\AMD\ROCm\6.2\lib" -ID:\sd\ComfyUI-Zluda\venv\Lib\site-packages\triton\backends\amd\include -ID:\idm-temp\temp\tmpg3bcjclq -IC:\Users\aslan\AppData\Local\Programs\Python\Python311\Include
In file included from D:\idm-temp\temp\tmpg3bcjclq\__triton_launcher.c:4:
In file included from C:\Users\aslan\AppData\Local\Programs\Python\Python311\Include\Python.h:87:
C:\Users\aslan\AppData\Local\Programs\Python\Python311\Include\cpython/pytime.h:184:12: warning: declaration of 'struct timeval' will not be visible outside of this function [-Wvisibility]
  184 |     struct timeval *tv,
      |            ^
C:\Users\aslan\AppData\Local\Programs\Python\Python311\Include\cpython/pytime.h:190:12: warning: declaration of 'struct timeval' will not be visible outside of this function [-Wvisibility]
  190 |     struct timeval *tv,
      |            ^
2 warnings generated.
  ::  Triton kernel test passed successfully
  ::  Triton initialized successfully
  ::  Patching ONNX Runtime for ZLUDA — disabling CUDA EP.
  ::  Using ZLUDA with device: AMD Radeon RX 6800 [ZLUDA]
  ::  Applying core ZLUDA patches...
  ::  Initializing Triton optimizations
  ::  Configuring Triton device properties...
  ::  Triton device properties configured
  ::  Flash attention components found
  ::  AMD flash attention enabled successfully
  ::  Configuring PyTorch backends...
  ::  Disabled CUDA flash attention
  ::  Enabled math attention fallback
  ::  ZLUDA initialization complete
  ::  ------------------------ ZLUDA -----------------------  ::

Total VRAM 16368 MB, total RAM 32690 MB
pytorch version: 2.7.0+cu118
Set vram state to: NORMAL_VRAM
Device: cuda:0 AMD Radeon RX 6800 [ZLUDA] : native

And here is a wan 2.1 image generation.

Sampling 1 frames at 1024x1536 with 4 steps
  0%|                                                                                                                                                                                                                 | 0/4 [00:00<?, ?it/s]%%% [debug] generate_ttir: removed invalid kernel arg: waves_per_eu=4  -- torch/_higher_order_ops/triton_kernel_wrap.py
%%% [debug] generate_ttir: removed invalid kernel arg: waves_per_eu=4  -- torch/_higher_order_ops/triton_kernel_wrap.py
%%% [debug] generate_ttir: removed invalid kernel arg: waves_per_eu=4  -- torch/_higher_order_ops/triton_kernel_wrap.py
%%% [debug] generate_ttir: removed invalid kernel arg: waves_per_eu=4  -- torch/_higher_order_ops/triton_kernel_wrap.py
%%% [fix] ignoring ASTSource::__init__ invalid constexpr: waves_per_eu=2 -- triton/compiler/compiler.py
%%% [fix] ignoring triton_heuristics::_get_args_with_constexprs: missing arg_name: waves_per_eu=2 -- torch/_inductor/runtime/triton_heuristics.py
Compilation is in progress. Please wait...

Other than the debug and info no errors. Just like the rock version I am using atm.
Tried wan 2.2 light_noise as well to check.

%% [warn] Your old Radeon is not supported for max_autotune_gemm mode.
%%% [debug] make_llir: adding llvm option for amdgpu-waves-per-eu=1  -- triton/backends/amd/compiler.py
%%% [debug] make_llir: adding llvm option for amdgpu-waves-per-eu=1  -- triton/backends/amd/compiler.py
%%% [debug] make_llir: adding llvm option for amdgpu-waves-per-eu=1  -- triton/backends/amd/compiler.py
%%% [debug] make_llir: adding llvm option for amdgpu-waves-per-eu=1  -- triton/backends/amd/compiler.py
%%% [debug] make_llir: adding llvm option for amdgpu-waves-per-eu=1  -- triton/backends/amd/compiler.py
%%% [debug] make_llir: adding llvm option for amdgpu-waves-per-eu=1  -- triton/backends/amd/compiler.py
%%% [debug] make_llir: adding llvm option for amdgpu-waves-per-eu=1  -- triton/backends/amd/compiler.py
%%% [debug] make_llir: adding llvm option for amdgpu-waves-per-eu=1  -- triton/backends/amd/compiler.py
%%% [debug] make_llir: adding llvm option for amdgpu-waves-per-eu=1  -- triton/backends/amd/compiler.py
%%% [debug] generate_ttir: removed invalid kernel arg: waves_per_eu=4  -- torch/_higher_order_ops/triton_kernel_wrap.py
%%% [debug] generate_ttir: removed invalid kernel arg: waves_per_eu=4  -- torch/_higher_order_ops/triton_kernel_wrap.py
%%% [debug] generate_ttir: removed invalid kernel arg: waves_per_eu=4  -- torch/_higher_order_ops/triton_kernel_wrap.py
%%% [debug] generate_ttir: removed invalid kernel arg: waves_per_eu=4  -- torch/_higher_order_ops/triton_kernel_wrap.py
%%% [debug] generate_ttir: removed invalid kernel arg: waves_per_eu=4  -- torch/_higher_order_ops/triton_kernel_wrap.py
%%% [debug] generate_ttir: removed invalid kernel arg: waves_per_eu=4  -- torch/_higher_order_ops/triton_kernel_wrap.py
%%% [debug] generate_ttir: removed invalid kernel arg: waves_per_eu=4  -- torch/_higher_order_ops/triton_kernel_wrap.py
%%% [debug] generate_ttir: removed invalid kernel arg: waves_per_eu=4  -- torch/_higher_order_ops/triton_kernel_wrap.py
%%% [debug] make_llir: adding llvm option for amdgpu-waves-per-eu=1  -- triton/backends/amd/compiler.py
%%% [debug] make_llir: adding llvm option for amdgpu-waves-per-eu=1  -- triton/backends/amd/compiler.py
%%% [debug] make_llir: adding llvm option for amdgpu-waves-per-eu=1  -- triton/backends/amd/compiler.py
%%% [debug] make_llir: adding llvm option for amdgpu-waves-per-eu=1  -- triton/backends/amd/compiler.py
%%% [debug] make_llir: adding llvm option for amdgpu-waves-per-eu=1  -- triton/backends/amd/compiler.py
%%% [debug] make_llir: adding llvm option for amdgpu-waves-per-eu=1  -- triton/backends/amd/compiler.py
%%% [debug] make_llir: adding llvm option for amdgpu-waves-per-eu=1  -- triton/backends/amd/compiler.py
%%% [debug] make_llir: adding llvm option for amdgpu-waves-per-eu=1  -- triton/backends/amd/compiler.py
%%% [debug] make_llir: adding llvm option for amdgpu-waves-per-eu=1  -- triton/backends/amd/compiler.py
%%% [debug] make_llir: adding llvm option for amdgpu-waves-per-eu=1  -- triton/backends/amd/compiler.py
%%% [debug] make_llir: adding llvm option for amdgpu-waves-per-eu=1  -- triton/backends/amd/compiler.py
%%% [debug] make_llir: adding llvm option for amdgpu-waves-per-eu=1  -- triton/backends/amd/compiler.py
%%% [debug] make_llir: adding llvm option for amdgpu-waves-per-eu=1  -- triton/backends/amd/compiler.py

So, doesn't this mean it is working ?

EDIT : Tested both zluda and rock versions with the same wan 2.1 wanwrapper 8 step image generation (1024x1536 , 1 frame generation) workflow.

(sage-attention is enabled always)

zluda version (install-n)
without torch compile / I just removed the compile node / ::: around 20.4 sec/it
with torch compile ::: around 18.6 sec/it

rock version
with torch compile 18.4 sec/it

So I think torch compile is definately working at least it IS making a difference and it is as good as rock version. /FOR MY RX 6800 at least/

[sfinktah]

Yeah, that all looks like it's working. I should note that I start with --use-sage-attention usually, or sometimes --use-flash-attention. Those can cause their own set of errors. Also, I am testing using a Chroma workflow that uses the native comfyui inductor compile node (seems to be faster and more reliable than Kijai's more complex one), and I seem to be having some kind of other issue that is causing the errors to continue even when I update to pytorch 2.8 again.

That's why I wanted you to check it.

The error that is getting me -- which you might have seen before -- is:

RuntimeError: CUDA error: CUBLAS_STATUS_NOT_SUPPORTED when calling `cublasGemmEx( handle, opa, opb, m, n, k, &falpha, a, CUDA_R_16BF, lda, b, CUDA_R_16BF, ldb, &fbeta, c, std::is_same_v<C_Dtype, float> ? CUDA_R_32F : CUDA_R_16BF, ldc, compute_type, CUBLAS_GEMM_DEFAULT_TENSOR_OP)`
Exception raised from gemm_internal_cublas_bfloat16_helper at C:\actions-runner\_work\pytorch\pytorch\pytorch\aten\src\ATen\cuda\CUDABlas.cpp:1234 (most recent call first):

I'll do some testing with my wan workflows, that is a good idea.

[sfinktah]

ack I didn't have the right comfy/customzluda/zluda.py file, so it wasn't doing those Triton checks at startup. No idea how that happened, but checking everything again now to see if pytorch 2.7 will work now.

[sfinktah]

okay, so that is sorted out. i also have stats. my chroma test render takes

• 105 seconds with pytorch2.70 + rocm 6.2 (full zluda triton patch from latest commit)
• 85 seconds with pytorch2.80 + rocm 6.2 (some non-triton zluda patch)
• 78 seconds with pytorch2.80 + rocm 6.5 (some non-triton zluda patch)

If memory serves, 14 seconds of that difference is in the prompt encoding. I will have to do proper breakdown graphs to be sure, and this "two versions of zluda.py" thing will have be checked -- one probably goes faster than the other.