gfx1151 on Windows: Pytorch with aotriton SDPA working. Awaiting review.

[scottt]

I've built TheRock without the math and machine learning libraries on Windows successfully following #244 (reply in thread) and am currently running into no ROCm-capable device is detected errors on the gfx1151 when running test programs.

• That seems within the realm of what could be expected
• I'll look into how "supported devices" are determined

My Windows kernel drivers are from:

Setup

1. Fresh Windows 11 Installation

2. Visual Studio 2022 Community Edition 17.13.6

   • During installation, select "Desktop development with C++".

3. Create a separate disk volume (D:) as a Dev Drive

   • git work trees, pip and ccache caches, and the build directory would be placed in this Dev Drive (D:)
   • Visual Studio and tools installed by winget are in C:

4. Windows Security / Microsoft Defender: exclude Dev Drive from virus scanning

   1. Open Windows Security (Press Win + S, search for "Windows Security").
   2. Go to Virus & threat protection.
   3. Scroll down and click Manage settings under Virus & threat protection settings.
   4. Find Exclusions and click Add or remove exclusions.
   5. Click Add an exclusion, then select Folder.
   6. Choose your Dev Drive (D:) to exclude it

5. Windows Search: exclude Dev Drive from indexing

   1. Open the Settings app (Win + I).
   2. Go to Privacy & security > Searching Windows.
   3. Scroll down to the Exclude folders from enhanced search section.
   4. Click Add an excluded folder.
   5. Choose your Dev Drive (D:) to exclude it

6. Enable Windows Developer Mode for symbolic links support

   1. Open the Settings app (Win + I).
   2. Go to System > For developers (or Privacy & security > For developers in older Win11 builds).
   3. Toggle the Developer Mode switch to On.
   4. Test symlink creation. Reference: Git Symbolic Links on Windows (Stack Overflow)

   echo test > test.txt
   mklink link_from_cmd.txt test.txt

7. Enable Windows Long Path Support

   1. Open Windows PowerShell as Administrator.
   2. Run the following command: New-ItemProperty -Path "HKLM:\SYSTEM\CurrentControlSet\Control\FileSystem" -Name "LongPathsEnabled" -Value 1 -PropertyType DWORD -Force
   3. Restart computer

8. Enable PowerShell Script Execution

   1. Open Windows PowerShell as Administrator.
   2. Run the following command:

   Set-ExecutionPolicy RemoteSigned -Scope CurrentUser -Force

9. Install Development Tools (using Winget)

   • Git:

   winget install -e --id Git.Git

   • Git LFS:

   winget install -e --id GitHub.GitLFS

   • Strawberry Perl:

   winget install -e --id StrawberryPerl.StrawberryPerl

   • uv (Python Package Manager):

   winget install --id=astral-sh.uv -e

   • Restart your terminal after these installations to ensure PATH changes take effect.

10. Configure Git

    • Enable symbolic link handling globally for Git (requires Developer Mode to be enabled):
    • Enable long path names on Windows

    git config --global core.symlinks true
    git config --global core.longpaths true

11. Clone amdgpu-windows-interop.git next to TheRock
    git clone https://github.com/nod-ai/amdgpu-windows-interop.git

12. Use patch stack recommended by @stellaraccident until they're merged: Vendor googletest. #404 Re-enable artifact generation on Windows. #408

13. Install Python 3.11 and Setup Virtual Environment
    uv venv --python 3.11 .venv

14. Activate Python Virtual Environment and Install Build Dependencies:

    • Start Developer PowerShell for VS 2022.
    • Activate the virtual environment:

      cd D:\work\TheRock
      .\.venv\Scripts\Activate.ps1

    • Install required Python packages using uv:

      uv pip install -r .\requirements.txt 'cmake<4'

    • (Edit: removed ninja from pip install and rely on ninja 1.12.0 from Strawberry Perl)

15. Fetch TheRock source tree

    python .\build_tools\fetch_sources.py

16. Configure and build with cmake
    python D:\work\therock-build.py -DAMD_GPU_FAMILIES=gfx1151 -DTHEROCK_ENABLE_ALL=OFF -DTHEROCK_ENABLE_HIP_RUNTIME=ON -DTHEROCK_ENABLE_PRIM=ON -DTHEROCK_ENABLE_RAND=ON

Where therock-build is essentially linux_portable_build_in_container.sh rewritten in Python:

import os
import sys
import subprocess
import platform
import time
import signal
from pathlib import Path

# --- Configuration ---
output_base_dir = Path(os.environ['THEROCK_OUTPUT_DIR']).resolve()
source_dir = Path(os.environ['THEROCK_SOURCE_DIR']).resolve()

CCACHE_EXECUTABLE = "ccache"

# --- Setup ---

build_dir = output_base_dir / "build"
caches_dir = output_base_dir / "caches"
ccache_dir = caches_dir / "ccache"
pip_cache_dir = caches_dir / "pip"

print(f"--- Configuration ---")
print(f"Source Directory: {source_dir}")
print(f"Output Base Directory: {output_base_dir}")
print(f"Build Directory: {build_dir}")
print(f"Cache Directory: {caches_dir}")
print(f"Script Arguments: {sys.argv[1:]}")
print(f"---------------------")

print(f"Ensuring directories exist...")
ccache_dir.mkdir(parents=True, exist_ok=True)
pip_cache_dir.mkdir(parents=True, exist_ok=True)
build_dir.mkdir(parents=True, exist_ok=True) # Also ensure build dir exists early

print("Setting environment variables...")
os.environ['CCACHE_DIR'] = str(ccache_dir.resolve())
os.environ['PIP_CACHE_DIR'] = str(pip_cache_dir.resolve())
# Check if ccache is desired/available before setting launchers
use_ccache = True # Set to False to disable ccache
if use_ccache:
    print(f"Configuring CMake to use ccache ('{CCACHE_EXECUTABLE}')...")
    os.environ['CMAKE_C_COMPILER_LAUNCHER'] = CCACHE_EXECUTABLE
    os.environ['CMAKE_CXX_COMPILER_LAUNCHER'] = CCACHE_EXECUTABLE
else:
    print("Skipping ccache configuration.")
    # Ensure they are unset if they existed before
    os.environ.pop('CMAKE_C_COMPILER_LAUNCHER', None)
    os.environ.pop('CMAKE_CXX_COMPILER_LAUNCHER', None)

print(f"CCACHE_DIR = {os.environ.get('CCACHE_DIR')}")
print(f"PIP_CACHE_DIR = {os.environ.get('PIP_CACHE_DIR')}")
print(f"CMAKE_C_COMPILER_LAUNCHER = {os.environ.get('CMAKE_C_COMPILER_LAUNCHER')}")
print(f"CMAKE_CXX_COMPILER_LAUNCHER = {os.environ.get('CMAKE_CXX_COMPILER_LAUNCHER')}")

def run_command(cmd_list, cwd=None):
    """Runs a command, prints it, times it, and checks for errors."""
    print(f"\n--- Executing: {' '.join(map(str, cmd_list))} ---", flush=True)
    start_time = time.monotonic()
    try:
        process = subprocess.run(cmd_list, cwd=cwd, check=True, text=True)
    except FileNotFoundError:
        print(f"ERROR: Command not found: {cmd_list[0]}. Is it installed and in PATH?", file=sys.stderr)
        sys.exit(1)
    except subprocess.CalledProcessError as e:
        print(f"ERROR: Command failed with exit code {e.returncode}", file=sys.stderr)
        sys.exit(e.returncode)
    except Exception as e:
        print(f"ERROR: An unexpected error occurred: {e}", file=sys.stderr)
        sys.exit(1)

    end_time = time.monotonic()
    print(f"--- Command finished successfully in {end_time - start_time:.2f} seconds ---", flush=True)
    return process

# 1. CMake Configure Step
# Pass through any extra arguments given to this python script ($@)
cmake_configure_cmd = [
    "cmake",
    "-G", "Ninja", # Make sure Ninja is installed and in PATH
    "-S", str(source_dir.resolve()),
    "-B", str(build_dir.resolve()),
]
cmake_configure_cmd.extend(sys.argv[1:]) # Add extra arguments from script call

run_command(cmake_configure_cmd)

# 2. CMake Build Step
cmake_build_cmd = [
    "cmake",
    "--build", str(build_dir.resolve())
]
run_command(cmake_build_cmd)

print("\nBuild script completed successfully.")

[stellaraccident]

We just got to that same point on gfx1151 this morning as well. We think there may be a driver update needed and will follow-up on Monday.

[stellaraccident]

@powderluv

[jammm]

clinfo (bundled with the driver) seems to work though.

[jammm]

A few TODOs AFAICS:

• Setup gfx1151 stack similar to @scottt above (except the dev drive stuff.. I want to try without it to make it easier for others to follow)
• Try to get triton built on Windows (its unsupported officially, but there have been CUDA wheels built for it on a fork , so we can try something similar with a corresponding HIP fork.
• Get aotriton building on Windows (needs triton)
• Help with hipBLAS/hipBLASLt/MIOpen building on Windows
• Finally, build pytorch with all of the above

the first 3 is something I plan to look at soon.. the others I will try as time permits. In 2-3 weeks I should hopefully have more time to dedicate. Will need the driver update to test as @stellaraccident pointed out, though it shouldn't be a blocker to get things built at the very least.

aotriton/triton is optional for getting pytorch running, though perf on anything that uses SDPA will be quite slow, and torch.compile won't work.

[stellaraccident]

(disambiguating, @ScottTodd is with AMD and on vacation. It's been a while since I've had two Scott's on the same thing)

[stellaraccident]

Looks like the GPU driver interop blob needs to be rebuilt with a few more device options enabled for gfx1150/1151/1200. We'll build and push a new version on Monday. We don't have an OSS plan for that piece yet, so we have to do it.

[scottt]

@jammm, your plan sounds sensible as usual. I'll finish up the Pytorch 2.7 PR on Linux then join you.

@stellaraccident, for the GPU driver interop blob, try to push an internal goal: "don't let the Windows GPU driver interop piece not being open source block community efforts to work on more hardware support".
If driver interface revisions or major hardware changes requires work that only AMD can do, that maybe unavoidable, but otherwise let outside contributors try to use amdhip64 on hardware not yet officially supported.

[stellaraccident]

@scottt we're on the same page on the goal. Ultimately we will end up with a windows userland that is completely accessible to outside developers. But that is going to take some time (there are things entangled there that cannot simply be released and we need to piece apart), so I opted for a shorter term stop gap with this blob to unblock progress, since this frees the vast majority of rocm dev and let's us parallel track the rest. My hope is that this stopgap gets us the current generation after we overcome a couple of bumps, and then the next generation can be backed by even more accessible code.

[scottt]

@stellaraccident, understood. I wrote the paragraph above so that you have an outside contributor saying this that you can quote if required 😉

[jammm]

I can reproduce the device not detected error that @scottt reported:

(.venv) PS C:\jam\TheRock\build\dist\rocm\bin> .\test_basic.exe
Running main() from C:\jam\TheRock\build\third-party\googletest\source\googletest\src\gtest_main.cc
[==========] Running 2 tests from 1 test suite.
[----------] Global test environment set-up.
[----------] 2 tests from RocprimBasicTests
[ RUN      ] RocprimBasicTests.GetVersion
[       OK ] RocprimBasicTests.GetVersion (0 ms)
[ RUN      ] RocprimBasicTests.GetVersionOnDevice
:1:paldevice.cpp            :906 : 11850983129 us:  Unsupported PAL device with ASIC revision #60
:3:comgrctx.cpp             :33  : 11850983515 us:  Loading COMGR library.
:3:comgrctx.cpp             :128 : 11850986955 us:  Loaded COMGR library version 3.0.
:3:hip_context.cpp          :54  : 11850987376 us:  Direct Dispatch: 0
:3:hip_device_runtime.cpp   :688 : 11850987515 us:   hipSetDevice ( 0 )
:3:hip_device_runtime.cpp   :700 : 11850987572 us:  hipSetDevice: Returned hipErrorNoDevice :
HIP error: no ROCm-capable device is detected file: C:/jam/TheRock/math-libs/rocPRIM/test/rocprim/test_basic.cpp line: 44

[scottt]

The Windows build of the hip compiler driver appears to have not quite correct assumptions about where to find tools, headers and libraries. I need to make a symlink to place clang.exe next to hipcc.exe and compile custom hip test programs like this:

\therock-output\build\dist\rocm\bin\hipcc -c -Og --driver-mode=g++ -D__HIP_PLATFORM_AMD__ --offload-arch=$AMD_GPU_FAMILY -I D:\therock-output-gfx1036\build\dist\rocm\include .\hello.cpp -o hello.o      
\therock-output\build\dist\rocm\lib\llvm\bin\clang.exe -Og --driver-mode=g++ --offload-arch=$AMD_GPU_FAMILY -D__HIP_PLATFORM_AMD__ -L D:\therock-output\build\core\clr\dist\lib .\hello.o -lamdhip64 -o hello.exe

If I build TheRock for gfx1036 instead of gfx1151 and target the iGPU in the Ryzen 9950X of my desktop, test_basic.exe, test_warp_sort.exe, and custom hello-gpu.exe could run successfully.

[stellaraccident]

Can you file an issue with the path repro things? I'm expecting there are several things like that which we will need to patch.

[scottt]

Filed as #410

[jammm]

Someone on discord shared this existing triton fork which they were able to get built and running on Windows w/ HIP SDK 6.2.4 https://github.com/lshqqytiger/triton?tab=readme-ov-file#build-for-windows

related: https://vladmandic.github.io/sdnext-docs/ZLUDA/#how-to-enable-triton

[jammm]

There has been a lot of pre-existing work done by the ZLUDA community to get hipBLASLt, triton etc. running on Windows.

[scottt]

@jammm, is this on the GPU MODE server?

[jammm]

@scottt it's the AMD | SHARK server https://discord.gg/DwjF9Rsv

[jammm]

Related: https://vladmandic.github.io/sdnext-docs/ZLUDA/#how-to-enable-triton

[jammm]

In step 13: Start Developer PowerShell for VS 2022. - it uses the x86 version of MSVC in the PATH. In order to use the x64 ones, you can either use Start -> x64 Native Tools Command Prompt for VS 2022 or Start > Terminal > click on arrow "⌄" > Developer PowerShell for VS2022, which uses the correct x64 version.

[stellaraccident]

I can't believe this is still the default in 2025. This cost me an hour last week because some of the launchers are broken and default incorrectly/silently. I ended up adding a hard fail in TheRock just to keep others from wasting time: https://github.com/ROCm/TheRock/blob/main/CMakeLists.txt#L264

[jammm]

Yeah! insane stuff. That made me facepalm like crazy.

[jammm]

Thank you for adding that warning @stellaraccident. That's how I caught this issue!

[powderluv]

I have a profile.ps1 that sets up x64 easily from here: https://github.com/nod-ai/SRT/wiki/SHARK-Runtime-Windows-Build
We can add a similar note in TheRock

[stellaraccident]

I just pushed this patch which adapts to a new directory layout in amdgpu-windows-interop: 5385753 and activates a new version. Note that there is a "latching" bug in the CMake var so it may cache the old directory. When upgrading, either use a fresh build or do ninja hip-clr+expunge to completely reset CLR.

I'm working on liberating a patch to clr that we have internally which enable gfx1150/1. Stand by. I won't be able to test these much until tomorrow because my dev gfx1150 laptop is being re-imaged.

[stellaraccident]

Alrighty - I don't have a theory as to why it wasn't working yesterday for you guys, but I'll take the progress!

Miscellaneous test errors, we'll catalog and test/fix on this side once we have the CI up. You guys are running a little bit ahead of what we're at, so thanks for being tolerant of the bumps.

I'll turn my attention to seeing what is required to enable the rest of the libraries.

[jammm]

Here's the test output after running all of the test_*.exe inside dist/rocm/bin folder...

out.txt

Some hipCUB tests failed with a missing file error when running test_hipcub_device_radix_sort.exe:
unknown file: error: C++ exception with description "resource deadlock would occur: resource deadlock would occur" thrown in the test body.

[==========] 126 tests from 32 test suites ran. (65228 ms total)
[  PASSED  ] 1 test.
[  FAILED  ] 125 tests, listed below:
[  FAILED  ] HipcubDeviceRadixSort.SortKeysLargeSizes
[  FAILED  ] Id54/HipcubDeviceRadixSort/0.SortKeys, where TypeParam = params<unsigned __int128,unsigned __int128,1,0,128,0,0>
[  FAILED  ] Id54/HipcubDeviceRadixSort/0.SortPairs, where TypeParam = params<unsigned __int128,unsigned __int128,1,0,128,0,0>
[  FAILED  ] Id54/HipcubDeviceRadixSort/0.SortKeysDoubleBuffer, where TypeParam = params<unsigned __int128,unsigned __int128,1,0,128,0,0>
[  FAILED  ] Id54/HipcubDeviceRadixSort/0.SortPairsDoubleBuffer, where TypeParam = params<unsigned __int128,unsigned __int128,1,0,128,0,0>
[  FAILED  ] Id55/HipcubDeviceRadixSort/0.SortKeys, where TypeParam = params<__int128,__int128,1,0,128,0,0>
[  FAILED  ] Id55/HipcubDeviceRadixSort/0.SortPairs, where TypeParam = params<__int128,__int128,1,0,128,0,0>
[  FAILED  ] Id55/HipcubDeviceRadixSort/0.SortKeysDoubleBuffer, where TypeParam = params<__int128,__int128,1,0,128,0,0>
[  FAILED  ] Id55/HipcubDeviceRadixSort/0.SortPairsDoubleBuffer, where TypeParam = params<__int128,__int128,1,0,128,0,0>
...

125 FAILED TESTS```

[jammm]

^ The above tests fail because HipcubDeviceRadixSort.SortKeysLargeSizes fails and that causes the rest of them to fall apart, likely due to a broken runtime or something. If you skip that particular test and run ./test_hipcub_device_radix_sort.exe --gtest_filter="-*LargeSizes*", then all of the tests except this one pass:

$ ./test_hipcub_device_radix_sort.exe --gtest_filter="-*LargeSizes*"
Running main() from D:\jam\TheRock\build\third-party\googletest\source\googletest\src\gtest_main.cc
Note: Google Test filter = -*LargeSizes*
[==========] Running 125 tests from 32 test suites.

...

[----------] Global test environment tear-down
[==========] 125 tests from 32 test suites ran. (44061 ms total)
[  PASSED  ] 125 tests.

It sounds like this particular test case was trying to run a kernel for too long, potentially causing a GPU hang. In my experience, Windows is known to not particularly run well with long-running GPU kernels for some reason.. I'm not sure why but it's probably driver related.

Other than that, every other test executable passes (with some skips), at least when compiled with -DTHEROCK_ENABLE_ALL=OFF -DTHEROCK_ENABLE_HIP_RUNTIME=ON -DTHEROCK_ENABLE_PRIM=ON -DTHEROCK_ENABLE_RAND=ON

[stellaraccident]

Good find, and that is likely tickling a legitimate bug that needs to be triaged/fixed in the runtime. Let's get it logged as an issue so we don't lose it.

[jammm]

Created #425

[stellaraccident]

Ok, I finally got switched over to developing on Strix Halo myself (I'm very hard to dislodge from my dev machine once I get in the groove -- people had been offering me a new system for a while), and everything built fine from the top and ran tests.

I'll shift my attention to getting the rest of the math libraries built on Windows tomorrow.

[jammm]

I came across a clue related to this boost issue - apparently bootstrap.bat doesn't parse --with-libraries argument but bootstrap.sh does. That could explain why it was trying to compile the python parts.

[jammm]

Adding the following two lines right after line 60 of https://github.com/boostorg/boost/blob/master/bootstrap.bat#L60 seems to have done the trick. The python stuff isn't compiled anymore and the build continues on to hipBLASLt.

ECHO libraries = --with-atomic --with-filesystem --with-system ; >> project-config.jam
ECHO. >> project-config.jam

hipBLASLt build fails as expected. Will investigate. Hopefully it's straightforward, because the Windows HIP SDK does include a hipblaslt.lib, so surely they would have it built internally at least.

[stellaraccident]

God, I really dislike boost.

I'm confident that some set of patches somewhere has at one time produced that library. And that's about as far as I care to project my optimism.

[jammm]

I feel the same! boost is weird.

[stellaraccident]

#440

[jammm]

hipBLASLt build fails while trying to build code using tensilelite:

[hipBLASLt] D:/jam/TheRock/math-libs/BLAS/hipBLASLt/tensilelite/rocisa/rocisa/include\helper.hpp:26:10: fatal error: 'sys/wait.h' file not found
[hipBLASLt]    26 | #include <sys/wait.h>
[hipBLASLt]       |          ^~~~~~~~~~~~
[hipBLASLt] 1 error generated.

To get to building this, I had to:

• Build with cmake -B build -GNinja -DTHEROCK_ENABLE_ALL=OFF -DTHEROCK_ENABLE_HIP_RUNTIME=ON -DTHEROCK_ENABLE_PRIM=ON -DTHEROCK_ENABLE_RAND=ON -DTHEROCK_ENABLE_MIOPEN=ON -DTHEROCK_BUILD_TESTING=OFF -DTHEROCK_AMDGPU_TARGETS=gfx1151 -DTHEROCK_VERBOSE=ON
• Attempt a build with cmake --build build, then once it fails at boost, Manually modify build/third_party/boost/source/third_party.bat gfx1151 on Windows: Pytorch with aotriton SDPA working. Awaiting review. #409 (reply in thread)
• Attempt a build again, after which boost compiles and you'll see the above hipBLASLt build errors.

Will investigate on how to fix. This one seems to be the bulk of the errors while building hipBLASLt.

[jammm]

Ok, or just put what you did in a gist - anywhere I can see it. Most likely I can fix it at the source where the package is set up vs at the consumer.

jammm@1c47e55 hope it helps!
On MIOpen itself, I hard-coded zstd and bzip2 stuff https://gist.github.com/jammm/f6d247bffa1c7bd936af1cd9784b75f5

[jammm]

There is an art to getting projects happy with portability patches that involves education, attention to detail, and trust. No one likes accepting mondo "make it work" patches that they can't reason about and without having someone to call when it breaks. Same thing I'm doing now with landing the big patches bit by bit into rocm proper.

If you find a triton fork you like, we can start from there and I'll think on a strategy.

Agreed completely. Upstreaming makes more sense. The triton fork can be a nice reference but ultimately we should go for upstreaming as much as possible. Having said that, It seems triton isn't committing to windows support yet (triton-lang/triton#1640) so we may have to maintain our own fork in the worst case..

[stellaraccident]

When we have an organized fork branch, I'll have some conversations and commit/donate support. Over the timeframe I care about, I'm sure we can find some good ways forward.

[stellaraccident]

I got a bit ahead of myself and ended up not getting to this yet. Presumably, you're busy traveling at this point. Will get it done and catch you on the flip side.

[jammm]

Thanks @stellaraccident ! Yeah i'm wrapping up here and moving countries in a couple days. Most likely be back online early next week, maybe monday-ish

[stellaraccident]

FYI - ROCm/hipBLASLt#1967

I wouldn't recommend using this yet. Builds on my machine with a rebased patch stack.

[stellaraccident]

Sorry I'm running behind. While that is turning, I'll fix the compression libs tomorrow. I also have Friday and Monday blocked off for work on this project.

[stellaraccident]

#464

bzip2 sysdep on Windows. Also taught it to create the executables and projects that depend on it should find them.

[stellaraccident]

@jammm you still need zstd, right?

[jammm]

Yes, MIOpen seems to require it https://gist.github.com/jammm/f6d247bffa1c7bd936af1cd9784b75f5#file-gistfile1-txt-L15-L29

[stellaraccident]

Very well. Will take care of it

[jammm]

Thank you! I'm eager to get back on next week and help getting pytorch building..

[stellaraccident]

zstd on windows: #466

Will get it landed before the weekend.

[scottt]

I'm working on building Triton on Windows using https://github.com/lshqqytiger/triton (linked from https://vladmandic.github.io/sdnext-docs/ZLUDA/#how-to-enable-triton).

I've done a first pass analysis of the difference with upstream Triton and will post findings here.

[jammm]

Thanks @scottt ! The jetlag hit me harder than I thought.. I also have to do some health checks tomorrow. I'll be with you in 24-48 hours on trying this version of triton.
BTW in parallel we should also try and build pytorch without aotriton and sort out any windows-specific issues there. That gives us two tasks we can distribute amongst us.

[stellaraccident]

This sounds like a good way to split it up.

I think I have all of the hipBLASLt Windows portability patches landed at head now. We also got the new build system landed. I'm going to flush TheRock's patch queue and try to make them meet up. Probably a few more things needed, but with any luck, this will just work.

[scottt]

I have Triton built on Windows following:

git clone [email protected]:llvm/llvm-project.git llvm-for-triton
cd llvm-for-triton
# Use llvm version specified in https://github.com/lshqqytiger/triton/blob/main/cmake/llvm-hash.txt
git checkout 1cec5fffd8fddd9d85b516f876093b0e3f0eec5f
cd ..
git clone [email protected]:lshqqytiger/triton.git triton-lshqqytiger
cd triton-lshqqytiger
uv venv --python=3.13
.venv\Script\Activate
python llvm-build.py
python triton-lshqqytiger-build.py # which calls triton-build.py

Where the llvm-build.py and triton-lshqqytiger-build.py scripts are posted here: https://gist.github.com/scottt/b85f07db488b3e37944a9f675370ed3c

[stellaraccident]

This rolls up and drops all hipBLASLt patches. Need to switch gears and test on Windows, but I believe this contains all of our roll ups and might just work. But it was a lot of stuff to land and I wouldn't be completely shocked if something slipped by: #476

[jammm]

Trying to build with -DTHEROCK_ENABLE_SOLVER=ON gives the following linker errors while trying to create cholmod.dll:

[therock-SuiteSparse] [276/280] Linking C shared library CHOLMOD\cholmod.dll
[therock-SuiteSparse] FAILED: CHOLMOD/cholmod.dll CHOLMOD/cholmod.lib
[therock-SuiteSparse] C:\WINDOWS\system32\cmd.exe /C "C:\WINDOWS\system32\cmd.exe /C "D:\jam\.venv\Lib\site-packages\cmake\data\bin\cmake.exe -E __create_def D:\jam\TheRock\build\third-party\SuiteSparse\build\CHOLMOD\CMakeFiles\CHOLMOD.dir\.\exports.def D:\jam\TheRock\build\third-party\SuiteSparse\build\CHOLMOD\CMakeFiles\CHOLMOD.dir\.\exports.def.objs && cd D:\jam\TheRock\build\third-party\SuiteSparse\build" && D:\jam\.venv\Lib\site-packages\cmake\data\bin\cmake.exe -E vs_link_dll --msvc-ver=1943 --intdir=CHOLMOD\CMakeFiles\CHOLMOD.dir --rc=C:\PROGRA~2\WI3CF2~1\10\bin\100226~1.0\x64\rc.exe --mt=C:\PROGRA~2\WI3CF2~1\10\bin\100226~1.0\x64\mt.exe --manifests  -- C:\PROGRA~1\MICROS~1\2022\COMMUN~1\VC\Tools\MSVC\1443~1.348\bin\Hostx64\x64\link.exe /nologo @CMakeFiles\CHOLMOD.rsp  /out:CHOLMOD\cholmod.dll /implib:CHOLMOD\cholmod.lib /pdb:CHOLMOD\cholmod.pdb /dll /version:5.3 /LIBPATH:D:/jam/TheRock/build/third-party/host-blas/stage/lib/host-math/lib /INCREMENTAL:NO  /DEF:CHOLMOD\CMakeFiles\CHOLMOD.dir\.\exports.def && cd ."
[therock-SuiteSparse] LINK: command "C:\PROGRA~1\MICROS~1\2022\COMMUN~1\VC\Tools\MSVC\1443~1.348\bin\Hostx64\x64\link.exe /nologo @CMakeFiles\CHOLMOD.rsp /out:CHOLMOD\cholmod.dll /implib:CHOLMOD\cholmod.lib /pdb:CHOLMOD\cholmod.pdb /dll /version:5.3 /LIBPATH:D:/jam/TheRock/build/third-party/host-blas/stage/lib/host-math/lib /INCREMENTAL:NO /DEF:CHOLMOD\CMakeFiles\CHOLMOD.dir\.\exports.def /MANIFEST:EMBED,ID=2" failed (exit code 1120) with the following output:
[therock-SuiteSparse]    Creating library CHOLMOD\cholmod.lib and object CHOLMOD\cholmod.exp
[therock-SuiteSparse] cholmod_l_super_numeric.c.obj : error LNK2019: unresolved external symbol dtrsm referenced in function rd_cholmod_super_numeric_worker
[therock-SuiteSparse] cholmod_l_super_solve.c.obj : error LNK2001: unresolved external symbol dtrsm
[therock-SuiteSparse] cholmod_super_numeric.c.obj : error LNK2001: unresolved external symbol dtrsm
[therock-SuiteSparse] cholmod_super_solve.c.obj : error LNK2001: unresolved external symbol dtrsm
[therock-SuiteSparse] cholmod_l_super_numeric.c.obj : error LNK2019: unresolved external symbol strsm referenced in function rs_cholmod_super_numeric_worker
[therock-SuiteSparse] cholmod_l_super_solve.c.obj : error LNK2001: unresolved external symbol strsm
[therock-SuiteSparse] cholmod_super_numeric.c.obj : error LNK2001: unresolved external symbol strsm
[therock-SuiteSparse] cholmod_super_solve.c.obj : error LNK2001: unresolved external symbol strsm
[therock-SuiteSparse] cholmod_l_super_numeric.c.obj : error LNK2019: unresolved external symbol ztrsm referenced in function cd_cholmod_super_numeric_worker
[therock-SuiteSparse] cholmod_l_super_solve.c.obj : error LNK2001: unresolved external symbol ztrsm
[therock-SuiteSparse] cholmod_super_numeric.c.obj : error LNK2001: unresolved external symbol ztrsm
[therock-SuiteSparse] cholmod_super_solve.c.obj : error LNK2001: unresolved external symbol ztrsm
[therock-SuiteSparse] cholmod_l_super_numeric.c.obj : error LNK2019: unresolved external symbol ctrsm referenced in function cs_cholmod_super_numeric_worker
[therock-SuiteSparse] cholmod_l_super_solve.c.obj : error LNK2001: unresolved external symbol ctrsm
[therock-SuiteSparse] cholmod_super_numeric.c.obj : error LNK2001: unresolved external symbol ctrsm
[therock-SuiteSparse] cholmod_super_solve.c.obj : error LNK2001: unresolved external symbol ctrsm
[therock-SuiteSparse] cholmod_l_super_numeric.c.obj : error LNK2019: unresolved external symbol dgemm referenced in function rd_cholmod_super_numeric_worker
[therock-SuiteSparse] cholmod_l_super_solve.c.obj : error LNK2001: unresolved external symbol dgemm
[therock-SuiteSparse] cholmod_super_numeric.c.obj : error LNK2001: unresolved external symbol dgemm
[therock-SuiteSparse] cholmod_super_solve.c.obj : error LNK2001: unresolved external symbol dgemm
[therock-SuiteSparse] cholmod_l_super_numeric.c.obj : error LNK2019: unresolved external symbol sgemm referenced in function rs_cholmod_super_numeric_worker
[therock-SuiteSparse] cholmod_l_super_solve.c.obj : error LNK2001: unresolved external symbol sgemm
[therock-SuiteSparse] cholmod_super_numeric.c.obj : error LNK2001: unresolved external symbol sgemm
[therock-SuiteSparse] cholmod_super_solve.c.obj : error LNK2001: unresolved external symbol sgemm
[therock-SuiteSparse] cholmod_l_super_numeric.c.obj : error LNK2019: unresolved external symbol zgemm referenced in function cd_cholmod_super_numeric_worker
[therock-SuiteSparse] cholmod_l_super_solve.c.obj : error LNK2001: unresolved external symbol zgemm
[therock-SuiteSparse] cholmod_super_numeric.c.obj : error LNK2001: unresolved external symbol zgemm
[therock-SuiteSparse] cholmod_super_solve.c.obj : error LNK2001: unresolved external symbol zgemm
[therock-SuiteSparse] cholmod_l_super_numeric.c.obj : error LNK2019: unresolved external symbol cgemm referenced in function cs_cholmod_super_numeric_worker
[therock-SuiteSparse] cholmod_l_super_solve.c.obj : error LNK2001: unresolved external symbol cgemm
[therock-SuiteSparse] cholmod_super_numeric.c.obj : error LNK2001: unresolved external symbol cgemm
[therock-SuiteSparse] cholmod_super_solve.c.obj : error LNK2001: unresolved external symbol cgemm
...

hipSOLVER is apparently a dependency of pytorch, so thought I'd try to build with this.

Release notes suggest BUILD_WITH_SPARSE=OFF for Windows https://github.com/ROCm/hipSOLVER/blob/bb8408f52a166c5ce587b204c61a51ae44bc1daa/CHANGELOG.md?plain=1#L101 but it seems to still use it as a runtime dependency, so will need to figure out how to fix this I think.

[stellaraccident]

Ok, will need to run it down. I can figure it out...

[jammm]

FWIW, This builds fine if I set BUILD_WITH_SPARSE=OFF. Though fixing those linker errors should help with keep it ON for Windows as well.

hipSPARSE itself builds fine with inclusion of <chrono> in a couple places.

[stellaraccident]

@ScottTodd and I will slog through that Thursday. Trying to get rocm build clean on Windows this week. Just taking time... Everything has to be run down individually. But now that I've just about got the hipblaslt monkey off my back, it should go better.

As an aside, I think we eliminated the need for bash on windows in hipBLASLt.

[jammm]

First attempt to building pytorch - some minor changes were needed in LoadHIP.cmake (disabling rccl and libhsa64, and adding another candidate dir for FindHIP.cmake as it's located in rocm/cmake instead of rocm/lib/cmake. I'll push my state soon to show those changes.

I also had to pip uninstall ninja because that ninja was an older version (1.11) and it didn't support path lengths greater than 260. Ninja 1.12.0 onwards supports it, and apparently visual studio 2022 comes bundled with ninja 1.12.1.

First compiler errors are around warnings not being recognized by MSVC:

FAILED: c10/hip/CMakeFiles/c10_hip.dir/driver_api.cpp.obj
C:\PROGRA~1\MICROS~1\2022\COMMUN~1\VC\Tools\MSVC\1443~1.348\bin\Hostx64\x64\cl.exe  /nologo /TP -DEXPORT_AOTI_FUNCTIONS -DMINIZ_DISABLE_ZIP_READER_CRC32_CHECKS -DNOMINMAX -DONNXIFI_ENABLE_EXT=1 -DONNX_ML=1 -DONNX_NAMESPACE=onnx_torch -DROCM_ON_WINDOWS -DROCM_VERSION=85675 -DTORCH_HIP_VERSION=605 -DUSE_EXTERNAL_MZCRC -DUSE_MIMALLOC -DUSE_PROF_API=1 -DWIN32_LEAN_AND_MEAN -D_CRT_SECURE_NO_DEPRECATE=1 -D_UCRT_LEGACY_INFINITY -Dc10_hip_EXPORTS -ID:\jam\TheRock\external-builds\pytorch\src\build\aten\src -ID:\jam\TheRock\external-builds\pytorch\src\aten\src -ID:\jam\TheRock\external-builds\pytorch\src\build -ID:\jam\TheRock\external-builds\pytorch\src -ID:\jam\TheRock\external-builds\pytorch\src\cmake\..\third_party\benchmark\include -ID:\jam\TheRock\external-builds\pytorch\src\third_party\onnx -ID:\jam\TheRock\external-builds\pytorch\src\build\third_party\onnx -ID:\jam\TheRock\external-builds\pytorch\src\nlohmann -ID:\jam\TheRock\external-builds\pytorch\src\third_party\mimalloc\include -ID:\jam\TheRock\external-builds\pytorch\src\c10\hip\..\.. -ID:\jam\TheRock\external-builds\pytorch\src\c10\.. -external:ID:\jam\TheRock\build\dist\rocm\include -external:ID:\jam\TheRock\external-builds\pytorch\src\cmake\..\third_party\googletest\googlemock\include -external:ID:\jam\TheRock\external-builds\pytorch\src\cmake\..\third_party\googletest\googletest\include -external:ID:\jam\TheRock\external-builds\pytorch\src\third_party\protobuf\src -external:ID:\jam\TheRock\external-builds\pytorch\src\third_party\XNNPACK\include -external:ID:\jam\TheRock\external-builds\pytorch\src\third_party\ittapi\include -external:ID:\jam\TheRock\external-builds\pytorch\src\cmake\..\third_party\eigen -external:ID:\jam\TheRock\external-builds\pytorch\src\third_party\ideep\mkl-dnn\include\oneapi\dnnl -external:ID:\jam\TheRock\external-builds\pytorch\src\third_party\ideep\include -external:ID:\jam\TheRock\external-builds\pytorch\src\INTERFACE -external:ID:\jam\TheRock\external-builds\pytorch\src\third_party\nlohmann\include -external:W0 /DWIN32 /D_WINDOWS /GR /EHsc /Zc:__cplusplus /bigobj /FS /utf-8 -DUSE_PTHREADPOOL -DNDEBUG -DUSE_FBGEMM -DUSE_XNNPACK -DSYMBOLICATE_MOBILE_DEBUG_HANDLE /wd4624 /wd4068 /wd4067 /wd4267 /wd4661 /wd4717 /wd4244 /wd4804 /wd4273 /O2 /Ob2 /DNDEBUG /bigobj -DNDEBUG -std:c++17 -MD -fms-runtime-lib=dll -D__HIP_PLATFORM_AMD__=1 -DCUDA_HAS_FP16=1 -DUSE_ROCM -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -DTORCH_HIP_VERSION=605 -Wno-shift-count-negative -Wno-shift-count-overflow -Wno-duplicate-decl-specifier -DCAFFE2_USE_MIOPEN -DTHRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_HIP -std=c++17 -DHIPBLAS_V2 -DC10_HIP_BUILD_MAIN_LIB /showIncludes /Foc10\hip\CMakeFiles\c10_hip.dir\driver_api.cpp.obj /Fdc10\hip\CMakeFiles\c10_hip.dir\ /FS -c D:\jam\TheRock\external-builds\pytorch\src\c10\hip\driver_api.cpp
cl : Command line error D8021 : invalid numeric argument '/Wno-shift-count-negative'

These warning options are inside of HIP_CLANG_FLAGS - but I've no idea who's appending -Wno-shift-count-negative in there, need to investigate.

[jammm]

I've a feeling I'm hitting a cmake bug here.. https://github.com/pytorch/pytorch/blob/main/c10/cuda/test/CMakeLists.txt#L17-L20
It's interfacing with the legacy FindHIP.cmake and it seems to do some cmake magic with the target properties.. and somehow that is messing with the paths involved in the CXX linkage.. the path used here is c10\hip\test\CMakeFiles\c10_hip_HIPAssertionsTest_multiple_writes_from_blocks_and_threads.dir\impl\c10_hip_....hip.obj but it should ideally be the absolute path or the filename itself c10_hip_....hip.obj right?

[jammm]

Wait.. this command fails too:

C:\PROGRA~1\MICROS~1\2022\COMMUN~1\VC\Tools\MSVC\1443~1.348\bin\Hostx64\x64\link.exe /nologo c10\hip\test\CMakeFiles\c10_hip_HIPAssertionsTest_catches_thread_and_block_and_device.dir\impl\c10_hip_HIPAssertionsTest_catches_thread_and_block_and_device_generated_HIPAssertionsTest_catches_thread_and_block_and_device.hip.obj /out:bin\c10_hip_HIPAssertionsTest_catches_thread_and_block_and_device.exe /implib:lib\c10_hip_HIPAssertionsTest_catches_thread_and_block_and_device.lib /pdb:bin\c10_hip_HIPAssertionsTest_catches_thread_and_block_and_device.pdb /version:0.0 /machine:x64 /ignore:4049 /ignore:4217 /ignore:4099 /INCREMENTAL:NO /subsystem:console lib\c10_hip.lib lib\c10.lib lib\gtest_main.lib D:\jam\TheRock\build\dist\rocm\lib\amdhip64.lib lib\gtest.lib kernel32.lib user32.lib gdi32.lib winspool.lib shell32.lib ole32.lib oleaut32.lib uuid.lib comdlg32.lib advapi32.lib /MANIFEST:EMBED,ID=1

LINK : fatal error LNK1181: cannot open input file 'c10\hip\test\CMakeFiles\c10_hip_HIPAssertionsTest_catches_thread_and_block_and_device.dir\impl\c10_hip_HIPAssertionsTest_catches_thread_and_block_and_device_generated_HIPAssertionsTest_catches_thread_and_block_and_device.hip.obj'

But I can cat c10\hip\test\CMakeFiles\c10_hip_HIPAssertionsTest_catches_thread_and_block_and_device.dir\impl\c10_hip_HIPAssertionsTest_catches_thread_and_block_and_device_generated_HIPAssertionsTest_catches_thread_and_block_and_device.hip.obj' and the file does exist in there. What could the issue be? max path length again?

[jammm]

Hmm. Switching to clang-cl seems to circumvent the above issue. Seems like MSVC doesn't work well with long paths!
With $env:CXX="clang-cl" and $env:CC="clang-cl", those linker errors disappear, and new ones come up that complain about triton and some attention stuff:

lld-link: error: undefined symbol: __declspec(dllimport) class std::tuple<class at::Tensor, class at::Tensor, class at::Tensor, class at::Tensor, class c10::SymInt, class c10::SymInt, class at::Tensor, class at::Tensor, class at::Tensor> __cdecl at::native::_cudnn_attention_forward(class at::Tensor const &, class at::Tensor const &, class at::Tensor const &, class std::optional<class at::Tensor> const &, class std::optional<class at::Tensor> const &, class std::optional<class at::Tensor> const &, __int64, __int64, bool, double, bool, bool, class std::optional<double>)
>>> referenced by caffe2/CMakeFiles/torch_hip.dir/__/aten/src/ATen/RegisterCUDA_0.cpp.obj:(class std::tuple<class at::Tensor, class at::Tensor, class at::Tensor, class at::Tensor, class c10::SymInt, class c10::SymInt, class at::Tensor, class at::Tensor, class at::Tensor> __cdecl at::`anonymous namespace'::`anonymous namespace'::wrapper_CUDA___cudnn_attention_forward(class 0xDB3EA737::Tensor const &, class 0xDB3EA737::Tensor const &, class 0xDB3EA737::Tensor const &, class tuple<class at::Tensor, class at::Tensor, class at::Tensor, class at::Tensor, class c10::SymInt, class c10::SymInt, class at::Tensor, class at::Tensor, class at::Tensor>::optional<class at::Tensor> const &, class tuple<class at::Tensor, class at::Tensor, class at::Tensor, class at::Tensor, class c10::SymInt, class c10::SymInt, class at::Tensor, class at::Tensor, class at::Tensor>::optional<class at::Tensor> const &, class tuple<class at::Tensor, class at::Tensor, class at::Tensor, class at::Tensor, class c10::SymInt, class c10::SymInt, class at::Tensor, class at::Tensor, class at::Tensor>::optional<class at::Tensor> const &, class c10::SymInt, class c10::SymInt, bool, double, bool, bool, class tuple<class at::Tensor, class at::Tensor, class at::Tensor, class at::Tensor, class c10::SymInt, class c10::SymInt, class at::Tensor, class at::Tensor, class at::Tensor>::optional<double>))

lld-link: error: undefined symbol: __declspec(dllimport) class at::Tensor __cdecl at::native::triton_scaled_dot_attention(class at::Tensor const &, class at::Tensor const &, class at::Tensor const &, double)
>>> referenced by caffe2/CMakeFiles/torch_hip.dir/__/aten/src/ATen/RegisterCUDA_0.cpp.obj:(class 0xDB3EA737::Tensor __cdecl at::`anonymous namespace'::`anonymous namespace'::wrapper_CUDA___triton_scaled_dot_attention(class 0xDB3EA737::at const &, class 0xDB3EA737::at const &, class 0xDB3EA737::at const &, double))

[ScottTodd]

I also see the same cl : Command line error D8021 : invalid numeric argument '/Wno-ignored-attributes' style errors when trying to use MSVC. clang-cl works, but I get 750MB of build output from these sorts of lines

D:/projects/TheRock/external-builds/pytorch/src/aten/src\ATen/core/Dimname.h:38:1: warning: __declspec attribute 'dllimport' is not supported [-Wignored-attributes]
   38 | TORCH_API std::ostream& operator<<(std::ostream& out, const Dimname& dimname);
      | ^~~~~~~~~
D:/projects/TheRock/external-builds/pytorch/src\c10/macros/Export.h:98:19: note: expanded from macro 'TORCH_API'
   98 | #define TORCH_API C10_IMPORT
      |                   ^~~~~~~~~~
D:/projects/TheRock/external-builds/pytorch/src\c10/macros/Export.h:51:31: note: expanded from macro 'C10_IMPORT'
   51 | #define C10_IMPORT __declspec(dllimport)
      |                               ^~~~~~~~~

Maybe we can patch that in https://github.com/pytorch/pytorch/blob/main/c10/macros/Export.h

[ScottTodd]

Filed #590 for all those warnings

[scottt]

I'm still working on Triton on Windows. I now have both the https://github.com/woct0rdho/triton-windows and https://github.com/lshqqytiger/triton forks built and have updated my scripts in https://gist.github.com/scottt/b85f07db488b3e37944a9f675370ed3c

In the gist link I've recorded the working toolchain and LLVM versions.

I've discovered that for upstream triton's main branch, both cl.exe and clang.exe as shipped in Visual Studio 2022 would produce errors on the C++ produced by mlir-tblgen so we may have some pain sending the changes upstream.

I've read over triton's mixed CMake and setuptools build system and will now work on aotriton.

[jammm]

@Sabrewarrior both me and @scottt gone through the rabbit hole of trying to get aotriton images to compile reliably on Windows and after a lot of pain, we ended up building those images in Linux instead, and copying over the .hsaco compressed .aks2 files onto Windows. Those hsaco files are cross-platform, so it shouldn't matter OS you build the images from. See #409 (reply in thread)

I did however try a workaround where I keep retrying the compile in case of permission errors, but it was quite hacky and honestly the aotriton build system here needs a quality-of-life overhaul in general. Even when I did get them building using multiple cores though, the backward kernels with larger tile sizes (>256) were taking forever. Each kernel was taking longer than an hour to compile, and I haven't been able to measure the actual build time. I suspect maybe this is a triton issue, or a compiler issue, or a mix of both, but I'm not sure as we're already in very unknown territory at this point. A deeper investigation needs to be done to fix the aotriton image builds on Windows. For the time being however, building the images on Linux and then bringing them over to Windows works very well and pain-free.

Re. the carat issue - I recall manually removing the carat and re-adding it, or modifying the code to do it automatically. It was hacky though, so I ended up scrapping it and getting the linux-built ones instead. Your solution of string(REPLACE "^" "^^" HSACO_OUTPATH_WIN32 ${HSACO}) could just work though, and might be worth a try!

[Sabrewarrior]

There are also 287 kernels that fail to build within 2 minutes. Getting stuck on py::class_<PassManager>.run in thirdparty/triton/python/src/ir.cc
Don't know much about MLIR to fix those.
https://gist.github.com/Sabrewarrior/cffa902c9688da352fb169b13fe5b7df#file-aotriton_build_failure-txt
MLIR dump is 600+MB so not going to be uploading that. But its for building llir for the first file.

os.environ['MLIR_ENABLE_DUMP'] = '1'
os.environ['MLIR_DUMP_PATH'] = 'C:\\w\\aotriton\\mlirllirdump.txt'

As for the ^ string replace worked very well!

[jammm]

Yup some of those fwd kernels take longer than 2 minutes. Some of the bwd ones may take an hour or longer

[Sabrewarrior]

Yup some of those fwd kernels take longer than 2 minutes. Some of the bwd ones may take an hour or longer

Does it take that long on linux?

[jammm]

Not really. It was a while ago but compiling a single arch seemed to run fine, where 2 minutes timeout should be ok

[jammm]

I'm facing test hipBLAS, hipBLASLt test regressions after rebuilding TheRock on the main as of yesterday.. something must be off.. it used to pass before on my older gfx1151-win branch #409 (reply in thread)

[jammm]

RelWthDebInfo isn't generating pdb files for hipblaslt, but it did generate for hipBLAS and rocBLAS. Weird.

[stellaraccident]

Let me mind meld with this today. I have developed a special relationship with hipblaslt.

Not to be confused with a good relationship. Just special

[jammm]

I might have found the issue right after attaching the debugger. libraryDir was set to "." https://github.com/ROCm/hipBLASLt/blob/f1c60c38592d03e30ff9730b86c6f9568c4b4647/tensilelite/Tensile/Source/lib/include/Tensile/PlaceholderLibrary.hpp#L143-L153

Copying all the tensile library files into the same directory as hipblaslt-test.exe made the tests run. Still running right now, so far so good...

EDIT: They all passed after the above workaround:

[----------] Global test environment tear-down
[==========] 15252 tests from 13 test suites ran. (951718 ms total)
[  PASSED  ] 15252 tests.
hipBLASLt version: 1500
hipBLASLt git version:

[jammm]

Had added this to the CMakeLists to get it to pick up CMAKE_BUILD_TYPE

diff --git a/next-cmake/CMakeLists.txt b/next-cmake/CMakeLists.txt
index 300b4f10..f508a7b7 100644
--- a/next-cmake/CMakeLists.txt
+++ b/next-cmake/CMakeLists.txt
@@ -35,6 +35,19 @@ include(GenerateExportHeader)
 include(CMakeDependentOption)
 include(hipblaslt_python)

+# Set a default build type if none was specified
+if(NOT CMAKE_BUILD_TYPE AND NOT CMAKE_CONFIGURATION_TYPES)
+  message(STATUS "Setting build type to 'Release' as none was specified.")
+  set(CMAKE_BUILD_TYPE "Release" CACHE STRING "Choose the type of build." FORCE)
+  set_property(CACHE CMAKE_BUILD_TYPE PROPERTY STRINGS "" "Debug" "Release" "MinSizeRel" "RelWithDebInfo")
+endif()
+
+message(STATUS "Setting CMAKE_BUILD_TYPE: ${CMAKE_BUILD_TYPE}")
+
+if(CMAKE_BUILD_TYPE STREQUAL "RelWithDebInfo")
+    set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -g")
+endif()
+
 set(HIPBLASLT_ENABLE_DEVICE ON CACHE BOOL "Build hipBLASLt device libraries.")
 set(HIPBLASLT_ENABLE_CLIENT ON CACHE BOOL "Build hipBLASLt client apps.")
 cmake_dependent_option(HIPBLASLT_ENABLE_HOST "Build hipBLASLt host library." ON "HIPBLASLT_ENABLE_CLIENT" OFF)

That -g might be unnecessary I guess.

[jammm]

The fix might be to handle windows paths here maybe
https://github.com/ROCm/hipBLASLt/blob/c937047ca3fac81e330e40c1f1476c9fe8169a46/tensilelite/Tensile/Source/lib/include/Tensile/Serialization/PlaceholderLibrary.hpp#L62-L68

Maybe the issue has to do with backslashes in the path?

[jammm]

(continuing from #409 (reply in thread))

I finally have pytorch building on Windows. I've pushed my state on this branch https://github.com/ROCm/TheRock/compare/main...jammm:TheRock:jam/gfx1151-win?expand=1

The meat of the pytorch related changes are in 673553d

pytorch was built with the following environment variables set:

USE_ROCM=ON
USE_KINETO=0
BUILD_TEST=0
USE_FLASH_ATTENTION=0
USE_MEM_EFF_ATTENTION=0
CMAKE_PREFIX_PATH="<path_to_therock>/dist/rocm"
HIP_CLANG_PATH="<path_to_therock>/dist/rocm/lib/llvm/bin"
CC=clang-cl
CXX=clang-cl
DISTUTILS_USE_SDK=1

TL;DR is that I had to switch from MSVC to clang-cl for this one. I also cherry-picked some commits that are in pytorch main but didn't land on 2.7.0. Some of my changes in that commit were specific to MSVC, but now that clang-cl works, maybe those MSVC specific changes aren't necessary.

At some point I also forced it to use lld as the linker when linking torch_hip.lib with target_link_options(torch_hip PRIVATE "-fuse-ld=lld"), but I'm not sure if it's needed anymore, though it works with it regardless.

Basic ops like torch.randn((10), device="cuda") work fine, but anything more complex and hipBLASLt crashes. It's likely related to the test regressions I reported in #409 (comment). Once hipBLASLt tests start passing again, hopefully so will pytorch.

[jammm]

Ah okay. I must have misremembered who was doing the jit compilation..
Re. test cases - will take a look soon. Haven't done smoke tests here before so might take some time

[ScottTodd]

Adding rocm/bin to my path did the trick. I was able to build a wheel and run import torch.

Got to the next error trying to run smoketests from https://github.com/ROCm/TheRock/blob/main/external-builds/pytorch/smoke-tests/pytorch_smoke_tests.py:

(.venv) λ python
Python 3.12.8 (tags/v3.12.8:2dc476b, Dec  3 2024, 19:30:04) [MSC v.1942 64 bit (AMD64)] on win32
Type "help", "copyright", "credits" or "license" for more information.
>>> import torch
>>> print(torch.cuda.is_available())
True
>>> matrix1 = torch.tensor([[1.0, 2.0, 3.0], [4.0, 5.0, 6.0]], device="cuda")
>>> batch_matrix1 = torch.ones(10, 2, 3, device="cuda")
>>> batch_matrix2 = torch.ones(10, 3, 4, device="cuda")
>>> result = torch.bmm(batch_matrix1, batch_matrix2)

rocBLAS error: Cannot read D:\projects\TheRock\build\dist\rocm\bin\/rocblas/library/TensileLibrary.dat: No such file or directory for GPU arch : gfx1100
 List of available TensileLibrary Files :

rocBLAS error: Could not initialize Tensile host:
directory_iterator::directory_iterator: The system cannot find the path specified.: "D:\projects\TheRock\build\dist\rocm\bin\/rocblas/library"

I'll try copying the TensileLibrary files from hipblaslt as suggested to fix that.

More notes: https://gist.github.com/ScottTodd/ea9da40039c69e42dc7ef1d3d977e2df. Not sure about the numpy error I was seeing before - that might just crop up after a rebuild with python setup.py develop, but now I'm trying to use python setup.py bdist_wheel instead to get something more portable to other machines.

[ScottTodd]

With

Forgot to mention an important detail - I had to manually copy over hipblaslt/library and rocblas/library into the dist/rocm/bin folder, and needed to have openblas.dll in PATH in order for the tests to run. hipblaslt/library was within dist/rocm/lib and rocblas/library was within the build/math-lbs/rocBLAS/stage/bin folder.

I can now run

import torch
batch_matrix1 = torch.ones(10, 2, 3, device="cuda")
batch_matrix2 = torch.ones(10, 3, 4, device="cuda")
result = torch.bmm(batch_matrix1, batch_matrix2)
print(result)
tensor([[[3., 3., 3., 3.],
         [3., 3., 3., 3.]],
         ...

[ScottTodd]

Here are my repro steps for a pytorch windows + rocm wheel build and running smoketests: https://gist.github.com/ScottTodd/cb4f9f5a697cce91084977ed3674a4bd. Could try https://github.com/ROCm/TheRock/blob/main/external-builds/pytorch/windows_patch_fat_wheel.py to make it self-contained. I'll try landing some of those patches and getting the steps more streamlined/scripted/documented.

[ScottTodd]

• Rename amd_comgr_3.dll to amd_comgr0605.dll (apparently HIP runtime looks for names postfixed with ${HIP_VERSION_MAJOR}${HIP_VERSION_MINOR.

That particular change wasn't needed when importing torch itself - IIRC it only triggered when running some ML ops like conv or attention etc., which would trigger MIOpen and/or hipBLASLt to compile stuff via. hiprtc which in turns calls this comgr dll.

Now I've hit that too while running conv ops. The same amd_comgr_3.dll to amd_comgr0605.dll renaming fixed for me. Here are the logs without the renaming/copying: https://gist.github.com/ScottTodd/3eb07ca3463f8baf7c85cefef3225125.

[stellaraccident]

@jammm shared a gdrive with me with his torch wheel and build/dist/rocm directory. Here is an easier recipe for getting it to work ootb:

• Copy build/dist/rocm to Lib/torch/lib in your Python venv (jam used 3.11)
• In Lib/torch/__init__.py _load_dll_libraries(), change the dll_paths = to this:

        rocm_dll_path = os.path.join(os.path.join(th_dll_path, "rocm", "bin"))
        dll_paths = [
            p
            for p in (rocm_dll_path, th_dll_path, py_dll_path, base_py_dll_path, usebase_path)
            if os.path.exists(p)
        ]

(note addition of rocm_dll_path)

Then seems to do things on my Strix Halo laptop with no further alterations. I didn't need to move files or anything beyond that (although I did delete the test_* files in bin/).

In theory, if we monkey patch the wheel and then rezip it, we can turn any of our standalone torch wheels into a fat torch+rocm wheel. There are better ways to do this for sure, but I'll take the easy button to get started. I'll write a monkey patch script.

[stellaraccident]

It was funny: my first few attempts did the whole exit with error code thing when I was iterating on python -c "import torch". It wasn't until I went to the trouble of firing up the VS debugger that I saw that the last edit exited cleanly without printing anything... because it didn't have an issue and worked.

[stellaraccident]

Handy dandy wheel patcher that should absolutely not be used in production, but will probably be still used ten years from now: https://github.com/ROCm/TheRock/blob/main/external-builds/pytorch/windows_patch_fat_wheel.py

[stellaraccident]

Getting a spurious (non fatal) comgr error. Need to run that down.

[jammm]

Getting a spurious (non fatal) comgr error. Need to run that down.

Yeah I’ve seen that too. Not sure why it’s there

[ScottTodd]

The wheel patcher script worked for me locally too. Got a 4GB torch-2.7.0a0+rocm_git5c0440f-cp312-cp312-win_amd64.whl that passes external-builds/pytorch/smoke-tests/pytorch_smoke_tests.py without needing any modifications to my PATH env var.

[jammm]

Here's a few things to keep in mind while building pytorch on Windows. This consolidates all the things I've said scattered around this discusison thread:

• Don't install ninja using pip. Instead use the one provided by MSVC or install 1.12.x one yourself (context - Should support filename longer than 260 characters ninja-build/ninja#1900 (comment))
• Copy over hipblaslt-bench.exe into the rocm/dist/bin folder gfx1151 on Windows: Pytorch with aotriton SDPA working. Awaiting review. #409 (reply in thread)
• Set HIP_CLANG_PATH to the dist/rocm/lib/llvm/bin directory.
• Rename amd_comgr_3.dll to amd_comgr0605.dll (or amd_comgr<hip_major><hip_minor>.dll - [Windows][comgr] Postfix the dll name with HIP major and minor versions #545 (comment)
• (when running pytorch) Manually copy over all the contents of dist/rocm/lib/hipBLASLt/library/ into the folder that contains the script you're trying to run. - gfx1151 on Windows: Pytorch with aotriton SDPA working. Awaiting review. #409 (reply in thread)
• (maybe optional) Use a symlink to <TheRock_root>/external-builds/pytorch/src to something like C:\src because it'll help avoid long path length issues.
• Use the default compilers while building TheRock itself, then once you want to build pytorch, set the env vars as mentioned in gfx1151 on Windows: Pytorch with aotriton SDPA working. Awaiting review. #409 (comment)
• Copying rocblas folder from build\math-libs\BLAS\rocBLAS\dist\bin to the build\rocm\dist\bin
• Copying build\dist\rocm\lib\host-math\bin\openblas.dll into build\rocm\dist\bin (or another place where appropriate)
• Setting -DMIOPEN_USE_HIPBLASLT=ON in ml-libs/CMakeLists.txt

I'll edit this post and add more notes as I remember them. For now this should do.

[stellaraccident]

Windows file locks are the worst... There is usually a way to do what you want but it can require some change of the algorithm. Rename across local directories is atomic, and so is making a hard link, iirc. Also newer versions of windows will let you delete a file that is in use with appropriate permissions, but it is better to avoid that since not everyone has the permission bit.

[jammm]

I've tried an approach of simply looping until it does the move (using MoveFileExW) and the same for the triton compilation pass (keep looping if compilation fails for whatever reason) <- this obviously isn't production-ready, but works for getting a wheel for the time being.

[jammm]

It's building now. Will report back how it goes

[scottt]

@jammm , the most informative reference I found while working on that is this abandoned PR: microsoft/STL#2062. In particular:

My plan was to use FileRenameInfoEx with FILE_RENAME_REPLACE_IF_EXISTS , FILE_RENAME_POSIX_SEMANTICS and FILE_RENAME_IGNORE_READONLY_ATTRIBUTE.

I suspect the permission settings on the directory containing the file Triton is calling os.replace() on may be significant.

[jammm]

I just rebased on the latest main as of today and it built out of the box which was great. A couple things still remain to be done:

• Copying rocblas folder from build\math-libs\BLAS\rocBLAS\dist\bin to the build\rocm\dist\bin
• Copying build\dist\rocm\lib\host-math\bin\openblas.dll into build\rocm\dist\bin (or another place where appropriate)
• Setting -DMIOPEN_USE_HIPBLASLT=ON in ml-libs/CMakeLists.txt

After I did these, pytorch just built and ran out of the box (after setting the env vars mentioned in #409 (comment).

@scottt will share the build steps for getting aotriton running as well.

[stellaraccident]

Fix for the amd_comgr version issue: #579

[stellaraccident]

FYI - I reverted this because in post submit testing, it breaks device enumeration on Windows (!). Sorry I missed that: if you picked up this commit, please don't waste the time I did and get the revert.

[jammm]

Didn't pick it up yet, no worries!

[jammm]

My bad as well - I kept copies of both the old and new dll in my bin folder. That would explain why I didn't come across this device enumeration issue :(

[stellaraccident]

I'll root cause it tomorrow and fix the patch. That plus the patch below should get most of the weird path handling at runtime taken care of.

[stellaraccident]

Ok, I promise this one doesn't break things worse: #580 -- fixes hipblaslt library finding, so no more need to copy the library files to bin.

Did we have more that had this problem (i.e. do I need to make the same basic change in rocblas)?

[jammm]

hopelully not - just have to make sure those hipblaslt/rocblas library files are both installed into rocm/dist/bin folder

[jammm]

We got aotriton running on PyTorch+Windows. This means flash attention runs on Windows PyTorch! -
SF3D:
Without aotriton (before) - 5-6 seconds
After aotriton - 1.7-1.8 seconds

SD3.5
Before: 14.8s/it
After: 5.2s/it

benchmark.py:

Before:

Benchmarking Scaled Dot-Product Attention (Flash) in FP16 ...
Average time: 111.14 ms
Average peak memory: 4728.06 MB
Average throughput: 1236.75 GFLOP/s

Benchmarking Conv2d in FP16 ...
Average time: 117.80 ms
Average peak memory: 6880.02 MB
Average throughput: 2679.74 GFLOP/s

After:

Benchmarking Scaled Dot-Product Attention (Flash) in FP16 ...
Average time: 10.86 ms
Average peak memory: 32.25 MB
Average throughput: 12673.09 GFLOP/s

Benchmarking Conv2d in FP16 ...
Average time: 116.93 ms
Average peak memory: 6848.02 MB
Average throughput: 2699.84 GFLOP/s

Some SD3.5 results (prompt: cute wallpaper art of a cat holding a chip that has the label 'gfx1151' on it:

[scottt]

Building Pytorch with AOTriton on Windows

• These instructions are for building the gfx1151-therock branch of this aotriton 0.9 based windows port: https://github.com/scottt/aotriton/commits/gfx1151-therock/

  • @scottt has the aotriton Windows portability patches prepared for the main (0.10) branch for upstreaming here but that's not what we're using

• We'll build aotriton in its AOTRITON_NOIMAGE_MODE=ON configuration, which means using pre-built GPU kernel code instead of compiling them from source using the Triton compiler

  • This defers the need of having a sable Triton build on Windows

Build Instructions

1. Download pre-built aotriton 0.9 GPU kernels

   1. @scottt's copy of the gfx1151 aotriton 0.9 GPU kernels is linked from the github repo releases page: aotriton-0.9.2b-images-gfx1151.tar.xz
   2. Note that aotriton 0.9 and 0.10 GPU kernels can use different filenames for its .aks2 files (Navi3x vs gfx1151)

2. (Alternatively) Build aotriton 0.9 on Linux

   1. git clone [email protected]:scottt/aotriton.git && cd aotriton && git checkout gfx1151-therock && python aotriton-build.py
   2. aotriton-build.py is a script like this where you'll likely need to edit the source and build paths

3. Build liblzma on Windows

   1. Used for aotriton GPU kernel decompression
   2. git clone https://github.com/tukaani-project/xz.git && cd xz && git checkout v5.8.1 && python lzma-build.py
   3. lzma-build.py is a script like this

4. Build dlfcn-win32 on Windows

   1. aotriton uses dladdr to locate the path of aotriton_v2.dll at run time which it then uses to locate the aotriton.images GPU kernel directory relative to its DLL path
   2. git clone [email protected]:dlfcn-win32/dlfcn-win32.git && cd dflcn-win32 && python dlfcn-build.py
   3. dlfcn-build.py is a script like this

5. Build aotriton in AOTRITON_NOIMAGE_MODE=on on Windows

   1. git clone [email protected]:scottt/aotriton.git && cd aotriton && git checkout gfx1151-therock && python aotriton-nokernel-build.py
   2. Note that aotriton-nokernel-build.py uses different build options than on Linux here
   3. After a successful build, copy the aotriton.images directory next to the aotriton_v2.dll in the install directory
   4. Copy liblzma.dll and dl.dll next to aotriton_v2.dll

6. Optional Sanity Test: aotriton C++ test program

   1. You can build attn-fwd-test to smoke test your aotriton build
   2. Set $env:PATH=<AOTRITON_DLL_DIR>
   3. Run attn-fwd-test.exe
   4. Verify that the attn-fwd-test finds the right kernel for your GPU, otherwise build packed_kernel.cc in aotriton with verbose logging to debug it
   5. The program currently produces a GPU memory access fault, which is not yet fixed because it was only used to debug GPU kernel decompression and loading

7. Build Pytorch with aotriton support on Windows

   1. Use the gfx1151-win-test branch of this repo https://github.com/scottt/rocm-TheRock/tree/gfx1151-win-test for the Pytorch aotriton patches
   2. python external-builds/pytorch/ptbuild.py --repo <PYTORCH_SRC> --pytorch-ref v2.7.0
   3. python pytorch-build.py
   4. where pytorch-build.py is a script like this
   5. This uses the AOTRITON_INSTALLED_PREFIX environment variable to locate the aotriton built previously

8. Smoke Test: Pytoch Scaled Dot Product Attention

   1. Install the pytorch wheel built in the last step
   2. Set $env:PATH=<AOTRITON_DLL_DIR>
   3. Run spda-test.py which compares output values to a naive implementation and spda-bench.py which shows throughput
   4. @jammm wrote spda-bench.py and @scottt wrote spda-test.py

[jammm]

aotriton will have to be modified and recompiled for gfx1030.

[Sabrewarrior]

Got it working!
Had to set the following two as well:

set 'HIP_CLANG_PATH' = '%ROCM_PATH%/lib/llvm/bin'
set 'CMAKE_LINKER_LINK' = '%ROCM_PATH%/lib/llvm/bin/lld-link.exe'

Will build it with more things working and build triton as well.

Passed sdpa-test.py and got the following on sdpa-bench.py:

Benchmarking Scaled Dot-Product Attention (Flash) in FP16 ...
Average time: 5.08 ms
Average peak memory: 32.25 MB
Average throughput: 27296.27 GFLOP/s

[jammm]

@Sabrewarrior Awesome! Yes you're indeed using aotriton given the numbers I see there.
Which GPU is it BTW?

[scottt]

@Sabrewarrior , nice! Which GPU are you using?

[Sabrewarrior]

Using RX 7900XTX

[maximilienleclei]

@ScottTodd In relation to #409 (reply in thread), I'm browsing the aotriton repo and see repeated messages about Navi32 not being officially supported on aotriton.

Could this be the source of my errors?

If so, this PR from @xinyazhang mentions

## Known problems
1. No official Navi32 support. Users may want to duplicate Navi31 tuning database entries to accomplish Navi 32 support in AOTriton.

[maximilienleclei]

@MaximilienLC are you facing this issue even after copying the aotriton.images folder as mentioned in #651 (comment) ? If yes, then try adding 'gfx1201' : 'gfx1100', in scott's aotriton branch, and rebuild aotriton https://github.com/scottt/aotriton/blob/gfx1151-therock/v2python/gpu_targets.py#L26. Then try again?

You might want to rebuild aotriton images though. For that, you'll need Linux. I'm not sure if WSL works or not, but might be worth the try if you have time.

Yep. Do you mean 1101: 1100?

[jammm]

Right forgot to edit - it’s 1101: 1100

[jammm]

Once you do that change you should build aotriton on Linux with the no image mode set to false, so it builds the gfx1101 images for you with the tuning reuse from gfx1100.

@scottt’s other script that builds aotriton from his gist should help with that

[scottt]

@MaximilienLC , @jammm and I analyzed the upstream aotriton 0.9 tar ball contents and verified that gfx1101 kernels need to built. I'll try doing that.

[scottt]

The pre-built aotriton GPU kernels for gfx110{0,1,2,3} is here: https://github.com/scottt/aotriton/releases/download/0.9.2b/aotriton-0.9.2-images-gfx110x.tar.xz

[dpredie]

trying to use the wheels
torch-2.7.0a0+rocm_git3f903c3-cp311-cp311-win_amd64.whl
torchaudio-2.7.0a0+52638ef-cp311-cp311-win_amd64.whl
torchvision-0.22.0+9eb57cd-cp311-cp311-win_amd64.whl

on a strix halo windows 11,

and gotten this c10.dll issue when trying to run comfyUI or Fooocus ( after pip install the 3 torch wheels in venv, and pip install comfy/fooocus' requirements and run the main.py)

OSError: [WinError 1114] A dynamic link library (DLL) initialization routine failed. Error loading "C:\RocM\rocm_venv\venv\Lib\site-packages\torch\lib\c10.dll" or one of its dependencies.

am i missing anything else?

[jammm]

can you run Dependency walker on the c10.dll ? It could give an idea of which dependent dll is missing.

usually it's a missing visual studio redistributable and/or vs2022 install with c++

[dpredie]

this is correct, new laptop issue, thank you

[0xDELUXA]

Hi @jammm and @scottt. Can I somehow use gfx1200 Pytorch with aotriton on Windows? I mean theoretically how can it be done

[0xDELUXA]

These wheels won't work on gfx1200. We're working on gfx1200 support.

Is there a certain repo that I should follow to see the advancements in gfx1200 support? Or it will just magically appear one day somewhere on the internet XD

[jammm]

This is the repo :)

[xinyazhang]

These wheels won't work on gfx1200. We're working on gfx1200 support.

Is there a certain repo that I should follow to see the advancements in gfx1200 support?

gfx1200 support is already released: https://github.com/ROCm/aotriton/blob/release/0.10/CMakeLists.txt#L61
I think what you are really looking for is Windows+gfx1200 support

[0xDELUXA]

Yes, its Windows support.
So TheRock now supports gfx1200 and aotriton supports it too experimentally. I just don't really understand the things that should be happening right now in order to have pytorch wheels for this card on windows like gfx1201 has. Think I'll just be patient and trust the developers

[Nem404]

@jammm is it normal to get the error 'RuntimeError: No HIP GPUs are available' when using the gfx1201 wheels on a gfx1200 system, or did I just install them incorrectly? I know you said they aren't compatible with this card but is this the error it should give for me?

[jammm]

Yes, that’s expected.

…

On Monday, June 30, 2025, Nem404 ***@***.***> wrote: @jammm <https://github.com/jammm> is it normal to get the error 'RuntimeError: No HIP GPUs are available' when using the gfx1201 wheels on a gfx1200 system, or did I just install them incorrectly? I know you said they aren't compatible with this card but is this the error it should give for me? — Reply to this email directly, view it on GitHub <#409 (reply in thread)>, or unsubscribe <https://github.com/notifications/unsubscribe-auth/AATCSOGNPBQVN6ADRQ2WGKL3F76VRAVCNFSM6AAAAAB3ASRN2WVHI2DSMVQWIX3LMV43URDJONRXK43TNFXW4Q3PNVWWK3TUHMYTGNRRGAYDONA> . You are receiving this because you were mentioned.Message ID: ***@***.***>

[Nem404]

One more thing Im curious about: @xinyazhang specifically said, quote: "Porting to Windows requires @scottt 's effort." Speaking about aotriton 0.10b that supports gfx1200 experimentally. But only on Linux, so this means I cant use it on Windows rn? Im just thinking if you devs can do smth about this too...

[jammm]

You can certainly build aotriton on windows but we don't officially support it just yet. You'll have to build @scottt's aotriton branch if you want it right now.

To get pytorch and comfyui running. aotriton is not required. It'll help as a perf enhancer, but getting pytorch itself running is more important. Getting @scottt's aotriton changes upstreamed is a separate effort that could be looked at after we land the initial gfx1200 support. There's some potential issues that needs to be ironed out, as building aotriton for windows is not straightforward and requires the aotriton.images prebuilt from linux as a dependency to workaround the current issues.

[stellaraccident]

It is likely that for the near future there continue to be "community special brew" releases on Windows that do a variety of difficult to automate/support things to get the best performing kernels enabled in the build. These are built mostly from our open-source footprint but with various procedures/patches/branches for secondary dependencies that we aren't in a position to support properly.

In terms of AMD official releases, we support this kind of scrappy effort to have the best options in the short term, and we are working on systematically enabling and doing upstream porting so that the out of the box PyTorch builds on Windows can be as good as they can be. A lot of this work takes a good deal of time to do since it requires coordinating with those projects, streamlining builds/releases, and making sure that people are responsible to support them. We'll get there, but in the meantime, I'm really thankful for the special brews that @scottt and others are doing to get the best performing options in people's hands as quickly as possible.

Basically three months ago, most of ROCm and PyTorch were in the category of not being able to be built by the community. Now, we've almost got that release pipeline worked for out of the box downloads from AMD official builds, but we're just getting started on consolidating the porting work needed on ancillary projects like *triton. It is a journey.

[Nem404]

What a great community AMD has. I'm actually following the developments towards PyTorch on Windows and see things like @ScottTodd and others making commits 24/7. Love to see it. Thank you all for your incredible efforts.

[V6ser]

It is likely that for the near future there continue to be "community special brew" releases on Windows that do a variety of difficult to automate/support things to get the best performing kernels enabled in the build. These are built mostly from our open-source footprint but with various procedures/patches/branches for secondary dependencies that we aren't in a position to support properly.

In terms of AMD official releases, we support this kind of scrappy effort to have the best options in the short term, and we are working on systematically enabling and doing upstream porting so that the out of the box PyTorch builds on Windows can be as good as they can be. A lot of this work takes a good deal of time to do since it requires coordinating with those projects, streamlining builds/releases, and making sure that people are responsible to support them. We'll get there, but in the meantime, I'm really thankful for the special brews that @scottt and others are doing to get the best performing options in people's hands as quickly as possible.

Basically three months ago, most of ROCm and PyTorch were in the category of not being able to be built by the community. Now, we've almost got that release pipeline worked for out of the box downloads from AMD official builds, but we're just getting started on consolidating the porting work needed on ancillary projects like *triton. It is a journey.

Hi Stella, is it possible to get gfx1031 MIOpen kernels?
I currently have a working Windows torch(+torchaudio/torchvision) build for this architecture thanks to all the work done here. As I understand gfx103X cards will be treated the same, so aotriton needs modifications for proper support. Is there anyone working on this currently? What can we expect challenges to be?

[Bazza-63]

Completely random and not really on topic but i thought I'd share. Yesterday, instead of launching the Visual Studio X64 Native environment I accidentally clicked on the OneAPI x64 one. It built the hip compiler/s fine and the build finished without any obvious breakage. Looking at the logs it was calling icx-cl, tbb, mpi at different times of the build.