Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
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
191 changes: 191 additions & 0 deletions benchmark/ops/bench_hipblaslt_algo_tuning.py

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This benchmark file can be removed. We should try to consolidate everything into the existing bench_gemm_turbo.py as much as possible.

Original file line number Diff line number Diff line change
@@ -0,0 +1,191 @@
###############################################################################
# Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
#
# See LICENSE for license information.
###############################################################################

"""hipBLASLt algorithm tuning benchmark.

Runs each dense-model GEMM shape (MBS=1) with level-2 AutoTune so
that the framework benchmarks all hipBLASLt heuristic algorithms and
prints the winner for each shape. The tuning logs come from the
AutoKernelDispatcher logger.

By default all dtype/granularity combinations are tested. Use --dtype
and --granularity to restrict to a single combination.

Usage:
python benchmark/ops/bench_hipblaslt_algo_tuning.py # all combos
python benchmark/ops/bench_hipblaslt_algo_tuning.py --dtype bf16 # bf16 only
python benchmark/ops/bench_hipblaslt_algo_tuning.py --dtype fp8 --granularity tensorwise
python benchmark/ops/bench_hipblaslt_algo_tuning.py --max-algos 5
"""

import argparse
import os
import sys
from datetime import datetime

os.environ.setdefault("PRIMUS_TURBO_LOG_LEVEL", "INFO")

import torch # noqa: E402
from config import ( # noqa: E402
DenseModelConfigs,
gen_gemm_test_cases,
get_platform_info,
)

import primus_turbo.pytorch as turbo # noqa: E402
from primus_turbo.pytorch.core.backend import GlobalBackendManager # noqa: E402
from primus_turbo.pytorch.core.low_precision import ( # noqa: E402
Float8QuantConfig,
Format,
ScaleDtype,
ScalingGranularity,
)

FP8_GRANULARITY_CONFIGS = {
"tensorwise": Float8QuantConfig(format=Format.E4M3, granularity=ScalingGranularity.TENSORWISE),
"rowwise": Float8QuantConfig(format=Format.E4M3, granularity=ScalingGranularity.ROWWISE),
"blockwise": Float8QuantConfig(
format=Format.E4M3,
granularity=ScalingGranularity.BLOCKWISE,
block_size=128,
),
}


def run_shapes(dtype_label, device, fp8_config=None):
"""Run all dense-model GEMM shapes for a given dtype/granularity."""
for model_name, model_config in DenseModelConfigs.items():
test_cases = gen_gemm_test_cases(model_config)
for shape in test_cases:
M, N, K = shape[0], shape[1], shape[2]

print(f"{'='*60}")
print(f"Case: {model_name}, M={M}, N={N}, K={K}, dtype={dtype_label}")
print(f"{'='*60}")

a = torch.randn((M, K), dtype=torch.bfloat16, device=device)
b = torch.randn((N, K), dtype=torch.bfloat16, device=device)

if fp8_config is not None:
turbo.ops.gemm_fp8(a, b, trans_b=True, config=fp8_config)
turbo.ops.gemm_fp8(a, b, trans_b=True, config=fp8_config)
else:
turbo.ops.gemm(a, b, trans_b=True)
turbo.ops.gemm(a, b, trans_b=True)

torch.cuda.synchronize()
print()


def main():
parser = argparse.ArgumentParser(description="hipBLASLt Algorithm Tuning Benchmark")
parser.add_argument(
"--dtype",
type=str,
choices=["bf16", "fp8", "all"],
default="all",
help="Data type to test (default: all)",
)
parser.add_argument(
"--granularity",
type=str,
choices=["tensorwise", "rowwise", "blockwise", "all"],
default="all",
help="FP8 scaling granularity (only used when dtype includes fp8, default: all)",
)
parser.add_argument(
"--max-algos",
type=int,
default=100,
help="Max hipBLASLt algorithms to try per shape (default: 10)",
)
parser.add_argument(
"--output",
"-o",
type=str,
default=None,
help="Output log file (default: auto-generated name). Use 'stdout' to skip file output.",
)
args = parser.parse_args()

os.environ["PRIMUS_TURBO_HIPBLASLT_TUNE_MAX_ALGOS"] = str(args.max_algos)
GlobalBackendManager.set_auto_tune(2)

platform, gpu_name = get_platform_info()
device = "cuda"

log_file = None
if args.output != "stdout":
filename = args.output or (
f"hipblaslt_algo_tuning_{args.dtype}_{args.granularity}"
f"_{datetime.now().strftime('%Y%m%d_%H%M%S')}_{gpu_name}.log"
)
log_file = open(filename, "w")

# Tee output to both stdout and file
class Tee:
def __init__(self, *streams):
self.streams = streams

def write(self, data):
for s in self.streams:
s.write(data)
s.flush()

def flush(self):
for s in self.streams:
s.flush()

sys.stdout = Tee(sys.__stdout__, log_file)
sys.stderr = Tee(sys.__stderr__, log_file)

# Point existing logger handlers at the new stderr so their
# output is captured in the log file too.
import logging

for handler in logging.getLogger("primus_turbo").handlers:
if isinstance(handler, logging.StreamHandler):
handler.stream = sys.stdout

print(f"Platform: {platform}, GPU: {gpu_name}")
print(f"AutoTune level: 2, hipBLASLt algo cap: {args.max_algos}")
print(f"Dtype: {args.dtype}, Granularity: {args.granularity}\n")

run_bf16 = args.dtype in ("bf16", "all")
run_fp8 = args.dtype in ("fp8", "all")

if run_bf16:
print(f"\n{'#'*60}")
print(f"# BF16 GEMM")
print(f"{'#'*60}\n")
run_shapes("bf16", device)
GlobalBackendManager.reset()
GlobalBackendManager.set_auto_tune(2)

if run_fp8:
granularities = (
list(FP8_GRANULARITY_CONFIGS.keys()) if args.granularity == "all" else [args.granularity]
)
for gran in granularities:
print(f"\n{'#'*60}")
print(f"# FP8 GEMM ({gran})")
print(f"{'#'*60}\n")
run_shapes(f"fp8-{gran}", device, fp8_config=FP8_GRANULARITY_CONFIGS[gran])
GlobalBackendManager.reset()
GlobalBackendManager.set_auto_tune(2)

GlobalBackendManager.reset()
print("Done.")

if log_file is not None:
sys.stdout = sys.__stdout__
sys.stderr = sys.__stderr__
log_file.close()
print(f"Log saved to {filename}")


if __name__ == "__main__":
main()
7 changes: 7 additions & 0 deletions benchmark/ops/benchmark_suite.yaml
Original file line number Diff line number Diff line change
Expand Up @@ -281,6 +281,13 @@ tasks:
env: {PRIMUS_TURBO_GROUPED_GEMM_BACKEND: TRITON}
output: grouped_gemm_fp8_blockwise_triton_benchmark.csv

# ── Online Tuning (level 2) ──────────────────────────────
- label: online_tuning_all
group: online_tuning
script: bench_hipblaslt_algo_tuning.py
args: [--max-algos, 100]
output: online_tuning_all_benchmark.log

# ── DeepEP (multi-GPU, placed last) ──────────────────────
- label: deep_ep_intranode
group: deepep
Expand Down
15 changes: 14 additions & 1 deletion csrc/include/primus_turbo/gemm.h
Original file line number Diff line number Diff line change
Expand Up @@ -42,6 +42,7 @@ inline size_t hipblaslt_dtype_bytes(hipDataType dtype) {

int64_t get_hipblaslt_workspace_size_in_byte();

// Run a GEMM using the heuristic algorithm at algo_index (0 = top-ranked).
void hipblaslt_gemm_impl(const void *A, const hipDataType A_type, const int64_t rows_a,
const int64_t cols_a, const int64_t lda, const void *scaleA_inv,
hipblasOperation_t transA, const void *B, const hipDataType B_type,
Expand All @@ -51,7 +52,19 @@ void hipblaslt_gemm_impl(const void *A, const hipDataType A_type, const int64_t
const int64_t ldd, void *workspace, const int64_t workspace_size,
const bool use_low_precision,
const hipblasLtMatmulMatrixScale_t scale_mode, hipblasLtHandle_t handle,
hipStream_t stream);
hipStream_t stream, const int algo_index = 0);

// Query how many heuristic algorithms are available for this problem.
int hipblaslt_gemm_get_algo_count(const void *A, const hipDataType A_type, const int64_t rows_a,
const int64_t cols_a, const int64_t lda, const void *scaleA_inv,
hipblasOperation_t transA, const void *B,
const hipDataType B_type, const int64_t rows_b,
const int64_t cols_b, const int64_t ldb, const void *scaleB_inv,
hipblasOperation_t transB, const hipDataType D_type,
const int64_t rows_d, const int64_t cols_d, const int64_t ldd,
const int64_t workspace_size, const bool use_low_precision,
hipblasLtMatmulMatrixScale_t scale_mode,
hipblasLtHandle_t handle);

//==================================================================
// CK GEMM
Expand Down
79 changes: 74 additions & 5 deletions csrc/kernels/gemm/hipblaslt_gemm.cu
Original file line number Diff line number Diff line change
Expand Up @@ -26,7 +26,7 @@ void hipblaslt_gemm_impl(const void *A, const hipDataType A_type, const int64_t
const hipDataType D_type, const int64_t rows_d, const int64_t cols_d,
const int64_t ldd, void *workspace, const int64_t workspace_size,
const bool use_low_precision, hipblasLtMatmulMatrixScale_t scale_mode,
hipblasLtHandle_t handle, hipStream_t stream) {
hipblasLtHandle_t handle, hipStream_t stream, const int algo_index) {
hipblasLtMatmulDesc_t operation_desc = nullptr;
hipblasLtMatrixLayout_t A_desc = nullptr, B_desc = nullptr, D_desc = nullptr;
hipblasLtMatmulPreference_t preference = nullptr;
Expand Down Expand Up @@ -69,7 +69,7 @@ void hipblaslt_gemm_impl(const void *A, const hipDataType A_type, const int64_t
operation_desc, scaleB_inv_ptr_desc, &scaleB_inv, sizeof(scaleB_inv)));
}

const int request_solutions = 1;
const int request_solutions = algo_index + 1;
std::vector<hipblasLtMatmulHeuristicResult_t> algos(request_solutions);
int returnedAlgoCount = 0;

Expand All @@ -81,8 +81,8 @@ void hipblaslt_gemm_impl(const void *A, const hipDataType A_type, const int64_t
PRIMUS_TURBO_CHECK_HIPBLAS(hipblasLtMatmulAlgoGetHeuristic(
handle, operation_desc, A_desc, B_desc, D_desc, D_desc, preference, request_solutions,
algos.data(), &returnedAlgoCount));
PRIMUS_TURBO_CHECK(returnedAlgoCount > 0,
"hipBLASLt: no valid algorithm found for current matmul config");
PRIMUS_TURBO_CHECK(returnedAlgoCount > algo_index,
"hipBLASLt: requested algo_index exceeds available algorithms");

const float alpha = 1.0;
const float beta = 0.0;
Expand All @@ -96,7 +96,7 @@ void hipblaslt_gemm_impl(const void *A, const hipDataType A_type, const int64_t
&beta,
D, D_desc,
D, D_desc,
&algos[0].algo,
&algos[algo_index].algo,
workspace, workspace_size,
stream));
// clang-format on
Expand All @@ -108,4 +108,73 @@ void hipblaslt_gemm_impl(const void *A, const hipDataType A_type, const int64_t
PRIMUS_TURBO_CHECK_HIPBLAS(hipblasLtMatmulPreferenceDestroy(preference));
}

int hipblaslt_gemm_get_algo_count(const void *A, const hipDataType A_type, const int64_t rows_a,
const int64_t cols_a, const int64_t lda, const void *scaleA_inv,
hipblasOperation_t transA, const void *B,
const hipDataType B_type, const int64_t rows_b,
const int64_t cols_b, const int64_t ldb, const void *scaleB_inv,
hipblasOperation_t transB, const hipDataType D_type,
const int64_t rows_d, const int64_t cols_d, const int64_t ldd,
const int64_t workspace_size, const bool use_low_precision,
hipblasLtMatmulMatrixScale_t scale_mode,
hipblasLtHandle_t handle) {
hipblasLtMatmulDesc_t operation_desc = nullptr;
hipblasLtMatrixLayout_t A_desc = nullptr, B_desc = nullptr, D_desc = nullptr;
hipblasLtMatmulPreference_t preference = nullptr;
hipblasLtEpilogue_t epilogue = HIPBLASLT_EPILOGUE_DEFAULT;
hipblasComputeType_t gemm_compute_type = HIPBLAS_COMPUTE_32F;

PRIMUS_TURBO_CHECK_HIPBLAS(hipblasLtMatrixLayoutCreate(&A_desc, A_type, rows_a, cols_a, lda));
PRIMUS_TURBO_CHECK_HIPBLAS(hipblasLtMatrixLayoutCreate(&B_desc, B_type, rows_b, cols_b, ldb));
PRIMUS_TURBO_CHECK_HIPBLAS(hipblasLtMatrixLayoutCreate(&D_desc, D_type, rows_d, cols_d, ldd));

PRIMUS_TURBO_CHECK_HIPBLAS(
hipblasLtMatmulDescCreate(&operation_desc, gemm_compute_type, HIP_R_32F));
PRIMUS_TURBO_CHECK_HIPBLAS(hipblasLtMatmulDescSetAttribute(
operation_desc, HIPBLASLT_MATMUL_DESC_TRANSA, &transA, sizeof(transA)));
PRIMUS_TURBO_CHECK_HIPBLAS(hipblasLtMatmulDescSetAttribute(
operation_desc, HIPBLASLT_MATMUL_DESC_TRANSB, &transB, sizeof(transB)));
PRIMUS_TURBO_CHECK_HIPBLAS(hipblasLtMatmulDescSetAttribute(
operation_desc, HIPBLASLT_MATMUL_DESC_EPILOGUE, &epilogue, sizeof(epilogue)));

if (use_low_precision) {
PRIMUS_TURBO_CHECK(scaleA_inv != nullptr);
PRIMUS_TURBO_CHECK(scaleB_inv != nullptr);

hipblasLtMatmulDescAttributes_t scaleA_inv_ptr_desc = HIPBLASLT_MATMUL_DESC_A_SCALE_POINTER;
hipblasLtMatmulDescAttributes_t scaleB_inv_ptr_desc = HIPBLASLT_MATMUL_DESC_B_SCALE_POINTER;

PRIMUS_TURBO_CHECK_HIPBLAS(hipblasLtMatmulDescSetAttribute(
operation_desc, HIPBLASLT_MATMUL_DESC_A_SCALE_MODE, &scale_mode, sizeof(scale_mode)));
PRIMUS_TURBO_CHECK_HIPBLAS(hipblasLtMatmulDescSetAttribute(
operation_desc, HIPBLASLT_MATMUL_DESC_B_SCALE_MODE, &scale_mode, sizeof(scale_mode)));

PRIMUS_TURBO_CHECK_HIPBLAS(hipblasLtMatmulDescSetAttribute(
operation_desc, scaleA_inv_ptr_desc, &scaleA_inv, sizeof(scaleA_inv)));
PRIMUS_TURBO_CHECK_HIPBLAS(hipblasLtMatmulDescSetAttribute(
operation_desc, scaleB_inv_ptr_desc, &scaleB_inv, sizeof(scaleB_inv)));
}

const int max_algos = 32768;
std::vector<hipblasLtMatmulHeuristicResult_t> algos(max_algos);
int returnedAlgoCount = 0;

PRIMUS_TURBO_CHECK_HIPBLAS(hipblasLtMatmulPreferenceCreate(&preference));
PRIMUS_TURBO_CHECK_HIPBLAS(
hipblasLtMatmulPreferenceSetAttribute(preference, HIPBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES,
&workspace_size, sizeof(workspace_size)));

PRIMUS_TURBO_CHECK_HIPBLAS(
hipblasLtMatmulAlgoGetHeuristic(handle, operation_desc, A_desc, B_desc, D_desc, D_desc,
preference, max_algos, algos.data(), &returnedAlgoCount));

PRIMUS_TURBO_CHECK_HIPBLAS(hipblasLtMatrixLayoutDestroy(D_desc));
PRIMUS_TURBO_CHECK_HIPBLAS(hipblasLtMatrixLayoutDestroy(B_desc));
PRIMUS_TURBO_CHECK_HIPBLAS(hipblasLtMatrixLayoutDestroy(A_desc));
PRIMUS_TURBO_CHECK_HIPBLAS(hipblasLtMatmulDescDestroy(operation_desc));
PRIMUS_TURBO_CHECK_HIPBLAS(hipblasLtMatmulPreferenceDestroy(preference));

return returnedAlgoCount;
}

} // namespace primus_turbo
24 changes: 17 additions & 7 deletions csrc/pytorch/bindings_pytorch.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -12,14 +12,21 @@ namespace primus_turbo::pytorch {

TORCH_LIBRARY(primus_turbo_cpp_extension, m) {
// ********* Gemm *********
m.def("hipblaslt_gemm(Tensor A, Tensor B, "
"ScalarType out_dtype, bool transA, bool transB, bool transC) -> Tensor");
m.def(
"hipblaslt_gemm_fp8(Tensor A, Tensor scaleA_inv, Tensor B, Tensor scaleB_inv,"
"ScalarType out_dtype, bool transA, bool transB, bool transC, str granularity) -> Tensor");
m.def(
"hipblaslt_gemm_fp4(Tensor A, Tensor scaleA_inv, Tensor B, Tensor scaleB_inv,"
"ScalarType out_dtype, bool transA, bool transB, bool transC, str granularity) -> Tensor");
"hipblaslt_gemm(Tensor A, Tensor B, "
"ScalarType out_dtype, bool transA, bool transB, bool transC, int algo_index=0) -> Tensor");
m.def("hipblaslt_gemm_fp8(Tensor A, Tensor scaleA_inv, Tensor B, Tensor scaleB_inv,"
"ScalarType out_dtype, bool transA, bool transB, bool transC, str granularity, int "
"algo_index=0) -> Tensor");
m.def("hipblaslt_gemm_fp4(Tensor A, Tensor scaleA_inv, Tensor B, Tensor scaleB_inv,"
"ScalarType out_dtype, bool transA, bool transB, bool transC, str granularity, int "
"algo_index=0) -> Tensor");
m.def("hipblaslt_gemm_algo_count(Tensor A, Tensor B, "
"ScalarType out_dtype, bool transA, bool transB, bool transC) -> int");
m.def("hipblaslt_gemm_fp8_algo_count(Tensor A, Tensor scaleA_inv, Tensor B, Tensor scaleB_inv,"
"ScalarType out_dtype, bool transA, bool transB, bool transC, str granularity) -> int");
m.def("hipblaslt_gemm_fp4_algo_count(Tensor A, Tensor scaleA_inv, Tensor B, Tensor scaleB_inv,"
"ScalarType out_dtype, bool transA, bool transB, bool transC, str granularity) -> int");
m.def("ck_gemm_fp8(Tensor a, Tensor b, Tensor a_scales, Tensor b_scales, bool transA,"
"bool transB, ScalarType out_dtype, str granularity) -> Tensor");

Expand Down Expand Up @@ -86,6 +93,9 @@ TORCH_LIBRARY_IMPL(primus_turbo_cpp_extension, CUDA, m) {
m.impl("hipblaslt_gemm", hipblaslt_gemm);
m.impl("hipblaslt_gemm_fp8", hipblaslt_gemm_fp8);
m.impl("hipblaslt_gemm_fp4", hipblaslt_gemm_fp4);
m.impl("hipblaslt_gemm_algo_count", hipblaslt_gemm_algo_count);
m.impl("hipblaslt_gemm_fp8_algo_count", hipblaslt_gemm_fp8_algo_count);
m.impl("hipblaslt_gemm_fp4_algo_count", hipblaslt_gemm_fp4_algo_count);
m.impl("ck_gemm_fp8", ck_gemm_fp8);
m.impl("turbo_gemm_fp8", turbo_gemm_fp8);
// ********* Quantization *********
Expand Down
Loading
Loading