From aaf60563fc821d090e29c6e6cccc6cd488b5b695 Mon Sep 17 00:00:00 2001 From: Seung-Hee Bae Date: Tue, 28 Apr 2026 22:19:01 +0000 Subject: [PATCH 01/27] minor fix of the type of an internal function name. --- projects/hipblaslt/library/src/amd_detail/hipblaslt.cpp | 2 +- .../amd_detail/rocblaslt/include/rocblaslt-auxiliary.h | 8 ++++---- .../library/src/amd_detail/rocblaslt/src/include/handle.h | 2 +- .../src/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp | 4 ++-- 4 files changed, 8 insertions(+), 8 deletions(-) diff --git a/projects/hipblaslt/library/src/amd_detail/hipblaslt.cpp b/projects/hipblaslt/library/src/amd_detail/hipblaslt.cpp index 7191a2a2bd13..eed7a873d66a 100644 --- a/projects/hipblaslt/library/src/amd_detail/hipblaslt.cpp +++ b/projects/hipblaslt/library/src/amd_detail/hipblaslt.cpp @@ -262,7 +262,7 @@ try { rocblaslt::Debug::Instance().markerStart("hipblasLtMatrixLayoutDestroy"); auto status = RocBlasLtStatusToHIPStatus( - rocblaslt_matrix_layout_destory((const rocblaslt_matrix_layout)descr)); + rocblaslt_matrix_layout_destroy((const rocblaslt_matrix_layout)descr)); rocblaslt::Debug::Instance().markerStop(); return status; } diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/include/rocblaslt-auxiliary.h b/projects/hipblaslt/library/src/amd_detail/rocblaslt/include/rocblaslt-auxiliary.h index ffa3191ef4c8..796fc86a2c6d 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/include/rocblaslt-auxiliary.h +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/include/rocblaslt-auxiliary.h @@ -117,7 +117,7 @@ rocblaslt_status rocblaslt_get_sm_count_target(rocblaslt_handle handle, * \brief Create a descriptor for matrix * \details * \p rocblaslt_matrix_layout_create creates a matrix descriptor It initializes - * It should be destroyed at the end using rocblaslt_matrix_layout_destory(). + * It should be destroyed at the end using rocblaslt_matrix_layout_destroy(). * * @param[out] * matDescr the pointer to the matrix descriptor @@ -136,7 +136,7 @@ rocblaslt_status rocblaslt_matrix_layout_create(rocblaslt_matrix_layout* matDesc * \brief Destroy a matrix descriptor * * \details - * \p rocblaslt_matrix_layout_destory destroys a matrix descriptor and releases + * \p rocblaslt_matrix_layout_destroy destroys a matrix descriptor and releases * all resources used by the descriptor * * @param[in] @@ -145,7 +145,7 @@ rocblaslt_status rocblaslt_matrix_layout_create(rocblaslt_matrix_layout* matDesc * \retval rocblaslt_status_success the operation completed successfully. * \retval rocblaslt_status_invalid_pointer \p descr is invalid. */ -rocblaslt_status rocblaslt_matrix_layout_destory(const rocblaslt_matrix_layout descr); +rocblaslt_status rocblaslt_matrix_layout_destroy(const rocblaslt_matrix_layout descr); rocblaslt_status rocblaslt_matrix_layout_set_attribute(rocblaslt_matrix_layout matLayout, rocblaslt_matrix_layout_attribute attr, @@ -187,7 +187,7 @@ rocblaslt_status rocblaslt_matmul_desc_create(rocblaslt_matmul_desc* matmulDesc, * \brief Destroy a matrix multiplication descriptor * * \details - * \p rocblaslt_matrix_layout_destory destroys a multiplication matrix descr. + * \p rocblaslt_matrix_layout_destroy destroys a multiplication matrix descr. * * @param[in] * descr the matrix multiplication descriptor diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/handle.h b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/handle.h index 32181747a0ba..859594812340 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/handle.h +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/handle.h @@ -133,7 +133,7 @@ struct _rocblaslt_handle * content. It must be initialized using rocblaslt_matrix_layout_create() * and the retured handle must be passed * to all subsequent library function calls that involve the matrix. - * It should be destroyed at the end using rocblaslt_matrix_layout_destory(). + * It should be destroyed at the end using rocblaslt_matrix_layout_destroy(). *******************************************************************************/ struct _rocblaslt_matrix_layout { diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp index 79dc03366528..697a0ed98d8f 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp @@ -634,7 +634,7 @@ rocblaslt_status rocblaslt_get_sm_count_target(rocblaslt_handle handle, * content. It must be initialized using rocblaslt_matrix_layout_create() * and the retured handle must be passed * to all subsequent library function calls that involve the matrix. - * It should be destroyed at the end using rocblaslt_matrix_layout_destory(). + * It should be destroyed at the end using rocblaslt_matrix_layout_destroy(). *******************************************************************************/ rocblaslt_status rocblaslt_matrix_layout_create(rocblaslt_matrix_layout* matDescr, hipDataType valueType, @@ -682,7 +682,7 @@ rocblaslt_status rocblaslt_matrix_layout_create(rocblaslt_matrix_layout* matDesc /******************************************************************************** * \brief destroy matrix descriptor *******************************************************************************/ -rocblaslt_status rocblaslt_matrix_layout_destory(const rocblaslt_matrix_layout matDescr) +rocblaslt_status rocblaslt_matrix_layout_destroy(const rocblaslt_matrix_layout matDescr) { if(matDescr == nullptr) { From c183e5bcda2c0e47102419a20284554f031967f7 Mon Sep 17 00:00:00 2001 From: Seung-Hee Bae Date: Tue, 12 May 2026 19:11:08 +0000 Subject: [PATCH 02/27] [hipblaslt] implement host-side code for the 64bit offset support in hipblaslt. --- .../library/include/hipblaslt/hipblaslt.h | 14 ++++++- .../rocblaslt/include/rocblaslt-types.h | 21 +++++++--- .../amd_detail/rocblaslt/src/include/handle.h | 1 + .../src/include/rocblaslt_mat_utils.hpp | 31 ++++++++++++-- .../rocblaslt/src/rocblaslt_auxiliary.cpp | 36 +++++++++++++++-- .../rocblaslt/src/rocblaslt_mat.cpp | 40 +++++++++++++++++-- .../amd_detail/rocblaslt/src/tensile_host.cpp | 8 ++++ .../src/amd_detail/rocblaslt/src/utility.cpp | 6 ++- 8 files changed, 136 insertions(+), 21 deletions(-) diff --git a/projects/hipblaslt/library/include/hipblaslt/hipblaslt.h b/projects/hipblaslt/library/include/hipblaslt/hipblaslt.h index dd391231c0f0..bcbb4631138e 100644 --- a/projects/hipblaslt/library/include/hipblaslt/hipblaslt.h +++ b/projects/hipblaslt/library/include/hipblaslt/hipblaslt.h @@ -163,15 +163,25 @@ typedef enum { * ``int64_t`` */ HIPBLASLT_MATRIX_LAYOUT_LD = 6, + /** Matrix Batch Mode. * Batched GEMM can be either: * 1. Strided Batch: Single contiguous memory allocation and stride between matrices in * the batch is specified in terms of number of elements. - * 2. General Batched: This uses pointer array with each pointer storing the base address + * 2. General Batched: This uses pointer array with each pointer storing the base address * of the matrices in the batch. * See hipblasLtBatchMode_t */ - HIPBLASLT_MATRIX_LAYOUT_BATCH_MODE = 7, + HIPBLASLT_MATRIX_LAYOUT_BATCH_MODE = 7, + + /** Matrix Offset. + * + * For ``General Batched GEMM``, we can support for users to access a sub-matrix of + * the original matrix by adding an ``offset`` value from the base address. + * Note that for non-batched or Strided Batch GEMM case, we can directly apply + * the offset value by using the strided-offset value. + */ + HIPBLASLT_MATRIX_LAYOUT_OFFSET = 8 } hipblasLtMatrixLayoutAttribute_t; /*! \ingroup types_module diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/include/rocblaslt-types.h b/projects/hipblaslt/library/src/amd_detail/rocblaslt/include/rocblaslt-types.h index 6c1f9cf7b727..7a983557332d 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/include/rocblaslt-types.h +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/include/rocblaslt-types.h @@ -337,13 +337,14 @@ typedef enum rocblaslt_matrix_layout_attribute_ ROCBLASLT_MATRIX_LAYOUT_STRIDED_BATCH_OFFSET = 1, /**< stride between consecutive matrices in a batch expressed in terms of matrix elements. */ - ROCBLASLT_MATRIX_LAYOUT_TYPE = 2, - ROCBLASLT_MATRIX_LAYOUT_ORDER = 3, - ROCBLASLT_MATRIX_LAYOUT_ROWS = 4, - ROCBLASLT_MATRIX_LAYOUT_COLS = 5, - ROCBLASLT_MATRIX_LAYOUT_LD = 6, + ROCBLASLT_MATRIX_LAYOUT_TYPE = 2, + ROCBLASLT_MATRIX_LAYOUT_ORDER = 3, + ROCBLASLT_MATRIX_LAYOUT_ROWS = 4, + ROCBLASLT_MATRIX_LAYOUT_COLS = 5, + ROCBLASLT_MATRIX_LAYOUT_LD = 6, ROCBLASLT_MATRIX_LAYOUT_BATCH_MODE = 7, - ROCBLASLT_MATRIX_LAYOUT_MAX = 8 + ROCBLASLT_MATRIX_LAYOUT_OFFSET = 8, + ROCBLASLT_MATRIX_LAYOUT_MAX = 9 } rocblaslt_matrix_layout_attribute; typedef enum @@ -523,6 +524,7 @@ struct RocblasltContractionProblem size_t row_stride_a; size_t col_stride_a; size_t batch_stride_a; + int64_t batch_offset_a; hipDataType b_type; const void* B; @@ -530,6 +532,7 @@ struct RocblasltContractionProblem size_t row_stride_b; size_t col_stride_b; size_t batch_stride_b; + int64_t batch_offset_b; const void* beta; @@ -539,6 +542,7 @@ struct RocblasltContractionProblem size_t row_stride_c; size_t col_stride_c; size_t batch_stride_c; + int64_t batch_offset_c; hipDataType d_type; void* D; @@ -546,6 +550,7 @@ struct RocblasltContractionProblem size_t row_stride_d; size_t col_stride_d; size_t batch_stride_d; + int64_t batch_offset_d; void* E; void* const* batch_E; @@ -612,22 +617,26 @@ struct RocblasltContractionProblem const void* const* batch_A, int64_t ld_a, int64_t batch_stride_a, + int64_t batch_offset_a, hipDataType b_type, const void* B, const void* const* batch_B, int64_t ld_b, int64_t batch_stride_b, + int64_t batch_offset_b, const void* beta, hipDataType c_type, const void* C, const void* const* batch_C, int64_t ld_c, int64_t batch_stride_c, + int64_t batch_offset_c, hipDataType d_type, void* D, void* const* batch_D, int64_t ld_d, int64_t batch_stride_d, + int64_t batch_offset_d, void* E, void* const* batch_E, int64_t ld_e, diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/handle.h b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/handle.h index 859594812340..130980ee9740 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/handle.h +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/handle.h @@ -152,6 +152,7 @@ struct _rocblaslt_matrix_layout hipDataType type; int32_t batch_count = 1; int64_t batch_stride = 0; + int64_t batch_offset = 0; hipblasLtOrder_t order = HIPBLASLT_ORDER_COL; // Batch Mode hipblasLtBatchMode_t batch_mode = HIPBLASLT_BATCH_MODE_STRIDED; diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/rocblaslt_mat_utils.hpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/rocblaslt_mat_utils.hpp index 6eca0b491472..cd4e3fe4aa70 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/rocblaslt_mat_utils.hpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/rocblaslt_mat_utils.hpp @@ -151,6 +151,10 @@ inline rocblaslt_status validateMatmulArgs(int64_t m, int64_t batch_stride_b = 0, int64_t batch_stride_c = 0, int64_t batch_stride_d = 0, + int64_t batch_offset_a = 0, + int64_t batch_offset_b = 0, + int64_t batch_offset_c = 0, + int64_t batch_offset_d = 0, const rocblaslt_pointer_mode& pointermode = rocblaslt_pointer_mode_host) { @@ -211,17 +215,26 @@ inline rocblaslt_status validateMatmulArgs(int64_t m, if(batch_stride_a < 0 || batch_stride_b < 0 || batch_stride_c < 0 || batch_stride_d < 0) { #ifndef CODE_COVERAGE - std::cerr << "matrix and stride size must be positive" << std::endl; + std::cerr << "matrix and stride size must be zero or positive" << std::endl; #endif return rocblaslt_status_invalid_size; } - // number of batches of matrics A,B,C,D must be the same and negative + // sizes must not be negative + if(batch_offset_a < 0 || batch_offset_b < 0 || batch_offset_c < 0 || batch_offset_d < 0) + { +#ifndef CODE_COVERAGE + std::cerr << "matrix offset size must be zero or positive" << std::endl; +#endif + return rocblaslt_status_invalid_size; + } + + // number of batches of matrics A,B,C,D must be the same and positive if(num_batches_a != num_batches_b || num_batches_a != num_batches_c || num_batches_a != num_batches_d || num_batches_a < 1) { #ifndef CODE_COVERAGE - std::cerr << " number of batches of matrics A,B,C,D must be the same and negative" + std::cerr << " number of batches of matrics A,B,C,D must be the same and positive" << std::endl; #endif return rocblaslt_status_invalid_size; @@ -329,15 +342,19 @@ inline rocblaslt_status rocblaslt_matmul_valid_args(const rocblaslt_matmul_desc hipDataType& a_type, int64_t& lda, int64_t& batch_stride_a, + int64_t& batch_offset_a, hipDataType& b_type, int64_t& ldb, int64_t& batch_stride_b, + int64_t& batch_offset_b, hipDataType& c_type, int64_t& ldc, int64_t& batch_stride_c, + int64_t& batch_offset_c, hipDataType& d_type, int64_t& ldd, int64_t& batch_stride_d, + int64_t& batch_offset_d, int64_t& lde, int64_t& batch_stride_e, void*& bias, @@ -367,18 +384,21 @@ inline rocblaslt_status rocblaslt_matmul_valid_args(const rocblaslt_matmul_desc a_type = matA->type; lda = matA->ld; batch_stride_a = matA->batch_stride; + batch_offset_a = matA->batch_offset; // matrix B int num_batches_b = matB->batch_count; b_type = matB->type; ldb = matB->ld; batch_stride_b = matB->batch_stride; + batch_offset_b = matB->batch_offset; // matrix C int num_batches_c = matC->batch_count; c_type = matC->type; ldc = matC->ld; batch_stride_c = matC->batch_stride; + batch_offset_c = matC->batch_offset; // matrix D int64_t num_rows_d = matD->m; @@ -387,6 +407,7 @@ inline rocblaslt_status rocblaslt_matmul_valid_args(const rocblaslt_matmul_desc d_type = matD->type; ldd = matD->ld; batch_stride_d = matD->batch_stride; + batch_offset_d = matD->batch_offset; compute_type = matmul_descr->compute_type; @@ -418,6 +439,10 @@ inline rocblaslt_status rocblaslt_matmul_valid_args(const rocblaslt_matmul_desc batch_stride_b, batch_stride_c, batch_stride_d, + batch_offset_a, + batch_offset_b, + batch_offset_c, + batch_offset_d, matmul_descr->pointermode); const void* alphaVecPtr = matmul_descr->pointermode ? alpha : nullptr; diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp index 697a0ed98d8f..37a6a1d2debc 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp @@ -367,8 +367,9 @@ RocblasltContractionProblem construct_rocblaslt_problem(rocblaslt_handle { int8_t dummy; const void* dummy_ptr = &dummy; - int64_t m, n, k, lda, ldb, ldc, ldd, lde, batch_stride_a, batch_stride_b, batch_stride_c, - batch_stride_d, batch_stride_e; + int64_t m, n, k, lda, ldb, ldc, ldd, lde; + int64_t batch_stride_a, batch_stride_b, batch_stride_c, batch_stride_d, batch_stride_e; + int64_t batch_offset_a, batch_offset_b, batch_offset_c, batch_offset_d; int32_t bias_stride = matmul_descr->bias_stride; hipDataType bias_type; hipDataType aux_type; @@ -396,15 +397,19 @@ RocblasltContractionProblem construct_rocblaslt_problem(rocblaslt_handle a_type, lda, batch_stride_a, + batch_offset_a, b_type, ldb, batch_stride_b, + batch_offset_b, c_type, ldc, batch_stride_c, + batch_offset_c, d_type, ldd, batch_stride_d, + batch_offset_d, lde, batch_stride_e, bias, @@ -460,22 +465,26 @@ RocblasltContractionProblem construct_rocblaslt_problem(rocblaslt_handle nullptr, lda, batch_stride_a, + batch_offset_a, b_type, nullptr, nullptr, ldb, batch_stride_b, + batch_offset_b, beta, c_type, nullptr, nullptr, ldc, batch_stride_c, + batch_offset_c, d_type, nullptr, nullptr, ldd, batch_stride_d, + batch_offset_d, e, nullptr, lde, @@ -804,8 +813,17 @@ rocblaslt_status rocblaslt_matrix_layout_set_attribute(rocblaslt_matrix_layout { log_error(__func__, "invalid buf size", sizeInBytes); return rocblaslt_status_invalid_value; - } - break; + } + break; + case ROCBLASLT_MATRIX_LAYOUT_OFFSET: + if(sizeof(int64_t) <= sizeInBytes) + memcpy(&matLayout->batch_offset, buf, sizeof(int64_t)); + else + { + log_error(__func__, "invalid buf size", sizeInBytes); + return rocblaslt_status_invalid_value; + } + break; default: log_error(__func__, "invalid attribute", attr); return rocblaslt_status_invalid_value; @@ -892,6 +910,16 @@ rocblaslt_status rocblaslt_matrix_layout_get_attribute(rocblaslt_matrix_layout } memcpy(buf, &matLayout->batch_mode, sizeof(int32_t)); break; + case ROCBLASLT_MATRIX_LAYOUT_OFFSET: + if(sizeWritten) + *sizeWritten = sizeof(int64_t); + if(sizeInBytes < sizeof(int64_t)) + { + log_error(__func__, "invalid buf size", sizeInBytes); + return rocblaslt_status_invalid_value; + } + memcpy(buf, &matLayout->batch_offset, sizeof(int64_t)); + break; default: log_error(__func__, "invalid attribute", attr); return rocblaslt_status_invalid_value; diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_mat.cpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_mat.cpp index 5f001b8ef88d..51461365ccd7 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_mat.cpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_mat.cpp @@ -58,8 +58,10 @@ rocblaslt_status rocblaslt_matmul_impl(const rocblaslt_handle handle, size_t workspaceSizeInBytes, hipStream_t stream) { - int64_t m, n, k, lda, ldb, ldc, ldd, lde, batch_stride_a, batch_stride_b, batch_stride_c, - batch_stride_d, batch_stride_e; + int64_t m, n, k, lda, ldb, ldc, ldd, lde; + int64_t batch_stride_a, batch_stride_b, batch_stride_c, batch_stride_d, batch_stride_e; + int64_t batch_offset_a, batch_offset_b, batch_offset_c, batch_offset_d; + hipDataType bias_type; hipDataType aux_type; hipDataType type_a, type_b, type_c, type_d; @@ -87,15 +89,19 @@ rocblaslt_status rocblaslt_matmul_impl(const rocblaslt_handle handle, type_a, lda, batch_stride_a, + batch_offset_a, type_b, ldb, batch_stride_b, + batch_offset_b, type_c, ldc, batch_stride_c, + batch_offset_c, type_d, ldd, batch_stride_d, + batch_offset_d, lde, batch_stride_e, bias, @@ -174,22 +180,26 @@ rocblaslt_status rocblaslt_matmul_impl(const rocblaslt_handle handle, nullptr, lda, batch_stride_a, + batch_offset_a, type_b, B, nullptr, ldb, batch_stride_b, + batch_offset_b, beta, type_c, C, nullptr, ldc, batch_stride_c, + batch_offset_c, type_d, D, nullptr, ldd, batch_stride_d, + batch_offset_d, E, nullptr, lde, @@ -265,8 +275,10 @@ rocblaslt_status rocblaslt_gemm_create_cpp_impl(const rocblaslt_handle std::shared_ptr& gemmData, size_t& gemmCount) { - int64_t m, n, k, lda, ldb, ldc, ldd, lde, batch_stride_a, batch_stride_b, batch_stride_c, - batch_stride_d, batch_stride_e; + int64_t m, n, k, lda, ldb, ldc, ldd, lde; + int64_t batch_stride_a, batch_stride_b, batch_stride_c, batch_stride_d, batch_stride_e; + int64_t batch_offset_a, batch_offset_b, batch_offset_c, batch_offset_d; + hipDataType bias_type; hipDataType aux_type; hipDataType type_a, type_b, type_c, type_d; @@ -293,15 +305,19 @@ rocblaslt_status rocblaslt_gemm_create_cpp_impl(const rocblaslt_handle type_a, lda, batch_stride_a, + batch_offset_a, type_b, ldb, batch_stride_b, + batch_offset_b, type_c, ldc, batch_stride_c, + batch_offset_c, type_d, ldd, batch_stride_d, + batch_offset_d, lde, batch_stride_e, bias, @@ -358,22 +374,26 @@ rocblaslt_status rocblaslt_gemm_create_cpp_impl(const rocblaslt_handle nullptr, lda, batch_stride_a, + batch_offset_a, type_b, B, nullptr, ldb, batch_stride_b, + batch_offset_b, beta, type_c, C, nullptr, ldc, batch_stride_c, + batch_offset_c, type_d, D, nullptr, ldd, batch_stride_d, + batch_offset_d, E, nullptr, lde, @@ -652,22 +672,26 @@ rocblaslt_status nullptr, lda_vec[i], batch_stride_a_vec[i], + 0, // batch_offset_a type_b, B_vec[i], nullptr, ldb_vec[i], batch_stride_b_vec[i], + 0, // batch_offset_b beta_vec[i], type_c, C_vec[i], nullptr, ldc_vec[i], batch_stride_c_vec[i], + 0, // batch_offset_c type_d, D_vec[i], nullptr, ldd_vec[i], batch_stride_d_vec[i], + 0, // batch_offset_d E_vec[i], nullptr, lde_vec[i], @@ -997,22 +1021,26 @@ rocblaslt_status rocblaslt_gemm_create_cpp_impl_2(const rocblaslt_handle handle, nullptr, lda, batch_stride_a, + 0, // batch_offset_a, type_b, B, nullptr, ldb, batch_stride_b, + 0, // batch_offset_b, beta, type_c, C, nullptr, ldc, batch_stride_c, + 0, // batch_offset_c, type_d, D, nullptr, ldd, batch_stride_d, + 0, // batch_offset_d, E, nullptr, lde, @@ -1317,22 +1345,26 @@ rocblaslt_status rocblaslt_groupedgemm_create_cpp_impl_2(const rocblaslt_handle nullptr, lda[i], strideA[i], + 0, // batch_offset_a type_b, B_vec[i], nullptr, ldb[i], strideB[i], + 0, // batch_offset_b beta_vec[i], type_c, C_vec[i], nullptr, ldc[i], strideC[i], + 0, // batch_offset_c type_d, D_vec[i], nullptr, ldd[i], strideD[i], + 0, // batch_offset_d E_vec[i], nullptr, lde_vec[i], diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp index 323e802910b2..acc061421449 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp @@ -89,22 +89,26 @@ RocblasltContractionProblem::RocblasltContractionProblem(hipblasOperation_t const void* const* batch_A, int64_t ld_a, int64_t batch_stride_a, + int64_t batch_offset_a, hipDataType b_type, const void* B, const void* const* batch_B, int64_t ld_b, int64_t batch_stride_b, + int64_t batch_offset_b, const void* beta, hipDataType c_type, const void* C, const void* const* batch_C, int64_t ld_c, int64_t batch_stride_c, + int64_t batch_offset_c, hipDataType d_type, void* D, void* const* batch_D, int64_t ld_d, int64_t batch_stride_d, + int64_t batch_offset_d, void* E, void* const* batch_E, int64_t ld_e, @@ -152,12 +156,14 @@ RocblasltContractionProblem::RocblasltContractionProblem(hipblasOperation_t , row_stride_a(1) , col_stride_a(ld_a) , batch_stride_a(batch_stride_a) + , batch_offset_a(batch_offset_a) , b_type(b_type) , B(B) , batch_B(batch_B) , row_stride_b(1) , col_stride_b(ld_b) , batch_stride_b(batch_stride_b) + , batch_offset_b(batch_offset_b) , beta(beta) , c_type(c_type) , C(C) @@ -165,12 +171,14 @@ RocblasltContractionProblem::RocblasltContractionProblem(hipblasOperation_t , row_stride_c(1) , col_stride_c(ld_c) , batch_stride_c(batch_stride_c) + , batch_offset_c(batch_offset_c) , d_type(d_type) , D(D) , batch_D(batch_D) , row_stride_d(1) , col_stride_d(ld_d) , batch_stride_d(batch_stride_d) + , batch_offset_d(batch_offset_d) , E(E) , batch_E(batch_E) , row_stride_e(1) diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/utility.cpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/utility.cpp index 57cfc325a651..d56ab2f29896 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/utility.cpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/utility.cpp @@ -202,9 +202,9 @@ const char* rocblaslt_matrix_layout_attributes_to_string(rocblaslt_matrix_layout switch(type) { case ROCBLASLT_MATRIX_LAYOUT_BATCH_COUNT: - return "MATRIX_LAYOUT_BATCH_COUNT"; + return "ROCBLASLT_MATRIX_LAYOUT_BATCH_COUNT"; case ROCBLASLT_MATRIX_LAYOUT_STRIDED_BATCH_OFFSET: - return "MATRIX_LAYOUT_STRIDED_BATCH_OFFSET"; + return "ROCBLASLT_MATRIX_LAYOUT_STRIDED_BATCH_OFFSET"; case ROCBLASLT_MATRIX_LAYOUT_TYPE: return "ROCBLASLT_MATRIX_LAYOUT_TYPE"; case ROCBLASLT_MATRIX_LAYOUT_ORDER: @@ -217,6 +217,8 @@ const char* rocblaslt_matrix_layout_attributes_to_string(rocblaslt_matrix_layout return "ROCBLASLT_MATRIX_LAYOUT_LD"; case ROCBLASLT_MATRIX_LAYOUT_BATCH_MODE: return "ROCBLASLT_MATRIX_LAYOUT_BATCH_MODE"; + case ROCBLASLT_MATRIX_LAYOUT_OFFSET: + return "ROCBLASLT_MATRIX_LAYOUT_OFFSET"; case ROCBLASLT_MATRIX_LAYOUT_MAX: return "ROCBLASLT_MATRIX_LAYOUT_MAX"; default: From e03ec4d016d6a6dfe935ff0af9a6d2db71113794 Mon Sep 17 00:00:00 2001 From: Seung-Hee Bae Date: Wed, 13 May 2026 01:13:52 +0000 Subject: [PATCH 03/27] [hipblaslt] pass batch offset values as kernel arguments. --- .../amd_detail/rocblaslt/src/tensile_host.cpp | 6 ++++ .../include/Tensile/ContractionProblem.hpp | 5 ++++ .../tensilelite/src/ContractionSolution.cpp | 28 +++++++++++++++++++ 3 files changed, 39 insertions(+) diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp index acc061421449..4b8c6542d796 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp @@ -2424,6 +2424,12 @@ namespace inputs.batchC = reinterpret_cast(prob.batch_C); inputs.batchD = reinterpret_cast(prob.batch_D); + // set the batched offset of A, B, C, and D. + inputs.batchOffsetA = prob.batch_offset_a; + inputs.batchOffsetB = prob.batch_offset_b; + inputs.batchOffsetC = prob.batch_offset_c; + inputs.batchOffsetD = prob.batch_offset_d; + // Set the GSU workspace inputs.ws = prob.workspace; inputs.workspaceSize = prob.workspaceSize; diff --git a/projects/hipblaslt/tensilelite/include/Tensile/ContractionProblem.hpp b/projects/hipblaslt/tensilelite/include/Tensile/ContractionProblem.hpp index 0e91852bb1be..051ce6acdded 100644 --- a/projects/hipblaslt/tensilelite/include/Tensile/ContractionProblem.hpp +++ b/projects/hipblaslt/tensilelite/include/Tensile/ContractionProblem.hpp @@ -1587,6 +1587,11 @@ namespace TensileLite unsigned char const* metadata = nullptr; void const* compressed = nullptr; + int64_t batchOffsetA = 0; + int64_t batchOffsetB = 0; + int64_t batchOffsetC = 0; + int64_t batchOffsetD = 0; + // Constants ConstantVariant alpha = static_cast(0); ConstantVariant beta = static_cast(0); diff --git a/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp b/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp index 9a60778d3580..dd1766bc5101 100644 --- a/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp +++ b/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp @@ -620,6 +620,12 @@ namespace TensileLite { args.template append("batchD", inputs.batchD); args.template append("batchC", inputs.batchC); + + if(problem.batchMode() == ContractionProblemGemm::BATCHMODE::POINTER_ARRAY) + { + args.template append("batchOffsetD", inputs.batchOffsetD); + args.template append("batchOffsetC", inputs.batchOffsetC); + } } if(problemType.stridedBatched) @@ -637,6 +643,12 @@ namespace TensileLite { args.template append("batchA", inputs.batchA); args.template append("batchB", inputs.batchB); + + if(problem.batchMode() == ContractionProblemGemm::BATCHMODE::POINTER_ARRAY) + { + args.template append("batchOffsetA", inputs.batchOffsetA); + args.template append("batchOffsetB", inputs.batchOffsetB); + } } if(problemType.sparse) @@ -2113,12 +2125,20 @@ namespace TensileLite else if(problemType.stridedBatched) rv.args.append("D", inputs.d); else + { rv.args.append("batchD", inputs.batchD); + if(problem.batchMode() == ContractionProblemGemm::BATCHMODE::POINTER_ARRAY) + rv.args.append("batchOffsetD", inputs.batchOffsetD); + } if(problemType.stridedBatched) rv.args.append("C", inputs.c); else + { rv.args.append("batchC", inputs.batchC); + if(problem.batchMode() == ContractionProblemGemm::BATCHMODE::POINTER_ARRAY) + rv.args.append("batchOffsetC", inputs.batchOffsetC); + } if(problemType.useBias && sizeMapping.globalAccumulation == 0 && (!problemType.useGradient)) { @@ -2280,14 +2300,22 @@ namespace TensileLite if(problemType.stridedBatched) args.template append("D", inputs.d); else + { args.template append("batchD", inputs.batchD); + if(problem.batchMode() == ContractionProblemGemm::BATCHMODE::POINTER_ARRAY) + args.template append("batchOffsetD", inputs.batchOffsetD); + } args.template append("WS", (uint8_t*)inputs.ws + workspaceOffsetInByte); if(problemType.stridedBatched) args.template append("C", inputs.c); else + { args.template append("batchC", inputs.batchC); + if(problem.batchMode() == ContractionProblemGemm::BATCHMODE::POINTER_ARRAY) + args.template append("batchOffsetC", inputs.batchOffsetC); + } bool useBias = false; if(problemType.useBias) From 763df2a7fe70625ab2488fb6e6cbeab7cbb20bdc Mon Sep 17 00:00:00 2001 From: Seung-Hee Bae Date: Tue, 19 May 2026 19:03:53 +0000 Subject: [PATCH 04/27] [hipblaslt] add offset parameters into the kernel signature. --- .../tensilelite/Tensile/Components/Signature.py | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/projects/hipblaslt/tensilelite/Tensile/Components/Signature.py b/projects/hipblaslt/tensilelite/Tensile/Components/Signature.py index 7d101e4ea4f5..a6fad12207fb 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Components/Signature.py +++ b/projects/hipblaslt/tensilelite/Tensile/Components/Signature.py @@ -163,6 +163,15 @@ def __call__(self, writer) -> SignatureBase: if kernel["ProblemType"]["Sparse"]: signature.addArg("MetaData", SVK.SIG_GLOBALBUFFER, "void" , "generic") + # Batch offset support for general batched mode (pointer array) + # Only for non-strided, non-grouped GEMM + if not kernel["ProblemType"]["StridedBatched"] and not kernel["ProblemType"]["GroupedGemm"]: + signature.addArg("batchOffsetD", SVK.SIG_VALUE, "u64") + signature.addArg("batchOffsetC", SVK.SIG_VALUE, "u64") + signature.addArg("batchOffsetA", SVK.SIG_VALUE, "u64") + signature.addArg("batchOffsetB", SVK.SIG_VALUE, "u64") + userArgumentsInfo.gemmArgumentSize += 32 # 4 offsets * 8 bytes each + if kernel["StreamK"] > 0 and kernel["StreamKAtomic"] == 0: signature.addArg("AddressWS", SVK.SIG_GLOBALBUFFER, cptValueType, "generic") signature.addArg("AddressFlags", SVK.SIG_GLOBALBUFFER, dstValueType, "generic") From f0bde2b25d8d75725a093c88c6e8f0655f9c7d91 Mon Sep 17 00:00:00 2001 From: Seung-Hee Bae Date: Wed, 20 May 2026 07:09:47 +0000 Subject: [PATCH 05/27] [hipblaslt] update kernel generation to use offsets in general batch mode. --- .../tensilelite/Tensile/KernelWriter.py | 7 ++++++ .../Tensile/KernelWriterAssembly.py | 23 +++++++++++++++++++ 2 files changed, 30 insertions(+) diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py index a63a57da1e7c..59f454f87825 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py @@ -9233,6 +9233,13 @@ def vgprAllocationImplSubtile(): if kernel["ProblemType"]["Sparse"]: self.defineSgpr("StridesMetadata", self.states.m.numSgprStrides) + # Batch offset support for general batched GEMM (pointer array mode) + if not kernel["ProblemType"]["StridedBatched"] and not kernel["ProblemType"]["GroupedGemm"]: + self.defineSgpr("BatchOffsetD", 2) # 64-bit offset requires 2 SGPRs + self.defineSgpr("BatchOffsetC", 2) # 64-bit offset requires 2 SGPRs + self.defineSgpr("BatchOffsetA", 2) # 64-bit offset requires 2 SGPRs + self.defineSgpr("BatchOffsetB", 2) # 64-bit offset requires 2 SGPRs + # for packed batches without stride restrictions need to do something different here assert sorted(kernel["PackedC0IdxChars"]+kernel["PackedC1IdxChars"]) == \ sorted(set(kernel["PackedC0IdxChars"]+kernel["PackedC1IdxChars"])) diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py index eb25ac4b68e4..877b1589b994 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py @@ -2130,6 +2130,9 @@ def loadBatchedAddress(self, kernel, Batch, tmpSgprResource: ContinuousRegister) if not isPackedIndex(kernel,idx): module.add(SMulI32(dst=sgpr(tmpSgpr), src0=sgpr(Batch), src1=0x8, comment="offset of global buffer address")) module.add(SLoadB64(dst=sgpr("AddressD", 2), base=sgpr("AddressD",2), soffset=sgpr(tmpSgpr), comment="load global buffer D address")) + # Apply batch offset to AddressD for general batched mode + if not kernel["ProblemType"]["StridedBatched"] and not kernel["ProblemType"]["GroupedGemm"]: + module.add(SAddU64(dst=sgpr("AddressD", 2), src0=sgpr("AddressD", 2), src1=sgpr("BatchOffsetD", 2), comment="add batch offset to D address")) # Only load C buffer address if Beta is used and potentially non-zero if kernel["ProblemType"]["UseBeta"]: @@ -2141,6 +2144,9 @@ def loadBatchedAddress(self, kernel, Batch, tmpSgprResource: ContinuousRegister) if not isPackedIndex(kernel,idx): module.add(SMulI32(dst=sgpr(tmpSgpr), src0=sgpr(Batch), src1=0x8, comment="offset of global buffer address")) module.add(SLoadB64(dst=sgpr("AddressC", 2), base=sgpr("AddressC",2), soffset=sgpr(tmpSgpr), comment="load global buffer C address")) + # Apply batch offset to AddressC for general batched mode + if not kernel["ProblemType"]["StridedBatched"] and not kernel["ProblemType"]["GroupedGemm"]: + module.add(SAddU64(dst=sgpr("AddressC", 2), src0=sgpr("AddressC", 2), src1=sgpr("BatchOffsetC", 2), comment="add batch offset to C address")) module.add(endCheckLabel) @@ -2159,6 +2165,10 @@ def loadBatchedAddress(self, kernel, Batch, tmpSgprResource: ContinuousRegister) if not isPackedIndex(kernel,idx): module.add(SLoadB64(dst=sgpr("AddressA", 2), base=sgpr("AddressA",2), soffset=sgpr(tmpSgpr), comment="load global buffer A address")) module.add(SLoadB64(dst=sgpr("AddressB", 2), base=sgpr("AddressB",2), soffset=sgpr(tmpSgpr), comment="load global buffer B address")) + # Apply batch offset to AddressA and AddressB for general batched mode + if not kernel["ProblemType"]["StridedBatched"] and not kernel["ProblemType"]["GroupedGemm"]: + module.add(SAddU64(dst=sgpr("AddressA", 2), src0=sgpr("AddressA", 2), src1=sgpr("BatchOffsetA", 2), comment="add batch offset to A address")) + module.add(SAddU64(dst=sgpr("AddressB", 2), src0=sgpr("AddressB", 2), src1=sgpr("BatchOffsetB", 2), comment="add batch offset to B address")) module.add(endCheckLabel) @@ -2184,6 +2194,19 @@ def getKernelArgLoadModule(self, kernel, sgprStartIdx, numsOfLoad, preloadNum): kernelArgs.add(self.argLoader.loadKernArg(self.states.esmRuntimeFlagSgpr, "KernArgAddress", sgprOffset=hex(sgprOffset), dword=1)) sgprOffset += 4 self.argLoader.setOffset(sgprOffset) + + # Load batch offset arguments for general batched mode + if not kernel["ProblemType"]["StridedBatched"] and not kernel["ProblemType"]["GroupedGemm"]: + sgprOffset = self.argLoader.getOffset() + kernelArgs.add(self.argLoader.loadKernArg("BatchOffsetD", "KernArgAddress", sgprOffset=hex(sgprOffset), dword=2)) + sgprOffset += 8 # 64-bit = 8 bytes + kernelArgs.add(self.argLoader.loadKernArg("BatchOffsetC", "KernArgAddress", sgprOffset=hex(sgprOffset), dword=2)) + sgprOffset += 8 + kernelArgs.add(self.argLoader.loadKernArg("BatchOffsetA", "KernArgAddress", sgprOffset=hex(sgprOffset), dword=2)) + sgprOffset += 8 + kernelArgs.add(self.argLoader.loadKernArg("BatchOffsetB", "KernArgAddress", sgprOffset=hex(sgprOffset), dword=2)) + sgprOffset += 8 + return kernelArgs def localReadAddresses(self, kernel, tPA, tPB, tPM): From 4ee679f2cbd072eb2ce23bd7de78b6b7fd94a451 Mon Sep 17 00:00:00 2001 From: Seung-Hee Bae Date: Wed, 20 May 2026 17:50:01 +0000 Subject: [PATCH 06/27] [hipblaslt] add waitcnt instruction after SLoadB64 before updating the offset. --- .../hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py index 877b1589b994..cd583a625db6 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py @@ -2132,6 +2132,7 @@ def loadBatchedAddress(self, kernel, Batch, tmpSgprResource: ContinuousRegister) module.add(SLoadB64(dst=sgpr("AddressD", 2), base=sgpr("AddressD",2), soffset=sgpr(tmpSgpr), comment="load global buffer D address")) # Apply batch offset to AddressD for general batched mode if not kernel["ProblemType"]["StridedBatched"] and not kernel["ProblemType"]["GroupedGemm"]: + module.add(SWaitCnt(kmcnt=0, comment="Wait for the Matrix Address Load from the Pointer Array before updating offset")) module.add(SAddU64(dst=sgpr("AddressD", 2), src0=sgpr("AddressD", 2), src1=sgpr("BatchOffsetD", 2), comment="add batch offset to D address")) # Only load C buffer address if Beta is used and potentially non-zero @@ -2146,6 +2147,7 @@ def loadBatchedAddress(self, kernel, Batch, tmpSgprResource: ContinuousRegister) module.add(SLoadB64(dst=sgpr("AddressC", 2), base=sgpr("AddressC",2), soffset=sgpr(tmpSgpr), comment="load global buffer C address")) # Apply batch offset to AddressC for general batched mode if not kernel["ProblemType"]["StridedBatched"] and not kernel["ProblemType"]["GroupedGemm"]: + module.add(SWaitCnt(kmcnt=0, comment="Wait for the Matrix Address Load from the Pointer Array before updating offset")) module.add(SAddU64(dst=sgpr("AddressC", 2), src0=sgpr("AddressC", 2), src1=sgpr("BatchOffsetC", 2), comment="add batch offset to C address")) module.add(endCheckLabel) @@ -2167,7 +2169,9 @@ def loadBatchedAddress(self, kernel, Batch, tmpSgprResource: ContinuousRegister) module.add(SLoadB64(dst=sgpr("AddressB", 2), base=sgpr("AddressB",2), soffset=sgpr(tmpSgpr), comment="load global buffer B address")) # Apply batch offset to AddressA and AddressB for general batched mode if not kernel["ProblemType"]["StridedBatched"] and not kernel["ProblemType"]["GroupedGemm"]: + module.add(SWaitCnt(kmcnt=1, comment="Wait for the AddressA Load from the Pointer Array before updating offset")) module.add(SAddU64(dst=sgpr("AddressA", 2), src0=sgpr("AddressA", 2), src1=sgpr("BatchOffsetA", 2), comment="add batch offset to A address")) + module.add(SWaitCnt(kmcnt=0, comment="Wait for the AddressB Load from the Pointer Array before updating offset")) module.add(SAddU64(dst=sgpr("AddressB", 2), src0=sgpr("AddressB", 2), src1=sgpr("BatchOffsetB", 2), comment="add batch offset to B address")) module.add(endCheckLabel) From cf9e3a9487b8c7c9c4dfe6e0a4e1985af9bea468 Mon Sep 17 00:00:00 2001 From: Seung-Hee Bae Date: Thu, 28 May 2026 23:35:19 +0000 Subject: [PATCH 07/27] [hipblaslt] implement tests for 64-bit offset support --- .../common/include/hipblaslt_arguments.hpp | 10 + .../include/testing_matmul_batch_offset.hpp | 411 ++++++++++++++++++ .../clients/tests/data/hipblaslt_common.yaml | 8 + .../clients/tests/data/matmul_gtest.yaml | 150 ++++++- .../clients/tests/src/matmul_gtest.cpp | 6 +- 5 files changed, 583 insertions(+), 2 deletions(-) create mode 100644 projects/hipblaslt/clients/common/include/testing_matmul_batch_offset.hpp diff --git a/projects/hipblaslt/clients/common/include/hipblaslt_arguments.hpp b/projects/hipblaslt/clients/common/include/hipblaslt_arguments.hpp index d3d2040686ad..ac50c9c48a21 100644 --- a/projects/hipblaslt/clients/common/include/hipblaslt_arguments.hpp +++ b/projects/hipblaslt/clients/common/include/hipblaslt_arguments.hpp @@ -112,6 +112,12 @@ struct Arguments int32_t batch_count; int32_t batch_mode; + // Batch offset support for general batched GEMM + int64_t batch_offset_a; + int64_t batch_offset_b; + int64_t batch_offset_c; + int64_t batch_offset_d; + int32_t iters; int32_t cold_iters; @@ -251,6 +257,10 @@ struct Arguments OPER(lde) SEP \ OPER(batch_count) SEP \ OPER(batch_mode) SEP \ + OPER(batch_offset_a) SEP \ + OPER(batch_offset_b) SEP \ + OPER(batch_offset_c) SEP \ + OPER(batch_offset_d) SEP \ OPER(iters) SEP \ OPER(cold_iters) SEP \ OPER(warmup_time) SEP \ diff --git a/projects/hipblaslt/clients/common/include/testing_matmul_batch_offset.hpp b/projects/hipblaslt/clients/common/include/testing_matmul_batch_offset.hpp new file mode 100644 index 000000000000..3637c9aeea08 --- /dev/null +++ b/projects/hipblaslt/clients/common/include/testing_matmul_batch_offset.hpp @@ -0,0 +1,411 @@ +/******************************************************************************* + * + * MIT License + * + * Copyright (C) 2022-2026 Advanced Micro Devices, Inc. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + * + *******************************************************************************/ + +#pragma once + +#include "cblas_interface.hpp" +#include "datatype_interface.hpp" +#include "flops.hpp" +#include "hipblaslt_datatype2string.hpp" +#include "hipblaslt_init.hpp" +#include "hipblaslt_math.hpp" +#include "hipblaslt_random.hpp" +#include "hipblaslt_test.hpp" +#include "hipblaslt_vector.hpp" +#include "near.hpp" +#include "norm.hpp" +#include "unit.hpp" +#include "utility.hpp" +#include + +/* ============================================================================================ */ +/*! \brief Test for 64-bit batch offset support in general batched GEMM */ + +template +void testing_matmul_batch_offset_impl(const Arguments& arg) +{ + hipblasOperation_t transA = char_to_hipblas_operation(arg.transA); + hipblasOperation_t transB = char_to_hipblas_operation(arg.transB); + + // Use first element from arrays (grouped GEMM uses arrays, we use first element) + int64_t M = arg.M[0]; + int64_t N = arg.N[0]; + int64_t K = arg.K[0]; + + int64_t lda = arg.lda[0]; + int64_t ldb = arg.ldb[0]; + int64_t ldc = arg.ldc[0]; + int64_t ldd = arg.ldd[0]; + + int32_t batch_count = arg.batch_count; + + // Batch offsets (from YAML - assumed to be in BYTES) + int64_t offset_a = arg.batch_offset_a; + int64_t offset_b = arg.batch_offset_b; + int64_t offset_c = arg.batch_offset_c; + int64_t offset_d = arg.batch_offset_d; + + // Only test general batched mode (pointer array) + if(arg.batch_mode != 1) // 1 = pointer array + { + GTEST_SKIP() << "Batch offset only supported for general batched mode (batch_mode=1)"; + } + + // Calculate matrix sizes + int64_t A_row = transA == HIPBLAS_OP_N ? M : K; + int64_t A_col = transA == HIPBLAS_OP_N ? K : M; + int64_t B_row = transB == HIPBLAS_OP_N ? K : N; + int64_t B_col = transB == HIPBLAS_OP_N ? N : K; + + // Sub-matrix sizes (actual GEMM size) in elements + size_t size_A_sub = size_t(lda) * size_t(A_col); + size_t size_B_sub = size_t(ldb) * size_t(B_col); + size_t size_C_sub = size_t(ldc) * size_t(N); + size_t size_D_sub = size_t(ldd) * size_t(N); + + // Convert byte offsets to element offsets + // Note: offsets must be aligned to element size + size_t offset_a_elem = offset_a / sizeof(Ti); + size_t offset_b_elem = offset_b / sizeof(Ti); + size_t offset_c_elem = offset_c / sizeof(To); + size_t offset_d_elem = offset_d / sizeof(To); + + // Full buffer sizes: [offset padding] + [matrix data] + // The offset padding comes FIRST because: + // - Pointer array points to buffer BASE + // - Kernel adds offset to BASE to get actual data location + // - So actual data is at BASE + offset + size_t size_A_full = offset_a_elem + size_A_sub; + size_t size_B_full = offset_b_elem + size_B_sub; + size_t size_C_full = offset_c_elem + size_C_sub; + size_t size_D_full = offset_d_elem + size_D_sub; + + // Allocate host memory for full buffers + host_vector h_A_full(size_A_full * batch_count); + host_vector h_B_full(size_B_full * batch_count); + host_vector h_C_full(size_C_full * batch_count); + host_vector h_D_full(size_D_full * batch_count); // GPU result with offset API + host_vector h_D_gold(size_D_sub * batch_count); // CPU reference + + // Initialize matrices with known pattern + // Data must be placed at (base + offset) because the offset API tells the kernel + // to start reading at that location + for(int b = 0; b < batch_count; b++) + { + // Initialize each batch's sub-matrix at the offset location + Ti* A_batch = h_A_full.data() + b * size_A_full + offset_a_elem; + Ti* B_batch = h_B_full.data() + b * size_B_full + offset_b_elem; + To* C_batch = h_C_full.data() + b * size_C_full + offset_c_elem; + + // Simple initialization: A and B with small integers + for(int64_t j = 0; j < A_col; j++) + for(int64_t i = 0; i < A_row; i++) + A_batch[i + j * lda] = Ti((i + j + b) % 7 + 1); + + for(int64_t j = 0; j < B_col; j++) + for(int64_t i = 0; i < B_row; i++) + B_batch[i + j * ldb] = Ti((i - j + b) % 5 + 1); + + for(int64_t j = 0; j < N; j++) + for(int64_t i = 0; i < M; i++) + C_batch[i + j * ldc] = To((i + j) % 3); + } + + // Allocate device memory + device_vector d_A_full(size_A_full * batch_count); + device_vector d_B_full(size_B_full * batch_count); + device_vector d_C_full(size_C_full * batch_count); + device_vector d_D_full(size_D_full * batch_count); + + // Copy to device + CHECK_HIP_ERROR( + hipMemcpy(d_A_full, h_A_full.data(), sizeof(Ti) * size_A_full * batch_count, hipMemcpyHostToDevice)); + CHECK_HIP_ERROR( + hipMemcpy(d_B_full, h_B_full.data(), sizeof(Ti) * size_B_full * batch_count, hipMemcpyHostToDevice)); + CHECK_HIP_ERROR( + hipMemcpy(d_C_full, h_C_full.data(), sizeof(To) * size_C_full * batch_count, hipMemcpyHostToDevice)); + + // Setup pointer arrays for base addresses + std::vector h_batch_A(batch_count); + std::vector h_batch_B(batch_count); + std::vector h_batch_C(batch_count); + std::vector h_batch_D(batch_count); + + for(int b = 0; b < batch_count; b++) + { + h_batch_A[b] = d_A_full + b * size_A_full; + h_batch_B[b] = d_B_full + b * size_B_full; + h_batch_C[b] = d_C_full + b * size_C_full; + h_batch_D[b] = d_D_full + b * size_D_full; + } + + // Allocate device memory for pointer arrays + Ti** d_batch_A; + Ti** d_batch_B; + To** d_batch_C; + To** d_batch_D; + + CHECK_HIP_ERROR(hipMalloc(&d_batch_A, sizeof(Ti*) * batch_count)); + CHECK_HIP_ERROR(hipMalloc(&d_batch_B, sizeof(Ti*) * batch_count)); + CHECK_HIP_ERROR(hipMalloc(&d_batch_C, sizeof(To*) * batch_count)); + CHECK_HIP_ERROR(hipMalloc(&d_batch_D, sizeof(To*) * batch_count)); + + CHECK_HIP_ERROR( + hipMemcpy(d_batch_A, h_batch_A.data(), sizeof(Ti*) * batch_count, hipMemcpyHostToDevice)); + CHECK_HIP_ERROR( + hipMemcpy(d_batch_B, h_batch_B.data(), sizeof(Ti*) * batch_count, hipMemcpyHostToDevice)); + CHECK_HIP_ERROR( + hipMemcpy(d_batch_C, h_batch_C.data(), sizeof(To*) * batch_count, hipMemcpyHostToDevice)); + CHECK_HIP_ERROR( + hipMemcpy(d_batch_D, h_batch_D.data(), sizeof(To*) * batch_count, hipMemcpyHostToDevice)); + + // Alpha and beta + Tc h_alpha = arg.get_alpha(); + Tc h_beta = arg.get_beta(); + + // Setup hipBLASLt + hipblasLtHandle_t handle; + CHECK_HIPBLASLT_ERROR(hipblasLtCreate(&handle)); + + hipblasLtMatmulDesc_t matmul_desc; + CHECK_HIPBLASLT_ERROR( + hipblasLtMatmulDescCreate(&matmul_desc, arg.compute_type, arg.scale_type)); + CHECK_HIPBLASLT_ERROR( + hipblasLtMatmulDescSetAttribute( + matmul_desc, HIPBLASLT_MATMUL_DESC_TRANSA, &transA, sizeof(transA))); + CHECK_HIPBLASLT_ERROR( + hipblasLtMatmulDescSetAttribute( + matmul_desc, HIPBLASLT_MATMUL_DESC_TRANSB, &transB, sizeof(transB))); + + // Create matrix layouts + hipblasLtMatrixLayout_t matA, matB, matC, matD; + CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutCreate(&matA, arg.a_type, A_row, A_col, lda)); + CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutCreate(&matB, arg.b_type, B_row, B_col, ldb)); + CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutCreate(&matC, arg.c_type, M, N, ldc)); + CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutCreate(&matD, arg.d_type, M, N, ldd)); + + // Set batch count and mode + CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutSetAttribute( + matA, HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT, &batch_count, sizeof(batch_count))); + CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutSetAttribute( + matB, HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT, &batch_count, sizeof(batch_count))); + CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutSetAttribute( + matC, HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT, &batch_count, sizeof(batch_count))); + CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutSetAttribute( + matD, HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT, &batch_count, sizeof(batch_count))); + + int32_t batch_mode = 1; // Pointer array + CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutSetAttribute( + matA, HIPBLASLT_MATRIX_LAYOUT_BATCH_MODE, &batch_mode, sizeof(batch_mode))); + CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutSetAttribute( + matB, HIPBLASLT_MATRIX_LAYOUT_BATCH_MODE, &batch_mode, sizeof(batch_mode))); + CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutSetAttribute( + matC, HIPBLASLT_MATRIX_LAYOUT_BATCH_MODE, &batch_mode, sizeof(batch_mode))); + CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutSetAttribute( + matD, HIPBLASLT_MATRIX_LAYOUT_BATCH_MODE, &batch_mode, sizeof(batch_mode))); + + // ======================================== + // GPU GEMM with offset API + // ======================================== + + // Set offsets for all matrices + CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutSetAttribute( + matA, HIPBLASLT_MATRIX_LAYOUT_OFFSET, &offset_a, sizeof(offset_a))); + CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutSetAttribute( + matB, HIPBLASLT_MATRIX_LAYOUT_OFFSET, &offset_b, sizeof(offset_b))); + CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutSetAttribute( + matC, HIPBLASLT_MATRIX_LAYOUT_OFFSET, &offset_c, sizeof(offset_c))); + CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutSetAttribute( + matD, HIPBLASLT_MATRIX_LAYOUT_OFFSET, &offset_d, sizeof(offset_d))); + + // Find algorithm + hipblasLtMatmulPreference_t pref; + CHECK_HIPBLASLT_ERROR(hipblasLtMatmulPreferenceCreate(&pref)); + size_t maxWorkspaceSize = 128 * 1024 * 1024; // 128 MB + CHECK_HIPBLASLT_ERROR(hipblasLtMatmulPreferenceSetAttribute( + pref, HIPBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES, &maxWorkspaceSize, sizeof(maxWorkspaceSize))); + + int numAlgos = 0; + const int requestedAlgos = 1; + hipblasLtMatmulHeuristicResult_t heuristicResult[requestedAlgos]; + CHECK_HIPBLASLT_ERROR(hipblasLtMatmulAlgoGetHeuristic( + handle, matmul_desc, matA, matB, matC, matD, pref, requestedAlgos, heuristicResult, &numAlgos)); + + if(numAlgos == 0) + { + CHECK_HIPBLASLT_ERROR(hipblasLtMatmulPreferenceDestroy(pref)); + GTEST_SKIP() << "No algorithm found for this configuration"; + } + + // Allocate workspace + size_t workspaceSize = heuristicResult[0].workspaceSize; + void* d_workspace = nullptr; + if(workspaceSize > 0) + { + CHECK_HIP_ERROR(hipMalloc(&d_workspace, workspaceSize)); + } + + CHECK_HIPBLASLT_ERROR(hipblasLtMatmul(handle, + matmul_desc, + &h_alpha, + d_batch_A, + matA, + d_batch_B, + matB, + &h_beta, + d_batch_C, + matC, + d_batch_D, + matD, + &heuristicResult[0].algo, + d_workspace, + workspaceSize, + 0)); + + CHECK_HIPBLASLT_ERROR(hipblasLtMatmulPreferenceDestroy(pref)); + + // Ensure kernel completes before reading result + CHECK_HIP_ERROR(hipDeviceSynchronize()); + + if(d_workspace) + { + CHECK_HIP_ERROR(hipFree(d_workspace)); + } + + // Copy GPU result + CHECK_HIP_ERROR(hipMemcpy(h_D_full.data(), + d_D_full, + sizeof(To) * size_D_full * batch_count, + hipMemcpyDeviceToHost)); + + // ======================================== + // CPU Reference: D = alpha * A * B + beta * C + // Simple manual GEMM for testing (avoids cblas_gemm complexity) + // ======================================== + for(int b = 0; b < batch_count; b++) + { + // Get pointers to sub-matrices (with element offset applied) + Ti* A_sub = h_A_full.data() + b * size_A_full + offset_a_elem; + Ti* B_sub = h_B_full.data() + b * size_B_full + offset_b_elem; + To* C_sub = h_C_full.data() + b * size_C_full + offset_c_elem; + To* D_sub = h_D_gold.data() + b * size_D_sub; + + // Simple GEMM: D = alpha * A * B + beta * C + for(int64_t i = 0; i < M; i++) + { + for(int64_t j = 0; j < N; j++) + { + Tc sum = 0; + for(int64_t k = 0; k < K; k++) + { + // A is A_row x A_col, B is B_row x B_col + // For transA=N: A(i,k) = A[i + k*lda] + // For transA=T: A(i,k) = A[k + i*lda] + // For transB=N: B(k,j) = B[k + j*ldb] + // For transB=T: B(k,j) = B[j + k*ldb] + Tc a_val = (transA == HIPBLAS_OP_N) + ? Tc(A_sub[i + k * lda]) + : Tc(A_sub[k + i * lda]); + Tc b_val = (transB == HIPBLAS_OP_N) + ? Tc(B_sub[k + j * ldb]) + : Tc(B_sub[j + k * ldb]); + sum += a_val * b_val; + } + D_sub[i + j * ldd] = To(h_alpha * sum + h_beta * Tc(C_sub[i + j * ldc])); + } + } + } + + // ======================================== + // VALIDATION: Compare GPU vs CPU + // ======================================== + double max_error = 0.0; + for(int b = 0; b < batch_count; b++) + { + // GPU result is at (base + offset) within each batch's buffer + To* result_gpu = h_D_full.data() + b * size_D_full + offset_d_elem; + To* result_cpu = h_D_gold.data() + b * size_D_sub; + + for(size_t i = 0; i < size_D_sub; i++) + { + double diff = std::abs(double(result_gpu[i]) - double(result_cpu[i])); + max_error = std::max(max_error, diff); + } + } + + // Cleanup + CHECK_HIP_ERROR(hipFree(d_batch_A)); + CHECK_HIP_ERROR(hipFree(d_batch_B)); + CHECK_HIP_ERROR(hipFree(d_batch_C)); + CHECK_HIP_ERROR(hipFree(d_batch_D)); + + CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutDestroy(matA)); + CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutDestroy(matB)); + CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutDestroy(matC)); + CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutDestroy(matD)); + CHECK_HIPBLASLT_ERROR(hipblasLtMatmulDescDestroy(matmul_desc)); + CHECK_HIPBLASLT_ERROR(hipblasLtDestroy(handle)); + + // Tolerance: epsilon * factor * K (accumulation over K elements) + double tol = std::numeric_limits::epsilon() * 100 * K; + + // Report results + if(arg.unit_check) + { + EXPECT_LT(max_error, tol) << "GPU vs CPU mismatch (error: " << max_error << ", tol: " << tol << ")"; + } + + if(arg.norm_check) + { + hipblaslt_cout << "GPU vs CPU max error: " << max_error << " (tol: " << tol << ")" << std::endl; + } +} + +// Type dispatcher based on Arguments +void testing_matmul_batch_offset(const Arguments& arg) +{ + // Dispatch based on data types in Arguments + // For now, support only f32, f16, bf16 with matching input/output types + if(arg.a_type == HIP_R_32F && arg.b_type == HIP_R_32F && + arg.c_type == HIP_R_32F && arg.d_type == HIP_R_32F) + { + testing_matmul_batch_offset_impl(arg); + } + else if(arg.a_type == HIP_R_16F && arg.b_type == HIP_R_16F && + arg.c_type == HIP_R_16F && arg.d_type == HIP_R_16F) + { + testing_matmul_batch_offset_impl(arg); + } + else if(arg.a_type == HIP_R_16BF && arg.b_type == HIP_R_16BF && + arg.c_type == HIP_R_16BF && arg.d_type == HIP_R_16BF) + { + testing_matmul_batch_offset_impl(arg); + } + else + { + GTEST_SKIP() << "Unsupported type combination for batch_offset test"; + } +} diff --git a/projects/hipblaslt/clients/tests/data/hipblaslt_common.yaml b/projects/hipblaslt/clients/tests/data/hipblaslt_common.yaml index ab4e28b116be..e4d4b63fdb66 100644 --- a/projects/hipblaslt/clients/tests/data/hipblaslt_common.yaml +++ b/projects/hipblaslt/clients/tests/data/hipblaslt_common.yaml @@ -597,6 +597,10 @@ Arguments: - lde: c_int64*32 - batch_count: c_int32 - batch_mode: c_int32 + - batch_offset_a: c_int64 + - batch_offset_b: c_int64 + - batch_offset_c: c_int64 + - batch_offset_d: c_int64 - iters: c_int32 - cold_iters: c_int32 - warmup_time: c_float @@ -723,6 +727,10 @@ Defaults: transB: '*' batch_count: 1 batch_mode: 0 + batch_offset_a: 0 + batch_offset_b: 0 + batch_offset_c: 0 + batch_offset_d: 0 HMM: false pad: 4096 threads: 0 diff --git a/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml b/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml index 915507cbb655..aa4b93442c4c 100755 --- a/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml +++ b/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml @@ -3211,4 +3211,152 @@ Tests: requested_solution_num: 10 gpu_arch: '950' -... +# ============================================================================== +# Batch Offset Tests - 64-bit offset support for general batched GEMM +# ============================================================================== + +# Quick validation test - verifies basic offset functionality +- name: matmul_batch_offset_quick + category: quick + function: matmul_batch_offset + precision: *real_precisions + transA: N + transB: N + M: 256 + N: 128 + K: 64 + lda: 256 + ldb: 64 + ldc: 256 + ldd: 256 + batch_mode: 1 # Pointer array mode + batch_count: 2 + batch_offset_a: 0 + batch_offset_b: 64 + batch_offset_c: 128 + batch_offset_d: 256 + alpha: 1.0 + beta: 0.0 + unit_check: 1 + norm_check: 1 + +# Offset variation test - various offset values +# Note: Uses M=256 to avoid known General Batched GEMM issue with larger sizes +- name: matmul_batch_offset_values + category: pre_checkin + function: matmul_batch_offset + precision: *real_precisions + transA: N + transB: N + M: 256 + N: 128 + K: 128 + lda: 256 + ldb: 128 + ldc: 256 + ldd: 256 + batch_mode: 1 + batch_count: 3 + batch_offset_a: [0, 64, 256, 512] + batch_offset_b: [0, 64, 256, 512] + batch_offset_c: [0, 64, 256, 512] + batch_offset_d: [0, 64, 256, 512] + alpha: 1.0 + beta: [0.0, 1.0] + unit_check: 1 + norm_check: 1 + +# Transpose with offset +# Note: lda/ldb must be valid for all transpose combinations +# For transA=N: lda >= M, for transA=T: lda >= K +# For transB=N: ldb >= K, for transB=T: ldb >= N +# Using max(M,K)=256 for lda and max(K,N)=256 for ldb to cover all cases +- name: matmul_batch_offset_transpose + category: pre_checkin + function: matmul_batch_offset + precision: *real_precisions + transA: [N, T] + transB: [N, T] + M: 256 + N: 256 + K: 128 + lda: 256 + ldb: 256 + ldc: 256 + ldd: 256 + batch_mode: 1 + batch_count: 2 + batch_offset_a: 128 + batch_offset_b: 128 + batch_offset_c: 128 + batch_offset_d: 128 + alpha: 1.0 + beta: 0.5 + unit_check: 1 + norm_check: 1 + +# Alpha/Beta edge cases with offset +- name: matmul_batch_offset_alpha_beta + category: pre_checkin + function: matmul_batch_offset + precision: *real_precisions + transA: N + transB: N + M: 256 + N: 128 + K: 64 + batch_mode: 1 + batch_count: 2 + batch_offset_a: 64 + batch_offset_b: 64 + batch_offset_c: 64 + batch_offset_d: 64 + alpha_beta: *alpha_beta_range + unit_check: 1 + norm_check: 1 + +# Debug test - 1024x1024 using regular matmul with general batched (batch_mode=1) +- name: matmul_general_batched_debug + category: quick + function: matmul + precision: *real_precisions + transA: N + transB: N + M: 1024 + N: 1024 + K: 256 + lda: 1024 + ldb: 256 + ldc: 1024 + ldd: 1024 + batch_mode: 1 + batch_count: 2 + alpha: 1.0 + beta: 1.0 + unit_check: 1 + norm_check: 1 + +# Large offset test +- name: matmul_batch_offset_large + category: nightly + function: matmul_batch_offset + precision: *real_precisions + transA: N + transB: N + M: 256 #1024 + N: 256 #1024 + K: 256 + lda: 1024 + ldb: 256 + ldc: 1024 + ldd: 1024 + batch_mode: 1 + batch_count: 4 + batch_offset_a: 4096 # [1024, 4096] + batch_offset_b: 2048 # [512, 2048] + batch_offset_c: 4096 # [1024, 4096] + batch_offset_d: 4096 # [1024, 4096] + alpha: 1.0 + beta: 1.0 + unit_check: 1 + norm_check: 1 diff --git a/projects/hipblaslt/clients/tests/src/matmul_gtest.cpp b/projects/hipblaslt/clients/tests/src/matmul_gtest.cpp index 5d278d315add..86f234388aef 100644 --- a/projects/hipblaslt/clients/tests/src/matmul_gtest.cpp +++ b/projects/hipblaslt/clients/tests/src/matmul_gtest.cpp @@ -27,6 +27,7 @@ #include "hipblaslt_datatype2string.hpp" #include "hipblaslt_test.hpp" #include "testing_matmul.hpp" +#include "testing_matmul_batch_offset.hpp" #include #include #include @@ -48,6 +49,8 @@ namespace testing_matmul(arg); else if(!strcmp(arg.function, "matmul_bad_arg")) testing_matmul_bad_arg(arg); + else if(!strcmp(arg.function, "matmul_batch_offset")) + testing_matmul_batch_offset(arg); else FAIL() << "Internal error: Test called with unknown function: " << arg.function; } @@ -64,7 +67,8 @@ namespace // Filter for which functions apply to this suite static bool function_filter(const Arguments& arg) { - return !strcmp(arg.function, "matmul") || !strcmp(arg.function, "matmul_bad_arg"); + return !strcmp(arg.function, "matmul") || !strcmp(arg.function, "matmul_bad_arg") + || !strcmp(arg.function, "matmul_batch_offset"); } // Google Test name suffix based on parameters From 842306adf1cc85f4e001cf4911298d0adf041b17 Mon Sep 17 00:00:00 2001 From: Seung-Hee Bae Date: Thu, 28 May 2026 23:47:20 +0000 Subject: [PATCH 08/27] [hipblaslt] update tensilelite code-gen part for 64bit offset support to add arguments appropriately and to use only two extra SGPRs --- .../rocblaslt/src/rocblaslt_mat.cpp | 4 + .../Tensile/Components/Signature.py | 19 ++--- .../tensilelite/Tensile/KernelWriter.py | 7 +- .../Tensile/KernelWriterAssembly.py | 80 ++++++++++++++----- .../tensilelite/src/ContractionSolution.cpp | 53 ++++++------ 5 files changed, 100 insertions(+), 63 deletions(-) diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_mat.cpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_mat.cpp index 51461365ccd7..679b19eb8779 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_mat.cpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_mat.cpp @@ -131,6 +131,8 @@ rocblaslt_status rocblaslt_matmul_impl(const rocblaslt_handle handle, hipDataType scale_type = matmul_descr->scale_type; // Others + // Use strided_batch=true for kernel selection (StridedBatched=true kernels with SupportUserArgs) + // The actual batch mode is tracked via problem.batchMode() for argument passing bool strided_batch = true; bool grouped_gemm = false; @@ -346,6 +348,8 @@ rocblaslt_status rocblaslt_gemm_create_cpp_impl(const rocblaslt_handle void* amaxD = matmul_descr->amaxD; // Others + // Use strided_batch=true for kernel selection (StridedBatched=true kernels with SupportUserArgs) + // The actual batch mode is tracked via problem.batchMode() for argument passing bool strided_batch = true; bool grouped_gemm = false; diff --git a/projects/hipblaslt/tensilelite/Tensile/Components/Signature.py b/projects/hipblaslt/tensilelite/Tensile/Components/Signature.py index a6fad12207fb..0ea36aa48ebc 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Components/Signature.py +++ b/projects/hipblaslt/tensilelite/Tensile/Components/Signature.py @@ -163,15 +163,6 @@ def __call__(self, writer) -> SignatureBase: if kernel["ProblemType"]["Sparse"]: signature.addArg("MetaData", SVK.SIG_GLOBALBUFFER, "void" , "generic") - # Batch offset support for general batched mode (pointer array) - # Only for non-strided, non-grouped GEMM - if not kernel["ProblemType"]["StridedBatched"] and not kernel["ProblemType"]["GroupedGemm"]: - signature.addArg("batchOffsetD", SVK.SIG_VALUE, "u64") - signature.addArg("batchOffsetC", SVK.SIG_VALUE, "u64") - signature.addArg("batchOffsetA", SVK.SIG_VALUE, "u64") - signature.addArg("batchOffsetB", SVK.SIG_VALUE, "u64") - userArgumentsInfo.gemmArgumentSize += 32 # 4 offsets * 8 bytes each - if kernel["StreamK"] > 0 and kernel["StreamKAtomic"] == 0: signature.addArg("AddressWS", SVK.SIG_GLOBALBUFFER, cptValueType, "generic") signature.addArg("AddressFlags", SVK.SIG_GLOBALBUFFER, dstValueType, "generic") @@ -263,6 +254,16 @@ def __call__(self, writer) -> SignatureBase: signature.addArg("skTiles", SVK.SIG_VALUE, "u32") userArgumentsInfo.gemmArgumentSize += 8 + # Batch offset support for general batched mode (pointer array) + # Placed after core GEMM args (strides, alpha/beta, StreamK) + # Enabled for all SupportUserArgs kernels (offsets are 0 when not in General Batched mode) + if kernel["ProblemType"]["SupportUserArgs"] and not kernel["ProblemType"]["GroupedGemm"]: + signature.addArg("batchOffsetD", SVK.SIG_VALUE, "u64") + signature.addArg("batchOffsetC", SVK.SIG_VALUE, "u64") + signature.addArg("batchOffsetA", SVK.SIG_VALUE, "u64") + signature.addArg("batchOffsetB", SVK.SIG_VALUE, "u64") + userArgumentsInfo.gemmArgumentSize += 32 # 4 offsets * 8 bytes each + if kernel["ProblemType"]["UseScaleAB"]: signature.addArg("AddressScaleA", SVK.SIG_GLOBALBUFFER, cptValueType, "generic") signature.addArg("AddressScaleB", SVK.SIG_GLOBALBUFFER, cptValueType, "generic") diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py index 59f454f87825..3da1a6e1def9 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py @@ -9234,11 +9234,8 @@ def vgprAllocationImplSubtile(): self.defineSgpr("StridesMetadata", self.states.m.numSgprStrides) # Batch offset support for general batched GEMM (pointer array mode) - if not kernel["ProblemType"]["StridedBatched"] and not kernel["ProblemType"]["GroupedGemm"]: - self.defineSgpr("BatchOffsetD", 2) # 64-bit offset requires 2 SGPRs - self.defineSgpr("BatchOffsetC", 2) # 64-bit offset requires 2 SGPRs - self.defineSgpr("BatchOffsetA", 2) # 64-bit offset requires 2 SGPRs - self.defineSgpr("BatchOffsetB", 2) # 64-bit offset requires 2 SGPRs + # Offsets are loaded on-demand from kernel arguments to avoid using persistent SGPRs + # No SGPR allocation here - offsets loaded directly when applying to addresses # for packed batches without stride restrictions need to do something different here assert sorted(kernel["PackedC0IdxChars"]+kernel["PackedC1IdxChars"]) == \ diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py index cd583a625db6..c7aee9d3d53d 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py @@ -2131,9 +2131,13 @@ def loadBatchedAddress(self, kernel, Batch, tmpSgprResource: ContinuousRegister) module.add(SMulI32(dst=sgpr(tmpSgpr), src0=sgpr(Batch), src1=0x8, comment="offset of global buffer address")) module.add(SLoadB64(dst=sgpr("AddressD", 2), base=sgpr("AddressD",2), soffset=sgpr(tmpSgpr), comment="load global buffer D address")) # Apply batch offset to AddressD for general batched mode - if not kernel["ProblemType"]["StridedBatched"] and not kernel["ProblemType"]["GroupedGemm"]: - module.add(SWaitCnt(kmcnt=0, comment="Wait for the Matrix Address Load from the Pointer Array before updating offset")) - module.add(SAddU64(dst=sgpr("AddressD", 2), src0=sgpr("AddressD", 2), src1=sgpr("BatchOffsetD", 2), comment="add batch offset to D address")) + # Load offset on-demand from kernel args (no persistent SGPRs needed) + if kernel["ProblemType"]["SupportUserArgs"] and not kernel["ProblemType"]["GroupedGemm"]: + with self.allocTmpSgpr(2) as tmpSgprOffset: + module.add(self.argLoader.loadKernArg(tmpSgprOffset.idx, "KernArgAddress", sgprOffset=hex(self.states.batchOffsetDKernArgOffset), dword=2)) + module.add(SWaitCnt(kmcnt=0, comment="Wait for the Matrix Address and Batch Offset Loads")) + module.add(SAddU32(dst=sgpr("AddressD+0"), src0=sgpr("AddressD+0"), src1=sgpr(tmpSgprOffset.idx), comment="add batch offset to D address (low)")) + module.add(SAddCU32(dst=sgpr("AddressD+1"), src0=sgpr("AddressD+1"), src1=sgpr(tmpSgprOffset.idx+1), comment="add batch offset to D address (high)")) # Only load C buffer address if Beta is used and potentially non-zero if kernel["ProblemType"]["UseBeta"]: @@ -2146,9 +2150,13 @@ def loadBatchedAddress(self, kernel, Batch, tmpSgprResource: ContinuousRegister) module.add(SMulI32(dst=sgpr(tmpSgpr), src0=sgpr(Batch), src1=0x8, comment="offset of global buffer address")) module.add(SLoadB64(dst=sgpr("AddressC", 2), base=sgpr("AddressC",2), soffset=sgpr(tmpSgpr), comment="load global buffer C address")) # Apply batch offset to AddressC for general batched mode - if not kernel["ProblemType"]["StridedBatched"] and not kernel["ProblemType"]["GroupedGemm"]: - module.add(SWaitCnt(kmcnt=0, comment="Wait for the Matrix Address Load from the Pointer Array before updating offset")) - module.add(SAddU64(dst=sgpr("AddressC", 2), src0=sgpr("AddressC", 2), src1=sgpr("BatchOffsetC", 2), comment="add batch offset to C address")) + # Load offset on-demand from kernel args (no persistent SGPRs needed) + if kernel["ProblemType"]["SupportUserArgs"] and not kernel["ProblemType"]["GroupedGemm"]: + with self.allocTmpSgpr(2) as tmpSgprOffset: + module.add(self.argLoader.loadKernArg(tmpSgprOffset.idx, "KernArgAddress", sgprOffset=hex(self.states.batchOffsetCKernArgOffset), dword=2)) + module.add(SWaitCnt(kmcnt=0, comment="Wait for the Matrix Address and Batch Offset Loads")) + module.add(SAddU32(dst=sgpr("AddressC+0"), src0=sgpr("AddressC+0"), src1=sgpr(tmpSgprOffset.idx), comment="add batch offset to C address (low)")) + module.add(SAddCU32(dst=sgpr("AddressC+1"), src0=sgpr("AddressC+1"), src1=sgpr(tmpSgprOffset.idx+1), comment="add batch offset to C address (high)")) module.add(endCheckLabel) @@ -2168,11 +2176,19 @@ def loadBatchedAddress(self, kernel, Batch, tmpSgprResource: ContinuousRegister) module.add(SLoadB64(dst=sgpr("AddressA", 2), base=sgpr("AddressA",2), soffset=sgpr(tmpSgpr), comment="load global buffer A address")) module.add(SLoadB64(dst=sgpr("AddressB", 2), base=sgpr("AddressB",2), soffset=sgpr(tmpSgpr), comment="load global buffer B address")) # Apply batch offset to AddressA and AddressB for general batched mode - if not kernel["ProblemType"]["StridedBatched"] and not kernel["ProblemType"]["GroupedGemm"]: - module.add(SWaitCnt(kmcnt=1, comment="Wait for the AddressA Load from the Pointer Array before updating offset")) - module.add(SAddU64(dst=sgpr("AddressA", 2), src0=sgpr("AddressA", 2), src1=sgpr("BatchOffsetA", 2), comment="add batch offset to A address")) - module.add(SWaitCnt(kmcnt=0, comment="Wait for the AddressB Load from the Pointer Array before updating offset")) - module.add(SAddU64(dst=sgpr("AddressB", 2), src0=sgpr("AddressB", 2), src1=sgpr("BatchOffsetB", 2), comment="add batch offset to B address")) + # Load offsets on-demand from kernel args (no persistent SGPRs needed) + if kernel["ProblemType"]["SupportUserArgs"] and not kernel["ProblemType"]["GroupedGemm"]: + with self.allocTmpSgpr(2) as tmpSgprOffset: + # Load and apply offset A + module.add(self.argLoader.loadKernArg(tmpSgprOffset.idx, "KernArgAddress", sgprOffset=hex(self.states.batchOffsetAKernArgOffset), dword=2)) + module.add(SWaitCnt(kmcnt=1, comment="Wait for the AddressA Load from the Pointer Array")) + module.add(SAddU32(dst=sgpr("AddressA+0"), src0=sgpr("AddressA+0"), src1=sgpr(tmpSgprOffset.idx), comment="add batch offset to A address (low)")) + module.add(SAddCU32(dst=sgpr("AddressA+1"), src0=sgpr("AddressA+1"), src1=sgpr(tmpSgprOffset.idx+1), comment="add batch offset to A address (high)")) + # Load and apply offset B (reuse same tmpSgprOffset) + module.add(self.argLoader.loadKernArg(tmpSgprOffset.idx, "KernArgAddress", sgprOffset=hex(self.states.batchOffsetBKernArgOffset), dword=2)) + module.add(SWaitCnt(kmcnt=0, comment="Wait for the AddressB and Batch Offset Loads")) + module.add(SAddU32(dst=sgpr("AddressB+0"), src0=sgpr("AddressB+0"), src1=sgpr(tmpSgprOffset.idx), comment="add batch offset to B address (low)")) + module.add(SAddCU32(dst=sgpr("AddressB+1"), src0=sgpr("AddressB+1"), src1=sgpr(tmpSgprOffset.idx+1), comment="add batch offset to B address (high)")) module.add(endCheckLabel) @@ -2199,17 +2215,22 @@ def getKernelArgLoadModule(self, kernel, sgprStartIdx, numsOfLoad, preloadNum): sgprOffset += 4 self.argLoader.setOffset(sgprOffset) - # Load batch offset arguments for general batched mode - if not kernel["ProblemType"]["StridedBatched"] and not kernel["ProblemType"]["GroupedGemm"]: + # Batch offset arguments for general batched mode + # Store kernarg offsets for on-demand loading (no persistent SGPRs needed) + if kernel["ProblemType"]["SupportUserArgs"] and not kernel["ProblemType"]["GroupedGemm"]: sgprOffset = self.argLoader.getOffset() - kernelArgs.add(self.argLoader.loadKernArg("BatchOffsetD", "KernArgAddress", sgprOffset=hex(sgprOffset), dword=2)) + self.states.batchOffsetDKernArgOffset = sgprOffset sgprOffset += 8 # 64-bit = 8 bytes - kernelArgs.add(self.argLoader.loadKernArg("BatchOffsetC", "KernArgAddress", sgprOffset=hex(sgprOffset), dword=2)) + self.states.batchOffsetCKernArgOffset = sgprOffset sgprOffset += 8 - kernelArgs.add(self.argLoader.loadKernArg("BatchOffsetA", "KernArgAddress", sgprOffset=hex(sgprOffset), dword=2)) + self.states.batchOffsetAKernArgOffset = sgprOffset sgprOffset += 8 - kernelArgs.add(self.argLoader.loadKernArg("BatchOffsetB", "KernArgAddress", sgprOffset=hex(sgprOffset), dword=2)) + self.states.batchOffsetBKernArgOffset = sgprOffset sgprOffset += 8 + self.argLoader.setOffset(sgprOffset) + # Debug: write kernel arg offsets to file for verification + with open("/tmp/kernel_batch_offset_debug.txt", "a") as f: + f.write(f"{self.states.kernelName}: D@{self.states.batchOffsetDKernArgOffset}, C@{self.states.batchOffsetCKernArgOffset}, A@{self.states.batchOffsetAKernArgOffset}, B@{self.states.batchOffsetBKernArgOffset}\n") return kernelArgs @@ -4665,7 +4686,15 @@ def computeLoadSrd(self, kernel, tP, tc, indices, bpe): moduleLoadGeneralBatch.add(SAddU32(dst=sgpr(stmp+0), src0=sgpr(stmp+0), src1=sgpr("Address%s+0"%tc), comment="Offsetting to the location [Lower half of address]")) moduleLoadGeneralBatch.add(SAddCU32(dst=sgpr(stmp+1), src0=sgpr("Address%s+1"%tc), src1=0, comment="Offsetting to the location [Higher half of address]")) moduleLoadGeneralBatch.add(SLoadB64(dst=sgpr("Srd%s"%tc, 2), base=sgpr(stmp, 2), soffset=0, comment="Load the Matrix Address in the Pointer Array")) - moduleLoadGeneralBatch.add(SWaitCnt(kmcnt=0, comment="Wait for the Matrix Address Load from the Pointer Array")) + # Load and apply batch offset for General Batched GEMM + if kernel["ProblemType"]["SupportUserArgs"] and not kernel["ProblemType"]["GroupedGemm"]: + batchOffsetKernArgOffset = self.states.batchOffsetAKernArgOffset if tc == "A" else self.states.batchOffsetBKernArgOffset + moduleLoadGeneralBatch.add(SLoadB64(dst=sgpr(stmp, 2), base=sgpr("KernArgAddress", 2), soffset=hex(batchOffsetKernArgOffset), comment="Load batchOffset%s from kernel args"%tc)) + moduleLoadGeneralBatch.add(SWaitCnt(kmcnt=0, comment="Wait for Matrix Address and Batch Offset Loads")) + moduleLoadGeneralBatch.add(SAddU32(dst=sgpr("Srd%s+0"%tc), src0=sgpr("Srd%s+0"%tc), src1=sgpr(stmp+0), comment="Add batch offset to %s address (low)"%tc)) + moduleLoadGeneralBatch.add(SAddCU32(dst=sgpr("Srd%s+1"%tc), src0=sgpr("Srd%s+1"%tc), src1=sgpr(stmp+1), comment="Add batch offset to %s address (high)"%tc)) + else: + moduleLoadGeneralBatch.add(SWaitCnt(kmcnt=0, comment="Wait for the Matrix Address Load from the Pointer Array")) if self.states.groOffsetInMacroTile and ((tc == "A" and not kernel["enableTDMA"]) or (tc == "B" and not kernel["enableTDMB"])): prePad1 = int(self.states.srdShiftLeft[tc] * tP["bpeGR"]) # leave room in case we have to pointer shift moduleLoadGeneralBatch.add(SSubU32(dst=sgpr("Srd%s+0"%tc), src0=sgpr("Srd%s+0"%tc), src1=prePad1, comment="pre-pad to make room for possible pointer shift")) @@ -13296,9 +13325,18 @@ def computeStoreSrdStart(self, kernel, srdTcList: list, sgprBpeList = [], useSiz module.add(SAddU32(dst=sgpr(tmpS0), src0=sgpr(tmpS0), src1=sgpr("Address%s+0"%mat), comment="Offsetting to the location [Lower half of address]")) module.add(SAddCU32(dst=sgpr(tmpS1), src0=sgpr("Address%s+1"%mat), src1=0, comment="Offsetting to the location [Higher half of address]")) module.add(SLoadB64(dst=sgpr(tmpS0, 2), base=sgpr(tmpS0, 2), soffset=0, comment="Load the Matrix Address in the Pointer Array")) - module.add(SWaitCnt(kmcnt=0, comment="Wait for the Matrix Address Load from the Pointer Array")) - module.add(SAddU32(dst=sgpr("Srd%s+0"%mat), src0=sgpr("Srd%s+0"%mat), src1=sgpr(tmpS0), comment="Offsetting within the Batch Matrix [Lower half of address]")) - module.add(SAddCU32(dst=sgpr("Srd%s+1"%mat), src0=sgpr("Srd%s+1"%mat), src1=sgpr(tmpS1), comment="Offsetting within the Batch Matrix [Higher half of address]")) + # Load and apply batch offset for General Batched GEMM (C or D matrix) + if kernel["ProblemType"]["SupportUserArgs"] and not kernel["ProblemType"]["GroupedGemm"]: + batchOffsetKernArgOffset = self.states.batchOffsetCKernArgOffset if mat == "C" else self.states.batchOffsetDKernArgOffset + module.add(SLoadB64(dst=sgpr("Srd%s"%mat, 2), base=sgpr("KernArgAddress", 2), soffset=hex(batchOffsetKernArgOffset), comment="Load batchOffset%s from kernel args"%mat)) + module.add(SWaitCnt(kmcnt=0, comment="Wait for Matrix Address and Batch Offset Loads")) + # Add loaded matrix address to SRD + module.add(SAddU32(dst=sgpr("Srd%s+0"%mat), src0=sgpr("Srd%s+0"%mat), src1=sgpr(tmpS0), comment="Add matrix address to SRD (low)")) + module.add(SAddCU32(dst=sgpr("Srd%s+1"%mat), src0=sgpr("Srd%s+1"%mat), src1=sgpr(tmpS1), comment="Add matrix address to SRD (high)")) + else: + module.add(SWaitCnt(kmcnt=0, comment="Wait for the Matrix Address Load from the Pointer Array")) + module.add(SAddU32(dst=sgpr("Srd%s+0"%mat), src0=sgpr("Srd%s+0"%mat), src1=sgpr(tmpS0), comment="Offsetting within the Batch Matrix [Lower half of address]")) + module.add(SAddCU32(dst=sgpr("Srd%s+1"%mat), src0=sgpr("Srd%s+1"%mat), src1=sgpr(tmpS1), comment="Offsetting within the Batch Matrix [Higher half of address]")) module.add(generalBatchedGemmLoad_End) module.addSpaceLine() diff --git a/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp b/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp index dd1766bc5101..917f769e715b 100644 --- a/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp +++ b/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp @@ -620,12 +620,6 @@ namespace TensileLite { args.template append("batchD", inputs.batchD); args.template append("batchC", inputs.batchC); - - if(problem.batchMode() == ContractionProblemGemm::BATCHMODE::POINTER_ARRAY) - { - args.template append("batchOffsetD", inputs.batchOffsetD); - args.template append("batchOffsetC", inputs.batchOffsetC); - } } if(problemType.stridedBatched) @@ -643,12 +637,6 @@ namespace TensileLite { args.template append("batchA", inputs.batchA); args.template append("batchB", inputs.batchB); - - if(problem.batchMode() == ContractionProblemGemm::BATCHMODE::POINTER_ARRAY) - { - args.template append("batchOffsetA", inputs.batchOffsetA); - args.template append("batchOffsetB", inputs.batchOffsetB); - } } if(problemType.sparse) @@ -1000,6 +988,16 @@ namespace TensileLite autoGsuVal, ntab); + // Batch offset support for General Batched GEMM (SupportUserArgs kernels) + // Placed after core GEMM args (strides, alpha/beta, StreamK) to match kernel signature + if(problemType.supportDeviceUserArguments && !problemType.groupedGemm) + { + args.template append("batchOffsetD", inputs.batchOffsetD); + args.template append("batchOffsetC", inputs.batchOffsetC); + args.template append("batchOffsetA", inputs.batchOffsetA); + args.template append("batchOffsetB", inputs.batchOffsetB); + } + // NOTE: an assumption here is A & B must be both MX data types or non-MX data types. // Mixing is not supported. if(!problemType.useScaleAB.empty()) @@ -1847,12 +1845,13 @@ namespace TensileLite rv.args.append("dstD", inputs.d); // MBSK: synchronizer address, MB: null address rv.args.append("Synchronizer", - gsuSettings.globalAccumulation == 3 - ? inputs.Synchronizer + gsuSettings.globalAccumulation == 3 + ? inputs.Synchronizer : NULL); rv.args.append("GSUSync", 0); } + if(problemType.stochasticRounding) { // generate seed from random generator @@ -2125,19 +2124,18 @@ namespace TensileLite else if(problemType.stridedBatched) rv.args.append("D", inputs.d); else - { rv.args.append("batchD", inputs.batchD); - if(problem.batchMode() == ContractionProblemGemm::BATCHMODE::POINTER_ARRAY) - rv.args.append("batchOffsetD", inputs.batchOffsetD); - } if(problemType.stridedBatched) rv.args.append("C", inputs.c); else - { rv.args.append("batchC", inputs.batchC); - if(problem.batchMode() == ContractionProblemGemm::BATCHMODE::POINTER_ARRAY) - rv.args.append("batchOffsetC", inputs.batchOffsetC); + + // Pass batch offsets when kernel expects them (SupportUserArgs=true, not GroupedGemm) + if(problemType.supportDeviceUserArguments && !problemType.groupedGemm) + { + rv.args.append("batchOffsetD", inputs.batchOffsetD); + rv.args.append("batchOffsetC", inputs.batchOffsetC); } if(problemType.useBias && sizeMapping.globalAccumulation == 0 && (!problemType.useGradient)) @@ -2300,21 +2298,20 @@ namespace TensileLite if(problemType.stridedBatched) args.template append("D", inputs.d); else - { args.template append("batchD", inputs.batchD); - if(problem.batchMode() == ContractionProblemGemm::BATCHMODE::POINTER_ARRAY) - args.template append("batchOffsetD", inputs.batchOffsetD); - } args.template append("WS", (uint8_t*)inputs.ws + workspaceOffsetInByte); if(problemType.stridedBatched) args.template append("C", inputs.c); else - { args.template append("batchC", inputs.batchC); - if(problem.batchMode() == ContractionProblemGemm::BATCHMODE::POINTER_ARRAY) - args.template append("batchOffsetC", inputs.batchOffsetC); + + // Pass batch offsets when kernel expects them (SupportUserArgs=true, not GroupedGemm) + if(problemType.supportDeviceUserArguments && !problemType.groupedGemm) + { + args.template append("batchOffsetD", inputs.batchOffsetD); + args.template append("batchOffsetC", inputs.batchOffsetC); } bool useBias = false; From 354afbbcb0a48349b9303c65cb1f8906b7a2e435 Mon Sep 17 00:00:00 2001 From: Seung-Hee Bae Date: Fri, 29 May 2026 21:15:33 +0000 Subject: [PATCH 09/27] [hipblaslt] fix bugs in computeStoreSrd() and remove unnecessary change. --- .../clients/tests/data/matmul_gtest.yaml | 16 ++++--- .../Tensile/KernelWriterAssembly.py | 42 +++---------------- 2 files changed, 16 insertions(+), 42 deletions(-) diff --git a/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml b/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml index aa4b93442c4c..a7b580934b5f 100755 --- a/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml +++ b/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml @@ -3343,8 +3343,8 @@ Tests: precision: *real_precisions transA: N transB: N - M: 256 #1024 - N: 256 #1024 + M: 1024 + N: 1024 K: 256 lda: 1024 ldb: 256 @@ -3352,10 +3352,14 @@ Tests: ldd: 1024 batch_mode: 1 batch_count: 4 - batch_offset_a: 4096 # [1024, 4096] - batch_offset_b: 2048 # [512, 2048] - batch_offset_c: 4096 # [1024, 4096] - batch_offset_d: 4096 # [1024, 4096] + batch_offset_a: [1024, 4096] + batch_offset_b: [512, 2048] + batch_offset_c: [1024, 4096] + batch_offset_d: [1024, 4096] + #batch_offset_a: 4096 + #batch_offset_b: 2048 + #batch_offset_c: 4096 + #batch_offset_d: 4096 alpha: 1.0 beta: 1.0 unit_check: 1 diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py index c7aee9d3d53d..c9e5e6598e5d 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py @@ -2130,14 +2130,6 @@ def loadBatchedAddress(self, kernel, Batch, tmpSgprResource: ContinuousRegister) if not isPackedIndex(kernel,idx): module.add(SMulI32(dst=sgpr(tmpSgpr), src0=sgpr(Batch), src1=0x8, comment="offset of global buffer address")) module.add(SLoadB64(dst=sgpr("AddressD", 2), base=sgpr("AddressD",2), soffset=sgpr(tmpSgpr), comment="load global buffer D address")) - # Apply batch offset to AddressD for general batched mode - # Load offset on-demand from kernel args (no persistent SGPRs needed) - if kernel["ProblemType"]["SupportUserArgs"] and not kernel["ProblemType"]["GroupedGemm"]: - with self.allocTmpSgpr(2) as tmpSgprOffset: - module.add(self.argLoader.loadKernArg(tmpSgprOffset.idx, "KernArgAddress", sgprOffset=hex(self.states.batchOffsetDKernArgOffset), dword=2)) - module.add(SWaitCnt(kmcnt=0, comment="Wait for the Matrix Address and Batch Offset Loads")) - module.add(SAddU32(dst=sgpr("AddressD+0"), src0=sgpr("AddressD+0"), src1=sgpr(tmpSgprOffset.idx), comment="add batch offset to D address (low)")) - module.add(SAddCU32(dst=sgpr("AddressD+1"), src0=sgpr("AddressD+1"), src1=sgpr(tmpSgprOffset.idx+1), comment="add batch offset to D address (high)")) # Only load C buffer address if Beta is used and potentially non-zero if kernel["ProblemType"]["UseBeta"]: @@ -2149,14 +2141,6 @@ def loadBatchedAddress(self, kernel, Batch, tmpSgprResource: ContinuousRegister) if not isPackedIndex(kernel,idx): module.add(SMulI32(dst=sgpr(tmpSgpr), src0=sgpr(Batch), src1=0x8, comment="offset of global buffer address")) module.add(SLoadB64(dst=sgpr("AddressC", 2), base=sgpr("AddressC",2), soffset=sgpr(tmpSgpr), comment="load global buffer C address")) - # Apply batch offset to AddressC for general batched mode - # Load offset on-demand from kernel args (no persistent SGPRs needed) - if kernel["ProblemType"]["SupportUserArgs"] and not kernel["ProblemType"]["GroupedGemm"]: - with self.allocTmpSgpr(2) as tmpSgprOffset: - module.add(self.argLoader.loadKernArg(tmpSgprOffset.idx, "KernArgAddress", sgprOffset=hex(self.states.batchOffsetCKernArgOffset), dword=2)) - module.add(SWaitCnt(kmcnt=0, comment="Wait for the Matrix Address and Batch Offset Loads")) - module.add(SAddU32(dst=sgpr("AddressC+0"), src0=sgpr("AddressC+0"), src1=sgpr(tmpSgprOffset.idx), comment="add batch offset to C address (low)")) - module.add(SAddCU32(dst=sgpr("AddressC+1"), src0=sgpr("AddressC+1"), src1=sgpr(tmpSgprOffset.idx+1), comment="add batch offset to C address (high)")) module.add(endCheckLabel) @@ -2175,20 +2159,6 @@ def loadBatchedAddress(self, kernel, Batch, tmpSgprResource: ContinuousRegister) if not isPackedIndex(kernel,idx): module.add(SLoadB64(dst=sgpr("AddressA", 2), base=sgpr("AddressA",2), soffset=sgpr(tmpSgpr), comment="load global buffer A address")) module.add(SLoadB64(dst=sgpr("AddressB", 2), base=sgpr("AddressB",2), soffset=sgpr(tmpSgpr), comment="load global buffer B address")) - # Apply batch offset to AddressA and AddressB for general batched mode - # Load offsets on-demand from kernel args (no persistent SGPRs needed) - if kernel["ProblemType"]["SupportUserArgs"] and not kernel["ProblemType"]["GroupedGemm"]: - with self.allocTmpSgpr(2) as tmpSgprOffset: - # Load and apply offset A - module.add(self.argLoader.loadKernArg(tmpSgprOffset.idx, "KernArgAddress", sgprOffset=hex(self.states.batchOffsetAKernArgOffset), dword=2)) - module.add(SWaitCnt(kmcnt=1, comment="Wait for the AddressA Load from the Pointer Array")) - module.add(SAddU32(dst=sgpr("AddressA+0"), src0=sgpr("AddressA+0"), src1=sgpr(tmpSgprOffset.idx), comment="add batch offset to A address (low)")) - module.add(SAddCU32(dst=sgpr("AddressA+1"), src0=sgpr("AddressA+1"), src1=sgpr(tmpSgprOffset.idx+1), comment="add batch offset to A address (high)")) - # Load and apply offset B (reuse same tmpSgprOffset) - module.add(self.argLoader.loadKernArg(tmpSgprOffset.idx, "KernArgAddress", sgprOffset=hex(self.states.batchOffsetBKernArgOffset), dword=2)) - module.add(SWaitCnt(kmcnt=0, comment="Wait for the AddressB and Batch Offset Loads")) - module.add(SAddU32(dst=sgpr("AddressB+0"), src0=sgpr("AddressB+0"), src1=sgpr(tmpSgprOffset.idx), comment="add batch offset to B address (low)")) - module.add(SAddCU32(dst=sgpr("AddressB+1"), src0=sgpr("AddressB+1"), src1=sgpr(tmpSgprOffset.idx+1), comment="add batch offset to B address (high)")) module.add(endCheckLabel) @@ -13325,18 +13295,18 @@ def computeStoreSrdStart(self, kernel, srdTcList: list, sgprBpeList = [], useSiz module.add(SAddU32(dst=sgpr(tmpS0), src0=sgpr(tmpS0), src1=sgpr("Address%s+0"%mat), comment="Offsetting to the location [Lower half of address]")) module.add(SAddCU32(dst=sgpr(tmpS1), src0=sgpr("Address%s+1"%mat), src1=0, comment="Offsetting to the location [Higher half of address]")) module.add(SLoadB64(dst=sgpr(tmpS0, 2), base=sgpr(tmpS0, 2), soffset=0, comment="Load the Matrix Address in the Pointer Array")) - # Load and apply batch offset for General Batched GEMM (C or D matrix) + module.add(SWaitCnt(kmcnt=0, comment="Wait for the Matrix Address Load from the Pointer Array")) + module.add(SAddU32(dst=sgpr("Srd%s+0"%mat), src0=sgpr("Srd%s+0"%mat), src1=sgpr(tmpS0), comment="Offsetting within the Batch Matrix [Lower half of address]")) + module.add(SAddCU32(dst=sgpr("Srd%s+1"%mat), src0=sgpr("Srd%s+1"%mat), src1=sgpr(tmpS1), comment="Offsetting within the Batch Matrix [Higher half of address]")) + # Now, we have starting matrix address of a specific batch in the corresponding Srd. + # Load and apply batch offset for General Batched GEMM (C or D matrix) as necessary. if kernel["ProblemType"]["SupportUserArgs"] and not kernel["ProblemType"]["GroupedGemm"]: batchOffsetKernArgOffset = self.states.batchOffsetCKernArgOffset if mat == "C" else self.states.batchOffsetDKernArgOffset - module.add(SLoadB64(dst=sgpr("Srd%s"%mat, 2), base=sgpr("KernArgAddress", 2), soffset=hex(batchOffsetKernArgOffset), comment="Load batchOffset%s from kernel args"%mat)) + module.add(SLoadB64(dst=sgpr(tmpS0, 2), base=sgpr("KernArgAddress", 2), soffset=hex(batchOffsetKernArgOffset), comment="Load batchOffset%s from kernel args"%mat)) module.add(SWaitCnt(kmcnt=0, comment="Wait for Matrix Address and Batch Offset Loads")) # Add loaded matrix address to SRD module.add(SAddU32(dst=sgpr("Srd%s+0"%mat), src0=sgpr("Srd%s+0"%mat), src1=sgpr(tmpS0), comment="Add matrix address to SRD (low)")) module.add(SAddCU32(dst=sgpr("Srd%s+1"%mat), src0=sgpr("Srd%s+1"%mat), src1=sgpr(tmpS1), comment="Add matrix address to SRD (high)")) - else: - module.add(SWaitCnt(kmcnt=0, comment="Wait for the Matrix Address Load from the Pointer Array")) - module.add(SAddU32(dst=sgpr("Srd%s+0"%mat), src0=sgpr("Srd%s+0"%mat), src1=sgpr(tmpS0), comment="Offsetting within the Batch Matrix [Lower half of address]")) - module.add(SAddCU32(dst=sgpr("Srd%s+1"%mat), src0=sgpr("Srd%s+1"%mat), src1=sgpr(tmpS1), comment="Offsetting within the Batch Matrix [Higher half of address]")) module.add(generalBatchedGemmLoad_End) module.addSpaceLine() From 46bd945437afffe7e2879b5f2dffc0f0b3d16b8a Mon Sep 17 00:00:00 2001 From: Seung-Hee Bae Date: Sat, 30 May 2026 07:04:38 +0000 Subject: [PATCH 10/27] [hipblaslt] update 64-bit offset support for BetaOnly and Conversion. --- .../tensilelite/Tensile/KernelWriterBetaOnly.py | 15 +++++++++++++-- .../tensilelite/Tensile/KernelWriterConversion.py | 15 +++++++++++++-- 2 files changed, 26 insertions(+), 4 deletions(-) diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriterBetaOnly.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriterBetaOnly.py index 5f1a31c432d6..3c55f403fbc2 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriterBetaOnly.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriterBetaOnly.py @@ -89,6 +89,11 @@ def functionSignature(self): batch = "" if isStridedBuffer else "Batch" kStr += " " + ptrStr + " const * " + batch + "C," + self.endLine + # Batch offsets for General Batched GEMM (SupportUserArgs kernels) + if self.state["ProblemType"]["SupportUserArgs"]: + kStr += " int64_t const batchOffsetD,%s" % (self.endLine) + kStr += " int64_t const batchOffsetC,%s" % (self.endLine) + # bias if self.state["ProblemType"]["BetaOnlyUseBias"]: biasPtrStr = self.state["ProblemType"]["BiasDataType"].toDevice(self.language) @@ -217,10 +222,16 @@ def kernelBodyBetaOnly(self): if not self.state["_GlobalAccumulation"]: ptrStr = self.state["ProblemType"]["DestDataType"].toDevice(self.language) - kStr += " " + ptrStr + " * D = BatchD[wg];" + self.endLine + if self.state["ProblemType"]["SupportUserArgs"]: + kStr += " " + ptrStr + " * D = BatchD[wg] + batchOffsetD;" + self.endLine + else: + kStr += " " + ptrStr + " * D = BatchD[wg];" + self.endLine ptrStr = self.state["ProblemType"]["DestDataType"].toDevice(self.language) zeroStr = self.state["ProblemType"]["ComputeDataType"].zeroString(self.language, 1) - kStr += " " + ptrStr + f" const* C = (beta == {zeroStr}) ? nullptr : BatchC[wg];" + self.endLine + if self.state["ProblemType"]["SupportUserArgs"]: + kStr += " " + ptrStr + f" const* C = (beta == {zeroStr}) ? nullptr : BatchC[wg] + batchOffsetC;" + self.endLine + else: + kStr += " " + ptrStr + f" const* C = (beta == {zeroStr}) ? nullptr : BatchC[wg];" + self.endLine kStr += self.endLine ######################################## diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriterConversion.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriterConversion.py index fa7fa8145f03..569da94f9e38 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriterConversion.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriterConversion.py @@ -110,6 +110,11 @@ def functionArgument(self): kStr += " " + self.datatype + " * W;" + self.endLine kStr += " " + ptrStr + " * " + bStr + "C;" + self.endLine + # Batch offsets for General Batched GEMM (SupportUserArgs kernels) + if self.state["ProblemType"]["SupportUserArgs"]: + kStr += " int64_t batchOffsetD;" + self.endLine + kStr += " int64_t batchOffsetC;" + self.endLine + # bias if self.state["ProblemType"]["UseBias"]: if (not self.state["ProblemType"]["Gradient"]): @@ -406,10 +411,16 @@ def kernelBody(self): ptrStr = self.state["ProblemType"]["DataTypeE"].toDevice(self.language) kStr += " " + ptrStr + " * arg.E = arg.BatchE[wg];" + self.endLine ptrStr = self.state["ProblemType"]["DestDataType"].toDevice(self.language) - kStr += " " + ptrStr + " * arg.D = arg.BatchD[wg];" + self.endLine + if self.state["ProblemType"]["SupportUserArgs"]: + kStr += " " + ptrStr + " * arg.D = arg.BatchD[wg] + arg.batchOffsetD;" + self.endLine + else: + kStr += " " + ptrStr + " * arg.D = arg.BatchD[wg];" + self.endLine ptrStr = self.state["ProblemType"]["DestDataType"].toDevice(self.language) zeroStr = self.state["ProblemType"]["ComputeDataType"].zeroString(self.language, 1) - kStr += " " + ptrStr + f" const* arg.C = (arg.beta == {zeroStr}) ? nullptr : arg.BatchC[wg];" + self.endLine + if self.state["ProblemType"]["SupportUserArgs"]: + kStr += " " + ptrStr + f" const* arg.C = (arg.beta == {zeroStr}) ? nullptr : arg.BatchC[wg] + arg.batchOffsetC;" + self.endLine + else: + kStr += " " + ptrStr + f" const* arg.C = (arg.beta == {zeroStr}) ? nullptr : arg.BatchC[wg];" + self.endLine ######################################## # D index From 5b53fdb21c19348102dc3d88c0dba87d7b44e91f Mon Sep 17 00:00:00 2001 From: Seung-Hee Bae Date: Mon, 1 Jun 2026 01:16:31 +0000 Subject: [PATCH 11/27] [hipblaslt] fix bugs in the changes of KernelWriterConversion.py --- .../Tensile/KernelWriterBetaOnly.py | 26 ++++++++++--------- .../Tensile/KernelWriterConversion.py | 23 +++++++++------- .../tensilelite/src/ContractionSolution.cpp | 14 ++++++---- 3 files changed, 36 insertions(+), 27 deletions(-) diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriterBetaOnly.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriterBetaOnly.py index 3c55f403fbc2..63a6d5f10f7f 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriterBetaOnly.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriterBetaOnly.py @@ -89,10 +89,10 @@ def functionSignature(self): batch = "" if isStridedBuffer else "Batch" kStr += " " + ptrStr + " const * " + batch + "C," + self.endLine - # Batch offsets for General Batched GEMM (SupportUserArgs kernels) - if self.state["ProblemType"]["SupportUserArgs"]: - kStr += " int64_t const batchOffsetD,%s" % (self.endLine) - kStr += " int64_t const batchOffsetC,%s" % (self.endLine) +# # Batch offsets for General Batched GEMM (SupportUserArgs kernels) +# if self.state["ProblemType"]["SupportUserArgs"]: +# kStr += " int64_t const batchOffsetD,%s" % (self.endLine) +# kStr += " int64_t const batchOffsetC,%s" % (self.endLine) # bias if self.state["ProblemType"]["BetaOnlyUseBias"]: @@ -222,16 +222,18 @@ def kernelBodyBetaOnly(self): if not self.state["_GlobalAccumulation"]: ptrStr = self.state["ProblemType"]["DestDataType"].toDevice(self.language) - if self.state["ProblemType"]["SupportUserArgs"]: - kStr += " " + ptrStr + " * D = BatchD[wg] + batchOffsetD;" + self.endLine - else: - kStr += " " + ptrStr + " * D = BatchD[wg];" + self.endLine + kStr += " " + ptrStr + " * D = BatchD[wg];" + self.endLine +# if self.state["ProblemType"]["SupportUserArgs"]: +# kStr += " " + ptrStr + " * D = BatchD[wg] + batchOffsetD;" + self.endLine +# else: +# kStr += " " + ptrStr + " * D = BatchD[wg];" + self.endLine ptrStr = self.state["ProblemType"]["DestDataType"].toDevice(self.language) zeroStr = self.state["ProblemType"]["ComputeDataType"].zeroString(self.language, 1) - if self.state["ProblemType"]["SupportUserArgs"]: - kStr += " " + ptrStr + f" const* C = (beta == {zeroStr}) ? nullptr : BatchC[wg] + batchOffsetC;" + self.endLine - else: - kStr += " " + ptrStr + f" const* C = (beta == {zeroStr}) ? nullptr : BatchC[wg];" + self.endLine + kStr += " " + ptrStr + f" const* C = (beta == {zeroStr}) ? nullptr : BatchC[wg];" + self.endLine +# if self.state["ProblemType"]["SupportUserArgs"]: +# kStr += " " + ptrStr + f" const* C = (beta == {zeroStr}) ? nullptr : BatchC[wg] + batchOffsetC;" + self.endLine +# else: +# kStr += " " + ptrStr + f" const* C = (beta == {zeroStr}) ? nullptr : BatchC[wg];" + self.endLine kStr += self.endLine ######################################## diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriterConversion.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriterConversion.py index 569da94f9e38..7be7425f3dfd 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriterConversion.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriterConversion.py @@ -111,7 +111,8 @@ def functionArgument(self): kStr += " " + ptrStr + " * " + bStr + "C;" + self.endLine # Batch offsets for General Batched GEMM (SupportUserArgs kernels) - if self.state["ProblemType"]["SupportUserArgs"]: + #if self.state["ProblemType"]["SupportUserArgs"]: + if not self.state["ProblemType"]["GroupedGemm"]: kStr += " int64_t batchOffsetD;" + self.endLine kStr += " int64_t batchOffsetC;" + self.endLine @@ -212,7 +213,7 @@ def functionSignature(self): # Additional argument batch_mode is added to distinguish between Strided Batch and General Batched GEMM # batch_mode will dictate how the GLOBAL_C and GLOBAL_D macros are defined and used in the kernel body # since the index calculation for Strided Batch and General Batch GEMM are different. - kStr += " argument_%s arg, uint32_t batch_mode, uint32_t additionalPaddingPerBatch)" % ( self.kernelName ) + self.endLine + kStr += " argument_%s arg, uint32_t batch_mode, uint32_t additionalPaddingPerBatch, int64_t batchOffsetD, int64_t batchOffsetC)" % ( self.kernelName ) + self.endLine return kStr @@ -411,16 +412,18 @@ def kernelBody(self): ptrStr = self.state["ProblemType"]["DataTypeE"].toDevice(self.language) kStr += " " + ptrStr + " * arg.E = arg.BatchE[wg];" + self.endLine ptrStr = self.state["ProblemType"]["DestDataType"].toDevice(self.language) - if self.state["ProblemType"]["SupportUserArgs"]: - kStr += " " + ptrStr + " * arg.D = arg.BatchD[wg] + arg.batchOffsetD;" + self.endLine - else: - kStr += " " + ptrStr + " * arg.D = arg.BatchD[wg];" + self.endLine + kStr += " " + ptrStr + " * arg.D = arg.BatchD[wg];" + self.endLine +# if self.state["ProblemType"]["SupportUserArgs"]: +# kStr += " " + ptrStr + " * arg.D = arg.BatchD[wg] + arg.batchOffsetD;" + self.endLine +# else: +# kStr += " " + ptrStr + " * arg.D = arg.BatchD[wg];" + self.endLine ptrStr = self.state["ProblemType"]["DestDataType"].toDevice(self.language) zeroStr = self.state["ProblemType"]["ComputeDataType"].zeroString(self.language, 1) - if self.state["ProblemType"]["SupportUserArgs"]: - kStr += " " + ptrStr + f" const* arg.C = (arg.beta == {zeroStr}) ? nullptr : arg.BatchC[wg] + arg.batchOffsetC;" + self.endLine - else: - kStr += " " + ptrStr + f" const* arg.C = (arg.beta == {zeroStr}) ? nullptr : arg.BatchC[wg];" + self.endLine + kStr += " " + ptrStr + f" const* arg.C = (arg.beta == {zeroStr}) ? nullptr : arg.BatchC[wg];" + self.endLine +# if self.state["ProblemType"]["SupportUserArgs"]: +# kStr += " " + ptrStr + f" const* arg.C = (arg.beta == {zeroStr}) ? nullptr : arg.BatchC[wg] + arg.batchOffsetC;" + self.endLine +# else: +# kStr += " " + ptrStr + f" const* arg.C = (arg.beta == {zeroStr}) ? nullptr : arg.BatchC[wg];" + self.endLine ######################################## # D index diff --git a/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp b/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp index 917f769e715b..233bb59ea46e 100644 --- a/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp +++ b/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp @@ -2308,11 +2308,12 @@ namespace TensileLite args.template append("batchC", inputs.batchC); // Pass batch offsets when kernel expects them (SupportUserArgs=true, not GroupedGemm) - if(problemType.supportDeviceUserArguments && !problemType.groupedGemm) - { - args.template append("batchOffsetD", inputs.batchOffsetD); - args.template append("batchOffsetC", inputs.batchOffsetC); - } + //if(problemType.supportDeviceUserArguments && !problemType.groupedGemm) +// if(!problemType.groupedGemm) +// { +// args.template append("batchOffsetD", inputs.batchOffsetD); +// args.template append("batchOffsetC", inputs.batchOffsetC); +// } bool useBias = false; if(problemType.useBias) @@ -2468,6 +2469,9 @@ namespace TensileLite ContractionProblemGemm::BATCHMODE batchMode = problem.batchMode(); args.template append("batchMode", static_cast(batchMode)); args.template append("additionalPaddingPerBatch", additionalPaddingPerBatchGeneralBatch); + + args.template append("batchOffsetD", inputs.batchOffsetD); + args.template append("batchOffsetC", inputs.batchOffsetC); } } From c19a56615447c9f6865a82b5b8875ad0e113a2ba Mon Sep 17 00:00:00 2001 From: Seung-Hee Bae Date: Tue, 2 Jun 2026 01:41:57 +0000 Subject: [PATCH 12/27] [hipblaslt] remove unnecessary if-condition and minor update. --- .../include/testing_matmul_batch_offset.hpp | 2 +- .../clients/tests/data/matmul_gtest.yaml | 2 ++ .../Tensile/KernelWriterAssembly.py | 8 ++++--- .../Tensile/KernelWriterConversion.py | 15 +++---------- .../tensilelite/src/ContractionSolution.cpp | 21 +++++++------------ 5 files changed, 18 insertions(+), 30 deletions(-) diff --git a/projects/hipblaslt/clients/common/include/testing_matmul_batch_offset.hpp b/projects/hipblaslt/clients/common/include/testing_matmul_batch_offset.hpp index 3637c9aeea08..1e9c17fedc8a 100644 --- a/projects/hipblaslt/clients/common/include/testing_matmul_batch_offset.hpp +++ b/projects/hipblaslt/clients/common/include/testing_matmul_batch_offset.hpp @@ -378,7 +378,7 @@ void testing_matmul_batch_offset_impl(const Arguments& arg) EXPECT_LT(max_error, tol) << "GPU vs CPU mismatch (error: " << max_error << ", tol: " << tol << ")"; } - if(arg.norm_check) + if(arg.norm_check && max_error >= tol) { hipblaslt_cout << "GPU vs CPU max error: " << max_error << " (tol: " << tol << ")" << std::endl; } diff --git a/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml b/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml index a7b580934b5f..1d7472c88b8f 100755 --- a/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml +++ b/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml @@ -3364,3 +3364,5 @@ Tests: beta: 1.0 unit_check: 1 norm_check: 1 + +... diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py index c9e5e6598e5d..4631355f9f56 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py @@ -2187,7 +2187,7 @@ def getKernelArgLoadModule(self, kernel, sgprStartIdx, numsOfLoad, preloadNum): # Batch offset arguments for general batched mode # Store kernarg offsets for on-demand loading (no persistent SGPRs needed) - if kernel["ProblemType"]["SupportUserArgs"] and not kernel["ProblemType"]["GroupedGemm"]: + if not kernel["ProblemType"]["GroupedGemm"]: sgprOffset = self.argLoader.getOffset() self.states.batchOffsetDKernArgOffset = sgprOffset sgprOffset += 8 # 64-bit = 8 bytes @@ -2198,6 +2198,7 @@ def getKernelArgLoadModule(self, kernel, sgprStartIdx, numsOfLoad, preloadNum): self.states.batchOffsetBKernArgOffset = sgprOffset sgprOffset += 8 self.argLoader.setOffset(sgprOffset) + # TODO: remove below before merging the PR! # Debug: write kernel arg offsets to file for verification with open("/tmp/kernel_batch_offset_debug.txt", "a") as f: f.write(f"{self.states.kernelName}: D@{self.states.batchOffsetDKernArgOffset}, C@{self.states.batchOffsetCKernArgOffset}, A@{self.states.batchOffsetAKernArgOffset}, B@{self.states.batchOffsetBKernArgOffset}\n") @@ -4657,7 +4658,7 @@ def computeLoadSrd(self, kernel, tP, tc, indices, bpe): moduleLoadGeneralBatch.add(SAddCU32(dst=sgpr(stmp+1), src0=sgpr("Address%s+1"%tc), src1=0, comment="Offsetting to the location [Higher half of address]")) moduleLoadGeneralBatch.add(SLoadB64(dst=sgpr("Srd%s"%tc, 2), base=sgpr(stmp, 2), soffset=0, comment="Load the Matrix Address in the Pointer Array")) # Load and apply batch offset for General Batched GEMM - if kernel["ProblemType"]["SupportUserArgs"] and not kernel["ProblemType"]["GroupedGemm"]: + if not kernel["ProblemType"]["GroupedGemm"]: batchOffsetKernArgOffset = self.states.batchOffsetAKernArgOffset if tc == "A" else self.states.batchOffsetBKernArgOffset moduleLoadGeneralBatch.add(SLoadB64(dst=sgpr(stmp, 2), base=sgpr("KernArgAddress", 2), soffset=hex(batchOffsetKernArgOffset), comment="Load batchOffset%s from kernel args"%tc)) moduleLoadGeneralBatch.add(SWaitCnt(kmcnt=0, comment="Wait for Matrix Address and Batch Offset Loads")) @@ -4665,6 +4666,7 @@ def computeLoadSrd(self, kernel, tP, tc, indices, bpe): moduleLoadGeneralBatch.add(SAddCU32(dst=sgpr("Srd%s+1"%tc), src0=sgpr("Srd%s+1"%tc), src1=sgpr(stmp+1), comment="Add batch offset to %s address (high)"%tc)) else: moduleLoadGeneralBatch.add(SWaitCnt(kmcnt=0, comment="Wait for the Matrix Address Load from the Pointer Array")) + if self.states.groOffsetInMacroTile and ((tc == "A" and not kernel["enableTDMA"]) or (tc == "B" and not kernel["enableTDMB"])): prePad1 = int(self.states.srdShiftLeft[tc] * tP["bpeGR"]) # leave room in case we have to pointer shift moduleLoadGeneralBatch.add(SSubU32(dst=sgpr("Srd%s+0"%tc), src0=sgpr("Srd%s+0"%tc), src1=prePad1, comment="pre-pad to make room for possible pointer shift")) @@ -13300,7 +13302,7 @@ def computeStoreSrdStart(self, kernel, srdTcList: list, sgprBpeList = [], useSiz module.add(SAddCU32(dst=sgpr("Srd%s+1"%mat), src0=sgpr("Srd%s+1"%mat), src1=sgpr(tmpS1), comment="Offsetting within the Batch Matrix [Higher half of address]")) # Now, we have starting matrix address of a specific batch in the corresponding Srd. # Load and apply batch offset for General Batched GEMM (C or D matrix) as necessary. - if kernel["ProblemType"]["SupportUserArgs"] and not kernel["ProblemType"]["GroupedGemm"]: + if not kernel["ProblemType"]["GroupedGemm"]: batchOffsetKernArgOffset = self.states.batchOffsetCKernArgOffset if mat == "C" else self.states.batchOffsetDKernArgOffset module.add(SLoadB64(dst=sgpr(tmpS0, 2), base=sgpr("KernArgAddress", 2), soffset=hex(batchOffsetKernArgOffset), comment="Load batchOffset%s from kernel args"%mat)) module.add(SWaitCnt(kmcnt=0, comment="Wait for Matrix Address and Batch Offset Loads")) diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriterConversion.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriterConversion.py index 7be7425f3dfd..1ff7291d8a79 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriterConversion.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriterConversion.py @@ -110,8 +110,7 @@ def functionArgument(self): kStr += " " + self.datatype + " * W;" + self.endLine kStr += " " + ptrStr + " * " + bStr + "C;" + self.endLine - # Batch offsets for General Batched GEMM (SupportUserArgs kernels) - #if self.state["ProblemType"]["SupportUserArgs"]: + # Batch offsets for General Batched GEMM if not self.state["ProblemType"]["GroupedGemm"]: kStr += " int64_t batchOffsetD;" + self.endLine kStr += " int64_t batchOffsetC;" + self.endLine @@ -413,17 +412,9 @@ def kernelBody(self): kStr += " " + ptrStr + " * arg.E = arg.BatchE[wg];" + self.endLine ptrStr = self.state["ProblemType"]["DestDataType"].toDevice(self.language) kStr += " " + ptrStr + " * arg.D = arg.BatchD[wg];" + self.endLine -# if self.state["ProblemType"]["SupportUserArgs"]: -# kStr += " " + ptrStr + " * arg.D = arg.BatchD[wg] + arg.batchOffsetD;" + self.endLine -# else: -# kStr += " " + ptrStr + " * arg.D = arg.BatchD[wg];" + self.endLine ptrStr = self.state["ProblemType"]["DestDataType"].toDevice(self.language) zeroStr = self.state["ProblemType"]["ComputeDataType"].zeroString(self.language, 1) kStr += " " + ptrStr + f" const* arg.C = (arg.beta == {zeroStr}) ? nullptr : arg.BatchC[wg];" + self.endLine -# if self.state["ProblemType"]["SupportUserArgs"]: -# kStr += " " + ptrStr + f" const* arg.C = (arg.beta == {zeroStr}) ? nullptr : arg.BatchC[wg] + arg.batchOffsetC;" + self.endLine -# else: -# kStr += " " + ptrStr + f" const* arg.C = (arg.beta == {zeroStr}) ? nullptr : arg.BatchC[wg];" + self.endLine ######################################## # D index @@ -740,7 +731,7 @@ def kernelBody(self): kStr += " }" + self.endLine kStr += " else" + self.endLine kStr += " {" + self.endLine - kStr += " %s *ptr = *(reinterpret_cast<%s **>(((char *)arg.C) + (8*id2)));" % (destTypeStr, destTypeStr) + self.endLine + kStr += " %s *ptr = *(reinterpret_cast<%s **>(((char *)arg.C) + (8*id2))) + batchOffsetC/sizeof(%s);" % (destTypeStr, destTypeStr, destTypeStr) + self.endLine for vIdx in range(self.num_dword_load): kStr += " %s[%d] += arg.beta * (%s)ptr[idxC+%d];%s" % (accumStr, vIdx, intermediateDataType, vIdx, self.endLine) kStr += " }" + self.endLine @@ -836,7 +827,7 @@ def kernelBody(self): kStr += " if(batch_mode == 0) {" + self.endLine kStr += " buffer_store<%s, sizeof(%s), CacheOperation::Kind::Always>(*(%s *)%s, arg.D, idxD * sizeof(%s), 0);%s" % (storeTypeStr, storeTypeStr, storeTypeStr, resultStr, destTypeStr, self.endLine) kStr += " } else {" + self.endLine - kStr += " %s *ptr = *(reinterpret_cast<%s **>(((char *)arg.D) + (8*id2)));" % (destTypeStr, destTypeStr) + self.endLine + kStr += " %s *ptr = *(reinterpret_cast<%s **>(((char *)arg.D) + (8*id2))) + batchOffsetD/sizeof(%s);" % (destTypeStr, destTypeStr, destTypeStr) + self.endLine kStr += " buffer_store<%s, sizeof(%s), CacheOperation::Kind::Always>(*(%s *)%s, ptr, idxD * sizeof(%s), 0);%s" % (storeTypeStr, storeTypeStr, storeTypeStr, resultStr, destTypeStr, self.endLine) kStr += " }" + self.endLine else: diff --git a/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp b/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp index 233bb59ea46e..082d0906083e 100644 --- a/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp +++ b/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp @@ -990,7 +990,8 @@ namespace TensileLite // Batch offset support for General Batched GEMM (SupportUserArgs kernels) // Placed after core GEMM args (strides, alpha/beta, StreamK) to match kernel signature - if(problemType.supportDeviceUserArguments && !problemType.groupedGemm) + //if(problemType.supportDeviceUserArguments && !problemType.groupedGemm) + if(!problemType.groupedGemm) { args.template append("batchOffsetD", inputs.batchOffsetD); args.template append("batchOffsetC", inputs.batchOffsetC); @@ -2132,11 +2133,11 @@ namespace TensileLite rv.args.append("batchC", inputs.batchC); // Pass batch offsets when kernel expects them (SupportUserArgs=true, not GroupedGemm) - if(problemType.supportDeviceUserArguments && !problemType.groupedGemm) - { - rv.args.append("batchOffsetD", inputs.batchOffsetD); - rv.args.append("batchOffsetC", inputs.batchOffsetC); - } +// if(problemType.supportDeviceUserArguments && !problemType.groupedGemm) +// { +// rv.args.append("batchOffsetD", inputs.batchOffsetD); +// rv.args.append("batchOffsetC", inputs.batchOffsetC); +// } if(problemType.useBias && sizeMapping.globalAccumulation == 0 && (!problemType.useGradient)) { @@ -2307,14 +2308,6 @@ namespace TensileLite else args.template append("batchC", inputs.batchC); - // Pass batch offsets when kernel expects them (SupportUserArgs=true, not GroupedGemm) - //if(problemType.supportDeviceUserArguments && !problemType.groupedGemm) -// if(!problemType.groupedGemm) -// { -// args.template append("batchOffsetD", inputs.batchOffsetD); -// args.template append("batchOffsetC", inputs.batchOffsetC); -// } - bool useBias = false; if(problemType.useBias) { From 343d2213407a5ea9baf2688f0855f61514e3b09d Mon Sep 17 00:00:00 2001 From: Seung-Hee Bae Date: Wed, 3 Jun 2026 00:07:37 +0000 Subject: [PATCH 13/27] [hipblaslt] fix CI test failures. --- .../tensilelite/Tensile/Components/Signature.py | 11 ++++++++++- .../tensilelite/Tensile/KernelWriterConversion.py | 10 +++++----- .../hipblaslt/tensilelite/src/ContractionSolution.cpp | 10 ++++++++++ 3 files changed, 25 insertions(+), 6 deletions(-) diff --git a/projects/hipblaslt/tensilelite/Tensile/Components/Signature.py b/projects/hipblaslt/tensilelite/Tensile/Components/Signature.py index 0ea36aa48ebc..0e829336919e 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Components/Signature.py +++ b/projects/hipblaslt/tensilelite/Tensile/Components/Signature.py @@ -257,7 +257,7 @@ def __call__(self, writer) -> SignatureBase: # Batch offset support for general batched mode (pointer array) # Placed after core GEMM args (strides, alpha/beta, StreamK) # Enabled for all SupportUserArgs kernels (offsets are 0 when not in General Batched mode) - if kernel["ProblemType"]["SupportUserArgs"] and not kernel["ProblemType"]["GroupedGemm"]: + if not kernel["ProblemType"]["GroupedGemm"]: signature.addArg("batchOffsetD", SVK.SIG_VALUE, "u64") signature.addArg("batchOffsetC", SVK.SIG_VALUE, "u64") signature.addArg("batchOffsetA", SVK.SIG_VALUE, "u64") @@ -316,6 +316,15 @@ def __call__(self, writer) -> SignatureBase: signature.addArg( "AmaxWS", SVK.SIG_GLOBALBUFFER, cptValueType, "generic") signature.addArg( "AmaxSync", SVK.SIG_GLOBALBUFFER, "u32", "generic") + # Batch offset support for general batched mode (pointer array) + # Enabled for all SupportUserArgs kernels (offsets are 0 when not in General Batched mode) +# if not kernel["ProblemType"]["GroupedGemm"]: +# signature.addArg("batchOffsetD", SVK.SIG_VALUE, "u64") +# signature.addArg("batchOffsetC", SVK.SIG_VALUE, "u64") +# signature.addArg("batchOffsetA", SVK.SIG_VALUE, "u64") +# signature.addArg("batchOffsetB", SVK.SIG_VALUE, "u64") +# userArgumentsInfo.gemmArgumentSize += 32 # 4 offsets * 8 bytes each + if (kernel["_GlobalAccumulation"] == "MultipleBufferSingleKernel" or kernel["AdaptiveGemmGSUA"] == 1): signature.addArg( "dstD", SVK.SIG_GLOBALBUFFER, dstValueType, "generic") signature.addArg( "Synchronizer", SVK.SIG_GLOBALBUFFER, cptValueType, "generic") diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriterConversion.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriterConversion.py index 1ff7291d8a79..ed1956b18bef 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriterConversion.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriterConversion.py @@ -110,11 +110,6 @@ def functionArgument(self): kStr += " " + self.datatype + " * W;" + self.endLine kStr += " " + ptrStr + " * " + bStr + "C;" + self.endLine - # Batch offsets for General Batched GEMM - if not self.state["ProblemType"]["GroupedGemm"]: - kStr += " int64_t batchOffsetD;" + self.endLine - kStr += " int64_t batchOffsetC;" + self.endLine - # bias if self.state["ProblemType"]["UseBias"]: if (not self.state["ProblemType"]["Gradient"]): @@ -190,6 +185,11 @@ def functionArgument(self): if enableFactorDim: kStr += " unsigned int factorDim;%s" % (self.endLine) + # Batch offsets for General Batched GEMM + if not self.state["ProblemType"]["GroupedGemm"]: + kStr += " int64_t batchOffsetD;" + self.endLine + kStr += " int64_t batchOffsetC;" + self.endLine + # argument structure end kStr += "};" + self.endLine diff --git a/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp b/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp index 082d0906083e..991728c7a0b9 100644 --- a/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp +++ b/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp @@ -1118,6 +1118,16 @@ namespace TensileLite (uint8_t*)inputs.ws + workspaceOffsetInByte); args.template append("AmaxSync", inputs.Synchronizer); } + + // Batch offset support for General Batched GEMM (SupportUserArgs kernels) + // Placed at the end of the parameter list +// if(!problemType.groupedGemm) +// { +// args.template append("batchOffsetD", inputs.batchOffsetD); +// args.template append("batchOffsetC", inputs.batchOffsetC); +// args.template append("batchOffsetA", inputs.batchOffsetA); +// args.template append("batchOffsetB", inputs.batchOffsetB); +// } } inline uint32_t getNumWorkGroups(const KernelInvocation& rv) From 366226bb20e379f433e742128c17b7c600479c2d Mon Sep 17 00:00:00 2001 From: Seung-Hee Bae Date: Wed, 3 Jun 2026 20:30:45 +0000 Subject: [PATCH 14/27] remove temporary debug implementation. --- .../hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py | 4 ---- 1 file changed, 4 deletions(-) diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py index 4631355f9f56..63c491f45d10 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py @@ -2198,10 +2198,6 @@ def getKernelArgLoadModule(self, kernel, sgprStartIdx, numsOfLoad, preloadNum): self.states.batchOffsetBKernArgOffset = sgprOffset sgprOffset += 8 self.argLoader.setOffset(sgprOffset) - # TODO: remove below before merging the PR! - # Debug: write kernel arg offsets to file for verification - with open("/tmp/kernel_batch_offset_debug.txt", "a") as f: - f.write(f"{self.states.kernelName}: D@{self.states.batchOffsetDKernArgOffset}, C@{self.states.batchOffsetCKernArgOffset}, A@{self.states.batchOffsetAKernArgOffset}, B@{self.states.batchOffsetBKernArgOffset}\n") return kernelArgs From 71a37030936652f065e799bc040891c577ad9a5f Mon Sep 17 00:00:00 2001 From: Seung-Hee Bae Date: Wed, 10 Jun 2026 18:28:23 -0400 Subject: [PATCH 15/27] [hipblaslt] fix the bug related to CI failures. --- .../hipblaslt/tensilelite/Tensile/KernelWriterConversion.py | 6 +----- 1 file changed, 1 insertion(+), 5 deletions(-) diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriterConversion.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriterConversion.py index ed1956b18bef..9f8691393abc 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriterConversion.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriterConversion.py @@ -185,11 +185,6 @@ def functionArgument(self): if enableFactorDim: kStr += " unsigned int factorDim;%s" % (self.endLine) - # Batch offsets for General Batched GEMM - if not self.state["ProblemType"]["GroupedGemm"]: - kStr += " int64_t batchOffsetD;" + self.endLine - kStr += " int64_t batchOffsetC;" + self.endLine - # argument structure end kStr += "};" + self.endLine @@ -212,6 +207,7 @@ def functionSignature(self): # Additional argument batch_mode is added to distinguish between Strided Batch and General Batched GEMM # batch_mode will dictate how the GLOBAL_C and GLOBAL_D macros are defined and used in the kernel body # since the index calculation for Strided Batch and General Batch GEMM are different. + # Also, batchOffsetD and batchOffsetC arguments added at the end. kStr += " argument_%s arg, uint32_t batch_mode, uint32_t additionalPaddingPerBatch, int64_t batchOffsetD, int64_t batchOffsetC)" % ( self.kernelName ) + self.endLine return kStr From b90f9985b49235e8a4c4e1626f869575cb7e2385 Mon Sep 17 00:00:00 2001 From: Seung-Hee Bae Date: Thu, 11 Jun 2026 19:13:07 -0400 Subject: [PATCH 16/27] [hipblaslt] place the batch offset argument at the tail of the kernarg buffer --- .../Tensile/Components/Signature.py | 41 ++++++++++--------- .../tensilelite/Tensile/KernelWriter.py | 6 +++ .../Tensile/KernelWriterAssembly.py | 17 ++------ .../Tensile/KernelWriterBetaOnly.py | 13 ------ .../tensilelite/rocisa/rocisa/src/code.cpp | 3 ++ .../tensilelite/src/ContractionSolution.cpp | 38 +++++------------ 6 files changed, 45 insertions(+), 73 deletions(-) diff --git a/projects/hipblaslt/tensilelite/Tensile/Components/Signature.py b/projects/hipblaslt/tensilelite/Tensile/Components/Signature.py index 0e829336919e..847a17e510e7 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Components/Signature.py +++ b/projects/hipblaslt/tensilelite/Tensile/Components/Signature.py @@ -254,16 +254,6 @@ def __call__(self, writer) -> SignatureBase: signature.addArg("skTiles", SVK.SIG_VALUE, "u32") userArgumentsInfo.gemmArgumentSize += 8 - # Batch offset support for general batched mode (pointer array) - # Placed after core GEMM args (strides, alpha/beta, StreamK) - # Enabled for all SupportUserArgs kernels (offsets are 0 when not in General Batched mode) - if not kernel["ProblemType"]["GroupedGemm"]: - signature.addArg("batchOffsetD", SVK.SIG_VALUE, "u64") - signature.addArg("batchOffsetC", SVK.SIG_VALUE, "u64") - signature.addArg("batchOffsetA", SVK.SIG_VALUE, "u64") - signature.addArg("batchOffsetB", SVK.SIG_VALUE, "u64") - userArgumentsInfo.gemmArgumentSize += 32 # 4 offsets * 8 bytes each - if kernel["ProblemType"]["UseScaleAB"]: signature.addArg("AddressScaleA", SVK.SIG_GLOBALBUFFER, cptValueType, "generic") signature.addArg("AddressScaleB", SVK.SIG_GLOBALBUFFER, cptValueType, "generic") @@ -316,20 +306,33 @@ def __call__(self, writer) -> SignatureBase: signature.addArg( "AmaxWS", SVK.SIG_GLOBALBUFFER, cptValueType, "generic") signature.addArg( "AmaxSync", SVK.SIG_GLOBALBUFFER, "u32", "generic") - # Batch offset support for general batched mode (pointer array) - # Enabled for all SupportUserArgs kernels (offsets are 0 when not in General Batched mode) -# if not kernel["ProblemType"]["GroupedGemm"]: -# signature.addArg("batchOffsetD", SVK.SIG_VALUE, "u64") -# signature.addArg("batchOffsetC", SVK.SIG_VALUE, "u64") -# signature.addArg("batchOffsetA", SVK.SIG_VALUE, "u64") -# signature.addArg("batchOffsetB", SVK.SIG_VALUE, "u64") -# userArgumentsInfo.gemmArgumentSize += 32 # 4 offsets * 8 bytes each - if (kernel["_GlobalAccumulation"] == "MultipleBufferSingleKernel" or kernel["AdaptiveGemmGSUA"] == 1): signature.addArg( "dstD", SVK.SIG_GLOBALBUFFER, dstValueType, "generic") signature.addArg( "Synchronizer", SVK.SIG_GLOBALBUFFER, cptValueType, "generic") signature.addArg( "GSUSync", SVK.SIG_VALUE, "u32") + # Batch offset support for general batched mode (pointer array). + # Placed at the tail of the kernarg buffer (after the dstD/Synchronizer block) + # so no later arg is shifted; the host appends them in the same position, + # after the dstD/Synchronizer/seed block. Record each arg's kernarg byte + # offset so the assembly loads them from the accurate position rather + # than re-deriving it. + # + # signature.offset counts from the very first arg including the common header. + # The assembly loads these args with KernArgAddress already advanced past + # that header by commonArgsSize, so subtract it. + if not kernel["ProblemType"]["GroupedGemm"]: + commonArgsSize = userArgumentsInfo.commonArgsSize + writer.states.batchOffsetDKernArgOffset = signature.offset - commonArgsSize + signature.addArg("batchOffsetD", SVK.SIG_VALUE, "u64") + writer.states.batchOffsetCKernArgOffset = signature.offset - commonArgsSize + signature.addArg("batchOffsetC", SVK.SIG_VALUE, "u64") + writer.states.batchOffsetAKernArgOffset = signature.offset - commonArgsSize + signature.addArg("batchOffsetA", SVK.SIG_VALUE, "u64") + writer.states.batchOffsetBKernArgOffset = signature.offset - commonArgsSize + signature.addArg("batchOffsetB", SVK.SIG_VALUE, "u64") + userArgumentsInfo.gemmArgumentSize += 32 # 4 offsets * 8 bytes each + activationType = ActivationType("all") for name in activationType.getAdditionalArgStringList(): userArgumentsInfo.activationSize += userArgumentsInfo.actMaxSize diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py index 3da1a6e1def9..d6926f0d24df 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py @@ -310,6 +310,12 @@ class StateValues: numSgprToLoad: int = 0 # For kernel args preloadGuard: List[int] = field(init=False) # For preload kernel args guard numSgprPreload: int = 0 # For kernel args + # Kernarg byte offsets of the general-batched offsets, set when the signature + # appends them at the tail (0 for grouped-gemm kernels that omit them) + batchOffsetDKernArgOffset: int = 0 + batchOffsetCKernArgOffset: int = 0 + batchOffsetAKernArgOffset: int = 0 + batchOffsetBKernArgOffset: int = 0 numSgprAlpha: int = 0 # For user arguments numSgprBeta: int = 0 # For user arguments numStoreSgprNames: List[str] = field(init=False) # For post-loop kernel args diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py index 63c491f45d10..7312c4610bae 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py @@ -2185,19 +2185,10 @@ def getKernelArgLoadModule(self, kernel, sgprStartIdx, numsOfLoad, preloadNum): sgprOffset += 4 self.argLoader.setOffset(sgprOffset) - # Batch offset arguments for general batched mode - # Store kernarg offsets for on-demand loading (no persistent SGPRs needed) - if not kernel["ProblemType"]["GroupedGemm"]: - sgprOffset = self.argLoader.getOffset() - self.states.batchOffsetDKernArgOffset = sgprOffset - sgprOffset += 8 # 64-bit = 8 bytes - self.states.batchOffsetCKernArgOffset = sgprOffset - sgprOffset += 8 - self.states.batchOffsetAKernArgOffset = sgprOffset - sgprOffset += 8 - self.states.batchOffsetBKernArgOffset = sgprOffset - sgprOffset += 8 - self.argLoader.setOffset(sgprOffset) + # Batch offset arguments for general batched mode are loaded on-demand from + # their kernarg byte offsets. Those offsets are recorded accurately by + # the signature builder (see Signature.py) into self.states.batchOffset*KernArgOffset, + # so nothing is computed here. return kernelArgs diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriterBetaOnly.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriterBetaOnly.py index 63a6d5f10f7f..5f1a31c432d6 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriterBetaOnly.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriterBetaOnly.py @@ -89,11 +89,6 @@ def functionSignature(self): batch = "" if isStridedBuffer else "Batch" kStr += " " + ptrStr + " const * " + batch + "C," + self.endLine -# # Batch offsets for General Batched GEMM (SupportUserArgs kernels) -# if self.state["ProblemType"]["SupportUserArgs"]: -# kStr += " int64_t const batchOffsetD,%s" % (self.endLine) -# kStr += " int64_t const batchOffsetC,%s" % (self.endLine) - # bias if self.state["ProblemType"]["BetaOnlyUseBias"]: biasPtrStr = self.state["ProblemType"]["BiasDataType"].toDevice(self.language) @@ -223,17 +218,9 @@ def kernelBodyBetaOnly(self): if not self.state["_GlobalAccumulation"]: ptrStr = self.state["ProblemType"]["DestDataType"].toDevice(self.language) kStr += " " + ptrStr + " * D = BatchD[wg];" + self.endLine -# if self.state["ProblemType"]["SupportUserArgs"]: -# kStr += " " + ptrStr + " * D = BatchD[wg] + batchOffsetD;" + self.endLine -# else: -# kStr += " " + ptrStr + " * D = BatchD[wg];" + self.endLine ptrStr = self.state["ProblemType"]["DestDataType"].toDevice(self.language) zeroStr = self.state["ProblemType"]["ComputeDataType"].zeroString(self.language, 1) kStr += " " + ptrStr + f" const* C = (beta == {zeroStr}) ? nullptr : BatchC[wg];" + self.endLine -# if self.state["ProblemType"]["SupportUserArgs"]: -# kStr += " " + ptrStr + f" const* C = (beta == {zeroStr}) ? nullptr : BatchC[wg] + batchOffsetC;" + self.endLine -# else: -# kStr += " " + ptrStr + f" const* C = (beta == {zeroStr}) ? nullptr : BatchC[wg];" + self.endLine kStr += self.endLine ######################################## diff --git a/projects/hipblaslt/tensilelite/rocisa/rocisa/src/code.cpp b/projects/hipblaslt/tensilelite/rocisa/rocisa/src/code.cpp index 9fdff89722ca..a7eafed0e3b3 100644 --- a/projects/hipblaslt/tensilelite/rocisa/rocisa/src/code.cpp +++ b/projects/hipblaslt/tensilelite/rocisa/rocisa/src/code.cpp @@ -367,6 +367,7 @@ void init_code(nb::module_ m) nb::arg("totalVgprs") = 0, nb::arg("totalSgprs") = 0) .def("setGprs", &rocisa::SignatureCodeMeta::setGprs) + .def_ro("offset", &rocisa::SignatureCodeMeta::offset) .def("addArg", &rocisa::SignatureCodeMeta::addArg, nb::arg("name"), @@ -406,6 +407,8 @@ void init_code(nb::module_ m) nb::arg("totalSgprs") = 0, nb::arg("numSgprPreload") = 0) .def("setGprs", &rocisa::SignatureBase::setGprs) + .def_prop_ro("offset", + [](const rocisa::SignatureBase& self) { return self.codeMeta.offset; }) .def("addArg", &rocisa::SignatureBase::addArg, nb::arg("name"), diff --git a/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp b/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp index 991728c7a0b9..da452edd3b48 100644 --- a/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp +++ b/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp @@ -988,17 +988,6 @@ namespace TensileLite autoGsuVal, ntab); - // Batch offset support for General Batched GEMM (SupportUserArgs kernels) - // Placed after core GEMM args (strides, alpha/beta, StreamK) to match kernel signature - //if(problemType.supportDeviceUserArguments && !problemType.groupedGemm) - if(!problemType.groupedGemm) - { - args.template append("batchOffsetD", inputs.batchOffsetD); - args.template append("batchOffsetC", inputs.batchOffsetC); - args.template append("batchOffsetA", inputs.batchOffsetA); - args.template append("batchOffsetB", inputs.batchOffsetB); - } - // NOTE: an assumption here is A & B must be both MX data types or non-MX data types. // Mixing is not supported. if(!problemType.useScaleAB.empty()) @@ -1118,16 +1107,6 @@ namespace TensileLite (uint8_t*)inputs.ws + workspaceOffsetInByte); args.template append("AmaxSync", inputs.Synchronizer); } - - // Batch offset support for General Batched GEMM (SupportUserArgs kernels) - // Placed at the end of the parameter list -// if(!problemType.groupedGemm) -// { -// args.template append("batchOffsetD", inputs.batchOffsetD); -// args.template append("batchOffsetC", inputs.batchOffsetC); -// args.template append("batchOffsetA", inputs.batchOffsetA); -// args.template append("batchOffsetB", inputs.batchOffsetB); -// } } inline uint32_t getNumWorkGroups(const KernelInvocation& rv) @@ -1862,6 +1841,16 @@ namespace TensileLite rv.args.append("GSUSync", 0); } + // Batch offset support for General Batched GEMM (SupportUserArgs kernels). + // Appended at the tail, after the dstD/Synchronizer block, to match the + // kernel signature order (see Signature.py). + if(!problemType.groupedGemm) + { + rv.args.append("batchOffsetD", inputs.batchOffsetD); + rv.args.append("batchOffsetC", inputs.batchOffsetC); + rv.args.append("batchOffsetA", inputs.batchOffsetA); + rv.args.append("batchOffsetB", inputs.batchOffsetB); + } if(problemType.stochasticRounding) { @@ -2142,13 +2131,6 @@ namespace TensileLite else rv.args.append("batchC", inputs.batchC); - // Pass batch offsets when kernel expects them (SupportUserArgs=true, not GroupedGemm) -// if(problemType.supportDeviceUserArguments && !problemType.groupedGemm) -// { -// rv.args.append("batchOffsetD", inputs.batchOffsetD); -// rv.args.append("batchOffsetC", inputs.batchOffsetC); -// } - if(problemType.useBias && sizeMapping.globalAccumulation == 0 && (!problemType.useGradient)) { if(problemType.stridedBatched) From 586c62b561fa35e8c399363632b26dd41ecd2dab Mon Sep 17 00:00:00 2001 From: Seung-Hee Bae Date: Thu, 11 Jun 2026 22:54:23 -0400 Subject: [PATCH 17/27] [hipblaslt] modify the batch offset value as in elements. --- .../include/testing_matmul_batch_offset.hpp | 32 ++++++++----------- .../library/include/hipblaslt/hipblaslt.h | 2 +- .../src/amd_detail/include/auxiliary.hpp | 16 ++++++++++ .../src/include/rocblaslt_mat_utils.hpp | 13 ++++++++ .../amd_detail/rocblaslt/src/tensile_host.cpp | 21 +++++++++--- 5 files changed, 59 insertions(+), 25 deletions(-) diff --git a/projects/hipblaslt/clients/common/include/testing_matmul_batch_offset.hpp b/projects/hipblaslt/clients/common/include/testing_matmul_batch_offset.hpp index 1e9c17fedc8a..8cb2ad8d46d5 100644 --- a/projects/hipblaslt/clients/common/include/testing_matmul_batch_offset.hpp +++ b/projects/hipblaslt/clients/common/include/testing_matmul_batch_offset.hpp @@ -62,7 +62,7 @@ void testing_matmul_batch_offset_impl(const Arguments& arg) int32_t batch_count = arg.batch_count; - // Batch offsets (from YAML - assumed to be in BYTES) + // Batch offsets (from YAML - in ELEMENTS) int64_t offset_a = arg.batch_offset_a; int64_t offset_b = arg.batch_offset_b; int64_t offset_c = arg.batch_offset_c; @@ -86,22 +86,16 @@ void testing_matmul_batch_offset_impl(const Arguments& arg) size_t size_C_sub = size_t(ldc) * size_t(N); size_t size_D_sub = size_t(ldd) * size_t(N); - // Convert byte offsets to element offsets - // Note: offsets must be aligned to element size - size_t offset_a_elem = offset_a / sizeof(Ti); - size_t offset_b_elem = offset_b / sizeof(Ti); - size_t offset_c_elem = offset_c / sizeof(To); - size_t offset_d_elem = offset_d / sizeof(To); - // Full buffer sizes: [offset padding] + [matrix data] // The offset padding comes FIRST because: // - Pointer array points to buffer BASE // - Kernel adds offset to BASE to get actual data location // - So actual data is at BASE + offset - size_t size_A_full = offset_a_elem + size_A_sub; - size_t size_B_full = offset_b_elem + size_B_sub; - size_t size_C_full = offset_c_elem + size_C_sub; - size_t size_D_full = offset_d_elem + size_D_sub; + // Offsets are already in elements (the API takes element offsets). + size_t size_A_full = offset_a + size_A_sub; + size_t size_B_full = offset_b + size_B_sub; + size_t size_C_full = offset_c + size_C_sub; + size_t size_D_full = offset_d + size_D_sub; // Allocate host memory for full buffers host_vector h_A_full(size_A_full * batch_count); @@ -116,9 +110,9 @@ void testing_matmul_batch_offset_impl(const Arguments& arg) for(int b = 0; b < batch_count; b++) { // Initialize each batch's sub-matrix at the offset location - Ti* A_batch = h_A_full.data() + b * size_A_full + offset_a_elem; - Ti* B_batch = h_B_full.data() + b * size_B_full + offset_b_elem; - To* C_batch = h_C_full.data() + b * size_C_full + offset_c_elem; + Ti* A_batch = h_A_full.data() + b * size_A_full + offset_a; + Ti* B_batch = h_B_full.data() + b * size_B_full + offset_b; + To* C_batch = h_C_full.data() + b * size_C_full + offset_c; // Simple initialization: A and B with small integers for(int64_t j = 0; j < A_col; j++) @@ -308,9 +302,9 @@ void testing_matmul_batch_offset_impl(const Arguments& arg) for(int b = 0; b < batch_count; b++) { // Get pointers to sub-matrices (with element offset applied) - Ti* A_sub = h_A_full.data() + b * size_A_full + offset_a_elem; - Ti* B_sub = h_B_full.data() + b * size_B_full + offset_b_elem; - To* C_sub = h_C_full.data() + b * size_C_full + offset_c_elem; + Ti* A_sub = h_A_full.data() + b * size_A_full + offset_a; + Ti* B_sub = h_B_full.data() + b * size_B_full + offset_b; + To* C_sub = h_C_full.data() + b * size_C_full + offset_c; To* D_sub = h_D_gold.data() + b * size_D_sub; // Simple GEMM: D = alpha * A * B + beta * C @@ -346,7 +340,7 @@ void testing_matmul_batch_offset_impl(const Arguments& arg) for(int b = 0; b < batch_count; b++) { // GPU result is at (base + offset) within each batch's buffer - To* result_gpu = h_D_full.data() + b * size_D_full + offset_d_elem; + To* result_gpu = h_D_full.data() + b * size_D_full + offset_d; To* result_cpu = h_D_gold.data() + b * size_D_sub; for(size_t i = 0; i < size_D_sub; i++) diff --git a/projects/hipblaslt/library/include/hipblaslt/hipblaslt.h b/projects/hipblaslt/library/include/hipblaslt/hipblaslt.h index bcbb4631138e..795e58661493 100644 --- a/projects/hipblaslt/library/include/hipblaslt/hipblaslt.h +++ b/projects/hipblaslt/library/include/hipblaslt/hipblaslt.h @@ -177,7 +177,7 @@ typedef enum { /** Matrix Offset. * * For ``General Batched GEMM``, we can support for users to access a sub-matrix of - * the original matrix by adding an ``offset`` value from the base address. + * the original matrix by adding an ``offset`` value (in elements) from the base address. * Note that for non-batched or Strided Batch GEMM case, we can directly apply * the offset value by using the strided-offset value. */ diff --git a/projects/hipblaslt/library/src/amd_detail/include/auxiliary.hpp b/projects/hipblaslt/library/src/amd_detail/include/auxiliary.hpp index c5dd621ad5f3..320de121e4ca 100644 --- a/projects/hipblaslt/library/src/amd_detail/include/auxiliary.hpp +++ b/projects/hipblaslt/library/src/amd_detail/include/auxiliary.hpp @@ -158,6 +158,22 @@ constexpr const char* hip_datatype_to_string(hipDataType type) return "invalid"; } +// Returns true for sub-byte MX-style data types (fp6/fp4). +// Used to reject features that require byte-addressable elements. +HIPBLASLT_EXPORT +constexpr bool hip_datatype_is_mxtype(hipDataType type) +{ + switch(type) + { + case HIP_R_6F_E2M3: + case HIP_R_6F_E3M2: + case HIP_R_4F_E2M1: + return true; + default: + return false; + } +} + // return precision string for hipDataType HIPBLASLT_EXPORT constexpr const char* hipblas_computetype_to_string(hipblasComputeType_t type) diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/rocblaslt_mat_utils.hpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/rocblaslt_mat_utils.hpp index cd4e3fe4aa70..eab1a2cca15b 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/rocblaslt_mat_utils.hpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/rocblaslt_mat_utils.hpp @@ -229,6 +229,19 @@ inline rocblaslt_status validateMatmulArgs(int64_t m, return rocblaslt_status_invalid_size; } + // Batch offsets are expressed in elements and converted to bytes; sub-byte + // (MX) data types are not byte-addressable, so a nonzero offset is unsupported. + if((batch_offset_a != 0 && hip_datatype_is_mxtype(type_a)) + || (batch_offset_b != 0 && hip_datatype_is_mxtype(type_b)) + || (batch_offset_c != 0 && hip_datatype_is_mxtype(type_c)) + || (batch_offset_d != 0 && hip_datatype_is_mxtype(type_d))) + { +#ifndef CODE_COVERAGE + std::cerr << "matrix offset is not supported for sub-byte (MX) data types" << std::endl; +#endif + return rocblaslt_status_not_implemented; + } + // number of batches of matrics A,B,C,D must be the same and positive if(num_batches_a != num_batches_b || num_batches_a != num_batches_c || num_batches_a != num_batches_d || num_batches_a < 1) diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp index 4b8c6542d796..a7bc70cf97d3 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp @@ -2424,11 +2424,22 @@ namespace inputs.batchC = reinterpret_cast(prob.batch_C); inputs.batchD = reinterpret_cast(prob.batch_D); - // set the batched offset of A, B, C, and D. - inputs.batchOffsetA = prob.batch_offset_a; - inputs.batchOffsetB = prob.batch_offset_b; - inputs.batchOffsetC = prob.batch_offset_c; - inputs.batchOffsetD = prob.batch_offset_d; + // The batch offsets are specified by the user in elements; convert them to + // bytes here so the kernel/assembly can add them straight to byte addresses. + // Only data types whose element size is at least one byte are supported + // (sub-byte types such as fp4/fp6 are rejected during argument validation). + inputs.batchOffsetA + = prob.batch_offset_a + * size_t(TensileLite::DataTypeInfo::Get(hip2TensileType(prob.a_type)).elementSize); + inputs.batchOffsetB + = prob.batch_offset_b + * size_t(TensileLite::DataTypeInfo::Get(hip2TensileType(prob.b_type)).elementSize); + inputs.batchOffsetC + = prob.batch_offset_c + * size_t(TensileLite::DataTypeInfo::Get(hip2TensileType(prob.c_type)).elementSize); + inputs.batchOffsetD + = prob.batch_offset_d + * size_t(TensileLite::DataTypeInfo::Get(hip2TensileType(prob.d_type)).elementSize); // Set the GSU workspace inputs.ws = prob.workspace; From b3bc4ab08eb02bc13d95e23d4f715b9991b596c1 Mon Sep 17 00:00:00 2001 From: Seung-Hee Bae Date: Fri, 12 Jun 2026 16:01:19 -0400 Subject: [PATCH 18/27] remove temporary debug test and update matmul_batch_offset_large test. --- .../clients/tests/data/matmul_gtest.yaml | 33 +++---------------- 1 file changed, 4 insertions(+), 29 deletions(-) diff --git a/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml b/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml index 1d7472c88b8f..b7f6c2ad8541 100755 --- a/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml +++ b/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml @@ -3315,27 +3315,6 @@ Tests: unit_check: 1 norm_check: 1 -# Debug test - 1024x1024 using regular matmul with general batched (batch_mode=1) -- name: matmul_general_batched_debug - category: quick - function: matmul - precision: *real_precisions - transA: N - transB: N - M: 1024 - N: 1024 - K: 256 - lda: 1024 - ldb: 256 - ldc: 1024 - ldd: 1024 - batch_mode: 1 - batch_count: 2 - alpha: 1.0 - beta: 1.0 - unit_check: 1 - norm_check: 1 - # Large offset test - name: matmul_batch_offset_large category: nightly @@ -3352,14 +3331,10 @@ Tests: ldd: 1024 batch_mode: 1 batch_count: 4 - batch_offset_a: [1024, 4096] - batch_offset_b: [512, 2048] - batch_offset_c: [1024, 4096] - batch_offset_d: [1024, 4096] - #batch_offset_a: 4096 - #batch_offset_b: 2048 - #batch_offset_c: 4096 - #batch_offset_d: 4096 + batch_offset_a: 4096 + batch_offset_b: 2048 + batch_offset_c: [4096, 1000000] + batch_offset_d: [4096, 500000] alpha: 1.0 beta: 1.0 unit_check: 1 From 713ac691a356d1ce5804f55f2b9b0c3a43138fa7 Mon Sep 17 00:00:00 2001 From: Seung-Hee Bae Date: Mon, 15 Jun 2026 21:42:58 +0000 Subject: [PATCH 19/27] [hipblaslt] modify matmul_batch_offset_large test with actual 64-bit offset inputs. --- .../hipblaslt/clients/tests/data/matmul_gtest.yaml | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml b/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml index b7f6c2ad8541..a96fe5b8d273 100755 --- a/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml +++ b/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml @@ -3319,7 +3319,7 @@ Tests: - name: matmul_batch_offset_large category: nightly function: matmul_batch_offset - precision: *real_precisions + precision: [*hpa_half_precision, *hpa_bf16_precision] transA: N transB: N M: 1024 @@ -3331,10 +3331,10 @@ Tests: ldd: 1024 batch_mode: 1 batch_count: 4 - batch_offset_a: 4096 - batch_offset_b: 2048 - batch_offset_c: [4096, 1000000] - batch_offset_d: [4096, 500000] + batch_offset_a: 4294967296 + batch_offset_b: 4294967297 + batch_offset_c: 4294967298 + batch_offset_d: 4097 alpha: 1.0 beta: 1.0 unit_check: 1 From e2be7df9a97923766396cf9fc71f841fd31c28fc Mon Sep 17 00:00:00 2001 From: Seung-Hee Bae Date: Mon, 15 Jun 2026 23:46:20 +0000 Subject: [PATCH 20/27] limit the tested gpu_arch for matmul_batch_offset_large tests. --- projects/hipblaslt/clients/tests/data/matmul_gtest.yaml | 1 + 1 file changed, 1 insertion(+) diff --git a/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml b/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml index a96fe5b8d273..dd4ca35262e8 100755 --- a/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml +++ b/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml @@ -3339,5 +3339,6 @@ Tests: beta: 1.0 unit_check: 1 norm_check: 1 + gpu_arch: '9(42|50)' ... From 7f13345658e53bc530a675d9e670256fad409573 Mon Sep 17 00:00:00 2001 From: Seung-Hee Bae Date: Wed, 17 Jun 2026 16:54:41 +0000 Subject: [PATCH 21/27] [hipblaslt] remove checking the positive offset value --- .../rocblaslt/src/include/rocblaslt_mat_utils.hpp | 9 --------- 1 file changed, 9 deletions(-) diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/rocblaslt_mat_utils.hpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/rocblaslt_mat_utils.hpp index eab1a2cca15b..90e93c005e5d 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/rocblaslt_mat_utils.hpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/rocblaslt_mat_utils.hpp @@ -220,15 +220,6 @@ inline rocblaslt_status validateMatmulArgs(int64_t m, return rocblaslt_status_invalid_size; } - // sizes must not be negative - if(batch_offset_a < 0 || batch_offset_b < 0 || batch_offset_c < 0 || batch_offset_d < 0) - { -#ifndef CODE_COVERAGE - std::cerr << "matrix offset size must be zero or positive" << std::endl; -#endif - return rocblaslt_status_invalid_size; - } - // Batch offsets are expressed in elements and converted to bytes; sub-byte // (MX) data types are not byte-addressable, so a nonzero offset is unsupported. if((batch_offset_a != 0 && hip_datatype_is_mxtype(type_a)) From 5187d51ad7e02e4134cc04e288740e45f6b83553 Mon Sep 17 00:00:00 2001 From: Seung-Hee Bae Date: Wed, 24 Jun 2026 21:16:41 +0000 Subject: [PATCH 22/27] [hipblaslt] add matmul_batch_offset_all_solutions test. --- .../include/testing_matmul_batch_offset.hpp | 172 +++++++++++------- .../clients/tests/data/matmul_gtest.yaml | 31 +++- 2 files changed, 133 insertions(+), 70 deletions(-) diff --git a/projects/hipblaslt/clients/common/include/testing_matmul_batch_offset.hpp b/projects/hipblaslt/clients/common/include/testing_matmul_batch_offset.hpp index 8cb2ad8d46d5..0b86f0e91f51 100644 --- a/projects/hipblaslt/clients/common/include/testing_matmul_batch_offset.hpp +++ b/projects/hipblaslt/clients/common/include/testing_matmul_batch_offset.hpp @@ -40,6 +40,7 @@ #include "unit.hpp" #include "utility.hpp" #include +#include /* ============================================================================================ */ /*! \brief Test for 64-bit batch offset support in general batched GEMM */ @@ -238,15 +239,21 @@ void testing_matmul_batch_offset_impl(const Arguments& arg) // Find algorithm hipblasLtMatmulPreference_t pref; CHECK_HIPBLASLT_ERROR(hipblasLtMatmulPreferenceCreate(&pref)); - size_t maxWorkspaceSize = 128 * 1024 * 1024; // 128 MB + size_t prefMaxWorkspaceSize = 128 * 1024 * 1024; // 128 MB CHECK_HIPBLASLT_ERROR(hipblasLtMatmulPreferenceSetAttribute( - pref, HIPBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES, &maxWorkspaceSize, sizeof(maxWorkspaceSize))); + pref, HIPBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES, &prefMaxWorkspaceSize, sizeof(prefMaxWorkspaceSize))); - int numAlgos = 0; - const int requestedAlgos = 1; - hipblasLtMatmulHeuristicResult_t heuristicResult[requestedAlgos]; + // Support testing multiple solutions via requested_solution_num parameter + // Default to 1 (test only best solution) for backward compatibility + // Use HIPBLASLT_MAX_REQUESTED_SOLUTION_NUM when -1 to get all available solutions + int32_t requestedAlgos = (arg.requested_solution_num < 0) ? HIPBLASLT_MAX_REQUESTED_SOLUTION_NUM + : (arg.requested_solution_num == 0) ? 1 + : arg.requested_solution_num; + + std::vector heuristicResult(requestedAlgos); + int numAlgos = 0; CHECK_HIPBLASLT_ERROR(hipblasLtMatmulAlgoGetHeuristic( - handle, matmul_desc, matA, matB, matC, matD, pref, requestedAlgos, heuristicResult, &numAlgos)); + handle, matmul_desc, matA, matB, matC, matD, pref, requestedAlgos, heuristicResult.data(), &numAlgos)); if(numAlgos == 0) { @@ -254,50 +261,23 @@ void testing_matmul_batch_offset_impl(const Arguments& arg) GTEST_SKIP() << "No algorithm found for this configuration"; } - // Allocate workspace - size_t workspaceSize = heuristicResult[0].workspaceSize; - void* d_workspace = nullptr; - if(workspaceSize > 0) + // Find maximum workspace size across all solutions + size_t maxWorkspaceSize = 0; + for(int i = 0; i < numAlgos; i++) { - CHECK_HIP_ERROR(hipMalloc(&d_workspace, workspaceSize)); + maxWorkspaceSize = std::max(maxWorkspaceSize, heuristicResult[i].workspaceSize); } - CHECK_HIPBLASLT_ERROR(hipblasLtMatmul(handle, - matmul_desc, - &h_alpha, - d_batch_A, - matA, - d_batch_B, - matB, - &h_beta, - d_batch_C, - matC, - d_batch_D, - matD, - &heuristicResult[0].algo, - d_workspace, - workspaceSize, - 0)); - - CHECK_HIPBLASLT_ERROR(hipblasLtMatmulPreferenceDestroy(pref)); - - // Ensure kernel completes before reading result - CHECK_HIP_ERROR(hipDeviceSynchronize()); - - if(d_workspace) + // Allocate workspace + void* d_workspace = nullptr; + if(maxWorkspaceSize > 0) { - CHECK_HIP_ERROR(hipFree(d_workspace)); + CHECK_HIP_ERROR(hipMalloc(&d_workspace, maxWorkspaceSize)); } - // Copy GPU result - CHECK_HIP_ERROR(hipMemcpy(h_D_full.data(), - d_D_full, - sizeof(To) * size_D_full * batch_count, - hipMemcpyDeviceToHost)); - // ======================================== // CPU Reference: D = alpha * A * B + beta * C - // Simple manual GEMM for testing (avoids cblas_gemm complexity) + // Computed once before testing any solutions // ======================================== for(int b = 0; b < batch_count; b++) { @@ -333,49 +313,107 @@ void testing_matmul_batch_offset_impl(const Arguments& arg) } } - // ======================================== - // VALIDATION: Compare GPU vs CPU - // ======================================== - double max_error = 0.0; - for(int b = 0; b < batch_count; b++) + // Tolerance: epsilon * factor * K (accumulation over K elements) + double tol = std::numeric_limits::epsilon() * 100 * K; + + // Track validation results across all solutions + int numPassed = 0; + int numFailed = 0; + + // Test each solution + for(int algoIdx = 0; algoIdx < numAlgos; algoIdx++) { - // GPU result is at (base + offset) within each batch's buffer - To* result_gpu = h_D_full.data() + b * size_D_full + offset_d; - To* result_cpu = h_D_gold.data() + b * size_D_sub; + // Reset only D output buffer to sentinel values before testing each solution + // A, B, C are inputs and don't change, so we can reuse them + CHECK_HIP_ERROR(hipMemcpy(d_D_full, + h_D_full.data(), + sizeof(To) * size_D_full * batch_count, + hipMemcpyHostToDevice)); + + CHECK_HIPBLASLT_ERROR(hipblasLtMatmul(handle, + matmul_desc, + &h_alpha, + d_batch_A, + matA, + d_batch_B, + matB, + &h_beta, + d_batch_C, + matC, + d_batch_D, + matD, + &heuristicResult[algoIdx].algo, + d_workspace, + heuristicResult[algoIdx].workspaceSize, + 0)); + + // Ensure kernel completes before reading result + CHECK_HIP_ERROR(hipDeviceSynchronize()); + + // Copy GPU result back for verification + CHECK_HIP_ERROR(hipMemcpy(h_D_full.data(), + d_D_full, + sizeof(To) * size_D_full * batch_count, + hipMemcpyDeviceToHost)); + + // ======================================== + // VALIDATION: Compare GPU vs CPU + // ======================================== + double max_error = 0.0; + for(int b = 0; b < batch_count; b++) + { + // GPU result is at (base + offset) within each batch's buffer + To* result_gpu = h_D_full.data() + b * size_D_full + offset_d; + To* result_cpu = h_D_gold.data() + b * size_D_sub; + + for(size_t i = 0; i < size_D_sub; i++) + { + double diff = std::abs(double(result_gpu[i]) - double(result_cpu[i])); + max_error = std::max(max_error, diff); + } + } - for(size_t i = 0; i < size_D_sub; i++) + // Check and count per-solution results + if(arg.unit_check) { - double diff = std::abs(double(result_gpu[i]) - double(result_cpu[i])); - max_error = std::max(max_error, diff); + if(max_error >= tol) + { + numFailed++; + EXPECT_LT(max_error, tol) + << "Solution " << algoIdx << "/" << numAlgos + << " FAILED (error: " << max_error << ", tol: " << tol << ")"; + } + else + { + numPassed++; + } } } + // Report summary when testing multiple solutions + if(numAlgos > 1 && arg.unit_check) + { + hipblaslt_cout << "Tested " << numAlgos << " solutions: " + << numPassed << " passed, " << numFailed << " failed" << std::endl; + } + // Cleanup + if(d_workspace) + { + CHECK_HIP_ERROR(hipFree(d_workspace)); + } CHECK_HIP_ERROR(hipFree(d_batch_A)); CHECK_HIP_ERROR(hipFree(d_batch_B)); CHECK_HIP_ERROR(hipFree(d_batch_C)); CHECK_HIP_ERROR(hipFree(d_batch_D)); + CHECK_HIPBLASLT_ERROR(hipblasLtMatmulPreferenceDestroy(pref)); CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutDestroy(matA)); CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutDestroy(matB)); CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutDestroy(matC)); CHECK_HIPBLASLT_ERROR(hipblasLtMatrixLayoutDestroy(matD)); CHECK_HIPBLASLT_ERROR(hipblasLtMatmulDescDestroy(matmul_desc)); CHECK_HIPBLASLT_ERROR(hipblasLtDestroy(handle)); - - // Tolerance: epsilon * factor * K (accumulation over K elements) - double tol = std::numeric_limits::epsilon() * 100 * K; - - // Report results - if(arg.unit_check) - { - EXPECT_LT(max_error, tol) << "GPU vs CPU mismatch (error: " << max_error << ", tol: " << tol << ")"; - } - - if(arg.norm_check && max_error >= tol) - { - hipblaslt_cout << "GPU vs CPU max error: " << max_error << " (tol: " << tol << ")" << std::endl; - } } // Type dispatcher based on Arguments diff --git a/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml b/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml index dd4ca35262e8..b1b071a1fa5d 100755 --- a/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml +++ b/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml @@ -3285,7 +3285,7 @@ Tests: ldc: 256 ldd: 256 batch_mode: 1 - batch_count: 2 + batch_count: 3 batch_offset_a: 128 batch_offset_b: 128 batch_offset_c: 128 @@ -3306,7 +3306,7 @@ Tests: N: 128 K: 64 batch_mode: 1 - batch_count: 2 + batch_count: 4 batch_offset_a: 64 batch_offset_b: 64 batch_offset_c: 64 @@ -3330,7 +3330,7 @@ Tests: ldc: 1024 ldd: 1024 batch_mode: 1 - batch_count: 4 + batch_count: 2 batch_offset_a: 4294967296 batch_offset_b: 4294967297 batch_offset_c: 4294967298 @@ -3341,4 +3341,29 @@ Tests: norm_check: 1 gpu_arch: '9(42|50)' +# Test all solutions with batch offsets and all transpose combinations +- name: matmul_batch_offset_all_solutions + category: nightly + function: matmul_batch_offset + precision: *real_precisions + transA_transB: *transA_transB_range + M: 256 + N: 256 + K: 128 + lda: 256 + ldb: 256 + ldc: 256 + ldd: 256 + batch_mode: 1 + batch_count: 2 + batch_offset_a: 256 + batch_offset_b: 128 + batch_offset_c: 512 + batch_offset_d: 256 + alpha: 1.0 + beta: 1.0 + algo_method: [0,1] + requested_solution_num: -1 + unit_check: 1 + ... From f939800e81a7e006b7b3af3eea81cdcd51f93499 Mon Sep 17 00:00:00 2001 From: Seung-Hee Bae Date: Wed, 24 Jun 2026 21:22:20 +0000 Subject: [PATCH 23/27] [hipblaslt] don't add batchOffset arguments to the Custom Kernels, and excludes custom kernels solutions for General batched GEMM. --- .../amd_detail/rocblaslt/src/tensile_host.cpp | 18 ++++++++++++++++++ .../tensilelite/src/ContractionSolution.cpp | 4 ++-- 2 files changed, 20 insertions(+), 2 deletions(-) diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp index a7bc70cf97d3..da394dcd3b5f 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp @@ -4481,6 +4481,24 @@ rocblaslt_status getAllSolutions(MyProblem& int duplicated_counts = 0; for(auto solution : solutions) { + // Custom kernels don't support general batched mode (pointer arrays) + // Only check for ContractionProblemGemm (grouped gemm doesn't use batchMode) + if constexpr(std::is_same::value) + { + if(prob.batchMode() == TensileLite::ContractionProblemGemm::BATCHMODE::POINTER_ARRAY + && !solution->sizeMapping.customKernelName.empty()) + { + if(get_logger_layer_mode() & rocblaslt_layer_mode_log_info) + { + std::ostringstream msg; + msg << "Skipping custom kernel " << solution->sizeMapping.customKernelName + << " - does not support batch_mode=POINTER_ARRAY" << std::endl; + log_info(__func__, msg.str()); + } + continue; + } + } + //workaround: findAllSolutions should get all solutions without duplications bool duplicated_sol = false; for(int j = 0; j < i; j++) diff --git a/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp b/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp index da452edd3b48..841aa279c166 100644 --- a/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp +++ b/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp @@ -1844,7 +1844,7 @@ namespace TensileLite // Batch offset support for General Batched GEMM (SupportUserArgs kernels). // Appended at the tail, after the dstD/Synchronizer block, to match the // kernel signature order (see Signature.py). - if(!problemType.groupedGemm) + if(!problemType.groupedGemm && sizeMapping.customKernelName.empty()) { rv.args.append("batchOffsetD", inputs.batchOffsetD); rv.args.append("batchOffsetC", inputs.batchOffsetC); @@ -2449,7 +2449,7 @@ namespace TensileLite } // Adding the batchmode kernel argument for post GSU kernel to determine // how to index the batch dimension in Strided Batch versus General Batched. - if(problemType.groupedGemm == false) + if(problemType.groupedGemm == false && sizeMapping.customKernelName.empty()) { ContractionProblemGemm::BATCHMODE batchMode = problem.batchMode(); args.template append("batchMode", static_cast(batchMode)); From b560fcc0505dd10aa1e7ad31b6807510b8897a9b Mon Sep 17 00:00:00 2001 From: Seung-Hee Bae Date: Wed, 24 Jun 2026 22:02:49 +0000 Subject: [PATCH 24/27] [hipblaslt] add tests for negative batch offset values. --- .../include/testing_matmul_batch_offset.hpp | 63 ++++++++++++------- .../clients/tests/data/matmul_gtest.yaml | 47 +++++++++++++- 2 files changed, 85 insertions(+), 25 deletions(-) diff --git a/projects/hipblaslt/clients/common/include/testing_matmul_batch_offset.hpp b/projects/hipblaslt/clients/common/include/testing_matmul_batch_offset.hpp index 0b86f0e91f51..4ca921f119b7 100644 --- a/projects/hipblaslt/clients/common/include/testing_matmul_batch_offset.hpp +++ b/projects/hipblaslt/clients/common/include/testing_matmul_batch_offset.hpp @@ -87,16 +87,21 @@ void testing_matmul_batch_offset_impl(const Arguments& arg) size_t size_C_sub = size_t(ldc) * size_t(N); size_t size_D_sub = size_t(ldd) * size_t(N); - // Full buffer sizes: [offset padding] + [matrix data] - // The offset padding comes FIRST because: - // - Pointer array points to buffer BASE - // - Kernel adds offset to BASE to get actual data location - // - So actual data is at BASE + offset - // Offsets are already in elements (the API takes element offsets). - size_t size_A_full = offset_a + size_A_sub; - size_t size_B_full = offset_b + size_B_sub; - size_t size_C_full = offset_c + size_C_sub; - size_t size_D_full = offset_d + size_D_sub; + // Calculate padding needed for negative offsets (in elements) + // If offset is negative, we need padding BEFORE the base pointer + // If offset is positive, padding is at the beginning (offset space) + size_t padding_a = (offset_a < 0) ? size_t(-offset_a) : 0; + size_t padding_b = (offset_b < 0) ? size_t(-offset_b) : 0; + size_t padding_c = (offset_c < 0) ? size_t(-offset_c) : 0; + size_t padding_d = (offset_d < 0) ? size_t(-offset_d) : 0; + + // Full buffer sizes: [padding for negative offsets] + [matrix data] + [positive offset space] + // Layout for negative offset: [padding|matrix_data] base points to start of matrix_data + // Layout for positive offset: [matrix_data|padding] base points to start of buffer + size_t size_A_full = padding_a + size_A_sub + (offset_a > 0 ? size_t(offset_a) : 0); + size_t size_B_full = padding_b + size_B_sub + (offset_b > 0 ? size_t(offset_b) : 0); + size_t size_C_full = padding_c + size_C_sub + (offset_c > 0 ? size_t(offset_c) : 0); + size_t size_D_full = padding_d + size_D_sub + (offset_d > 0 ? size_t(offset_d) : 0); // Allocate host memory for full buffers host_vector h_A_full(size_A_full * batch_count); @@ -106,14 +111,21 @@ void testing_matmul_batch_offset_impl(const Arguments& arg) host_vector h_D_gold(size_D_sub * batch_count); // CPU reference // Initialize matrices with known pattern - // Data must be placed at (base + offset) because the offset API tells the kernel - // to start reading at that location + // For negative offsets: base points to (padding + 0), data at (base + negative_offset) = padding region + // For positive offsets: base points to 0, data at (base + positive_offset) = offset region for(int b = 0; b < batch_count; b++) { - // Initialize each batch's sub-matrix at the offset location - Ti* A_batch = h_A_full.data() + b * size_A_full + offset_a; - Ti* B_batch = h_B_full.data() + b * size_B_full + offset_b; - To* C_batch = h_C_full.data() + b * size_C_full + offset_c; + // Calculate base pointer for this batch (after padding for negative offsets) + Ti* A_base = h_A_full.data() + b * size_A_full + padding_a; + Ti* B_base = h_B_full.data() + b * size_B_full + padding_b; + To* C_base = h_C_full.data() + b * size_C_full + padding_c; + + // Data location is at (base + offset) + // For negative offset: base + (-N) points backward into padding region + // For positive offset: base + (+N) points forward into buffer + Ti* A_batch = A_base + offset_a; + Ti* B_batch = B_base + offset_b; + To* C_batch = C_base + offset_c; // Simple initialization: A and B with small integers for(int64_t j = 0; j < A_col; j++) @@ -144,6 +156,7 @@ void testing_matmul_batch_offset_impl(const Arguments& arg) hipMemcpy(d_C_full, h_C_full.data(), sizeof(To) * size_C_full * batch_count, hipMemcpyHostToDevice)); // Setup pointer arrays for base addresses + // Base pointers must point AFTER the padding (for negative offset support) std::vector h_batch_A(batch_count); std::vector h_batch_B(batch_count); std::vector h_batch_C(batch_count); @@ -151,10 +164,10 @@ void testing_matmul_batch_offset_impl(const Arguments& arg) for(int b = 0; b < batch_count; b++) { - h_batch_A[b] = d_A_full + b * size_A_full; - h_batch_B[b] = d_B_full + b * size_B_full; - h_batch_C[b] = d_C_full + b * size_C_full; - h_batch_D[b] = d_D_full + b * size_D_full; + h_batch_A[b] = d_A_full + b * size_A_full + padding_a; + h_batch_B[b] = d_B_full + b * size_B_full + padding_b; + h_batch_C[b] = d_C_full + b * size_C_full + padding_c; + h_batch_D[b] = d_D_full + b * size_D_full + padding_d; } // Allocate device memory for pointer arrays @@ -282,9 +295,10 @@ void testing_matmul_batch_offset_impl(const Arguments& arg) for(int b = 0; b < batch_count; b++) { // Get pointers to sub-matrices (with element offset applied) - Ti* A_sub = h_A_full.data() + b * size_A_full + offset_a; - Ti* B_sub = h_B_full.data() + b * size_B_full + offset_b; - To* C_sub = h_C_full.data() + b * size_C_full + offset_c; + // Base is at padding, then add offset (which may be negative) + Ti* A_sub = h_A_full.data() + b * size_A_full + padding_a + offset_a; + Ti* B_sub = h_B_full.data() + b * size_B_full + padding_b + offset_b; + To* C_sub = h_C_full.data() + b * size_C_full + padding_c + offset_c; To* D_sub = h_D_gold.data() + b * size_D_sub; // Simple GEMM: D = alpha * A * B + beta * C @@ -363,7 +377,8 @@ void testing_matmul_batch_offset_impl(const Arguments& arg) for(int b = 0; b < batch_count; b++) { // GPU result is at (base + offset) within each batch's buffer - To* result_gpu = h_D_full.data() + b * size_D_full + offset_d; + // Base is at padding_d, then add offset_d (which may be negative) + To* result_gpu = h_D_full.data() + b * size_D_full + padding_d + offset_d; To* result_cpu = h_D_gold.data() + b * size_D_sub; for(size_t i = 0; i < size_D_sub; i++) diff --git a/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml b/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml index b1b071a1fa5d..5f0f5815dac6 100755 --- a/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml +++ b/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml @@ -3362,8 +3362,53 @@ Tests: batch_offset_d: 256 alpha: 1.0 beta: 1.0 - algo_method: [0,1] requested_solution_num: -1 unit_check: 1 +# Test with negative batch offsets to verify proper memory layout handling +- name: matmul_batch_offset_negative + category: nightly + function: matmul_batch_offset + precision: *real_precisions + transA_transB: *transA_transB_range + M: 256 + N: 256 + K: 128 + lda: 256 + ldb: 256 + ldc: 256 + ldd: 256 + batch_mode: 1 + batch_count: 2 + batch_offset_a: -128 + batch_offset_b: -64 + batch_offset_c: -256 + batch_offset_d: -128 + alpha: 1.0 + beta: 1.0 + unit_check: 1 + +# Test with mixed positive and negative offsets +- name: matmul_batch_offset_mixed + category: nightly + function: matmul_batch_offset + precision: *real_precisions + transA_transB: *transA_transB_range + M: 256 + N: 256 + K: 128 + lda: 256 + ldb: 256 + ldc: 256 + ldd: 256 + batch_mode: 1 + batch_count: 2 + batch_offset_a: -64 + batch_offset_b: 128 + batch_offset_c: -128 + batch_offset_d: 256 + alpha: 1.0 + beta: 1.0 + unit_check: 1 + ... From b92d9b9839224d7cdff1abb89571f101fb3e5357 Mon Sep 17 00:00:00 2001 From: Seung-Hee Bae Date: Thu, 25 Jun 2026 00:13:19 +0000 Subject: [PATCH 25/27] [hipblaslt] check validity with non-zero offsets with POINTER_ARRAY mode only. --- .../src/include/rocblaslt_mat_utils.hpp | 26 +++++++++++++++++++ 1 file changed, 26 insertions(+) diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/rocblaslt_mat_utils.hpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/rocblaslt_mat_utils.hpp index 90e93c005e5d..90db7c08ac0f 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/rocblaslt_mat_utils.hpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/rocblaslt_mat_utils.hpp @@ -419,6 +419,32 @@ inline rocblaslt_status rocblaslt_matmul_valid_args(const rocblaslt_matmul_desc n = num_cols_d; k = (opA == HIPBLAS_OP_N) ? num_cols_a : num_rows_a; + // Validate: batch offsets are only valid with POINTER_ARRAY mode (general batched) + bool hasNonZeroOffset = (batch_offset_a != 0) || (batch_offset_b != 0) + || (batch_offset_c != 0) || (batch_offset_d != 0); + + if(hasNonZeroOffset) + { + // Check that all matrices are using POINTER_ARRAY mode + if(matA->batch_mode != HIPBLASLT_BATCH_MODE_POINTER_ARRAY + || matB->batch_mode != HIPBLASLT_BATCH_MODE_POINTER_ARRAY + || matC->batch_mode != HIPBLASLT_BATCH_MODE_POINTER_ARRAY + || matD->batch_mode != HIPBLASLT_BATCH_MODE_POINTER_ARRAY) + { + log_error(__func__, + "Batch offsets require all matrices to use batch_mode=POINTER_ARRAY. ", + "Current modes: A=", matA->batch_mode, + ", B=", matB->batch_mode, + ", C=", matC->batch_mode, + ", D=", matD->batch_mode, + ". Offsets: A=", batch_offset_a, + ", B=", batch_offset_b, + ", C=", batch_offset_c, + ", D=", batch_offset_d); + return rocblaslt_status_invalid_value; + } + } + auto matmul_status = validateMatmulArgs(m, n, k, From 5918c7510f96a2691568a438303f6432cfe9a1d5 Mon Sep 17 00:00:00 2001 From: Seung-Hee Bae Date: Wed, 1 Jul 2026 03:46:24 -0400 Subject: [PATCH 26/27] [hipblaslt] fix batch offset test failure with GSU + MBSK. --- .../tensilelite/Tensile/KernelWriterAssembly.py | 11 +++++++++++ 1 file changed, 11 insertions(+) diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py index 7312c4610bae..62e1d7d18a11 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py @@ -13289,6 +13289,11 @@ def computeStoreSrdStart(self, kernel, srdTcList: list, sgprBpeList = [], useSiz module.add(SAddCU32(dst=sgpr("Srd%s+1"%mat), src0=sgpr("Srd%s+1"%mat), src1=sgpr(tmpS1), comment="Offsetting within the Batch Matrix [Higher half of address]")) # Now, we have starting matrix address of a specific batch in the corresponding Srd. # Load and apply batch offset for General Batched GEMM (C or D matrix) as necessary. + # This block sits inside the generalBatchedGemmLoad label, which the GSU routing + # only reaches at runtime GSU==1 -- where SrdC/SrdD are the real user pointer arrays + # (not the GSU workspace), so the offset is correctly applied here. When GSU>1 the + # routing branches to the strided/workspace path and skips this block; the PostGSU + # conversion kernel applies the offset when it writes the final result to C/D. if not kernel["ProblemType"]["GroupedGemm"]: batchOffsetKernArgOffset = self.states.batchOffsetCKernArgOffset if mat == "C" else self.states.batchOffsetDKernArgOffset module.add(SLoadB64(dst=sgpr(tmpS0, 2), base=sgpr("KernArgAddress", 2), soffset=hex(batchOffsetKernArgOffset), comment="Load batchOffset%s from kernel args"%mat)) @@ -13388,6 +13393,12 @@ def SrdTDInit(self, kernel): module.add(SAddCU32(dst=sgpr(tmpspgr2+1), src0=sgpr("AddressTD+1"), src1=0, comment="Offsetting to the location [Higher half of address]")) module.add(SLoadB64(dst=sgpr("SrdTD", 2), base=sgpr(tmpspgr2, 2), soffset=0, comment="Load the Matrix Address in the Pointer Array")) module.add(SWaitCnt(kmcnt=0, comment="Wait for the Matrix Address Load from the Pointer Array")) + # Load and apply batch offset for General Batched GEMM (dstD) as necessary. + if not kernel["ProblemType"]["GroupedGemm"]: + module.add(SLoadB64(dst=sgpr(tmpspgr2, 2), base=sgpr("KernArgAddress", 2), soffset=hex(self.states.batchOffsetDKernArgOffset), comment="Load batchOffsetD from kernel args")) + module.add(SWaitCnt(kmcnt=0, comment="Wait for batchOffsetD Load")) + module.add(SAddU32(dst=sgpr("SrdTD+0"), src0=sgpr("SrdTD+0"), src1=sgpr(tmpspgr2+0), comment="Apply batchOffsetD to SrdTD (low)")) + module.add(SAddCU32(dst=sgpr("SrdTD+1"), src0=sgpr("SrdTD+1"), src1=sgpr(tmpspgr2+1), comment="Apply batchOffsetD to SrdTD (high)")) module.add(SrdTDGeneralBatched_End) module.add(SAddU32(dst=sgpr("SrdTD+0"), src0=sgpr("SrdTD+0"), src1=sgpr(tmpspgr1+0), comment="add lo to SRTD" )) module.add(SAddCU32(dst=sgpr("SrdTD+1"), src0=sgpr("SrdTD+1"), src1=sgpr(tmpspgr1+1), comment="add hi to SRTD" )) From 05fff7deb57838cee6c63a7c08186a3c0d0a475c Mon Sep 17 00:00:00 2001 From: Seung-Hee Bae Date: Wed, 1 Jul 2026 03:54:59 -0400 Subject: [PATCH 27/27] [hipblaslt] fix conversion kernel arguments alignment issue with batch offset args. --- .../Tensile/KernelWriterConversion.py | 24 +++++++++++++------ .../include/Tensile/ContractionSolution.hpp | 2 +- .../include/Tensile/KernelArguments.hpp | 14 +++++++++++ .../tensilelite/src/ContractionSolution.cpp | 16 +++++++++---- 4 files changed, 43 insertions(+), 13 deletions(-) diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriterConversion.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriterConversion.py index 9f8691393abc..d685abe2a73e 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriterConversion.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriterConversion.py @@ -350,9 +350,13 @@ def kernelBody(self): ######################################## # kernel start kStr += self.endLine + # Declare batchIdx for indexing pointer arrays in general batched mode + if not self.state["ProblemType"]["GroupedGemm"]: + kStr += " uint64_t batchIdx = 0;%s" % self.endLine + if not self.state["ProblemType"]["GroupedGemm"]: kStr += " if(batch_mode == 0)" + self.endLine - kStr += " {" + self.endLine + kStr += " {" + self.endLine kStr += " if (id*NUM_ELEMENT_LOAD >= (arg.size%s" % self.indexChars[0] for i in range(1, problemType["NumIndicesC"]): kStr += " * arg.size%s" % self.indexChars[i] @@ -362,14 +366,14 @@ def kernelBody(self): kStr += " }" + self.endLine kStr += " else" + self.endLine kStr += " {" + self.endLine - kStr += " uint64_t index2 = ((id*NUM_ELEMENT_LOAD) / (arg.size%s" % self.indexChars[0] + kStr += " batchIdx = ((id*NUM_ELEMENT_LOAD) / (arg.size%s" % self.indexChars[0] for i in range(1, problemType["NumIndicesC"]-1): kStr += " * arg.size%s" % self.indexChars[i] kStr += " + additionalPaddingPerBatch));%s" % self.endLine - kStr += " if (id*NUM_ELEMENT_LOAD >= ((index2+1) * (arg.size%s * arg.size%s)) + index2 * additionalPaddingPerBatch)%s" % (self.indexChars[0], self.indexChars[1], self.endLine) + kStr += " if (id*NUM_ELEMENT_LOAD >= ((batchIdx+1) * (arg.size%s * arg.size%s)) + batchIdx * additionalPaddingPerBatch)%s" % (self.indexChars[0], self.indexChars[1], self.endLine) kStr += " return;%s" % self.endLine - kStr += " if(index2 > 0)%s" % self.endLine - kStr += " id = id - (index2 * additionalPaddingPerBatch) / NUM_ELEMENT_LOAD;%s" % self.endLine + kStr += " if(batchIdx > 0)%s" % self.endLine + kStr += " id = id - (batchIdx * additionalPaddingPerBatch) / NUM_ELEMENT_LOAD;%s" % self.endLine kStr += " }" + self.endLine kStr += self.endLine @@ -386,6 +390,10 @@ def kernelBody(self): kStr += " id%d = id %% arg.size%s;%s" % (i, self.indexChars[i], self.endLine) kStr += " id = id / arg.size%s;%s" % (self.indexChars[i], self.endLine) + # Set batchIdx = id2 for strided batched mode (batch_mode == 0) + if not self.state["ProblemType"]["GroupedGemm"]: + kStr += " if(batch_mode == 0) batchIdx = id2;%s" % self.endLine + nonTileFreeIndices = [] ######################################## @@ -727,7 +735,8 @@ def kernelBody(self): kStr += " }" + self.endLine kStr += " else" + self.endLine kStr += " {" + self.endLine - kStr += " %s *ptr = *(reinterpret_cast<%s **>(((char *)arg.C) + (8*id2))) + batchOffsetC/sizeof(%s);" % (destTypeStr, destTypeStr, destTypeStr) + self.endLine + # Dereference C pointer array and apply the batch offset (in bytes). + kStr += " %s *ptr = *(reinterpret_cast<%s **>(((char *)arg.C) + (8*batchIdx))) + batchOffsetC/sizeof(%s);" % (destTypeStr, destTypeStr, destTypeStr) + self.endLine for vIdx in range(self.num_dword_load): kStr += " %s[%d] += arg.beta * (%s)ptr[idxC+%d];%s" % (accumStr, vIdx, intermediateDataType, vIdx, self.endLine) kStr += " }" + self.endLine @@ -823,7 +832,8 @@ def kernelBody(self): kStr += " if(batch_mode == 0) {" + self.endLine kStr += " buffer_store<%s, sizeof(%s), CacheOperation::Kind::Always>(*(%s *)%s, arg.D, idxD * sizeof(%s), 0);%s" % (storeTypeStr, storeTypeStr, storeTypeStr, resultStr, destTypeStr, self.endLine) kStr += " } else {" + self.endLine - kStr += " %s *ptr = *(reinterpret_cast<%s **>(((char *)arg.D) + (8*id2))) + batchOffsetD/sizeof(%s);" % (destTypeStr, destTypeStr, destTypeStr) + self.endLine + # Dereference D pointer array and apply the batch offset (in bytes). + kStr += " %s *ptr = *(reinterpret_cast<%s **>(((char *)arg.D) + (8*batchIdx))) + batchOffsetD/sizeof(%s);" % (destTypeStr, destTypeStr, destTypeStr) + self.endLine kStr += " buffer_store<%s, sizeof(%s), CacheOperation::Kind::Always>(*(%s *)%s, ptr, idxD * sizeof(%s), 0);%s" % (storeTypeStr, storeTypeStr, storeTypeStr, resultStr, destTypeStr, self.endLine) kStr += " }" + self.endLine else: diff --git a/projects/hipblaslt/tensilelite/include/Tensile/ContractionSolution.hpp b/projects/hipblaslt/tensilelite/include/Tensile/ContractionSolution.hpp index f2fc1d8fa5c3..0c68cca2d1fd 100644 --- a/projects/hipblaslt/tensilelite/include/Tensile/ContractionSolution.hpp +++ b/projects/hipblaslt/tensilelite/include/Tensile/ContractionSolution.hpp @@ -529,7 +529,7 @@ namespace TensileLite KA& args, StreamKSettings const& sk, uint32_t autoGsuVal, - uint32_t additionalPaddingPerBatchGeneralBatch=0) const; + uint32_t additionalPaddingPerBatchGeneralBatch=0) const; template inline void calculateConversionCallWorkGroupItems( diff --git a/projects/hipblaslt/tensilelite/include/Tensile/KernelArguments.hpp b/projects/hipblaslt/tensilelite/include/Tensile/KernelArguments.hpp index c4a1518bf2dc..210ee40563cb 100644 --- a/projects/hipblaslt/tensilelite/include/Tensile/KernelArguments.hpp +++ b/projects/hipblaslt/tensilelite/include/Tensile/KernelArguments.hpp @@ -738,6 +738,20 @@ namespace TensileLite counter += sizeof(value); } + void alignTo(size_t alignment) + { + size_t extraElements = counter % alignment; + size_t padding = (alignment - extraElements) % alignment; + counter += padding; + } + + template + inline void appendAligned(std::string const& name, T value) + { + alignTo(alignof(T)); + append(name, value); + } + template inline void appendUnbound(std::string const& name) { diff --git a/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp b/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp index 841aa279c166..02689c9b6b5f 100644 --- a/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp +++ b/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp @@ -2274,7 +2274,7 @@ namespace TensileLite KA& args, StreamKSettings const& sk, uint32_t autoGsuVal, - uint32_t additionalPaddingPerBatchGeneralBatch) const + uint32_t additionalPaddingPerBatchGeneralBatch) const { TensorDescriptor const& c = problem.c(); TensorDescriptor const& d = problem.d(); @@ -2453,10 +2453,16 @@ namespace TensileLite { ContractionProblemGemm::BATCHMODE batchMode = problem.batchMode(); args.template append("batchMode", static_cast(batchMode)); - args.template append("additionalPaddingPerBatch", additionalPaddingPerBatchGeneralBatch); - - args.template append("batchOffsetD", inputs.batchOffsetD); - args.template append("batchOffsetC", inputs.batchOffsetC); + args.template append("additionalPaddingPerBatch", additionalPaddingPerBatchGeneralBatch); + + // The HIP-compiled conversion kernel lays out these int64_t params on + // 8-byte-aligned kernarg slots. Match that alignment on the host so the + // bytes line up; a bare append() leaves them 4-byte-shifted when the + // preceding args don't end on an 8-byte boundary (e.g. the non-HAS + // variant), causing the kernel to read the neighboring offset into the + // high dword of the address and fault. + args.template appendAligned("batchOffsetD", inputs.batchOffsetD); + args.template appendAligned("batchOffsetC", inputs.batchOffsetC); } }