add the support of rocm arch detecting #661
Conversation
|
👋 Hi! Thank you for contributing to the TileLang project. Please remember to run We appreciate you taking this step! Our team will review your contribution, and we look forward to your awesome work! 🚀 |
![gemini-code-assist[bot]](https://avatars.githubusercontent.com/in/956858?s=60&v=4){kind=small}
There was a problem hiding this comment.
Summary of Changes
Hello @zhangnju, I'm Gemini Code Assist1! I'm currently reviewing this pull request and will post my feedback shortly. In the meantime, here's a summary to help you and other reviewers quickly get up to speed!
This pull request significantly improves the tilelang framework's hardware compatibility by adding robust support for AMD GPUs via ROCm (HIP). It refactors existing code to dynamically detect the underlying GPU architecture, moving away from hardcoded NVIDIA CUDA references. This change broadens the framework's usability and ensures it can leverage performance optimizations on both NVIDIA and AMD platforms.
Highlights
- ROCm/HIP Support: Introduced conditional logic across various benchmark and example scripts to detect the presence of
torch.version.hip. If detected, the system will now utilize theCDNAarchitecture (for AMD GPUs) instead of hardcodingCUDA. - Dynamic Architecture Inference: Enhanced the
auto_infer_current_archfunction intilelang/carver/arch/__init__.pyto intelligently select the appropriate GPU architecture. It now prioritizes ROCm (HIP) if available, falls back to CUDA iftorch.cuda.is_available(), and defaults to LLVM if neither is found. - Code Portability: Applied these architecture detection and selection changes consistently across multiple
tilelangcomponents, includingmatmulbenchmarks,convandgemmanalysis, and autotuning examples, improving the framework's adaptability to different GPU hardware.
Using Gemini Code Assist
The full guide for Gemini Code Assist can be found on our documentation page, here are some quick tips.
Invoking Gemini
You can request assistance from Gemini at any point in your pull request via creating an issue comment (i.e. comment on the pull request page) using either /gemini <command> or @gemini-code-assist <command>. Below is a summary of the supported commands.
| Feature | Command | Description |
|---|---|---|
| Code Review | /gemini review |
Performs a code review for the current pull request in its current state. |
| Pull Request Summary | /gemini summary |
Provides a summary of the current pull request in its current state. |
| Comment | @gemini-code-assist | Responds in comments when explicitly tagged, both in issue comments and review comments. |
| Help | /gemini help |
Displays a list of available commands. |
Customization
To customize Gemini Code Assist for GitHub experience, repository maintainers can create a configuration file and/or provide a custom code review style guide (such as PEP-8 for Python) by creating and adding files to a .gemini/ folder in the base of the repository. Detailed instructions can be found here.
Limitations & Feedback
Gemini Code Assist may make mistakes. Please leave feedback on any instances where its feedback is incorrect or counter productive. You can react with 👍 and 👎 on @gemini-code-assist comments or fill out our survey to provide feedback.
You can also get AI-powered code generation, chat, as well as code reviews directly in the IDE at no cost with the Gemini Code Assist IDE Extension.
Footnotes
-
Review the Privacy Notices, Generative AI Prohibited Use Policy, Terms of Service, and learn how to configure Gemini Code Assist in GitHub here. Gemini can make mistakes, so double check it and use code with caution. ↩
There was a problem hiding this comment.
Code Review
This pull request adds support for detecting ROCm/HIP environments to select the appropriate architecture (CDNA) instead of hardcoding CUDA. The implementation in tilelang/carver/arch/__init__.py is good. However, the detection logic is duplicated across numerous benchmark and example files. My main feedback is to refactor this duplicated code by using the newly improved auto_infer_current_arch helper function, which will significantly improve code maintainability.
| if torch.version.hip is not None: | ||
| arch=CDNA("hip") | ||
| else: | ||
| arch = CUDA("cuda") |
There was a problem hiding this comment.
This architecture detection logic is duplicated across several files. It's better to use the auto_infer_current_arch function you've enhanced in tilelang/carver/arch/__init__.py to avoid code duplication and improve maintainability. You'll also need to adjust the imports.
Consider replacing:
from tilelang.carver.arch import CUDA
from tilelang.carver.arch import CDNA
import torchwith:
from tilelang.carver.arch import auto_infer_current_archarch = auto_infer_current_arch()| if torch.version.hip is not None: | ||
| arch=CDNA("hip") | ||
| else: | ||
| arch = CUDA("cuda") |
There was a problem hiding this comment.
This architecture detection logic is duplicated across several files. Using the auto_infer_current_arch function from tilelang/carver/arch/__init__.py would centralize this logic and improve maintainability. You'll also need to adjust the imports.
Consider replacing:
from tilelang.carver.arch import CUDA
from tilelang.carver.arch import CDNA
import torchwith:
from tilelang.carver.arch import auto_infer_current_archarch = auto_infer_current_arch()| if torch.version.hip is not None: | ||
| arch=CDNA("hip") | ||
| else: | ||
| arch = CUDA("cuda") |
There was a problem hiding this comment.
This architecture detection logic is duplicated across several files. Using the auto_infer_current_arch function from tilelang/carver/arch/__init__.py would centralize this logic and improve maintainability. You'll also need to adjust the imports.
Consider replacing:
from tilelang.carver.arch import CUDA
from tilelang.carver.arch import CDNA
import torchwith:
from tilelang.carver.arch import auto_infer_current_archarch = auto_infer_current_arch()| if torch.version.hip is not None: | ||
| cuda_device=CDNA("hip") | ||
| else: | ||
| cuda_device = CUDA("cuda") |
There was a problem hiding this comment.
This architecture detection logic is duplicated. Using the auto_infer_current_arch function from tilelang.carver.arch would centralize this logic and improve maintainability.
Consider replacing:
from tilelang.carver.arch import CUDA
from tilelang.carver.arch import CDNA
import torchwith:
from tilelang.carver.arch import auto_infer_current_archcuda_device = auto_infer_current_arch()| if torch.version.hip is not None: | ||
| cuda_device=CDNA("hip") | ||
| else: | ||
| cuda_device = CUDA("cuda") |
There was a problem hiding this comment.
This architecture detection logic is duplicated. Using the auto_infer_current_arch function from tilelang.carver.arch would centralize this logic and improve maintainability.
Consider replacing:
from tilelang.carver.arch import CUDA
from tilelang.carver.arch import CDNA
import torchwith:
from tilelang.carver.arch import auto_infer_current_archcuda_device = auto_infer_current_arch()| if torch.version.hip is not None: | ||
| arch=CDNA("hip") | ||
| else: | ||
| arch = CUDA("cuda") |
There was a problem hiding this comment.
This architecture detection logic is duplicated across several files. It's better to use the auto_infer_current_arch function you've enhanced in tilelang/carver/arch/__init__.py to avoid code duplication and improve maintainability. You'll also need to adjust the imports.
Consider replacing:
from tilelang.carver.arch import CUDA
from tilelang.carver.arch import CDNA
import torchwith:
from tilelang.carver.arch import auto_infer_current_archarch = auto_infer_current_arch()| if torch.version.hip is not None: | ||
| arch=CDNA("hip") | ||
| else: | ||
| arch = CUDA("cuda") |
There was a problem hiding this comment.
This architecture detection logic is duplicated across several files. It's better to use the auto_infer_current_arch function you've enhanced in tilelang/carver/arch/__init__.py to avoid code duplication and improve maintainability. You'll also need to adjust the imports.
Consider replacing:
from tilelang.carver.arch import CUDA
from tilelang.carver.arch import CDNA
import torchwith:
from tilelang.carver.arch import auto_infer_current_archarch = auto_infer_current_arch()
|
Thanks @zhangnju , merged :) |
Co-authored-by: zhangnju <ningzhan@SMC-SC-DI08-33.dh144.dcgpu>
2 participants
HI @LeiWang1999
when reading the codes of tilelang, I found some codes had been hardcoded as "CUDA" arch. so I added some codes to detect CDNA GPU arch. please check it.
Thanks