From 54c95bca321260ecad1e7a856e505789d199eead Mon Sep 17 00:00:00 2001 From: tomchengchitang Date: Wed, 25 Jun 2025 07:14:24 +0000 Subject: [PATCH 1/6] Increase the coverage via covering rocblaslt_auxiliary.cpp and auxiliary.hpp --- projects/hipblaslt/CMakeLists.txt | 4 +- .../hipblaslt/clients/gtest/CMakeLists.txt | 12 +- .../clients/gtest/auxiliary_gtest.cpp | 8 +- .../clients/gtest/auxiliary_gtest.yaml | 10 + .../clients/include/testing_auxiliary.hpp | 366 +++++++++++++++++- .../src/amd_detail/hipblaslt_internal.hpp | 1 - .../rocblaslt/src/include/utility.hpp | 43 -- .../library/src/include/auxiliary.hpp | 29 +- 8 files changed, 388 insertions(+), 85 deletions(-) diff --git a/projects/hipblaslt/CMakeLists.txt b/projects/hipblaslt/CMakeLists.txt index 7e9badd2260a..0d2d42807d2c 100644 --- a/projects/hipblaslt/CMakeLists.txt +++ b/projects/hipblaslt/CMakeLists.txt @@ -466,12 +466,12 @@ if(BUILD_CODE_COVERAGE) COMMAND ${LLVM_COV_EXECUTABLE} report - -object ./library/libhipblaslt-d.so.0.15 + -object ./library/libhipblaslt-d.so.${hipblaslt_SOVERSION} -instr-profile=./coverage_test.profdata --ignore-filename-regex=.*/build/debug/.* COMMAND ${LLVM_COV_EXECUTABLE} show - -object ./library/libhipblaslt-d.so.0.15 + -object ./library/libhipblaslt-d.so.${hipblaslt_SOVERSION} -instr-profile=./coverage_test.profdata -format=html -output-dir=coverage-report --ignore-filename-regex=.*/build/debug/.* diff --git a/projects/hipblaslt/clients/gtest/CMakeLists.txt b/projects/hipblaslt/clients/gtest/CMakeLists.txt index ac7df61cdd03..b5b2fe6b6c5f 100644 --- a/projects/hipblaslt/clients/gtest/CMakeLists.txt +++ b/projects/hipblaslt/clients/gtest/CMakeLists.txt @@ -29,7 +29,7 @@ find_package (Python3 REQUIRED COMPONENTS Interpreter) set(hipblaslt_test_source hipblaslt_gtest_main.cpp hipblaslt_test.cpp - matmul_gtest.cpp + # matmul_gtest.cpp auxiliary_gtest.cpp matrix_transform_gtest.cpp hipblaslt_gtest_ext_op.cpp @@ -89,6 +89,16 @@ target_compile_definitions( hipblaslt-test PRIVATE ROCM_USE_FLOAT16 HIPBLASLT_IN target_compile_options(hipblaslt-test PRIVATE $<$:${COMMON_CXX_OPTIONS}>) # target_compile_options does not go to linker like CMAKE_CXX_FLAGS does, so manually add +# We need to cover the inline functions in library when they are used in tests. +# If we don't do this, the coverage instrumentation in libhipblaslt-d.so +# does not record execution because the functions are inlined and not compiled separately. +if(BUILD_CODE_COVERAGE) + target_compile_options(hipblaslt-test PRIVATE + -fprofile-instr-generate + -fcoverage-mapping + ) +endif() + if (BUILD_CUDA) target_link_libraries( hipblaslt-test PRIVATE -llapack -lcblas ) else() diff --git a/projects/hipblaslt/clients/gtest/auxiliary_gtest.cpp b/projects/hipblaslt/clients/gtest/auxiliary_gtest.cpp index fef39d67b112..c19c62a8960c 100644 --- a/projects/hipblaslt/clients/gtest/auxiliary_gtest.cpp +++ b/projects/hipblaslt/clients/gtest/auxiliary_gtest.cpp @@ -92,6 +92,10 @@ namespace testing_aux_matmul_pref_init(arg); else if(!strcmp(arg.function, "aux_matmul_alg_null_matmul")) testing_aux_matmul_alg_null_matmul(arg); + else if(!strcmp(arg.function, "aux_mat_copy")) + testing_aux_mat_copy(arg); + else if(!strcmp(arg.function, "aux_auxiliary_func")) + testing_aux_auxiliary_func(arg); else FAIL() << "Internal error: Test called with unknown function: " << arg.function; } @@ -132,7 +136,9 @@ namespace || !strcmp(arg.function, "aux_matmul_plan_init") || !strcmp(arg.function, "aux_matmul_alg_null_matmul") || !strcmp(arg.function, "aux_matmul_pref_get_attr_bad_arg") - || !strcmp(arg.function, "aux_matmul_pref_get_attr"); + || !strcmp(arg.function, "aux_matmul_pref_get_attr") + || !strcmp(arg.function, "aux_mat_copy") + || !strcmp(arg.function, "aux_auxiliary_func"); } // Google Test name suffix based on parameters diff --git a/projects/hipblaslt/clients/gtest/auxiliary_gtest.yaml b/projects/hipblaslt/clients/gtest/auxiliary_gtest.yaml index c64574d8d6c0..8b8869d20e61 100644 --- a/projects/hipblaslt/clients/gtest/auxiliary_gtest.yaml +++ b/projects/hipblaslt/clients/gtest/auxiliary_gtest.yaml @@ -132,4 +132,14 @@ Tests: transB: N alpha: 1 beta: 0 + +- name: aux_mat_copy + category: pre_checkin + function: + - aux_mat_copy: *hpa_half_precision + +- name: aux_auxiliary_func + category: pre_checkin + function: + - aux_auxiliary_func: *hpa_half_precision ... diff --git a/projects/hipblaslt/clients/include/testing_auxiliary.hpp b/projects/hipblaslt/clients/include/testing_auxiliary.hpp index 9317e2eb3711..a40668d12ddd 100644 --- a/projects/hipblaslt/clients/include/testing_auxiliary.hpp +++ b/projects/hipblaslt/clients/include/testing_auxiliary.hpp @@ -89,30 +89,126 @@ void testing_aux_mat_set_attr_bad_arg(const Arguments& arg) const int64_t col = 128; const int64_t ld = 128; - int data; + int32_t data32; int64_t data64; + uint32_t udata32; + uint64_t udata64; hipblaslt_local_matrix_layout mat(row, col, ld, arg.a_type); EXPECT_HIPBLAS_STATUS(mat.status(), HIPBLAS_STATUS_SUCCESS); + // Test with null matrix layout + EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( + nullptr, HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT, &data32, sizeof(int32_t)), + HIPBLAS_STATUS_NOT_INITIALIZED); EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( - mat, HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT, nullptr, sizeof(int)), - HIPBLAS_STATUS_INVALID_VALUE); + mat, HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT, nullptr, sizeof(int32_t)), + HIPBLAS_STATUS_INVALID_VALUE); + - data = 1; - EXPECT_HIPBLAS_STATUS( - hipblasLtMatrixLayoutSetAttribute(mat, HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT, &data, 1), + // Test with zero size + EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( + mat, HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT, &data32, 0), + HIPBLAS_STATUS_INVALID_VALUE); + + // Test ROCBLASLT_MATRIX_LAYOUT_BATCH_COUNT with insufficient buffer size + data32 = 1; + EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( + mat, HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT, &data32, sizeof(int32_t)/2), HIPBLAS_STATUS_INVALID_VALUE); - EXPECT_HIPBLAS_STATUS( - hipblasLtMatrixLayoutSetAttribute( - mat, HIPBLASLT_MATRIX_LAYOUT_STRIDED_BATCH_OFFSET, nullptr, sizeof(int64_t)), + EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( + mat, HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT, &data32, 1), HIPBLAS_STATUS_INVALID_VALUE); + // Test ROCBLASLT_MATRIX_LAYOUT_STRIDED_BATCH_OFFSET with insufficient buffer size data64 = ld * col; EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( - mat, HIPBLASLT_MATRIX_LAYOUT_STRIDED_BATCH_OFFSET, &data64, 1), - HIPBLAS_STATUS_INVALID_VALUE); + mat, HIPBLASLT_MATRIX_LAYOUT_STRIDED_BATCH_OFFSET, &data64, sizeof(int64_t)/2), + HIPBLAS_STATUS_INVALID_VALUE); + + EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( + mat, HIPBLASLT_MATRIX_LAYOUT_STRIDED_BATCH_OFFSET, &data64, 1), + HIPBLAS_STATUS_INVALID_VALUE); + + EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( + mat, HIPBLASLT_MATRIX_LAYOUT_STRIDED_BATCH_OFFSET, nullptr, sizeof(int64_t)), + HIPBLAS_STATUS_INVALID_VALUE); + + // Test ROCBLASLT_MATRIX_LAYOUT_TYPE with insufficient buffer size + udata32 = static_cast(arg.a_type); + EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( + mat, HIPBLASLT_MATRIX_LAYOUT_TYPE, &udata32, sizeof(uint32_t)/2), + HIPBLAS_STATUS_INVALID_VALUE); + + EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( + mat, HIPBLASLT_MATRIX_LAYOUT_TYPE, &udata32, 1), + HIPBLAS_STATUS_INVALID_VALUE); + + EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( + mat, HIPBLASLT_MATRIX_LAYOUT_TYPE, nullptr, sizeof(uint32_t)), + HIPBLAS_STATUS_INVALID_VALUE); + + // Test ROCBLASLT_MATRIX_LAYOUT_ORDER with insufficient buffer size + data32 = HIPBLASLT_ORDER_COL; + EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( + mat, HIPBLASLT_MATRIX_LAYOUT_ORDER, &data32, sizeof(int32_t)/2), + HIPBLAS_STATUS_INVALID_VALUE); + + EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( + mat, HIPBLASLT_MATRIX_LAYOUT_ORDER, &data32, 1), + HIPBLAS_STATUS_INVALID_VALUE); + + EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( + mat, HIPBLASLT_MATRIX_LAYOUT_ORDER, nullptr, sizeof(int32_t)), + HIPBLAS_STATUS_INVALID_VALUE); + + // Test ROCBLASLT_MATRIX_LAYOUT_ROWS with insufficient buffer size + udata64 = static_cast(row); + EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( + mat, HIPBLASLT_MATRIX_LAYOUT_ROWS, &udata64, sizeof(uint64_t)/2), + HIPBLAS_STATUS_INVALID_VALUE); + + EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( + mat, HIPBLASLT_MATRIX_LAYOUT_ROWS, &udata64, 1), + HIPBLAS_STATUS_INVALID_VALUE); + + EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( + mat, HIPBLASLT_MATRIX_LAYOUT_ROWS, nullptr, sizeof(uint64_t)), + HIPBLAS_STATUS_INVALID_VALUE); + + // Test ROCBLASLT_MATRIX_LAYOUT_COLS with insufficient buffer size + udata64 = static_cast(col); + EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( + mat, HIPBLASLT_MATRIX_LAYOUT_COLS, &udata64, sizeof(uint64_t)/2), + HIPBLAS_STATUS_INVALID_VALUE); + + EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( + mat, HIPBLASLT_MATRIX_LAYOUT_COLS, &udata64, 1), + HIPBLAS_STATUS_INVALID_VALUE); + + EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( + mat, HIPBLASLT_MATRIX_LAYOUT_COLS, nullptr, sizeof(uint64_t)), + HIPBLAS_STATUS_INVALID_VALUE); + + // Test ROCBLASLT_MATRIX_LAYOUT_LD with insufficient buffer size + data64 = ld; + EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( + mat, HIPBLASLT_MATRIX_LAYOUT_LD, &data64, sizeof(int64_t)/2), + HIPBLAS_STATUS_INVALID_VALUE); + + EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( + mat, HIPBLASLT_MATRIX_LAYOUT_LD, &data64, 1), + HIPBLAS_STATUS_INVALID_VALUE); + + EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( + mat, HIPBLASLT_MATRIX_LAYOUT_LD, nullptr, sizeof(int64_t)), + HIPBLAS_STATUS_INVALID_VALUE); + + // Test with invalid attribute (assuming there's a MAX value for bounds checking) + EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( + mat, static_cast(999), &data32, sizeof(int32_t)), + HIPBLAS_STATUS_INVALID_VALUE); } void testing_aux_mat_get_attr_bad_arg(const Arguments& arg) @@ -137,6 +233,12 @@ void testing_aux_mat_get_attr_bad_arg(const Arguments& arg) hipblasLtMatrixLayoutGetAttribute( mat, HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT, nullptr, sizeof(int), &sizeWritten), HIPBLAS_STATUS_INVALID_VALUE); + + EXPECT_HIPBLAS_STATUS( + hipblasLtMatrixLayoutGetAttribute( + mat, HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT, nullptr, 0, nullptr), + HIPBLAS_STATUS_INVALID_VALUE); + EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutGetAttribute( mat, HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT, &data, 1, &sizeWritten), HIPBLAS_STATUS_INVALID_VALUE); @@ -357,6 +459,10 @@ void testing_aux_matmul_set_get_attr(const Arguments& arg) void* d_bias_r; CHECK_HIP_ERROR(hipMalloc(&d_bias, k * sizeof(hipblasLtHalf))); + EXPECT_HIPBLAS_STATUS(hipblasLtMatmulDescSetAttribute( + matmul, HIPBLASLT_MATMUL_DESC_BIAS_POINTER, &d_bias, sizeof(void*)/2), + HIPBLAS_STATUS_INVALID_VALUE); + EXPECT_HIPBLAS_STATUS(hipblasLtMatmulDescSetAttribute( matmul, HIPBLASLT_MATMUL_DESC_BIAS_POINTER, &d_bias, sizeof(void*)), HIPBLAS_STATUS_SUCCESS); @@ -367,8 +473,8 @@ void testing_aux_matmul_set_get_attr(const Arguments& arg) ASSERT_TRUE(d_bias_r == d_bias); - // for HIPBLASLT_MATMUL_DESC_A_SCALE_POINTER - // We create a new matmul_descr becasue we don't dont want to set matmulDesc->scaleAType as Scalar in advance. + // for ROCBLASLT_MATMUL_DESC_A_SCALE_POINTER & ROCBLASLT_MATMUL_DESC_A_SCALE_POINTER_VEC_EXT + // We create a new matmul_descr becasue we don't want to set matmulDesc->scaleAType as Scalar in advance. // If we set it as Scalar, it will skip some cases, which is not what we desire for. const hipblasOperation_t opA = HIPBLAS_OP_T; const hipblasOperation_t opB = HIPBLAS_OP_N; @@ -381,6 +487,10 @@ void testing_aux_matmul_set_get_attr(const Arguments& arg) float* d_scale_a_r; CHECK_HIP_ERROR(hipMalloc(&d_scale_a, sizeof(float))); CHECK_HIP_ERROR(hipMemcpy(d_scale_a, &h_scale_a_for_desc_a_scale_pointer, sizeof(float), hipMemcpyHostToDevice)); + // EXPECT_HIPBLAS_STATUS(hipblasLtMatmulDescSetAttribute( + // matmul_descr, HIPBLASLT_MATMUL_MATRIX_SCALE_OUTER_VEC_32F, &d_scale_a, sizeof(void*)), + // HIPBLAS_STATUS_SUCCESS); // this is necessary to cover in set + EXPECT_HIPBLAS_STATUS(hipblasLtMatmulDescSetAttribute( matmul_descr, HIPBLASLT_MATMUL_DESC_A_SCALE_POINTER, &d_scale_a, sizeof(float*)/2), HIPBLAS_STATUS_INVALID_VALUE); @@ -525,13 +635,21 @@ void testing_aux_matmul_set_get_attr(const Arguments& arg) hipStream_t stream; CHECK_HIP_ERROR(hipStreamCreate(&stream)); - // for HIPBLASLT_MATMUL_DESC_B_SCALE_POINTER + // for HIPBLASLT_MATMUL_DESC_B_SCALE_POINTER_VEC_EXT & ROCBLASLT_MATMUL_DESC_B_SCALE_POINTER float h_scale_b = 3.f; float* d_scale_b; float* d_scale_b_r; CHECK_HIP_ERROR(hipMalloc(&d_scale_b, sizeof(float))); CHECK_HIP_ERROR(hipMemcpyAsync(d_scale_b, &h_scale_b, sizeof(float), hipMemcpyHostToDevice, stream)); + // EXPECT_HIPBLAS_STATUS(hipblasLtMatmulDescSetAttribute( + // matmul, HIPBLASLT_MATMUL_MATRIX_SCALE_OUTER_VEC_32F, &d_scale_b, sizeof(float*)), + // HIPBLAS_STATUS_SUCCESS); + + // EXPECT_HIPBLAS_STATUS(hipblasLtMatmulDescSetAttribute( + // matmul, HIPBLASLT_MATMUL_MATRIX_SCALE_OUTER_VEC_32F, &d_scale_b, sizeof(float*)/2), + // HIPBLAS_STATUS_INVALID_VALUE); + EXPECT_HIPBLAS_STATUS(hipblasLtMatmulDescSetAttribute( matmul, HIPBLASLT_MATMUL_DESC_B_SCALE_POINTER, &d_scale_b, sizeof(float*)/2), HIPBLAS_STATUS_INVALID_VALUE); @@ -820,6 +938,28 @@ void testing_aux_matmul_pref_get_attr(const Arguments& arg) uint64_t data_get = 0; size_t sizeWritten; + // Test edge cases + EXPECT_HIPBLAS_STATUS( + hipblasLtMatmulPreferenceSetAttribute(nullptr, + HIPBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES, + &data_set, + sizeof(data_set)), + HIPBLAS_STATUS_INVALID_VALUE); + + EXPECT_HIPBLAS_STATUS( + hipblasLtMatmulPreferenceSetAttribute(pref, + HIPBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES, + &data_set, + 0), + HIPBLAS_STATUS_INVALID_VALUE); + + + EXPECT_HIPBLAS_STATUS(hipblasLtMatmulPreferenceGetAttribute(pref, + HIPBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES, + nullptr, + 0, + &sizeWritten), + HIPBLAS_STATUS_INVALID_VALUE); // Set the attribute EXPECT_HIPBLAS_STATUS( hipblasLtMatmulPreferenceSetAttribute( @@ -840,13 +980,35 @@ void testing_aux_matmul_pref_get_attr(const Arguments& arg) ASSERT_TRUE(sizeWritten == sizeof(data_get)); // Test getting other attributes (assuming they have default values) - int32_t search_mode; + int32_t search_mode = 0; + int32_t search_mode_r = 1; + EXPECT_HIPBLAS_STATUS(hipblasLtMatmulPreferenceSetAttribute(pref, + HIPBLASLT_MATMUL_PREF_SEARCH_MODE, + &search_mode, + sizeof(search_mode)), + HIPBLAS_STATUS_SUCCESS); + EXPECT_HIPBLAS_STATUS(hipblasLtMatmulPreferenceGetAttribute(pref, HIPBLASLT_MATMUL_PREF_SEARCH_MODE, + &search_mode_r, + sizeof(search_mode_r), + &sizeWritten), + HIPBLAS_STATUS_SUCCESS); + ASSERT_TRUE(data_get == data_set); + + // Test default value + EXPECT_HIPBLAS_STATUS(hipblasLtMatmulPreferenceSetAttribute(pref, + HIPBLASLT_MATMUL_PREF_MAX, &search_mode, - sizeof(search_mode), + sizeof(search_mode)), + HIPBLAS_STATUS_INVALID_VALUE); + + EXPECT_HIPBLAS_STATUS(hipblasLtMatmulPreferenceGetAttribute(pref, + HIPBLASLT_MATMUL_PREF_MAX, + &search_mode_r, + sizeof(search_mode_r), &sizeWritten), - HIPBLAS_STATUS_SUCCESS); + HIPBLAS_STATUS_INVALID_VALUE); // You might want to add more attribute tests here } @@ -1219,3 +1381,173 @@ void testing_aux_matmul_pref_init(const Arguments& arg) hipblaslt_local_preference pref; EXPECT_HIPBLAS_STATUS(pref.status(), HIPBLAS_STATUS_SUCCESS); } + +void testing_aux_mat_copy(const Arguments& arg) +{ + hipblasLtMatmulDesc_t matmul_src; + CHECK_HIPBLASLT_ERROR(hipblasLtMatmulDescCreate(&matmul_src, arg.compute_type, arg.scale_type)); + CHECK_HIPBLASLT_ERROR(hipblasLtMatmulDescCreate(&matmul_src, arg.compute_type, arg.scale_type)); + hipblasLtMatmulDesc_t matmul_dest; + CHECK_HIPBLASLT_ERROR(hipblasLtMatmulDescCreate(&matmul_dest, HIPBLAS_COMPUTE_32F, HIP_R_32F)); + + EXPECT_HIPBLAS_STATUS(hipblaslt_ext::copyMatmul(nullptr, matmul_dest), HIPBLAS_STATUS_INVALID_VALUE); + EXPECT_HIPBLAS_STATUS(hipblaslt_ext::copyMatmul(matmul_src, nullptr), HIPBLAS_STATUS_INVALID_VALUE); + EXPECT_HIPBLAS_STATUS(hipblaslt_ext::copyMatmul(matmul_src, matmul_dest), HIPBLAS_STATUS_SUCCESS); +} + +void testing_aux_auxiliary_func(const Arguments& arg) +{ + // Test gpu_arch_match + int deviceId; + hipDeviceProp_t deviceProperties; + static_cast(hipGetDevice(&deviceId)); + static_cast(hipGetDeviceProperties(&deviceProperties, deviceId)); + ASSERT_TRUE(gpu_arch_match(deviceProperties.gcnArchName, "")); + ASSERT_TRUE(gpu_arch_match(deviceProperties.gcnArchName, "\\d")); + + // Test hipblas_status_to_string + ASSERT_TRUE(hipblas_status_to_string(HIPBLAS_STATUS_SUCCESS) == "HIPBLAS_STATUS_SUCCESS"); + ASSERT_TRUE(hipblas_status_to_string(HIPBLAS_STATUS_NOT_INITIALIZED) == "HIPBLAS_STATUS_NOT_INITIALIZED"); + ASSERT_TRUE(hipblas_status_to_string(HIPBLAS_STATUS_ALLOC_FAILED) == "HIPBLAS_STATUS_ALLOC_FAILED"); + ASSERT_TRUE(hipblas_status_to_string(HIPBLAS_STATUS_INVALID_VALUE) == "HIPBLAS_STATUS_INVALID_VALUE"); + ASSERT_TRUE(hipblas_status_to_string(HIPBLAS_STATUS_MAPPING_ERROR) == "HIPBLAS_STATUS_MAPPING_ERROR"); + ASSERT_TRUE(hipblas_status_to_string(HIPBLAS_STATUS_EXECUTION_FAILED) == "HIPBLAS_STATUS_EXECUTION_FAILED"); + ASSERT_TRUE(hipblas_status_to_string(HIPBLAS_STATUS_INTERNAL_ERROR) == "HIPBLAS_STATUS_INTERNAL_ERROR"); + ASSERT_TRUE(hipblas_status_to_string(HIPBLAS_STATUS_NOT_SUPPORTED) == "HIPBLAS_STATUS_NOT_SUPPORTED"); + ASSERT_TRUE(hipblas_status_to_string(HIPBLAS_STATUS_ARCH_MISMATCH) == "HIPBLAS_STATUS_ARCH_MISMATCH"); + ASSERT_TRUE(hipblas_status_to_string(HIPBLAS_STATUS_INVALID_ENUM) == "HIPBLAS_STATUS_INVALID_ENUM"); + ASSERT_TRUE(hipblas_status_to_string(HIPBLAS_STATUS_UNKNOWN) == "HIPBLAS_STATUS_UNKNOWN"); + ASSERT_TRUE(hipblas_status_to_string(static_cast(12)) == ""); + + // Test hipblas_operation_to_string + + ASSERT_TRUE(hipblas_operation_to_string(HIPBLAS_OP_N) == "N"); + ASSERT_TRUE(hipblas_operation_to_string(HIPBLAS_OP_T) == "T"); + ASSERT_TRUE(hipblas_operation_to_string(HIPBLAS_OP_C) == "C"); + ASSERT_TRUE(hipblas_operation_to_string(static_cast(114)) == "invalid"); + + // Test char_to_hipblas_operation + ASSERT_TRUE(char_to_hipblas_operation('N') == HIPBLAS_OP_N); + ASSERT_TRUE(char_to_hipblas_operation('n') == HIPBLAS_OP_N); + ASSERT_TRUE(char_to_hipblas_operation('T') == HIPBLAS_OP_T); + ASSERT_TRUE(char_to_hipblas_operation('t') == HIPBLAS_OP_T); + ASSERT_TRUE(char_to_hipblas_operation('C') == HIPBLAS_OP_C); + ASSERT_TRUE(char_to_hipblas_operation('c') == HIPBLAS_OP_C); + ASSERT_TRUE(char_to_hipblas_operation('X') == HIPBLASLT_OPERATION_INVALID); + + + // Test char_to_hipblas_operation + ASSERT_TRUE(char_to_hipblas_operation('N') == HIPBLAS_OP_N); + ASSERT_TRUE(char_to_hipblas_operation('n') == HIPBLAS_OP_N); + ASSERT_TRUE(char_to_hipblas_operation('T') == HIPBLAS_OP_T); + ASSERT_TRUE(char_to_hipblas_operation('t') == HIPBLAS_OP_T); + ASSERT_TRUE(char_to_hipblas_operation('C') == HIPBLAS_OP_C); + ASSERT_TRUE(char_to_hipblas_operation('c') == HIPBLAS_OP_C); + + // Test hip_datatype_to_string + ASSERT_TRUE(hip_datatype_to_string(HIP_R_32F) == "f32_r"); + ASSERT_TRUE(hip_datatype_to_string(HIP_R_64F) == "f64_r"); + ASSERT_TRUE(hip_datatype_to_string(HIP_R_16F) == "f16_r"); + ASSERT_TRUE(hip_datatype_to_string(HIP_R_16BF) == "bf16_r"); + ASSERT_TRUE(hip_datatype_to_string(HIP_R_8I) == "i8_r"); + ASSERT_TRUE(hip_datatype_to_string(HIP_R_32I) == "i32_r"); + ASSERT_TRUE(hip_datatype_to_string(HIP_R_8F_E4M3_FNUZ) == "f8_fnuz_r"); + ASSERT_TRUE(hip_datatype_to_string(HIP_R_8F_E5M2_FNUZ) == "bf8_fnuz_r"); + ASSERT_TRUE(hip_datatype_to_string(HIP_R_8F_E4M3) == "f8_r"); + ASSERT_TRUE(hip_datatype_to_string(HIP_R_8F_E5M2) == "bf8_r"); + ASSERT_TRUE(hip_datatype_to_string(static_cast(HIP_R_6F_E2M3_EXT)) == "f6_r"); + ASSERT_TRUE(hip_datatype_to_string(static_cast(HIP_R_6F_E3M2_EXT)) == "bf6_r"); + ASSERT_TRUE(hip_datatype_to_string(static_cast(HIP_R_4F_E2M1_EXT)) == "f4_r"); + + + // Test hipblas_computetype_to_string + hipblas_computetype_to_string(HIPBLAS_COMPUTE_16F); + ASSERT_TRUE(hipblas_computetype_to_string(HIPBLAS_COMPUTE_16F) == "f16_r"); + ASSERT_TRUE(hipblas_computetype_to_string(HIPBLAS_COMPUTE_32F) == "f32_r"); + ASSERT_TRUE(hipblas_computetype_to_string(HIPBLAS_COMPUTE_32F_FAST_TF32) == "xf32_r"); + ASSERT_TRUE(hipblas_computetype_to_string(HIPBLAS_COMPUTE_64F) == "f64_r"); + ASSERT_TRUE(hipblas_computetype_to_string(HIPBLAS_COMPUTE_32I) == "i32_r"); + ASSERT_TRUE(hipblas_computetype_to_string(HIPBLAS_COMPUTE_32F_FAST_16F) == "f32_f16_r"); + ASSERT_TRUE(hipblas_computetype_to_string(HIPBLAS_COMPUTE_32F_FAST_16BF) == "f32_bf16_r"); + ASSERT_TRUE(hipblas_computetype_to_string(HIPBLAS_COMPUTE_64F_PEDANTIC) == "non-supported compute type"); + + // Test string_to_hip_datatype + ASSERT_TRUE(string_to_hip_datatype("f8_fnuz_r") == HIP_R_8F_E4M3_FNUZ); + ASSERT_TRUE(string_to_hip_datatype("bf8_fnuz_r") == HIP_R_8F_E5M2_FNUZ); + ASSERT_TRUE(string_to_hip_datatype("f8_r") == HIP_R_8F_E4M3); + ASSERT_TRUE(string_to_hip_datatype("bf8_r") == HIP_R_8F_E5M2); + ASSERT_TRUE(string_to_hip_datatype("f32_r") == HIP_R_32F); + ASSERT_TRUE(string_to_hip_datatype("f64_r") == HIP_R_64F); + ASSERT_TRUE(string_to_hip_datatype("f16_r") == HIP_R_16F); + ASSERT_TRUE(string_to_hip_datatype("bf16_r") == HIP_R_16BF); + ASSERT_TRUE(string_to_hip_datatype("i8_r") == HIP_R_8I ); + ASSERT_TRUE(string_to_hip_datatype("f6_r") == static_cast(HIP_R_6F_E2M3_EXT)); + ASSERT_TRUE(string_to_hip_datatype("bf6_r") == static_cast(HIP_R_6F_E3M2_EXT)); + ASSERT_TRUE(string_to_hip_datatype("f4_r") == static_cast(HIP_R_4F_E2M1_EXT)); + ASSERT_TRUE(string_to_hip_datatype("i32_r") == HIP_R_32I); + ASSERT_TRUE(string_to_hip_datatype("") == HIPBLASLT_DATATYPE_INVALID); + + // Test string_to_hip_datatype_assert + string_to_hip_datatype_assert("f8_fnuz_r"); + + // Test string_to_hipblas_computetype + ASSERT_TRUE(string_to_hipblas_computetype("f32_r") == HIPBLAS_COMPUTE_32F); + ASSERT_TRUE(string_to_hipblas_computetype("xf32_r") == HIPBLAS_COMPUTE_32F_FAST_TF32); + ASSERT_TRUE(string_to_hipblas_computetype("f64_r") == HIPBLAS_COMPUTE_64F); + ASSERT_TRUE(string_to_hipblas_computetype("i32_r") == HIPBLAS_COMPUTE_32I); + ASSERT_TRUE(string_to_hipblas_computetype("f32_f16_r") == HIPBLAS_COMPUTE_32F_FAST_16F); + ASSERT_TRUE(string_to_hipblas_computetype("f32_bf16_r") == HIPBLAS_COMPUTE_32F_FAST_16BF); + ASSERT_TRUE(string_to_hipblas_computetype("") == HIPBLASLT_COMPUTE_TYPE_INVALID); + + // Test string_to_hipblas_computetype_assert + string_to_hipblas_computetype_assert("f32_r"); + + // Test string_to_epilogue_type + ASSERT_TRUE(string_to_epilogue_type("HIPBLASLT_EPILOGUE_RELU") == HIPBLASLT_EPILOGUE_RELU ); + ASSERT_TRUE(string_to_epilogue_type("HIPBLASLT_EPILOGUE_BIAS") == HIPBLASLT_EPILOGUE_BIAS); + ASSERT_TRUE(string_to_epilogue_type("HIPBLASLT_EPILOGUE_RELU_BIAS") == HIPBLASLT_EPILOGUE_RELU_BIAS); + ASSERT_TRUE(string_to_epilogue_type("HIPBLASLT_EPILOGUE_GELU") == HIPBLASLT_EPILOGUE_GELU); + ASSERT_TRUE(string_to_epilogue_type("HIPBLASLT_EPILOGUE_GELU_BIAS") == HIPBLASLT_EPILOGUE_GELU_BIAS); + ASSERT_TRUE(string_to_epilogue_type("HIPBLASLT_EPILOGUE_GELU_AUX") == HIPBLASLT_EPILOGUE_GELU_AUX); + ASSERT_TRUE(string_to_epilogue_type("HIPBLASLT_EPILOGUE_GELU_AUX_BIAS") == HIPBLASLT_EPILOGUE_GELU_AUX_BIAS); + ASSERT_TRUE(string_to_epilogue_type("HIPBLASLT_EPILOGUE_DGELU") == HIPBLASLT_EPILOGUE_DGELU); + ASSERT_TRUE(string_to_epilogue_type("HIPBLASLT_EPILOGUE_DGELU_BGRAD") == HIPBLASLT_EPILOGUE_DGELU_BGRAD); + ASSERT_TRUE(string_to_epilogue_type("HIPBLASLT_EPILOGUE_BGRADA") == HIPBLASLT_EPILOGUE_BGRADA); + ASSERT_TRUE(string_to_epilogue_type("HIPBLASLT_EPILOGUE_BGRADB") == HIPBLASLT_EPILOGUE_BGRADB); + ASSERT_TRUE(string_to_epilogue_type("HIPBLASLT_EPILOGUE_SWISH_EXT") == HIPBLASLT_EPILOGUE_SWISH_EXT); + ASSERT_TRUE(string_to_epilogue_type("HIPBLASLT_EPILOGUE_SWISH_BIAS_EXT") == HIPBLASLT_EPILOGUE_SWISH_BIAS_EXT); + ASSERT_TRUE(string_to_epilogue_type("HIPBLASLT_EPILOGUE_DEFAULT") == HIPBLASLT_EPILOGUE_DEFAULT); + ASSERT_TRUE(string_to_epilogue_type("test") == static_cast(0)); + + // Test string_to_epilogue_type_assert + ASSERT_TRUE(string_to_epilogue_type_assert("HIPBLASLT_EPILOGUE_RELU") == HIPBLASLT_EPILOGUE_RELU); + + // Test hipblaslt_isnan + uint32_t other_type_value = 5; + ASSERT_FALSE(hipblaslt_isnan(other_type_value)); + + float non_integral_nan = std::numeric_limits::quiet_NaN(); + ASSERT_TRUE(hipblaslt_isnan(non_integral_nan)); + + hipblaslt_f8_fnuz arg_hipblaslt_f8_fnuz; + arg_hipblaslt_f8_fnuz.__x = 0x80; + ASSERT_TRUE(hipblaslt_isnan(arg_hipblaslt_f8_fnuz)); + + hipblaslt_bf8_fnuz arg_hipblaslt_bf8_fnuz; + arg_hipblaslt_bf8_fnuz.__x = 0x80; + ASSERT_TRUE(hipblaslt_isnan(arg_hipblaslt_bf8_fnuz)); + + hipblaslt_f8 arg_hipblaslt_f8; + arg_hipblaslt_f8.__x = 0x80; + ASSERT_TRUE(hipblaslt_isnan(arg_hipblaslt_f8)); + + hipblaslt_bf8 arg_hipblaslt_bf8; + arg_hipblaslt_bf8.__x = 0x7d; + ASSERT_TRUE(hipblaslt_isnan(arg_hipblaslt_bf8)); + + hipblasLtHalf pos_inf = static_cast(INFINITY); + ASSERT_TRUE(hipblaslt_isinf(pos_inf)); + + hipblasLtHalf nan = static_cast(NAN); + ASSERT_TRUE(hipblaslt_isnan(nan)); +} \ No newline at end of file diff --git a/projects/hipblaslt/library/src/amd_detail/hipblaslt_internal.hpp b/projects/hipblaslt/library/src/amd_detail/hipblaslt_internal.hpp index 4570d5b91105..fbf60e73f72c 100644 --- a/projects/hipblaslt/library/src/amd_detail/hipblaslt_internal.hpp +++ b/projects/hipblaslt/library/src/amd_detail/hipblaslt_internal.hpp @@ -33,6 +33,5 @@ #endif #include #include -hipblasStatus_t hipErrorToHIPBLASStatus(hipError_t status); hipblasStatus_t RocBlasLtStatusToHIPStatus(rocblaslt_status_ status); diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/utility.hpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/utility.hpp index 6e8ccc070f74..b2f551332d5c 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/utility.hpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/utility.hpp @@ -105,49 +105,6 @@ constexpr const char* rocblaslt_compute_type_string(rocblaslt_compute_type type) } } -constexpr const char* rocblaslt_transpose_letter(hipblasOperation_t op) -{ - switch(op) - { - case HIPBLAS_OP_N: - return "N"; - case HIPBLAS_OP_T: - return "T"; - case HIPBLAS_OP_C: - return "C"; - default: - return "invalidTranspose"; - } -} -// Convert rocblaslt_status to string -constexpr const char* rocblaslt_status_to_string(rocblaslt_status status) -{ -#define CASE(x) \ - case x: \ - return #x - switch(status) - { - CASE(rocblaslt_status_success); - CASE(rocblaslt_status_invalid_handle); - CASE(rocblaslt_status_not_implemented); - CASE(rocblaslt_status_invalid_pointer); - CASE(rocblaslt_status_invalid_size); - CASE(rocblaslt_status_memory_error); - CASE(rocblaslt_status_internal_error); - CASE(rocblaslt_status_invalid_value); - CASE(rocblaslt_status_arch_mismatch); - CASE(rocblaslt_status_zero_pivot); - CASE(rocblaslt_status_not_initialized); - CASE(rocblaslt_status_type_mismatch); - CASE(rocblaslt_status_requires_sorted_storage); - CASE(rocblaslt_status_continue); - } -#undef CASE - // We don't use default: so that the compiler warns us if any valid enums are - // missing from our switch. If the value is not a valid rocblaslt_status, we - // return this string. - return ""; -} template static constexpr char rocblaslt_precision_string[] = "invalid"; template <> diff --git a/projects/hipblaslt/library/src/include/auxiliary.hpp b/projects/hipblaslt/library/src/include/auxiliary.hpp index 827815b2243a..be08ba9047ad 100644 --- a/projects/hipblaslt/library/src/include/auxiliary.hpp +++ b/projects/hipblaslt/library/src/include/auxiliary.hpp @@ -151,26 +151,15 @@ constexpr const char* hip_datatype_to_string(hipDataType type) HIPBLASLT_EXPORT constexpr const char* hipblas_computetype_to_string(hipblasComputeType_t type) { - switch(type) - { - case HIPBLAS_COMPUTE_16F: - return "f16_r"; - case HIPBLAS_COMPUTE_32F: - return "f32_r"; - case HIPBLAS_COMPUTE_32F_FAST_TF32: - return "xf32_r"; - case HIPBLAS_COMPUTE_64F: - return "f64_r"; - case HIPBLAS_COMPUTE_32I: - return "i32_r"; - case HIPBLAS_COMPUTE_32F_FAST_16F: - return "f32_f16_r"; - case HIPBLAS_COMPUTE_32F_FAST_16BF: - return "f32_bf16_r"; - default: - return "non-supported compute type"; - } - return "invalid"; + return + type == HIPBLAS_COMPUTE_16F ? "f16_r" : + type == HIPBLAS_COMPUTE_32F ? "f32_r" : + type == HIPBLAS_COMPUTE_32F_FAST_TF32 ? "xf32_r" : + type == HIPBLAS_COMPUTE_64F ? "f64_r" : + type == HIPBLAS_COMPUTE_32I ? "i32_r" : + type == HIPBLAS_COMPUTE_32F_FAST_16F ? "f32_f16_r" : + type == HIPBLAS_COMPUTE_32F_FAST_16BF ? "f32_bf16_r" : + "non-supported compute type"; } // clang-format off From 1d890e5a2e3b339dbff1e2a345f6f5a939ce8f90 Mon Sep 17 00:00:00 2001 From: tomchengchitang Date: Wed, 25 Jun 2025 07:55:58 +0000 Subject: [PATCH 2/6] Fix typo --- projects/hipblaslt/clients/gtest/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hipblaslt/clients/gtest/CMakeLists.txt b/projects/hipblaslt/clients/gtest/CMakeLists.txt index b5b2fe6b6c5f..c12172917643 100644 --- a/projects/hipblaslt/clients/gtest/CMakeLists.txt +++ b/projects/hipblaslt/clients/gtest/CMakeLists.txt @@ -29,7 +29,7 @@ find_package (Python3 REQUIRED COMPONENTS Interpreter) set(hipblaslt_test_source hipblaslt_gtest_main.cpp hipblaslt_test.cpp - # matmul_gtest.cpp + matmul_gtest.cpp auxiliary_gtest.cpp matrix_transform_gtest.cpp hipblaslt_gtest_ext_op.cpp From 70c8dd29d3c63054b8974ef1fa32099961ff0e70 Mon Sep 17 00:00:00 2001 From: tomchengchitang Date: Fri, 27 Jun 2025 09:22:04 +0000 Subject: [PATCH 3/6] Delete ROCBLASLT_MATMUL_DESC_A_SCALE_POINTER_VEC_EXT and unused code --- .../clients/include/testing_auxiliary.hpp | 15 ++------------- 1 file changed, 2 insertions(+), 13 deletions(-) diff --git a/projects/hipblaslt/clients/include/testing_auxiliary.hpp b/projects/hipblaslt/clients/include/testing_auxiliary.hpp index a40668d12ddd..fef3510fbbe1 100644 --- a/projects/hipblaslt/clients/include/testing_auxiliary.hpp +++ b/projects/hipblaslt/clients/include/testing_auxiliary.hpp @@ -473,7 +473,7 @@ void testing_aux_matmul_set_get_attr(const Arguments& arg) ASSERT_TRUE(d_bias_r == d_bias); - // for ROCBLASLT_MATMUL_DESC_A_SCALE_POINTER & ROCBLASLT_MATMUL_DESC_A_SCALE_POINTER_VEC_EXT + // for ROCBLASLT_MATMUL_DESC_A_SCALE_POINTER // We create a new matmul_descr becasue we don't want to set matmulDesc->scaleAType as Scalar in advance. // If we set it as Scalar, it will skip some cases, which is not what we desire for. const hipblasOperation_t opA = HIPBLAS_OP_T; @@ -487,9 +487,6 @@ void testing_aux_matmul_set_get_attr(const Arguments& arg) float* d_scale_a_r; CHECK_HIP_ERROR(hipMalloc(&d_scale_a, sizeof(float))); CHECK_HIP_ERROR(hipMemcpy(d_scale_a, &h_scale_a_for_desc_a_scale_pointer, sizeof(float), hipMemcpyHostToDevice)); - // EXPECT_HIPBLAS_STATUS(hipblasLtMatmulDescSetAttribute( - // matmul_descr, HIPBLASLT_MATMUL_MATRIX_SCALE_OUTER_VEC_32F, &d_scale_a, sizeof(void*)), - // HIPBLAS_STATUS_SUCCESS); // this is necessary to cover in set EXPECT_HIPBLAS_STATUS(hipblasLtMatmulDescSetAttribute( matmul_descr, HIPBLASLT_MATMUL_DESC_A_SCALE_POINTER, &d_scale_a, sizeof(float*)/2), @@ -635,21 +632,13 @@ void testing_aux_matmul_set_get_attr(const Arguments& arg) hipStream_t stream; CHECK_HIP_ERROR(hipStreamCreate(&stream)); - // for HIPBLASLT_MATMUL_DESC_B_SCALE_POINTER_VEC_EXT & ROCBLASLT_MATMUL_DESC_B_SCALE_POINTER + // for ROCBLASLT_MATMUL_DESC_B_SCALE_POINTER float h_scale_b = 3.f; float* d_scale_b; float* d_scale_b_r; CHECK_HIP_ERROR(hipMalloc(&d_scale_b, sizeof(float))); CHECK_HIP_ERROR(hipMemcpyAsync(d_scale_b, &h_scale_b, sizeof(float), hipMemcpyHostToDevice, stream)); - // EXPECT_HIPBLAS_STATUS(hipblasLtMatmulDescSetAttribute( - // matmul, HIPBLASLT_MATMUL_MATRIX_SCALE_OUTER_VEC_32F, &d_scale_b, sizeof(float*)), - // HIPBLAS_STATUS_SUCCESS); - - // EXPECT_HIPBLAS_STATUS(hipblasLtMatmulDescSetAttribute( - // matmul, HIPBLASLT_MATMUL_MATRIX_SCALE_OUTER_VEC_32F, &d_scale_b, sizeof(float*)/2), - // HIPBLAS_STATUS_INVALID_VALUE); - EXPECT_HIPBLAS_STATUS(hipblasLtMatmulDescSetAttribute( matmul, HIPBLASLT_MATMUL_DESC_B_SCALE_POINTER, &d_scale_b, sizeof(float*)/2), HIPBLAS_STATUS_INVALID_VALUE); From 1e91b676f16c4f423cfe62a3f6b8614a9e2c656b Mon Sep 17 00:00:00 2001 From: tomchengchitang Date: Fri, 27 Jun 2025 09:32:42 +0000 Subject: [PATCH 4/6] Merge develop From 9efe56af29a934d716eae8f1d6eba1dede3a5b80 Mon Sep 17 00:00:00 2001 From: tomchengchitang Date: Tue, 1 Jul 2025 03:16:47 +0000 Subject: [PATCH 5/6] Destroy handles and change some comments to HIPBLASLT --- .../clients/include/testing_auxiliary.hpp | 55 ++++++++++--------- 1 file changed, 29 insertions(+), 26 deletions(-) diff --git a/projects/hipblaslt/clients/include/testing_auxiliary.hpp b/projects/hipblaslt/clients/include/testing_auxiliary.hpp index fef3510fbbe1..25b7be341b03 100644 --- a/projects/hipblaslt/clients/include/testing_auxiliary.hpp +++ b/projects/hipblaslt/clients/include/testing_auxiliary.hpp @@ -111,7 +111,7 @@ void testing_aux_mat_set_attr_bad_arg(const Arguments& arg) mat, HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT, &data32, 0), HIPBLAS_STATUS_INVALID_VALUE); - // Test ROCBLASLT_MATRIX_LAYOUT_BATCH_COUNT with insufficient buffer size + // Test HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT with insufficient buffer size data32 = 1; EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( mat, HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT, &data32, sizeof(int32_t)/2), @@ -121,7 +121,7 @@ void testing_aux_mat_set_attr_bad_arg(const Arguments& arg) mat, HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT, &data32, 1), HIPBLAS_STATUS_INVALID_VALUE); - // Test ROCBLASLT_MATRIX_LAYOUT_STRIDED_BATCH_OFFSET with insufficient buffer size + // Test HIPBLASLT_MATRIX_LAYOUT_STRIDED_BATCH_OFFSET with insufficient buffer size data64 = ld * col; EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( mat, HIPBLASLT_MATRIX_LAYOUT_STRIDED_BATCH_OFFSET, &data64, sizeof(int64_t)/2), @@ -135,7 +135,7 @@ void testing_aux_mat_set_attr_bad_arg(const Arguments& arg) mat, HIPBLASLT_MATRIX_LAYOUT_STRIDED_BATCH_OFFSET, nullptr, sizeof(int64_t)), HIPBLAS_STATUS_INVALID_VALUE); - // Test ROCBLASLT_MATRIX_LAYOUT_TYPE with insufficient buffer size + // Test HIPBLASLT_MATRIX_LAYOUT_TYPE with insufficient buffer size udata32 = static_cast(arg.a_type); EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( mat, HIPBLASLT_MATRIX_LAYOUT_TYPE, &udata32, sizeof(uint32_t)/2), @@ -149,7 +149,7 @@ void testing_aux_mat_set_attr_bad_arg(const Arguments& arg) mat, HIPBLASLT_MATRIX_LAYOUT_TYPE, nullptr, sizeof(uint32_t)), HIPBLAS_STATUS_INVALID_VALUE); - // Test ROCBLASLT_MATRIX_LAYOUT_ORDER with insufficient buffer size + // Test HIPBLASLT_MATRIX_LAYOUT_ORDER with insufficient buffer size data32 = HIPBLASLT_ORDER_COL; EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( mat, HIPBLASLT_MATRIX_LAYOUT_ORDER, &data32, sizeof(int32_t)/2), @@ -163,7 +163,7 @@ void testing_aux_mat_set_attr_bad_arg(const Arguments& arg) mat, HIPBLASLT_MATRIX_LAYOUT_ORDER, nullptr, sizeof(int32_t)), HIPBLAS_STATUS_INVALID_VALUE); - // Test ROCBLASLT_MATRIX_LAYOUT_ROWS with insufficient buffer size + // Test HIPBLASLT_MATRIX_LAYOUT_ROWS with insufficient buffer size udata64 = static_cast(row); EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( mat, HIPBLASLT_MATRIX_LAYOUT_ROWS, &udata64, sizeof(uint64_t)/2), @@ -177,7 +177,7 @@ void testing_aux_mat_set_attr_bad_arg(const Arguments& arg) mat, HIPBLASLT_MATRIX_LAYOUT_ROWS, nullptr, sizeof(uint64_t)), HIPBLAS_STATUS_INVALID_VALUE); - // Test ROCBLASLT_MATRIX_LAYOUT_COLS with insufficient buffer size + // Test HIPBLASLT_MATRIX_LAYOUT_COLS with insufficient buffer size udata64 = static_cast(col); EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( mat, HIPBLASLT_MATRIX_LAYOUT_COLS, &udata64, sizeof(uint64_t)/2), @@ -191,7 +191,7 @@ void testing_aux_mat_set_attr_bad_arg(const Arguments& arg) mat, HIPBLASLT_MATRIX_LAYOUT_COLS, nullptr, sizeof(uint64_t)), HIPBLAS_STATUS_INVALID_VALUE); - // Test ROCBLASLT_MATRIX_LAYOUT_LD with insufficient buffer size + // Test HIPBLASLT_MATRIX_LAYOUT_LD with insufficient buffer size data64 = ld; EXPECT_HIPBLAS_STATUS(hipblasLtMatrixLayoutSetAttribute( mat, HIPBLASLT_MATRIX_LAYOUT_LD, &data64, sizeof(int64_t)/2), @@ -386,7 +386,7 @@ void testing_aux_matmul_set_get_attr(const Arguments& arg) int64_t n = 1; int64_t k = 1; - // For ROCBLASLT_MATMUL_DESC_TRANSA + // For HIPBLASLT_MATMUL_DESC_TRANSA hipblasOperation_t transA = HIPBLAS_OP_T; hipblasOperation_t transA_r = HIPBLAS_OP_N; @@ -419,7 +419,7 @@ void testing_aux_matmul_set_get_attr(const Arguments& arg) ASSERT_TRUE(transA_r == transA); // validate - // For ROCBLASLT_MATMUL_DESC_TRANSB + // For HIPBLASLT_MATMUL_DESC_TRANSB hipblasOperation_t transB = HIPBLAS_OP_N; hipblasOperation_t transB_r = HIPBLAS_OP_T; @@ -442,7 +442,7 @@ void testing_aux_matmul_set_get_attr(const Arguments& arg) ASSERT_TRUE(transB_r == transB); // validate - // for ROCBLASLT_MATMUL_DESC_EPILOGUE + // for HIPBLASLT_MATMUL_DESC_EPILOGUE hipblasLtEpilogue_t data = HIPBLASLT_EPILOGUE_DEFAULT; hipblasLtEpilogue_t data_r = HIPBLASLT_EPILOGUE_RELU; EXPECT_HIPBLAS_STATUS(hipblasLtMatmulDescSetAttribute( @@ -454,7 +454,7 @@ void testing_aux_matmul_set_get_attr(const Arguments& arg) ASSERT_TRUE(data_r == data); - // for ROCBLASLT_MATMUL_DESC_BIAS_POINTER set and get + // for HIPBLASLT_MATMUL_DESC_BIAS_POINTER set and get void* d_bias; void* d_bias_r; CHECK_HIP_ERROR(hipMalloc(&d_bias, k * sizeof(hipblasLtHalf))); @@ -473,7 +473,7 @@ void testing_aux_matmul_set_get_attr(const Arguments& arg) ASSERT_TRUE(d_bias_r == d_bias); - // for ROCBLASLT_MATMUL_DESC_A_SCALE_POINTER + // for HIPBLASLT_MATMUL_DESC_A_SCALE_POINTER // We create a new matmul_descr becasue we don't want to set matmulDesc->scaleAType as Scalar in advance. // If we set it as Scalar, it will skip some cases, which is not what we desire for. const hipblasOperation_t opA = HIPBLAS_OP_T; @@ -505,7 +505,7 @@ void testing_aux_matmul_set_get_attr(const Arguments& arg) HIPBLAS_STATUS_SUCCESS); ASSERT_TRUE(d_scale_a_r == d_scale_a); // validate - // ROCBLASLT_MATMUL_DESC_A_SCALE_MODE + // HIPBLASLT_MATMUL_DESC_A_SCALE_MODE hipblasLtMatmulMatrixScale_t scale_mode_a = HIPBLASLT_MATMUL_MATRIX_SCALE_VEC32_UE8M0; hipblasLtMatmulMatrixScale_t scale_mode_a_r = HIPBLASLT_MATMUL_MATRIX_SCALE_END; @@ -532,7 +532,7 @@ void testing_aux_matmul_set_get_attr(const Arguments& arg) ASSERT_TRUE(scale_mode_a_r == scale_mode_a); // validate - // ROCBLASLT_MATMUL_DESC_B_SCALE_MODE + // HIPBLASLT_MATMUL_DESC_B_SCALE_MODE hipblasLtMatmulMatrixScale_t scale_mode_b = HIPBLASLT_MATMUL_MATRIX_SCALE_VEC32_UE8M0; hipblasLtMatmulMatrixScale_t scale_mode_b_r = HIPBLASLT_MATMUL_MATRIX_SCALE_END; @@ -632,7 +632,7 @@ void testing_aux_matmul_set_get_attr(const Arguments& arg) hipStream_t stream; CHECK_HIP_ERROR(hipStreamCreate(&stream)); - // for ROCBLASLT_MATMUL_DESC_B_SCALE_POINTER + // for HIPBLASLT_MATMUL_DESC_B_SCALE_POINTER float h_scale_b = 3.f; float* d_scale_b; float* d_scale_b_r; @@ -657,7 +657,7 @@ void testing_aux_matmul_set_get_attr(const Arguments& arg) ASSERT_TRUE(d_scale_b == d_scale_b_r); // validate - // for ROCBLASLT_MATMUL_DESC_C_SCALE_POINTER & ROCBLASLT_MATMUL_DESC_D_SCALE_POINTER + // for HIPBLASLT_MATMUL_DESC_C_SCALE_POINTER & HIPBLASLT_MATMUL_DESC_D_SCALE_POINTER float h_scale_c = 3.f; float h_scale_d = 3.f; float* d_scale_c; @@ -684,7 +684,7 @@ void testing_aux_matmul_set_get_attr(const Arguments& arg) HIPBLAS_STATUS_SUCCESS); - // for ROCBLASLT_MATMUL_DESC_EPILOGUE_AUX_SCALE_POINTER + // for HIPBLASLT_MATMUL_DESC_EPILOGUE_AUX_SCALE_POINTER float h_scale_e = 3.f; float* d_scale_e; CHECK_HIP_ERROR(hipMalloc(&d_scale_e, sizeof(float))); @@ -698,7 +698,7 @@ void testing_aux_matmul_set_get_attr(const Arguments& arg) matmul, HIPBLASLT_MATMUL_DESC_EPILOGUE_AUX_SCALE_POINTER, &d_scale_e, sizeof(float*)), HIPBLAS_STATUS_SUCCESS); - // For ROCBLASLT_MATMUL_DESC_POINTER_MODE + // For HIPBLASLT_MATMUL_DESC_POINTER_MODE hipblasLtPointerMode_t pMode = HIPBLASLT_POINTER_MODE_ALPHA_DEVICE_VECTOR_BETA_HOST; hipblasLtPointerMode_t pMode_r = HIPBLASLT_POINTER_MODE_HOST; EXPECT_HIPBLAS_STATUS(hipblasLtMatmulDescSetAttribute( @@ -716,7 +716,7 @@ void testing_aux_matmul_set_get_attr(const Arguments& arg) ASSERT_TRUE(pMode_r == pMode); // validate - // For ROCBLASLT_MATMUL_DESC_BIAS_DATA_TYPE Set Desc Bias Data Type + // For HIPBLASLT_MATMUL_DESC_BIAS_DATA_TYPE Set Desc Bias Data Type int32_t bias_data_type = HIP_R_16F; int32_t bias_data_type_r = HIP_R_32F; @@ -740,7 +740,7 @@ void testing_aux_matmul_set_get_attr(const Arguments& arg) ASSERT_TRUE(bias_data_type_r == bias_data_type); // validate - // For ROCBLASLT_MATMUL_DESC_EPILOGUE_AUX_POINTER Set auxiliary buffer + // For HIPBLASLT_MATMUL_DESC_EPILOGUE_AUX_POINTER Set auxiliary buffer void* d_aux_buffer; CHECK_HIP_ERROR(hipMalloc(&d_aux_buffer, m * n * sizeof(hipblasLtHalf))); EXPECT_HIPBLAS_STATUS(hipblasLtMatmulDescSetAttribute( @@ -751,7 +751,7 @@ void testing_aux_matmul_set_get_attr(const Arguments& arg) HIPBLAS_STATUS_SUCCESS); - // For ROCBLASLT_MATMUL_DESC_EPILOGUE_AUX_POINTER Set auxiliary leading dimension (ld) + // For HIPBLASLT_MATMUL_DESC_EPILOGUE_AUX_LD Set auxiliary leading dimension (ld) const int64_t aux_ld = m; EXPECT_HIPBLAS_STATUS(hipblasLtMatmulDescSetAttribute( matmul, HIPBLASLT_MATMUL_DESC_EPILOGUE_AUX_LD, &aux_ld, sizeof(aux_ld)/2), @@ -761,17 +761,17 @@ void testing_aux_matmul_set_get_attr(const Arguments& arg) matmul, HIPBLASLT_MATMUL_DESC_EPILOGUE_AUX_LD, &aux_ld, sizeof(aux_ld)), HIPBLAS_STATUS_SUCCESS); - // for ROCBLASLT_MATMUL_DESC_EPILOGUE_AUX_BATCH_STRIDE & Set Epilogue Aux Batch Stride + // for HIPBLASLT_MATMUL_DESC_EPILOGUE_AUX_BATCH_STRIDE & Set Epilogue Aux Batch Stride const int64_t aux_batch_stride = m * n; EXPECT_HIPBLAS_STATUS( hipblasLtMatmulDescSetAttribute( - matmul,HIPBLASLT_MATMUL_DESC_EPILOGUE_AUX_BATCH_STRIDE, &aux_batch_stride, sizeof(aux_batch_stride)/2), + matmul, HIPBLASLT_MATMUL_DESC_EPILOGUE_AUX_BATCH_STRIDE, &aux_batch_stride, sizeof(aux_batch_stride)/2), HIPBLAS_STATUS_INVALID_VALUE); EXPECT_HIPBLAS_STATUS(hipblasLtMatmulDescSetAttribute( matmul, HIPBLASLT_MATMUL_DESC_EPILOGUE_AUX_BATCH_STRIDE, &aux_batch_stride, sizeof(aux_batch_stride)), HIPBLAS_STATUS_SUCCESS); - // for ROCBLASLT_MATMUL_DESC_AMAX_D_POINTER + // for HIPBLASLT_MATMUL_DESC_AMAX_D_POINTER void *d_out_amax; void *d_out_amax_r; CHECK_HIP_ERROR(hipMalloc(&d_out_amax, 1 * sizeof(float))); @@ -794,7 +794,7 @@ void testing_aux_matmul_set_get_attr(const Arguments& arg) ASSERT_TRUE(d_out_amax == d_out_amax_r); // validate - // for ROCBLASLT_MATMUL_DESC_EPILOGUE_AUX_DATA_TYPE + // for HIPBLASLT_MATMUL_DESC_EPILOGUE_AUX_DATA_TYPE hipDataType aux_type_r; EXPECT_HIPBLAS_STATUS(hipblasLtMatmulDescSetAttribute( matmul, HIPBLASLT_MATMUL_DESC_EPILOGUE_AUX_DATA_TYPE, &arg.aux_type, sizeof(hipDataType)/2), @@ -814,7 +814,7 @@ void testing_aux_matmul_set_get_attr(const Arguments& arg) ASSERT_TRUE(arg.aux_type == aux_type_r); // validate - // for ROCBLASLT_MATMUL_DESC_COMPUTE_INPUT_TYPE_A_EXT & ROCBLASLT_MATMUL_DESC_COMPUTE_INPUT_TYPE_B_EXT + // for HIPBLASLT_MATMUL_DESC_COMPUTE_INPUT_TYPE_A_EXT & HIPBLASLT_MATMUL_DESC_COMPUTE_INPUT_TYPE_B_EXT hipDataType computeTypeA = HIP_R_16F; hipDataType computeTypeA_r = HIP_R_32F; hipDataType computeTypeB = HIP_R_16F; @@ -1382,6 +1382,9 @@ void testing_aux_mat_copy(const Arguments& arg) EXPECT_HIPBLAS_STATUS(hipblaslt_ext::copyMatmul(nullptr, matmul_dest), HIPBLAS_STATUS_INVALID_VALUE); EXPECT_HIPBLAS_STATUS(hipblaslt_ext::copyMatmul(matmul_src, nullptr), HIPBLAS_STATUS_INVALID_VALUE); EXPECT_HIPBLAS_STATUS(hipblaslt_ext::copyMatmul(matmul_src, matmul_dest), HIPBLAS_STATUS_SUCCESS); + + CHECK_HIPBLASLT_ERROR(hipblasLtMatmulDescDestroy(matmul_src)); + CHECK_HIPBLASLT_ERROR(hipblasLtMatmulDescDestroy(matmul_dest)); } void testing_aux_auxiliary_func(const Arguments& arg) From 4cd906dee4479417f647a64a5dad60dfcf9e8705 Mon Sep 17 00:00:00 2001 From: tomchengchitang Date: Tue, 1 Jul 2025 03:29:51 +0000 Subject: [PATCH 6/6] delete duplicate code --- projects/hipblaslt/clients/include/testing_auxiliary.hpp | 1 - 1 file changed, 1 deletion(-) diff --git a/projects/hipblaslt/clients/include/testing_auxiliary.hpp b/projects/hipblaslt/clients/include/testing_auxiliary.hpp index 25b7be341b03..0c10a1c61f04 100644 --- a/projects/hipblaslt/clients/include/testing_auxiliary.hpp +++ b/projects/hipblaslt/clients/include/testing_auxiliary.hpp @@ -1375,7 +1375,6 @@ void testing_aux_mat_copy(const Arguments& arg) { hipblasLtMatmulDesc_t matmul_src; CHECK_HIPBLASLT_ERROR(hipblasLtMatmulDescCreate(&matmul_src, arg.compute_type, arg.scale_type)); - CHECK_HIPBLASLT_ERROR(hipblasLtMatmulDescCreate(&matmul_src, arg.compute_type, arg.scale_type)); hipblasLtMatmulDesc_t matmul_dest; CHECK_HIPBLASLT_ERROR(hipblasLtMatmulDescCreate(&matmul_dest, HIPBLAS_COMPUTE_32F, HIP_R_32F));