Skip to content

feat: support build on gfx1250 (Part2) #423

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
xiaobochen-amd merged 1 commit into from
Jul 21, 2026
Merged

feat: support build on gfx1250 (Part2) #423

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
49 changes: 4 additions & 45 deletions csrc/kernels/gemm/turbo_gemm.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,11 +3,8 @@
// See LICENSE for license information.

#include "primus_turbo/gemm.h"
#include <hip/hip_runtime.h>

#ifdef BUILD_TURBO_BACKEND

#include "turbo/turbo_gemm_mxfp8_kernel.h"
#include <hip/hip_runtime.h>
Comment thread
RuibinCheung marked this conversation as resolved.

namespace primus_turbo {

Expand All @@ -29,6 +26,9 @@ void turbo_gemm_mxfp8_impl(const AType *a_ptr, const BType *b_ptr,
const dtype::float8_e8m0 *b_scale_ptr, CType *c_ptr, int32_t m,
int32_t n, int32_t k, void *workspace, size_t workspace_size,
hipStream_t stream) {
PRIMUS_TURBO_CHECK(!primus_turbo::is_gfx1250(),
"turbo_gemm_fp8 is unavailable: the turbo backend is not built "
"on this architecture (gfx1250 is unsupported).");
constexpr int32_t MX_BLOCK_SIZE = 32;
const int32_t scale_cols = (k + MX_BLOCK_SIZE - 1) / MX_BLOCK_SIZE;
const size_t a_scale_bytes = (size_t) m * scale_cols * sizeof(uint32_t);
Expand Down Expand Up @@ -73,44 +73,3 @@ INSTANTIATE_TURBO_GEMM(dtype::float8_e5m2, dtype::float8_e4m3, dtype::bfloat16)
#undef INSTANTIATE_TURBO_GEMM

} // namespace primus_turbo

#else // !BUILD_TURBO_BACKEND : turbo (MFMA) GEMM not compiled on gfx1250.

#include "primus_turbo/macros.h"

namespace primus_turbo {

// Stub symbols so host bindings link uniformly. The real kernels are unavailable on
// gfx1250; callers must gate on is_gfx1250() before invoking the impl.

size_t turbo_gemm_mxfp8_workspace_size(int32_t, int32_t, int32_t) {
return 0;
}

template <typename AType, typename BType, typename CType>
void turbo_gemm_mxfp8_impl(const AType *, const BType *, const dtype::float8_e8m0 *,
const dtype::float8_e8m0 *, CType *, int32_t, int32_t, int32_t, void *,
size_t, hipStream_t) {
PRIMUS_TURBO_ERROR("turbo_gemm_mxfp8 is unavailable: the turbo backend is not built "
"on this architecture (gfx1250 is unsupported).");
}

#define INSTANTIATE_TURBO_GEMM(A, B, C) \
template void turbo_gemm_mxfp8_impl<A, B, C>(const A *, const B *, const dtype::float8_e8m0 *, \
const dtype::float8_e8m0 *, C *, int32_t, \
int32_t, int32_t, void *, size_t, hipStream_t);

INSTANTIATE_TURBO_GEMM(dtype::float8_e4m3, dtype::float8_e4m3, dtype::float16)
INSTANTIATE_TURBO_GEMM(dtype::float8_e4m3, dtype::float8_e4m3, dtype::bfloat16)
INSTANTIATE_TURBO_GEMM(dtype::float8_e5m2, dtype::float8_e5m2, dtype::float16)
INSTANTIATE_TURBO_GEMM(dtype::float8_e5m2, dtype::float8_e5m2, dtype::bfloat16)
INSTANTIATE_TURBO_GEMM(dtype::float8_e4m3, dtype::float8_e5m2, dtype::float16)
INSTANTIATE_TURBO_GEMM(dtype::float8_e4m3, dtype::float8_e5m2, dtype::bfloat16)
INSTANTIATE_TURBO_GEMM(dtype::float8_e5m2, dtype::float8_e4m3, dtype::float16)
INSTANTIATE_TURBO_GEMM(dtype::float8_e5m2, dtype::float8_e4m3, dtype::bfloat16)

#undef INSTANTIATE_TURBO_GEMM

} // namespace primus_turbo

#endif // BUILD_TURBO_BACKEND
14 changes: 10 additions & 4 deletions csrc/kernels/quantization/dequantization_mxfp4.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -46,6 +46,16 @@ __device__ __forceinline__ void cvt_fp4x2_scaled(uint32_t fp4, float scale, OTyp
: "v"(fp4), "v"(scale));
*reinterpret_cast<uint64_t *>(out) = r;
}
#elif defined(__gfx1250__)
// gfx1250 (CDNA5) only exposes the PK8 scale converter for de-quantization, and
// its E8M0 scale is selected cross-lane (lanes 16..31 read lanes 0..15). To keep
// each thread's scale independent we run the converter as a pure FP4->F32 unpack
// (scale byte = 0x7f == 2^0, opsel 0) and apply the per-thread ``scale`` in
// registers. Only the low 2 nibbles of ``fp4`` are meaningful here.
typedef float float32x8_t __attribute__((ext_vector_type(8)));
const float32x8_t r = __builtin_amdgcn_cvt_scale_pk8_f32_fp4(fp4, 0x7f7f7f7fu, 0);
out[0] = static_cast<OType>(r[0] * scale);
out[1] = static_cast<OType>(r[1] * scale);
#else
__builtin_trap();
#endif
Expand Down Expand Up @@ -154,10 +164,6 @@ void dequantize_mxfp4_impl(const uint8_t *x, OType *y, const int64_t stride_x_ro
const int64_t stride_scale_col, const int scale_n_rows,
const int scale_n_cols, const int block_size, const bool use_rowwise,
hipStream_t stream) {
// TODO(ruibin): add kernel support for gfx1250
if (is_gfx1250()) {
PRIMUS_TURBO_ERROR("dequantize_mxfp4 is not implemented on gfx1250.");
}
(void) stride_x_col; // packed input is contiguous along columns (stride == 1)
if (n_rows == 0 || n_cols == 0)
return;
Expand Down
8 changes: 0 additions & 8 deletions csrc/kernels/quantization/dequantization_mxfp8.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -258,10 +258,6 @@ void grouped_dequantize_mxfp8_impl(const QType *x, OType *y, const int64_t strid
const int scale_n_rows, const int scale_n_cols,
const int64_t *group_offs, const int64_t *group_offs_padded,
int G, int block_size, bool use_rowwise, hipStream_t stream) {
// TODO(ruibin): add kernel support for gfx1250
if (is_gfx1250()) {
PRIMUS_TURBO_ERROR("grouped_dequantize_mxfp8 is not implemented on gfx1250.");
}
if (total_M == 0 || n_cols == 0)
return;

Expand Down Expand Up @@ -289,10 +285,6 @@ void dequantize_mxfp8_impl(const QType *x, OType *y, const int64_t stride_x_row,
const int64_t stride_scale_col, const int scale_n_rows,
const int scale_n_cols, const int block_size, const bool use_rowwise,
hipStream_t stream) {
// TODO(ruibin): add kernel support for gfx1250
if (is_gfx1250()) {
PRIMUS_TURBO_ERROR("dequantize_mxfp8 is not implemented on gfx1250.");
}
(void) stride_x_col; // input is contiguous along columns (stride == 1)
if (n_rows == 0 || n_cols == 0)
return;
Expand Down
Loading
Loading