Skip to content

Commit 22411a3

Browse files
add Hygon HCU backend support on master.
Port HCU backend alongside existing GCU/MACA backends after upstream master refactor (examples-based layout). Includes review fixes: validate shared memory only via hipModuleLaunchKernel path, and gate librt.so stub creation behind HCU_CREATE_LIBRT_STUB (default OFF). Co-authored-by: Cursor <cursoragent@cursor.com>
1 parent 434c7b7 commit 22411a3

12 files changed

Lines changed: 386 additions & 7 deletions

File tree

CMakeLists.txt

Lines changed: 12 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -3,11 +3,11 @@ cmake_minimum_required(VERSION 3.25 FATAL_ERROR)
33
# ==============================================================================
44
# Backend Selection (must be before project())
55
# ==============================================================================
6-
set(BACKEND "CUDA" CACHE STRING "Target backend: CUDA, IX, MUSA, MACA, NPU, or GCU")
7-
set_property(CACHE BACKEND PROPERTY STRINGS "CUDA" "IX" "MUSA" "MACA" "NPU" "GCU")
6+
set(BACKEND "CUDA" CACHE STRING "Target backend: CUDA, IX, MUSA, MACA, NPU, GCU, or HCU")
7+
set_property(CACHE BACKEND PROPERTY STRINGS "CUDA" "IX" "MUSA" "MACA" "NPU" "GCU" "HCU")
88

9-
if(NOT BACKEND MATCHES "^(CUDA|IX|MUSA|MACA|NPU|GCU)$")
10-
message(FATAL_ERROR "Invalid BACKEND: ${BACKEND}. Must be CUDA, IX, MUSA, MACA, NPU, or GCU")
9+
if(NOT BACKEND MATCHES "^(CUDA|IX|MUSA|MACA|NPU|GCU|HCU)$")
10+
message(FATAL_ERROR "Invalid BACKEND: ${BACKEND}. Must be CUDA, IX, MUSA, MACA, NPU, GCU, or HCU")
1111
endif()
1212
message(STATUS "Building with backend: ${BACKEND}")
1313

@@ -24,7 +24,7 @@ if(BACKEND STREQUAL "MACA")
2424
set(USE_MACA ON)
2525
endif()
2626

27-
# Project definition (NPU, MUSA, and MACA don't need CUDA language)
27+
# Project definition (NPU, MUSA, MACA, GCU, and HCU don't need CUDA language)
2828
if(BACKEND STREQUAL "NPU")
2929
project(TritonJIT LANGUAGES CXX VERSION 0.1.0)
3030
elseif(BACKEND STREQUAL "MUSA")
@@ -33,6 +33,8 @@ elseif(BACKEND STREQUAL "MACA")
3333
project(TritonJIT LANGUAGES CXX VERSION 0.1.0)
3434
elseif(BACKEND STREQUAL "GCU")
3535
project(TritonJIT LANGUAGES CXX VERSION 0.1.0)
36+
elseif(BACKEND STREQUAL "HCU")
37+
project(TritonJIT LANGUAGES CXX VERSION 0.1.0)
3638
else()
3739
project(TritonJIT LANGUAGES CUDA CXX VERSION 0.1.0)
3840
endif()
@@ -54,6 +56,8 @@ elseif(BACKEND STREQUAL "MACA")
5456
add_compile_definitions(BACKEND_MACA USE_MACA)
5557
elseif(BACKEND STREQUAL "GCU")
5658
add_compile_definitions(BACKEND_GCU)
59+
elseif(BACKEND STREQUAL "HCU")
60+
add_compile_definitions(BACKEND_HCU)
5761
endif()
5862

5963
# String macro for Python runtime
@@ -71,7 +75,7 @@ add_compile_definitions(BACKEND_NAME="${TRITON_BACKEND_NAME}")
7175
set(CMAKE_CXX_STANDARD 20)
7276
set(CMAKE_CXX_STANDARD_REQUIRED ON)
7377
set(CMAKE_CXX_EXTENSIONS OFF)
74-
if(NOT (BACKEND STREQUAL "NPU" OR BACKEND STREQUAL "MUSA" OR BACKEND STREQUAL "GCU"))
78+
if(NOT (BACKEND STREQUAL "NPU" OR BACKEND STREQUAL "MUSA" OR BACKEND STREQUAL "GCU" OR BACKEND STREQUAL "HCU"))
7579
set(CMAKE_CUDA_STANDARD 17)
7680
endif()
7781
set(CMAKE_EXPORT_COMPILE_COMMANDS ON)
@@ -172,6 +176,8 @@ elseif(BACKEND STREQUAL "MACA")
172176
include(BackendMACA)
173177
elseif(BACKEND STREQUAL "GCU")
174178
include(BackendGCU)
179+
elseif(BACKEND STREQUAL "HCU")
180+
include(BackendHCU)
175181
endif()
176182

177183
# ==============================================================================

README.md

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -13,6 +13,7 @@ It supports multiple hardware backends through a compile-time backend policy des
1313
- **MUSA**: Moore Threads GPUs (warp size 32)
1414
- **NPU**: Ascend/Huawei (ACL API)
1515
- **IX**: Tianshu GPUs (warp size 64)
16+
- **HCU**: Hygon GPUs (warp size 64, HIP-compatible)
1617

1718
The project aims to reduce the inevitable Python overhead when using Triton in Python code.
1819
For many kernels, the execution time of the kernel is much shorter than the CPU overhead.
@@ -225,8 +226,19 @@ export PATH=$CUCC_PATH/tools:$PATH
225226
export CUCC_CMAKE_ENTRY=2
226227
cmake_maca -S . -B build/ -DPython_ROOT="$(which python)/../.." -DBACKEND=MACA
227228
make_maca -C build/ -j2
229+
230+
# HCU (Hygon)
231+
cmake -S . -B build/ -DPython_ROOT="$(which python)/../.." -DBACKEND=HCU
232+
```
233+
234+
On Ubuntu 22.04+ with DTK, if linking fails due to missing `librt.so`, create the stub once:
235+
236+
```shell
237+
sudo ln -s /usr/lib/x86_64-linux-gnu/librt.so.1 /usr/lib/x86_64-linux-gnu/librt.so
228238
```
229239

240+
Alternatively, pass `-DHCU_CREATE_LIBRT_STUB=ON` at configure time (requires write access to `/usr/lib`).
241+
230242
You can also specify build type via `-DCMAKE_BUILD_TYPE` and the install prefix using `-DCMAKE_INSTALL_PREFIX`.
231243

232244
### Build

cmake/BackendHCU.cmake

Lines changed: 96 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,96 @@
1+
# ==============================================================================
2+
# HCU Backend Configuration
3+
# ==============================================================================
4+
5+
message(STATUS "Configuring HCU backend...")
6+
7+
# HCU uses HIP runtime
8+
# Find HIP package - typically provided by DTK installation
9+
list(APPEND CMAKE_PREFIX_PATH
10+
"$ENV{ROCM_PATH}"
11+
"$ENV{HIP_PATH}"
12+
"/opt/rocm"
13+
"/opt/rocm/hip"
14+
"/opt/dtk"
15+
"/opt/dtk/hip"
16+
)
17+
18+
# DTK's MIOpen depends on /usr/lib/x86_64-linux-gnu/librt.so in its CMake target,
19+
# but on Ubuntu 22.04+ (glibc 2.34) librt was merged into libc and the stub
20+
# symlink librt.so is no longer shipped. Create it manually if needed:
21+
# sudo ln -s /usr/lib/x86_64-linux-gnu/librt.so.1 /usr/lib/x86_64-linux-gnu/librt.so
22+
# Or opt in to automatic creation at configure time (requires write access to /usr/lib):
23+
option(HCU_CREATE_LIBRT_STUB
24+
"Create /usr/lib/x86_64-linux-gnu/librt.so symlink for DTK MIOpen (requires root)"
25+
OFF)
26+
if(HCU_CREATE_LIBRT_STUB
27+
AND NOT EXISTS "/usr/lib/x86_64-linux-gnu/librt.so"
28+
AND EXISTS "/usr/lib/x86_64-linux-gnu/librt.so.1")
29+
message(STATUS "HCU_CREATE_LIBRT_STUB=ON: creating librt.so symlink (required by DTK MIOpen)")
30+
execute_process(
31+
COMMAND ${CMAKE_COMMAND} -E create_symlink
32+
/usr/lib/x86_64-linux-gnu/librt.so.1
33+
/usr/lib/x86_64-linux-gnu/librt.so
34+
RESULT_VARIABLE _librt_symlink_result
35+
)
36+
if(_librt_symlink_result EQUAL 0)
37+
message(STATUS "Created /usr/lib/x86_64-linux-gnu/librt.so -> librt.so.1")
38+
else()
39+
message(WARNING
40+
"Failed to create /usr/lib/x86_64-linux-gnu/librt.so -> librt.so.1 "
41+
"(exit code ${_librt_symlink_result}). "
42+
"Build may fail with 'No rule to make target librt.so'. "
43+
"Fix manually with: "
44+
"sudo ln -s /usr/lib/x86_64-linux-gnu/librt.so.1 /usr/lib/x86_64-linux-gnu/librt.so")
45+
endif()
46+
elseif(NOT EXISTS "/usr/lib/x86_64-linux-gnu/librt.so"
47+
AND EXISTS "/usr/lib/x86_64-linux-gnu/librt.so.1")
48+
message(WARNING
49+
"Missing /usr/lib/x86_64-linux-gnu/librt.so (required by DTK MIOpen on glibc 2.34+). "
50+
"Create it manually with: "
51+
"sudo ln -s /usr/lib/x86_64-linux-gnu/librt.so.1 /usr/lib/x86_64-linux-gnu/librt.so "
52+
"or reconfigure with -DHCU_CREATE_LIBRT_STUB=ON")
53+
endif()
54+
55+
find_package(hip REQUIRED)
56+
message(STATUS "Found HIP: ${hip_VERSION}")
57+
58+
# The HIP runtime (libgalaxyhip.so) carries RUNPATH=/opt/dtk/lib which
59+
# contains a bundled libunwind.so.8 (from hipprof_utils). CMake warns
60+
# about an RPATH conflict with the identical-SONAME system libunwind in
61+
# /usr/lib/x86_64-linux-gnu. The DTK compiler already treats several
62+
# /opt/dtk subdirectories as implicit link dirs; extend this to the parent
63+
# /opt/dtk/lib so CMake omits it from RPATH (the HIP library's own RUNPATH
64+
# still provides it at runtime).
65+
if(DEFINED ENV{ROCM_PATH})
66+
file(REAL_PATH "$ENV{ROCM_PATH}/lib" _rocm_lib_real)
67+
list(APPEND CMAKE_CXX_IMPLICIT_LINK_DIRECTORIES
68+
"$ENV{ROCM_PATH}/lib" "${_rocm_lib_real}")
69+
endif()
70+
71+
get_target_property(_torch_hip_opts torch_hip INTERFACE_COMPILE_OPTIONS)
72+
if(_torch_hip_opts)
73+
# torch_hip sets -std=c++17 in INTERFACE_COMPILE_OPTIONS, which overrides
74+
# our C++20 standard. Strip it so C++20 concepts work.
75+
# See Caffe2Targets.cmake for more details.
76+
list(FILTER _torch_hip_opts EXCLUDE REGEX "-std=c\\+\\+17")
77+
# -Wno-duplicate-decl-specifier is valid for C/ObjC only; GCC warns when
78+
# passed to C++ compilation. Strip it from the interface flags.
79+
# See Caffe2Targets.cmake for more details.
80+
list(FILTER _torch_hip_opts EXCLUDE REGEX "-Wno-duplicate-decl-specifier")
81+
set_property(TARGET torch_hip PROPERTY INTERFACE_COMPILE_OPTIONS ${_torch_hip_opts})
82+
endif()
83+
84+
# PyTorch pip wheels bundle libibverbs inside torch.libs/ with hashed SONAMEs
85+
# (e.g. libnl-3-04364822.so.200.26.0) that are not on the linker search path.
86+
# The linker follows libibverbs's DT_NEEDED chain and fails to resolve them at
87+
# link time, even though libtorch_hip.so's $ORIGIN/../../torch.libs RPATH
88+
# resolves them correctly at runtime. Suppress the link-time error.
89+
add_link_options(-Wl,--allow-shlib-undefined)
90+
91+
# ------------------------------- Helper Function ------------------------------
92+
function(target_link_hcu_libraries target_name)
93+
target_link_libraries(${target_name} PRIVATE hip::host)
94+
endfunction()
95+
96+
message(STATUS "HCU backend configuration complete")

examples/common/backend_ops.h

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -22,6 +22,9 @@
2222
#include "c10/cuda/CUDAStream.h"
2323
#elif defined(BACKEND_GCU)
2424
#include <tops_runtime_api.h>
25+
#elif defined(BACKEND_HCU)
26+
#include <hip/hip_runtime.h>
27+
#include "c10/hip/HIPStream.h"
2528
#else
2629
#include "c10/cuda/CUDAStream.h"
2730
#endif
@@ -37,6 +40,8 @@ using RawStream = musaStream_t;
3740
using RawStream = mcStream_t;
3841
#elif defined(BACKEND_GCU)
3942
using RawStream = topsStream_t;
43+
#elif defined(BACKEND_HCU)
44+
using RawStream = hipStream_t;
4045
#else
4146
using RawStream = CUstream;
4247
#endif
@@ -55,6 +60,8 @@ inline RawStream get_device_stream([[maybe_unused]] const at::Tensor& t) {
5560
return reinterpret_cast<mcStream_t>(c10::cuda::getCurrentCUDAStream(t.device().index()).stream());
5661
#elif defined(BACKEND_GCU)
5762
return nullptr;
63+
#elif defined(BACKEND_HCU)
64+
return static_cast<hipStream_t>(c10::hip::getCurrentHIPStream(t.device().index()).stream());
5865
#else
5966
return static_cast<CUstream>(c10::cuda::getCurrentCUDAStream(t.device().index()).stream());
6067
#endif

include/triton_jit/backend_config.h

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10,6 +10,8 @@
1010
#include "triton_jit/backends/ix_backend.h"
1111
#elif defined(BACKEND_GCU)
1212
#include "triton_jit/backends/gcu_backend.h"
13+
#elif defined(BACKEND_HCU)
14+
#include "triton_jit/backends/hcu_backend.h"
1315
#else
1416
#include "triton_jit/backends/cuda_backend.h"
1517
#endif
@@ -40,6 +42,10 @@ using DefaultBackend = IxBackend;
4042
/// Default backend for GCU (Enflame)
4143
using DefaultBackend = GcuBackend;
4244

45+
#elif defined(BACKEND_HCU)
46+
/// Default backend for HCU (Hygon, HIP-compatible)
47+
using DefaultBackend = HcuBackend;
48+
4349
#else
4450
// Default to CUDA if no backend specified
4551
#warning "No backend specified, defaulting to CUDA. Use -DBACKEND=CUDA explicitly."

0 commit comments

Comments
 (0)