Triton fork for Windows support

Based on andreigh, wkpark, mantaionut, eaplatanios, anmyachev, and more development in the community. Thank you all!

What's supported

• triton.jit and torch.compile just work
• All unit tests passed
• It's as fast as on Linux on the same GPU
• Windows 10 and 11 are supported
• Recent Nvidia and AMD GPUs are supported
  • For Intel XPU, see https://github.com/intel/intel-xpu-backend-for-triton . They have Windows support
• Proton and GSan are not actively maintained in this repo. If you want to try them, you can build from source
• We welcome quality-of-life improvements for local AI users, and PRs to add support for older GPUs. triton-windows does not regress Linux support, so these improvements also work on Linux

Installation

Triton accelerates your AI model by compiling things on your computer. You need to install it in the correct environment.

1. GPU

Check your GPU model. Also make sure you have the latest GPU driver.

• Blackwell only works with Triton >= 3.3, PyTorch >= 2.7, and CUDA >= 12.8
• Support for Volta and Turing was dropped since Triton 3.3, see https://github.com/triton-lang/triton/pull/5066
• If you want to help maintaining Turing in the latest Triton, see https://github.com/triton-lang/triton-windows/issues/28
• If you want to help with Pascal, see https://github.com/woct0rdho/triton-windows/issues/133 for the previous discussions, and open a new issue in this repo. It may be handled in a way similar to RDNA2
• If you want to help with RDNA2, see https://github.com/triton-lang/triton-windows/issues/41

Besides the overall Triton support, also check hardware-supported data types:

• 'Y' means it's possible to achieve the theoretical speed in matmul: fp16/bf16 is 2x fp32 or faster, fp8/int8 is 2x fp16, and fp4/int4 is 4x fp16, not considering sparsity
  • Without hardware support, it's also possible to work by silently upcasting to a supported data type in matmul, but it will not achieve the theoretical speed
  • The official Triton does not support the silent upcast of fp8 on Ampere and older GPUs. It's supported since triton-windows post21
  • Even with hardware support, the speed still depends on how well the code is optimized, rather than simply achieves the theoretical speed
• Consumer Pascal does not have tensor cores, but it has DP4A for fast int8. Similarly, RDNA2 does not have matrix cores, but it has V_DOT4_I32_I8/V_DOT8_I32_I4 for fast int8/int4
• Volta does not have fast int8/int4. It has DP4A, but that's slower than fp16 tensor core
• RDNA3/3.5's int8 matrix core is only 1x fp16, and int4 matrix core is only 2x fp16
• SageAttention requires int8

2. Python environment

Check how your Python is installed. Either of the following environments is supported:

• Embeded: You use an all-in-one AI software package such as ComfyUI
  • There should be a folder python_embeded in the ComfyUI installation folder
  • Other AI software may put this folder at a different path
  • In this case, don't directly run python, but use the full path C:\path\to\python_embeded\python.exe
  • Also, don't directly run pip, but instead run C:\path\to\python_embeded\python.exe -m pip
  • In the folder python_embeded, by default pip is not installed as a direct executable pip.exe, but a Python package. If you directly run pip, you're actually running a pip.exe installed somewhere else on your computer
  • It's ok to first cd to python_embeded, then run .\python.exe, but remember to add .\ to run an executable in the current folder. In PowerShell, without .\, you're still running a python.exe installed somewhere else on your computer
• System-wide: You install Python at a location like C:\Python312\ or C:\Program Files\Python312\ and directly use it
• User-wide: You install Python at a location like C:\Users\<your username>\AppData\Local\Programs\Python\Python312\ and directly use it
• Python venv: You create a virtual environment using venv or virtualenv
• uv: You create a virtual environment using uv venv
• conda: You create a virtual environment using conda

I don't recommend installing Python from Windows Store, because it's complicated to interact with a 'packaged' Windows app.

For other environment managers like poetry, if you find problems, please open an issue.

Make sure what environment you're using. You can run Get-Command -All python in PowerShell (or where python in cmd) to see the installation path of Python, and python --version to see its version. If you see multiple Python installations, make sure that you install and run everything from the first one.

• For example, if you think you're using Python 3.12, but pip downloads a wheel with cp311 in its name, then it means you're not using the Python environment you think

Don't mix two environments, unless you know them very well.

• If you're using ComfyUI with embeded Python, then don't use conda or venv
• If you're already using conda, then always create a new env using conda, and don't use Python venv

3. PyTorch

Although technically Triton can be used alone, in the following let's assume you use it with PyTorch. Each PyTorch minor version is only guaranteed to work with a specific Triton minor version:

PyTorch 2.3 and older are not supported in this repo.

If you have to use Triton 3.2 because you're using an old GPU, then you can try to use Triton 3.2 with PyTorch >= 2.7, but it's not guaranteed to always work.

For AMD GPU, you may install ROCm and PyTorch from TheRock to get the latest features, see https://github.com/ROCm/TheRock/blob/main/RELEASES.md

4. CUDA

You can skip this.

5. C compiler

You can skip this.

&"C:\Program Files (x86)\Microsoft Visual Studio\2022\BuildTools\Common7\Tools\Launch-VsDevShell.ps1" -Arch amd64

"C:\Program Files (x86)\Microsoft Visual Studio\2022\BuildTools\Common7\Tools\VsDevCmd.bat" -arch=amd64

6. vcredist

vcredist is required (also known as 'Visual C++ Redistributable for Visual Studio 2015-2022', msvcp140.dll, vcruntime140.dll), because libtriton.pyd is compiled by MSVC. Install it from https://aka.ms/vs/17/release/vc_redist.x64.exe

7. Triton

Since triton-windows post11, the wheels are published to https://pypi.org/project/triton-windows/ , so you don't need to manually download a wheel from GitHub releases, and pip will automatically find it.

If you've installed an old version of triton, first uninstall it:

pip uninstall triton

Now you can install triton-windows 3.8, or upgrade the already installed version. To prevent breaking with your installed PyTorch when a new version of Triton is released in future, you can limit the version to be < 3.9:

pip install -U "triton-windows<3.9"

Note again that if you're using the embeded Python, then instead of directly run pip, you need:

C:\path\to\python_embeded\python.exe -m pip install -U "triton-windows<3.9"

Or if you want triton-windows 3.2, then run:

pip install -U "triton-windows<3.3"

8. Special notes for embeded Python

• There should be a Python folder
  • For ComfyUI, it's python_embeded in the ComfyUI installation folder
  • Other AI software may put the Python folder at a different path
  • If you created a venv, depending on how you created it, the Python folder may be just the venv folder or the venv\Scripts folder
  • If you're not sure, you can run os.path.dirname(sysconfig.get_paths()["include"]) to find the Python folder, see py_include_dir
• You need to put two folders include and libs into the Python folder to make Triton work
  • Be careful: It is 'libs', not 'lib'. There may already be a folder Lib in the Python folder, containing things like site-packages or __future__.py. You should not modify the Lib folder
  • If you're using ComfyUI_windows_portable >= 0.3.50 with Python 3.13, then download the two folders here: python_3.13.2_include_libs.zip
  • The minor version (3.9/3.10 ...) must be correct, but the patch version (3.10.6/3.10.7 ...) can be different
  • If you're using another Python version, you can find the two folders at https://github.com/woct0rdho/triton-windows/releases/v3.0.0-windows.post1/
• (For developers: This is equivalent to python-dev on Linux, and you can obtain the two folders from nuget when bundling Python in your app, see https://github.com/comfyanonymous/ComfyUI/pull/7200 )

Test if it works

Before using Triton in your project like ComfyUI, please run the following script to test if Triton itself works.

• You need to save the code in a file, such as test_triton.py, then run python test_triton.py
• When you open an issue, please show the command you use to run this test, and the full error log

import torch
import triton
import triton.language as tl

@triton.jit
def add_kernel(x_ptr, y_ptr, output_ptr, n_elements, BLOCK_SIZE: tl.constexpr):
    pid = tl.program_id(axis=0)
    block_start = pid * BLOCK_SIZE
    offsets = block_start + tl.arange(0, BLOCK_SIZE)
    mask = offsets < n_elements
    x = tl.load(x_ptr + offsets, mask=mask)
    y = tl.load(y_ptr + offsets, mask=mask)
    output = x + y
    tl.store(output_ptr + offsets, output, mask=mask)

def add(x: torch.Tensor, y: torch.Tensor):
    output = torch.empty_like(x)
    n_elements = output.numel()
    grid = lambda meta: (triton.cdiv(n_elements, meta["BLOCK_SIZE"]),)
    add_kernel[grid](x, y, output, n_elements, BLOCK_SIZE=1024)
    return output

a = torch.rand(3, device="cuda")
b = a + a
b_compiled = add(a, a)
print(b_compiled - b)
print("If you see tensor([0., 0., 0.], device='cuda:0'), then it works")

Troubleshoot the test above

ModuleNotFoundError: No module named 'triton.language'; 'triton' is not a package

Don't name the test script triton.py. Also, check if there is a folder named triton in your current directory. If so, Python will think it's the 'triton' package and fail to import.

AttributeError: module 'pkgutil' has no attribute 'ImpImporter'. Did you mean: 'zipimporter'

This is because your setuptools is outdated. Run the following and try again:

python -m ensurepip -U
python -m pip install -U pip
python -m pip install -U setuptools

PermissionError: [WinError 5] Access is denied: 'C:\\Users\\<your username>\\.triton'

This is because of the permission settings of your user folder, see https://github.com/lllyasviel/FramePack/issues/221

ImportError: DLL load failed while importing libtriton

This is usually because your vcredist DLLs are too old.

If you're using conda, then you may try:

conda install -c conda-forge vc14_runtime

If you're not using conda, then you need to find the vcredist DLLs (vcruntime140.dll, vcruntime140_1.dll) in your Python installation folder:

import sysconfig
print(sysconfig.get_paths())

After finding the DLLs in the Python installation folder, you can install the latest vcredist, then copy the DLLs msvcp140.dll, vcruntime140.dll, vcruntime140_1.dll from C:\Windows\System32\ to the Python installation folder, and replace the existing ones.

You can right-click the DLL -> Properties -> Details to see its version. A new enough version, such as 14.42, is required by my Triton wheels.

ImportError: DLL load failed while importing cuda_utils

1. Delete the cache folders:

   C:\Users\<your username>\.triton\cache\
   C:\Users\<your username>\AppData\Local\Temp\torchinductor_<your username>\
   

   • You may also need to delete these cache folders when you change the Python version, install another version of Triton, or change the C compiler or CUDA
   • It's ok if these folders do not exist on your computer. The first folder exists only if you have used triton.jit (which is used by packages like SageAttention), and the second folder exists only if you have used torch.compile
2. Double check your Python version: You can run Get-Command -All python in PowerShell (or where python in cmd) to see the installation path of Python, and python --version to see its version. If you see multiple Python installations, make sure that you install and run everything from the first one
3. If you're using ComfyUI with embeded Python, make sure that you copy-pasted the folders include and libs from the correct version of Python

SystemError: PY_SSIZE_T_CLEAN macro must be defined for '#' formats

You also need to delete the cache folders above.

This should not happen if you upgrade to Python 3.13, see https://github.com/python/cpython/issues/104922 . If the error still exists, you may try to debug following https://github.com/woct0rdho/triton-windows/issues/163

dlltracer

If the above still doesn't work, you may try:

• Install dlltracer in the same Python environment
• In an administrator PowerShell, run the following script:

import sys
import dlltracer

print("import torch")
with dlltracer.Trace(out=sys.stdout):
    import torch

print("import triton")
with dlltracer.Trace(out=sys.stdout):
    import triton

print("begin definition")
with dlltracer.Trace(out=sys.stdout):
    import triton.language as tl

    @triton.jit
    def add_kernel(x_ptr, y_ptr, output_ptr, n_elements, BLOCK_SIZE: tl.constexpr):
        pid = tl.program_id(axis=0)
        block_start = pid * BLOCK_SIZE
        offsets = block_start + tl.arange(0, BLOCK_SIZE)
        mask = offsets < n_elements
        x = tl.load(x_ptr + offsets, mask=mask)
        y = tl.load(y_ptr + offsets, mask=mask)
        output = x + y
        tl.store(output_ptr + offsets, output, mask=mask)

    def add(x: torch.Tensor, y: torch.Tensor):
        output = torch.empty_like(x)
        n_elements = output.numel()
        grid = lambda meta: (triton.cdiv(n_elements, meta["BLOCK_SIZE"]),)
        add_kernel[grid](x, y, output, n_elements, BLOCK_SIZE=1024)
        return output

print("begin torch add")
with dlltracer.Trace(out=sys.stdout):
    a = torch.rand(3, device="cuda")
    b = a + a

print("begin jit add")
with dlltracer.Trace(out=sys.stdout):
    b_compiled = add(a, a)

print(b_compiled - b)
print("If you see tensor([0., 0., 0.], device='cuda:0'), then it works")

• Open an issue. Please show the command you use to run this test, and the full error log

If it shows PermissionError: [WinError 5] failed to start trace (0x00000005), then you need to make sure to run it as administrator.

(Security reminder: You don't need the administrator privilege to run Triton and other usual Python code. It's only dlltracer that needs it.)

If it shows Failed \Device\...\cuda_utils.pyd, please also:

• Find cuda_utils.pyd at this location
• Use DependenciesGui (or similar tools) to check what DLLs this cuda_utils.pyd depends on, and send a screenshot (or other related information) in the issue

Known issues

Windows file path length limit (260) causes compilation failure

torch.compile may create temp files with very long filenames, causing errors like:

  File "C:\...\Lib\site-packages\torch\_inductor\runtime\triton_heuristics.py", line 537, in _precompile_config
    binary = triton.compile(*compile_args, **compile_kwargs)
             ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "C:\...\Lib\site-packages\triton\compiler\compiler.py", line 288, in compile
    metadata_group[ir_filename] = fn_cache_manager.put(next_module, ir_filename)
                                  ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "C:\...\Lib\site-packages\triton\runtime\cache.py", line 122, in put
    with open(temp_path, mode) as f:
         ^^^^^^^^^^^^^^^^^^^^^
torch._inductor.exc.InductorError: FileNotFoundError: [Errno 2] No such file or directory: 'C:\\Users\\<your username>\\AppData\\Local\\Temp\\torchinductor_<your username>\\triton\\0\\...LONG...FILE...NAME...'

Or errors like:

[WinError 206] The filename or extension is too long

The solution is to enable Windows' long path support. A reboot is required after the modification.

Since triton-windows post26, it should not happen.

fp8 is not supported on RTX 30xx and older GPUs

If you see errors like

torch._dynamo.exc.BackendCompilerFailed: backend='inductor' raised:
CompilationError: at 8:11:
def triton_(in_ptr0, out_ptr0, xnumel, XBLOCK : tl.constexpr):
    xnumel = 196608
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = tl.full([XBLOCK], True, tl.int1)
    x0 = xindex
    tmp0 = tl.load(in_ptr0 + (x0), None)
    tmp1 = tmp0.to(tl.float32)
           ^

and in the full error log you find

AssertionError: fp8e4nv data type is not supported on CUDA arch < 89

then it's because in the official Triton, fp8 only works on Nvidia GPUs with sm >= 89, such as RTX 40xx and newer.

Since triton-windows post21, fp8 on RTX 30xx is supported.

Error with os.rename

If you see errors like

FileExistsError: [WinError 183] Cannot create a file when that file already exists: ...

then you need: https://github.com/pytorch/pytorch/issues/138211

This has been fixed since PyTorch 2.6 .

Error with CPU tensor

If see errors like

ValueError: Pointer argument (at 0) cannot be accessed from Triton (cpu tensor?)

It means some tensor should be on GPU but it's not there. A possible solution is adding --gpu-only when launching ComfyUI to disable CPU offloading, see https://github.com/woct0rdho/triton-windows/issues/61 . If you actually need CPU offloading or multi-GPU, please open an issue.

No module named 'triton.ops'

triton.ops was removed in Triton 3.1, and this is because some of your Python package is outdated (most likely bitsandbytes), see https://github.com/woct0rdho/triton-windows/issues/65

Exception Code: 0x80000003

If you see Exception Code: 0x80000003 in libtriton.pyd in a function like registerImplicitTypeID, that may not actually be an issue in Triton, but a CUDA error happened earlier, see https://github.com/thu-ml/SageAttention/issues/270 . If you want to help debug, you may set the environment variable CUDA_LAUNCH_BLOCKING=1 and run again.

Build from source

See BUILD.md. This is for developers.