feat: support build on gfx1250 (Part1)#374
Open
RuibinCheung wants to merge 5 commits into
Open
Conversation
![cursor[bot]](https://avatars.githubusercontent.com/in/1210556?s=60&v=4){kind=small}
There was a problem hiding this comment.
安全审查完成：未发现本 PR 引入或暴露的 medium/high/critical 漏洞。
检查范围包括新增/修改的 gfx1250 build 支持、可选 GEMM backend 选择、ROCM 路径探测、pybind stub、grouped_gemm_compute_offs 独立实现，以及相关 GPU kernel/inline asm 变更。未发现攻击者可控输入到注入、权限绕过、SSRF/XSS/路径穿越、不安全反序列化、秘密泄露或新增供应链 sink 的可利用路径。
既有自动化安全审查线程：未发现需要复核的未解决线程。
Sent by Cursor Automation: Find vulnerabilities
Contributor
There was a problem hiding this comment.
Pull request overview
Adds build support for the gfx1250 (CDNA5) architecture by making optional GEMM backends selectable at build time and porting device-side waitcnt usage to gfx12’s split wait counters, while consolidating grouped-GEMM helper code so it remains available even when CK/turbo backends are disabled.
Changes:
- Adds gfx1250 to supported arch list and introduces
PRIMUS_TURBO_BUILD_BACKENDto selectively build CK and/or turbo GEMM backends (with guarded headers and stubbed host bindings when disabled). - Ports device inline-asm wait logic for gfx12 using split counters and shared macros.
- Consolidates
compute_group_offsinto an always-built, backend-agnostic kernel and moves the PyTorch host binding to an always-built utility TU.
Reviewed changes
Copilot reviewed 15 out of 15 changed files in this pull request and generated 12 comments.
Show a summary per file
| File | Description |
|---|---|
| tools/build_utils.py | Improves ROCm discovery and adds backend-selection helpers; introduces _join_rocm_home. |
| tools/build_ext.py | Reuses _join_rocm_home from build utils. |
| setup.py | Adds gfx1250 arch, filters sources based on selected backends, and defines backend guard macros in compile flags. |
| csrc/pytorch/type_traits.h | Guards CK-only type mapping behind PRIMUS_TURBO_BUILD_CK_BACKEND. |
| csrc/pytorch/grouped_gemm/utils/grouped_gemm_utils.cpp | New always-built PyTorch grouped-GEMM utility binding for offset computation. |
| csrc/pytorch/grouped_gemm/ck_grouped_gemm.cpp | CK implementation now guarded; adds CK-disabled stubs. |
| csrc/pytorch/gemm/turbo_gemm.cpp | Adds turbo-disabled stub guarded by PRIMUS_TURBO_BUILD_TURBO_BACKEND. |
| csrc/pytorch/gemm/ck_gemm.cpp | Adds CK-disabled stub guarded by PRIMUS_TURBO_BUILD_CK_BACKEND. |
| csrc/kernels/grouped_gemm/compute_group_offs.cu | New standalone always-built device implementation of compute_group_offs. |
| csrc/kernels/grouped_gemm/ck_grouped_gemm.cu | Removes duplicated compute_group_offs implementation/instantiation. |
| csrc/kernels/deep_ep/utils.cuh | Introduces gfx12 split-counter wait macro and uses it in fences. |
| csrc/include/primus_turbo/grouped_gemm.h | Guards CK-only includes/types/APIs behind PRIMUS_TURBO_BUILD_CK_BACKEND. |
| csrc/include/primus_turbo/gemm.h | Guards CK/turbo-only declarations behind backend macros. |
| csrc/include/primus_turbo/device/reduce.cuh | Uses a gfx12 LDS-wait macro for ds swizzles/reductions. |
| csrc/include/primus_turbo/device/memory.cuh | Switches wait helpers to gfx12 split counters for gfx1250. |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
Comment on lines
+110
to
+111
| if rocm_home and torch.version.hip is None: | ||
| logger.warning("No ROCm runtime is found, using ROCM_HOME='%s'", rocm_home) |
Comment on lines
+115
to
120
| def _join_rocm_home(*paths) -> str: | ||
| return os.path.join(find_rocm_home(), *paths) | ||
|
|
||
|
|
||
| HIP_HOME: Path = find_hip_home() | ||
| HIP_HOME: Path = _join_rocm_home("hip") | ||
| HIP_LIBRARY_PATH = os.path.join(HIP_HOME, "lib") |
Comment on lines
+40
to
+46
| # csrc/kernels path fragments that depend on CK, skipped when the CK backend is disabled. | ||
| CK_DISABLED_SOURCE_MARKERS = ( | ||
| "/gemm/ck/", | ||
| "/gemm/ck_gemm.cu", | ||
| "/grouped_gemm/ck/", | ||
| "/grouped_gemm/ck_grouped_gemm.cu", | ||
| ) |
Comment on lines
+13
to
+20
| // gfx12 (gfx1250 / gfx1251, CDNA5) removed the unified `s_waitcnt`; the combined | ||
| // LDS + VMEM wait is expressed via the split counters. | ||
| #if defined(__gfx1250__) | ||
| #define PRIMUS_TURBO_WAIT_ALL_STR \ | ||
| "s_wait_dscnt 0\n\ts_wait_kmcnt 0\n\ts_wait_loadcnt 0\n\ts_wait_storecnt 0" | ||
| #else | ||
| #define PRIMUS_TURBO_WAIT_ALL_STR "s_waitcnt lgkmcnt(0) vmcnt(0)" | ||
| #endif |
Comment on lines
+10
to
+16
| // gfx12 (gfx1250 / gfx1251, CDNA5) removed the unified `s_waitcnt`; LDS waits use | ||
| // the split `s_wait_dscnt` counter instead. | ||
| #if defined(__gfx1250__) | ||
| #define PRIMUS_TURBO_WAIT_DS_STR "s_wait_dscnt 0" | ||
| #else | ||
| #define PRIMUS_TURBO_WAIT_DS_STR "s_waitcnt lgkmcnt(0)" | ||
| #endif |
Comment on lines
+88
to
+93
| at::Tensor turbo_gemm_fp8(at::Tensor A, at::Tensor scaleA_inv, at::Tensor B, at::Tensor scaleB_inv, | ||
| const at::ScalarType out_dtype, bool transA, bool transB, bool transC, | ||
| const std::string &granularity) { | ||
| PRIMUS_TURBO_ERROR("turbo_gemm_fp8 is unavailable: turbo backend not built. " | ||
| "Rebuild with PRIMUS_TURBO_BUILD_BACKEND including 'turbo'."); | ||
| } |
Comment on lines
+310
to
+315
| at::Tensor ck_grouped_gemm(at::Tensor &a, at::Tensor &b, at::Tensor &group_lens, | ||
| at::Tensor &group_offs, const bool transA, const bool transB, | ||
| c10::optional<int64_t> num_cu) { | ||
| PRIMUS_TURBO_ERROR("ck_grouped_gemm is unavailable: CK backend not built. " | ||
| "Rebuild with PRIMUS_TURBO_BUILD_BACKEND including 'ck'."); | ||
| } |
Comment on lines
+317
to
+322
| at::Tensor ck_grouped_gemm_variable_k(at::Tensor &a, at::Tensor &b, at::Tensor &group_lens, | ||
| at::Tensor &group_offs, const bool transA, const bool transB, | ||
| c10::optional<int64_t> num_cu) { | ||
| PRIMUS_TURBO_ERROR("ck_grouped_gemm_variable_k is unavailable: CK backend not built. " | ||
| "Rebuild with PRIMUS_TURBO_BUILD_BACKEND including 'ck'."); | ||
| } |
Comment on lines
+324
to
+330
| at::Tensor ck_grouped_gemm_fp8(at::Tensor &a, at::Tensor &b, at::Tensor &a_scales, | ||
| at::Tensor &b_scales, at::Tensor &group_lens, at::Tensor &group_offs, | ||
| const bool transA, const bool transB, at::ScalarType out_dtype, | ||
| const std::string &granularity, c10::optional<int64_t> num_cu) { | ||
| PRIMUS_TURBO_ERROR("ck_grouped_gemm_fp8 is unavailable: CK backend not built. " | ||
| "Rebuild with PRIMUS_TURBO_BUILD_BACKEND including 'ck'."); | ||
| } |
Comment on lines
+332
to
+340
| at::Tensor ck_grouped_gemm_fp8_variable_k(at::Tensor &a, at::Tensor &b, at::Tensor &a_scales, | ||
| at::Tensor &b_scales, at::Tensor &group_lens, | ||
| at::Tensor &group_offs, const bool transA, | ||
| const bool transB, at::ScalarType out_dtype, | ||
| const std::string &granularity, | ||
| c10::optional<int64_t> num_cu) { | ||
| PRIMUS_TURBO_ERROR("ck_grouped_gemm_fp8_variable_k is unavailable: CK backend not built. " | ||
| "Rebuild with PRIMUS_TURBO_BUILD_BACKEND including 'ck'."); | ||
| } |
Comment on lines
+131
to
+132
| def _join_rocm_home(*paths) -> str: | ||
| return os.path.join(find_rocm_home(), *paths) |
Comment on lines
+125
to
+127
| if rocm_home is None: | ||
| logger.warning("No ROCm runtime is found, using ROCM_HOME='%s'", rocm_home) | ||
|
|
Comment on lines
+104
to
+109
| at::Tensor ck_gemm_fp8(at::Tensor &a, at::Tensor &b, at::Tensor &a_scales, at::Tensor &b_scales, | ||
| const bool transA, const bool transB, at::ScalarType out_dtype, | ||
| const std::string &granularity) { | ||
| PRIMUS_TURBO_ERROR("ck_gemm_fp8 is unavailable: CK backend not built. " | ||
| "Rebuild with PRIMUS_TURBO_BUILD_BACKEND including 'ck'."); | ||
| } |
Comment on lines
+88
to
+93
| at::Tensor turbo_gemm_fp8(at::Tensor A, at::Tensor scaleA_inv, at::Tensor B, at::Tensor scaleB_inv, | ||
| const at::ScalarType out_dtype, bool transA, bool transB, bool transC, | ||
| const std::string &granularity) { | ||
| PRIMUS_TURBO_ERROR("turbo_gemm_fp8 is unavailable: turbo backend not built. " | ||
| "Rebuild with PRIMUS_TURBO_BUILD_BACKEND including 'turbo'."); | ||
| } |
Comment on lines
+316
to
+321
| at::Tensor ck_grouped_gemm(at::Tensor &a, at::Tensor &b, at::Tensor &group_lens, | ||
| at::Tensor &group_offs, const bool transA, const bool transB, | ||
| c10::optional<int64_t> num_cu) { | ||
| PRIMUS_TURBO_ERROR("ck_grouped_gemm is unavailable: CK backend not built. " | ||
| "Rebuild with PRIMUS_TURBO_BUILD_BACKEND including 'ck'."); | ||
| } |
Comment on lines
+338
to
+346
| at::Tensor ck_grouped_gemm_fp8_variable_k(at::Tensor &a, at::Tensor &b, at::Tensor &a_scales, | ||
| at::Tensor &b_scales, at::Tensor &group_lens, | ||
| at::Tensor &group_offs, const bool transA, | ||
| const bool transB, at::ScalarType out_dtype, | ||
| const std::string &granularity, | ||
| c10::optional<int64_t> num_cu) { | ||
| PRIMUS_TURBO_ERROR("ck_grouped_gemm_fp8_variable_k is unavailable: CK backend not built. " | ||
| "Rebuild with PRIMUS_TURBO_BUILD_BACKEND including 'ck'."); | ||
| } |
Comment on lines
+13
to
+17
| // gfx12 (gfx1250 / gfx1251, CDNA5) removed the unified `s_waitcnt`; the combined | ||
| // LDS + VMEM wait is expressed via the split counters. | ||
| #if defined(__gfx1250__) | ||
| #define PRIMUS_TURBO_WAIT_ALL_STR \ | ||
| "s_wait_dscnt 0\n\ts_wait_kmcnt 0\n\ts_wait_loadcnt 0\n\ts_wait_storecnt 0" |
Comment on lines
+10
to
+13
| // gfx12 (gfx1250 / gfx1251, CDNA5) removed the unified `s_waitcnt`; LDS waits use | ||
| // the split `s_wait_dscnt` counter instead. | ||
| #if defined(__gfx1250__) | ||
| #define PRIMUS_TURBO_WAIT_DS_STR "s_wait_dscnt 0" |
Comment on lines
18
to
24
| template <int CNT> __device__ __forceinline__ void wait_lgkmcnt() { | ||
| #if defined(__gfx1250__) | ||
| asm volatile("s_wait_dscnt %0" : : "n"(CNT) : "memory"); | ||
| #else | ||
| asm volatile("s_waitcnt lgkmcnt(%0)" : : "n"(CNT) : "memory"); | ||
| #endif | ||
| } |
Comment on lines
26
to
32
| template <int CNT> __device__ __forceinline__ void wait_vmcnt() { | ||
| #if defined(__gfx1250__) | ||
| asm volatile("s_wait_loadcnt %0" : : "n"(CNT) : "memory"); | ||
| #else | ||
| asm volatile("s_waitcnt vmcnt(%0)" : : "n"(CNT) : "memory"); | ||
| #endif | ||
| } |
RuibinCheung
changed the title
RuibinCheung
marked this pull request as draft
June 15, 2026 02:41
RuibinCheung
force-pushed
the
dev/zhangrb/gfx1250_build_support
branch
from
July 15, 2026 08:55
a7a196c to
7d2cf3b
Compare
RuibinCheung
force-pushed
the
dev/zhangrb/gfx1250_build_support
branch
from
July 15, 2026 10:25
7d2cf3b to
7cb6cf5
Compare
RuibinCheung
marked this pull request as ready for review
July 15, 2026 10:39
Comment on lines
+159
to
+160
| if rocm_home and torch.version.hip is None: | ||
| logger.warning("No ROCm runtime is found, using ROCM_HOME='%s'", rocm_home) |
Comment on lines
+164
to
+168
| def _join_rocm_home(*paths) -> str: | ||
| return os.path.join(find_rocm_home(), *paths) | ||
|
|
||
|
|
||
| HIP_HOME: Path = find_hip_home() | ||
| HIP_HOME: Path = _join_rocm_home("hip") |
Comment on lines
+36
to
+41
| const int threads_per_block = 256; | ||
| const int blocks = static_cast<int>((group_num + threads_per_block - 1) / threads_per_block); | ||
|
|
||
| compute_group_offs_device<IndexType><<<blocks, threads_per_block, 0, stream>>>( | ||
| group_lens_ptr, group_offs_ptr, static_cast<int>(group_num)); | ||
| } |
Comment on lines
+25
to
+29
| IndexType cumsum = 0; | ||
| for (int i = 0; i < idx; i++) { | ||
| cumsum += group_lens_ptr[i]; | ||
| } | ||
| group_offs_ptr[idx + 1] = cumsum + group_lens_ptr[idx]; |
Comment on lines
60
to
+62
| PRIMUS_TURBO_CHECK( | ||
| is_gfx950(), | ||
| "The HIPBLASLT_MATMUL_MATRIX_SCALE_VEC32_UE8M0 only support on gfx950."); | ||
| !is_gfx942(), | ||
| "The HIPBLASLT_MATMUL_MATRIX_SCALE_VEC32_UE8M0 not support on gfx942."); |
RuibinCheung
changed the title
Comment on lines
+159
to
+160
| if rocm_home: | ||
| logger.warning("No ROCm runtime is found, using ROCM_HOME='%s'", rocm_home) |
Comment on lines
+164
to
+168
| def _join_rocm_home(*paths) -> str: | ||
| return os.path.join(find_rocm_home(), *paths) | ||
|
|
||
|
|
||
| HIP_HOME: Path = find_hip_home() | ||
| HIP_HOME: Path = _join_rocm_home("hip") |
Comment on lines
+52
to
+59
| # TODO(ruibin): gfx1250 is not yet fully supported by the Primus-Turbo. | ||
| from primus_turbo.pytorch.core.utils import is_gfx1250 | ||
|
|
||
| if is_gfx1250(): | ||
| skip_gfx1250 = pytest.mark.skip(reason="Not yet supported on gfx1250") | ||
| for item in items: | ||
| item.add_marker(skip_gfx1250) | ||
| return |
Comment on lines
+31
to
+34
| # mega_moe_fused (the flydsl mega path) is not supported on gfx1250. Skip the whole | ||
| # module before importing the flydsl-backed modules below | ||
| if is_gfx1250(): | ||
| pytest.skip("mega_moe_fused is not supported on gfx1250", allow_module_level=True) |
Comment on lines
+14
to
+31
| template <typename IndexType> | ||
| __global__ void compute_group_offs_device(const IndexType *group_lens_ptr, | ||
| IndexType *group_offs_ptr, const int group_num) { | ||
| const int64_t idx = blockIdx.x * blockDim.x + threadIdx.x; | ||
|
|
||
| if (idx == 0) { | ||
| group_offs_ptr[0] = 0; | ||
| } | ||
|
|
||
| if (idx < group_num) { | ||
| // Exclusive-prefix-sum to produce group offsets. | ||
| IndexType cumsum = 0; | ||
| for (int i = 0; i < idx; i++) { | ||
| cumsum += group_lens_ptr[i]; | ||
| } | ||
| group_offs_ptr[idx + 1] = cumsum + group_lens_ptr[idx]; | ||
| } | ||
| } |
Comment on lines
+159
to
161
| if rocm_home: | ||
| logger.warning("No ROCm runtime is found, using ROCM_HOME='%s'", rocm_home) | ||
| return rocm_home |
Comment on lines
+164
to
+165
| def _join_rocm_home(*paths) -> str: | ||
| return os.path.join(find_rocm_home(), *paths) |
Comment on lines
+52
to
+59
| # TODO(ruibin): gfx1250 is not yet fully supported by the Primus-Turbo. | ||
| from primus_turbo.pytorch.core.utils import is_gfx1250 | ||
|
|
||
| if is_gfx1250(): | ||
| skip_gfx1250 = pytest.mark.skip(reason="Not yet supported on gfx1250") | ||
| for item in items: | ||
| item.add_marker(skip_gfx1250) | ||
| return |
Comment on lines
+33
to
+41
| template <typename IndexType> | ||
| void compute_group_offs(const IndexType *group_lens_ptr, IndexType *group_offs_ptr, | ||
| const int64_t group_num, hipStream_t stream) { | ||
| const int threads_per_block = 256; | ||
| const int blocks = static_cast<int>((group_num + threads_per_block - 1) / threads_per_block); | ||
|
|
||
| compute_group_offs_device<IndexType><<<blocks, threads_per_block, 0, stream>>>( | ||
| group_lens_ptr, group_offs_ptr, static_cast<int>(group_num)); | ||
| } |
Comment on lines
+24
to
+28
| // Exclusive-prefix-sum to produce group offsets. | ||
| IndexType cumsum = 0; | ||
| for (int i = 0; i < idx; i++) { | ||
| cumsum += group_lens_ptr[i]; | ||
| } |
RuibinCheung
changed the title
Comment on lines
+164
to
+165
| def _join_rocm_home(*paths) -> str: | ||
| return os.path.join(find_rocm_home(), *paths) |
Comment on lines
+159
to
+160
| if rocm_home: | ||
| logger.warning("No ROCm runtime is found, using ROCM_HOME='%s'", rocm_home) |
Comment on lines
+233
to
+238
| # TODO(ruibin): add CK and turbo support for gfx1250 | ||
| if "gfx1250" in arch_list: | ||
| if build_ck_backend() or build_turbo_backend(): | ||
| raise ValueError( | ||
| "The CK or turbo backend is not supported on gfx1250. Please set PRIMUS_TURBO_BUILD_BACKEND='' to disable build the CK and turbo backends." | ||
| ) |
Comment on lines
+14
to
+31
| template <typename IndexType> | ||
| __global__ void compute_group_offs_device(const IndexType *group_lens_ptr, | ||
| IndexType *group_offs_ptr, const int group_num) { | ||
| const int64_t idx = blockIdx.x * blockDim.x + threadIdx.x; | ||
|
|
||
| if (idx == 0) { | ||
| group_offs_ptr[0] = 0; | ||
| } | ||
|
|
||
| if (idx < group_num) { | ||
| // Exclusive-prefix-sum to produce group offsets. | ||
| IndexType cumsum = 0; | ||
| for (int i = 0; i < idx; i++) { | ||
| cumsum += group_lens_ptr[i]; | ||
| } | ||
| group_offs_ptr[idx + 1] = cumsum + group_lens_ptr[idx]; | ||
| } | ||
| } |
Comment on lines
+33
to
+41
| template <typename IndexType> | ||
| void compute_group_offs(const IndexType *group_lens_ptr, IndexType *group_offs_ptr, | ||
| const int64_t group_num, hipStream_t stream) { | ||
| const int threads_per_block = 256; | ||
| const int blocks = static_cast<int>((group_num + threads_per_block - 1) / threads_per_block); | ||
|
|
||
| compute_group_offs_device<IndexType><<<blocks, threads_per_block, 0, stream>>>( | ||
| group_lens_ptr, group_offs_ptr, static_cast<int>(group_num)); | ||
| } |
Comment on lines
59
to
63
| if (scale_mode == HIPBLASLT_MATMUL_MATRIX_SCALE_VEC32_UE8M0) { | ||
| PRIMUS_TURBO_CHECK( | ||
| is_gfx950(), | ||
| "The HIPBLASLT_MATMUL_MATRIX_SCALE_VEC32_UE8M0 only support on gfx950."); | ||
| !is_gfx942(), | ||
| "The HIPBLASLT_MATMUL_MATRIX_SCALE_VEC32_UE8M0 not support on gfx942."); | ||
| } |
zhenhuang12
self-requested a review
July 17, 2026 05:28
Comment on lines
+159
to
161
| if rocm_home: | ||
| logger.warning("No ROCm runtime is found, using ROCM_HOME='%s'", rocm_home) | ||
| return rocm_home |
Comment on lines
+164
to
+165
| def _join_rocm_home(*paths) -> str: | ||
| return os.path.join(find_rocm_home(), *paths) |
Comment on lines
+40
to
+46
| // Device pass: rely on the per-arch predefined macros so fat-binary builds | ||
| // (multiple --offload-arch) still get the correct value for each arch. | ||
| #if defined(__gfx942__) || defined(__gfx950__) | ||
| constexpr int THREADS_PER_WARP = 64; | ||
| #else | ||
| constexpr int THREADS_PER_WARP = 32; | ||
| #endif |
Comment on lines
+23
to
+27
| if (idx < group_num) { | ||
| // Exclusive-prefix-sum to produce group offsets. | ||
| IndexType cumsum = 0; | ||
| for (int i = 0; i < idx; i++) { | ||
| cumsum += group_lens_ptr[i]; |
Comment on lines
60
to
+62
| PRIMUS_TURBO_CHECK( | ||
| is_gfx950(), | ||
| "The HIPBLASLT_MATMUL_MATRIX_SCALE_VEC32_UE8M0 only support on gfx950."); | ||
| !is_gfx942(), | ||
| "The HIPBLASLT_MATMUL_MATRIX_SCALE_VEC32_UE8M0 not support on gfx942."); |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Description
This PR is the first part of enabling Primus-Turbo on the AMD
gfx1250(CDNA5) architecture.It teaches the whole stack — build system, host bindings, and device kernels — to recognize
gfx1250and to cope with the two hardware differences that break the existinggfx942/gfx950assumptions: a 32-lane wavefront (instead of 64) and the new splits_wait_*counters that replace the unifieds_waitcnt.Because the CK (
ck_tile) and turbo (MFMA) GEMM backends, Triton/flydsl, and several quantization kernels are not yet available ongfx1250, this part also makes those backends optional and selectable at build time so that agfx1250build can be produced today while the remaining kernel support lands in follow-up PRs.Fixes # (issue)
Type of change
Changes
Please list the changes introduced in this PR:
Add
gfx1250architecture detection on both host and device.GPUArch::GFX1250,is_gfx1250(), and awarp_size()helper incsrc/include/primus_turbo/arch.h.is_gfx1250()inprimus_turbo/pytorch/core/utils.py.gfx1250toSUPPORTED_GPU_ARCHSand the--offload-archfiltering pipeline insetup.py.Handle the 32-lane wavefront on
gfx1250.THREADS_PER_WARParch-aware inplatform.h(64 forgfx942/gfx950, 32 otherwise) so fat-binary builds get the correct value per arch.THREADS_PER_WARPusages with the runtimewarp_size()in the reduce and normalization host launchers.warp_reduce_max_64_dppwhen the wavefront is already 32 lanes.Support the
gfx12split wait counters.s_wait_dscnt/s_wait_loadcntinstead ofs_waitcnt lgkmcnt/vmcntongfx1250indevice/memory.cuhanddevice/reduce.cuh.Clarify FP8 encoding selection by architecture.
float8.hsogfx942uses FNUZ and all other archs (gfx950,gfx1250) use OCP, instead of special-casing onlygfx950.Make the CK and turbo GEMM backends optional and build-time selectable.
PRIMUS_TURBO_BUILD_BACKENDenv var (comma/semicolon separated, e.g.ck,turbo) parsed byenabled_backends()/build_ck_backend()/build_turbo_backend()intools/build_utils.py; unset builds all backends, empty string builds none.-DBUILD_CK_BACKEND/-DBUILD_TURBO_BACKENDonly for enabled backends, guard the CK/turbo declarations and host bindings behind them, and fall back to loud runtime-error stubs when a backend is disabled.primus_turbo/_build_info.py, and exposebuild_ck()/build_turbo()incore/utils.pyso kernel dispatchers only select a backend that was actually compiled in.Add a standalone, CK-free
compute_group_offskernel (csrc/kernels/grouped_gemm/compute_group_offs.cu) that is always built so the symbol is available to both CK and non-CK host bindings.Guard unsupported paths on
gfx1250.gfx1250.triton==3.7.0andflydsl>=0.2.0when building forgfx1250.gfx1250(kernel support pending).mega_moe_fusedtest module ongfx1250.Improve the build tooling.
find_rocm_home()to honorROCM_HOME/ROCM_PATHand support TheRockrocm-sdk-core, and deriveHIP_HOMEfrom it.ninjaavailability before building extensions and remove dead debug code intools/build_ext.py._build_info.pywhen git is unavailable.Checklist: