/* ************************************************************************ * Copyright (c) 2020 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 #ifndef TESTING_BSRMM_HPP #define TESTING_BSRMM_HPP #include "hipsparse.hpp" #include "hipsparse_test_unique_ptr.hpp" #include "unit.hpp" #include "utility.hpp" #include #include #include using namespace hipsparse; using namespace hipsparse_test; template void testing_bsrmm_bad_arg(void) { #ifdef __HIP_PLATFORM_NVIDIA__ // do not test for bad args return; #endif int mb = 100; int n = 100; int kb = 100; int nnzb = 100; int block_dim = 100; int ldb = 100; int ldc = 100; int safe_size = 100; T alpha = 0.6; T beta = 0.2; hipsparseDirection_t dirA = HIPSPARSE_DIRECTION_ROW; hipsparseOperation_t transA = HIPSPARSE_OPERATION_NON_TRANSPOSE; hipsparseOperation_t transB = HIPSPARSE_OPERATION_NON_TRANSPOSE; hipsparseStatus_t status; std::unique_ptr unique_ptr_handle(new handle_struct); hipsparseHandle_t handle = unique_ptr_handle->handle; std::unique_ptr unique_ptr_descr(new descr_struct); hipsparseMatDescr_t descr = unique_ptr_descr->descr; auto dbsr_row_ptr_managed = hipsparse_unique_ptr{device_malloc(sizeof(int) * safe_size), device_free}; auto dbsr_col_ind_managed = hipsparse_unique_ptr{device_malloc(sizeof(int) * safe_size), device_free}; auto dbsr_val_managed = hipsparse_unique_ptr{device_malloc(sizeof(T) * safe_size), device_free}; auto dB_managed = hipsparse_unique_ptr{device_malloc(sizeof(T) * safe_size), device_free}; auto dC_managed = hipsparse_unique_ptr{device_malloc(sizeof(T) * safe_size), device_free}; int* dbsr_row_ptr = (int*)dbsr_row_ptr_managed.get(); int* dbsr_col_ind = (int*)dbsr_col_ind_managed.get(); T* dbsr_val = (T*)dbsr_val_managed.get(); T* dB = (T*)dB_managed.get(); T* dC = (T*)dC_managed.get(); if(!dbsr_val || !dbsr_row_ptr || !dbsr_col_ind || !dB || !dC) { PRINT_IF_HIP_ERROR(hipErrorOutOfMemory); return; } // Testing hipsparseXbsrmm() // Test invalid handle status = hipsparseXbsrmm(nullptr, dirA, transA, transB, mb, n, kb, nnzb, &alpha, descr, dbsr_val, dbsr_row_ptr, dbsr_col_ind, block_dim, dB, ldb, &beta, dC, ldc); verify_hipsparse_status_invalid_handle(status); // Test invalid pointers status = hipsparseXbsrmm(handle, dirA, transA, transB, mb, n, kb, nnzb, (T*)nullptr, descr, dbsr_val, dbsr_row_ptr, dbsr_col_ind, block_dim, dB, ldb, &beta, dC, ldc); verify_hipsparse_status_invalid_pointer(status, "Error: alpha is nullptr"); status = hipsparseXbsrmm(handle, dirA, transA, transB, mb, n, kb, nnzb, &alpha, nullptr, dbsr_val, dbsr_row_ptr, dbsr_col_ind, block_dim, dB, ldb, &beta, dC, ldc); verify_hipsparse_status_invalid_pointer(status, "Error: descr is nullptr"); status = hipsparseXbsrmm(handle, dirA, transA, transB, mb, n, kb, nnzb, &alpha, descr, (T*)nullptr, dbsr_row_ptr, dbsr_col_ind, block_dim, dB, ldb, &beta, dC, ldc); verify_hipsparse_status_invalid_pointer(status, "Error: dbsr_val is nullptr"); status = hipsparseXbsrmm(handle, dirA, transA, transB, mb, n, kb, nnzb, &alpha, descr, dbsr_val, nullptr, dbsr_col_ind, block_dim, dB, ldb, &beta, dC, ldc); verify_hipsparse_status_invalid_pointer(status, "Error: dbsr_row_ptr is nullptr"); status = hipsparseXbsrmm(handle, dirA, transA, transB, mb, n, kb, nnzb, &alpha, descr, dbsr_val, dbsr_row_ptr, nullptr, block_dim, dB, ldb, &beta, dC, ldc); verify_hipsparse_status_invalid_pointer(status, "Error: dbsr_col_ind is nullptr"); status = hipsparseXbsrmm(handle, dirA, transA, transB, mb, n, kb, nnzb, &alpha, descr, dbsr_val, dbsr_row_ptr, dbsr_col_ind, block_dim, (T*)nullptr, ldb, &beta, dC, ldc); verify_hipsparse_status_invalid_pointer(status, "Error: dB is nullptr"); status = hipsparseXbsrmm(handle, dirA, transA, transB, mb, n, kb, nnzb, &alpha, descr, dbsr_val, dbsr_row_ptr, dbsr_col_ind, block_dim, dB, ldb, (T*)nullptr, dC, ldc); verify_hipsparse_status_invalid_pointer(status, "Error: beta is nullptr"); status = hipsparseXbsrmm(handle, dirA, transA, transB, mb, n, kb, nnzb, &alpha, descr, dbsr_val, dbsr_row_ptr, dbsr_col_ind, block_dim, dB, ldb, &beta, (T*)nullptr, ldc); verify_hipsparse_status_invalid_pointer(status, "Error: dC is nullptr"); // Test invalid sizes status = hipsparseXbsrmm(handle, dirA, transA, transB, -1, n, kb, nnzb, &alpha, descr, dbsr_val, dbsr_row_ptr, dbsr_col_ind, block_dim, dB, ldb, &beta, dC, ldc); verify_hipsparse_status_invalid_size(status, "Error: mb is invalid"); status = hipsparseXbsrmm(handle, dirA, transA, transB, mb, -1, kb, nnzb, &alpha, descr, dbsr_val, dbsr_row_ptr, dbsr_col_ind, block_dim, dB, ldb, &beta, dC, ldc); verify_hipsparse_status_invalid_size(status, "Error: n is invalid"); status = hipsparseXbsrmm(handle, dirA, transA, transB, mb, n, -1, nnzb, &alpha, descr, dbsr_val, dbsr_row_ptr, dbsr_col_ind, block_dim, dB, ldb, &beta, dC, ldc); verify_hipsparse_status_invalid_size(status, "Error: kb is invalid"); status = hipsparseXbsrmm(handle, dirA, transA, transB, mb, n, kb, nnzb, &alpha, descr, dbsr_val, dbsr_row_ptr, dbsr_col_ind, 0, dB, ldb, &beta, dC, ldc); verify_hipsparse_status_invalid_size(status, "Error: block_dim is invalid"); // Test not implemented (mapped to hiparse internal error) status = hipsparseXbsrmm(handle, dirA, HIPSPARSE_OPERATION_TRANSPOSE, transB, mb, n, kb, nnzb, &alpha, descr, dbsr_val, dbsr_row_ptr, dbsr_col_ind, block_dim, dB, ldb, &beta, dC, ldc); verify_hipsparse_status_not_supported(status, "Error: Passed value for transA is not supported"); status = hipsparseXbsrmm(handle, dirA, HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE, transB, mb, n, kb, nnzb, &alpha, descr, dbsr_val, dbsr_row_ptr, dbsr_col_ind, block_dim, dB, ldb, &beta, dC, ldc); verify_hipsparse_status_not_supported(status, "Error: Passed value for transA is not supported"); status = hipsparseXbsrmm(handle, dirA, transA, HIPSPARSE_OPERATION_CONJUGATE_TRANSPOSE, mb, n, kb, nnzb, &alpha, descr, dbsr_val, dbsr_row_ptr, dbsr_col_ind, block_dim, dB, ldb, &beta, dC, ldc); verify_hipsparse_status_not_supported(status, "Error: Passed value for transB is not supported"); } template hipsparseStatus_t testing_bsrmm(Arguments argus) { int safe_size = 100; int m = argus.M; int n = argus.N; int k = argus.K; int block_dim = argus.block_dim; int ldb = argus.ldb; int ldc = argus.ldc; T h_alpha = make_DataType(argus.alpha); T h_beta = make_DataType(argus.beta); hipsparseDirection_t dirA = argus.dirA; hipsparseOperation_t transA = argus.transA; hipsparseOperation_t transB = argus.transB; hipsparseIndexBase_t idx_base = argus.idx_base; std::string binfile = ""; std::string filename = ""; hipsparseStatus_t status; // When in testing mode, m == n == -99 indicates that we are testing with a real // matrix from cise.ufl.edu if(m == -99 && k == -99 && argus.timing == 0) { binfile = argus.filename; m = k = safe_size; } if(argus.timing == 1) { filename = argus.filename; } int mb = -1; int kb = -1; if(block_dim > 0) { mb = (m + block_dim - 1) / block_dim; kb = (k + block_dim - 1) / block_dim; } std::unique_ptr test_handle(new handle_struct); hipsparseHandle_t handle = test_handle->handle; std::unique_ptr test_descr(new descr_struct); hipsparseMatDescr_t descr = test_descr->descr; // Set matrix index base CHECK_HIPSPARSE_ERROR(hipsparseSetMatIndexBase(descr, idx_base)); // Argument sanity check before allocating invalid memory if(mb <= 0 || n <= 0 || kb <= 0 || block_dim <= 0) { auto dbsr_row_ptr_managed = hipsparse_unique_ptr{device_malloc(sizeof(int) * safe_size), device_free}; auto dbsr_col_ind_managed = hipsparse_unique_ptr{device_malloc(sizeof(int) * safe_size), device_free}; auto dbsr_val_managed = hipsparse_unique_ptr{device_malloc(sizeof(T) * safe_size), device_free}; auto dB_managed = hipsparse_unique_ptr{device_malloc(sizeof(T) * safe_size), device_free}; auto dC_managed = hipsparse_unique_ptr{device_malloc(sizeof(T) * safe_size), device_free}; int* dbsr_row_ptr = (int*)dbsr_row_ptr_managed.get(); int* dbsr_col_ind = (int*)dbsr_col_ind_managed.get(); T* dbsr_val = (T*)dbsr_val_managed.get(); T* dB = (T*)dB_managed.get(); T* dC = (T*)dC_managed.get(); if(!dbsr_val || !dbsr_row_ptr || !dbsr_col_ind || !dB || !dC) { verify_hipsparse_status_success( HIPSPARSE_STATUS_ALLOC_FAILED, "!dbsr_row_ptr || !dbsr_col_ind || !dbsr_val || !dB || !dC"); return HIPSPARSE_STATUS_ALLOC_FAILED; } CHECK_HIPSPARSE_ERROR(hipsparseSetPointerMode(handle, HIPSPARSE_POINTER_MODE_HOST)); status = hipsparseXbsrmm(handle, dirA, transA, transB, mb, n, kb, safe_size, &h_alpha, descr, dbsr_val, dbsr_row_ptr, dbsr_col_ind, block_dim, dB, ldb, &h_beta, dC, ldc); if(mb < 0 || n < 0 || kb < 0 || block_dim <= 0) { verify_hipsparse_status_invalid_size( status, "Error: mb < 0 || n < 0 || kb < 0 || block_dim <= 0"); } else { verify_hipsparse_status_success(status, "mb >= 0 && n >= 0 && kb >= 0 && block_dim > 0"); } return HIPSPARSE_STATUS_SUCCESS; } // Read or construct CSR matrix std::vector csr_row_ptr; std::vector csr_col_ind; std::vector csr_val; int nnz; srand(12345ULL); if(binfile != "") { if(read_bin_matrix(binfile.c_str(), m, k, nnz, csr_row_ptr, csr_col_ind, csr_val, idx_base) != 0) { fprintf(stderr, "Cannot open [read] %s\n", binfile.c_str()); return HIPSPARSE_STATUS_INTERNAL_ERROR; } } else if(argus.laplacian) { m = k = gen_2d_laplacian(argus.laplacian, csr_row_ptr, csr_col_ind, csr_val, idx_base); nnz = csr_row_ptr[m]; } else { std::vector coo_row_ind; if(filename != "") { if(read_mtx_matrix( filename.c_str(), m, k, nnz, coo_row_ind, csr_col_ind, csr_val, idx_base) != 0) { fprintf(stderr, "Cannot open [read] %s\n", filename.c_str()); return HIPSPARSE_STATUS_INTERNAL_ERROR; } } else { double scale = 0.02; if(m > 1000 || k > 1000) { scale = 2.0 / std::max(m, k); } nnz = m * scale * k; gen_matrix_coo(m, k, nnz, coo_row_ind, csr_col_ind, csr_val, idx_base); } // Convert COO to CSR csr_row_ptr.resize(m + 1, 0); for(int i = 0; i < nnz; ++i) { ++csr_row_ptr[coo_row_ind[i] + 1 - idx_base]; } csr_row_ptr[0] = idx_base; for(int i = 0; i < m; ++i) { csr_row_ptr[i + 1] += csr_row_ptr[i]; } } // m and k can be modifed if we read in a matrix from a file mb = (m + block_dim - 1) / block_dim; kb = (k + block_dim - 1) / block_dim; // Allocate memory on device for CSR matrix and BSR row pointer array auto dcsr_row_ptrA_managed = hipsparse_unique_ptr{device_malloc(sizeof(int) * (m + 1)), device_free}; auto dcsr_col_indA_managed = hipsparse_unique_ptr{device_malloc(sizeof(int) * nnz), device_free}; auto dcsr_valA_managed = hipsparse_unique_ptr{device_malloc(sizeof(T) * nnz), device_free}; auto dbsr_row_ptrA_managed = hipsparse_unique_ptr{device_malloc(sizeof(int) * (mb + 1)), device_free}; int* dcsr_row_ptrA = (int*)dcsr_row_ptrA_managed.get(); int* dcsr_col_indA = (int*)dcsr_col_indA_managed.get(); T* dcsr_valA = (T*)dcsr_valA_managed.get(); int* dbsr_row_ptrA = (int*)dbsr_row_ptrA_managed.get(); if(!dcsr_row_ptrA || !dcsr_col_indA || !dcsr_valA || !dbsr_row_ptrA) { verify_hipsparse_status_success( HIPSPARSE_STATUS_ALLOC_FAILED, "!dcsr_row_ptrA || !dcsr_col_indA || !dcsr_valA || !dbsr_row_ptrA"); return HIPSPARSE_STATUS_ALLOC_FAILED; } // Copy CSR from host to device CHECK_HIP_ERROR( hipMemcpy(dcsr_row_ptrA, csr_row_ptr.data(), sizeof(int) * (m + 1), hipMemcpyHostToDevice)); CHECK_HIP_ERROR( hipMemcpy(dcsr_col_indA, csr_col_ind.data(), sizeof(int) * nnz, hipMemcpyHostToDevice)); CHECK_HIP_ERROR(hipMemcpy(dcsr_valA, csr_val.data(), sizeof(T) * nnz, hipMemcpyHostToDevice)); // Convert CSR to BSR int nnzb; CHECK_HIPSPARSE_ERROR(hipsparseSetPointerMode(handle, HIPSPARSE_POINTER_MODE_HOST)); CHECK_HIPSPARSE_ERROR(hipsparseXcsr2bsrNnz(handle, dirA, m, k, descr, dcsr_row_ptrA, dcsr_col_indA, block_dim, descr, dbsr_row_ptrA, &nnzb)); auto dbsr_col_indA_managed = hipsparse_unique_ptr{device_malloc(sizeof(int) * nnzb), device_free}; auto dbsr_valA_managed = hipsparse_unique_ptr{ device_malloc(sizeof(T) * nnzb * block_dim * block_dim), device_free}; int* dbsr_col_indA = (int*)dbsr_col_indA_managed.get(); T* dbsr_valA = (T*)dbsr_valA_managed.get(); if(!dbsr_valA || !dbsr_col_indA) { verify_hipsparse_status_success(HIPSPARSE_STATUS_ALLOC_FAILED, "!dbsr_valA || !dbsr_col_indA"); return HIPSPARSE_STATUS_ALLOC_FAILED; } CHECK_HIPSPARSE_ERROR(hipsparseXcsr2bsr(handle, dirA, m, k, descr, dcsr_valA, dcsr_row_ptrA, dcsr_col_indA, block_dim, descr, dbsr_valA, dbsr_row_ptrA, dbsr_col_indA)); // Host BSR matrix std::vector hbsr_row_ptrA(mb + 1); std::vector hbsr_col_indA(nnzb); std::vector hbsr_valA(nnzb * block_dim * block_dim); // Copy BSR matrix to host CHECK_HIP_ERROR(hipMemcpy( hbsr_row_ptrA.data(), dbsr_row_ptrA, sizeof(int) * (mb + 1), hipMemcpyDeviceToHost)); CHECK_HIP_ERROR( hipMemcpy(hbsr_col_indA.data(), dbsr_col_indA, sizeof(int) * nnzb, hipMemcpyDeviceToHost)); CHECK_HIP_ERROR(hipMemcpy(hbsr_valA.data(), dbsr_valA, sizeof(T) * nnzb * block_dim * block_dim, hipMemcpyDeviceToHost)); m = mb * block_dim; k = kb * block_dim; // Some matrix properties ldb = (transB == HIPSPARSE_OPERATION_NON_TRANSPOSE) ? k : n; ldc = m; int ncol_B = (transB == HIPSPARSE_OPERATION_NON_TRANSPOSE ? n : k); int nnz_B = ldb * ncol_B; int nnz_C = ldc * n; // Allocate host memory for dense matrices std::vector hB(nnz_B); std::vector hC_1(nnz_C); std::vector hC_2(nnz_C); std::vector hC_gold(nnz_C); hipsparseInit(hB, ldb, ncol_B); hipsparseInit(hC_1, ldc, n); // copy vector is easy in STL; hC_gold = hC_1: save a copy in hy_gold which will be output of // CPU hC_gold = hC_1; hC_2 = hC_1; // allocate memory on device auto dB_managed = hipsparse_unique_ptr{device_malloc(sizeof(T) * nnz_B), device_free}; auto dC_1_managed = hipsparse_unique_ptr{device_malloc(sizeof(T) * nnz_C), device_free}; auto dC_2_managed = hipsparse_unique_ptr{device_malloc(sizeof(T) * nnz_C), device_free}; auto d_alpha_managed = hipsparse_unique_ptr{device_malloc(sizeof(T)), device_free}; auto d_beta_managed = hipsparse_unique_ptr{device_malloc(sizeof(T)), device_free}; T* dB = (T*)dB_managed.get(); T* dC_1 = (T*)dC_1_managed.get(); T* dC_2 = (T*)dC_2_managed.get(); T* d_alpha = (T*)d_alpha_managed.get(); T* d_beta = (T*)d_beta_managed.get(); if(!dB || !dC_1 || !dC_2 || !d_alpha || !d_beta) { verify_hipsparse_status_success(HIPSPARSE_STATUS_ALLOC_FAILED, "!dB || !dC_1 || !dC_2 || !d_alpha || !d_beta"); return HIPSPARSE_STATUS_ALLOC_FAILED; } // Copy data from host to device CHECK_HIP_ERROR(hipMemcpy(dB, hB.data(), sizeof(T) * nnz_B, hipMemcpyHostToDevice)); CHECK_HIP_ERROR(hipMemcpy(dC_1, hC_1.data(), sizeof(T) * nnz_C, hipMemcpyHostToDevice)); CHECK_HIP_ERROR(hipMemcpy(dC_2, hC_2.data(), sizeof(T) * nnz_C, hipMemcpyHostToDevice)); CHECK_HIP_ERROR(hipMemcpy(d_alpha, &h_alpha, sizeof(T), hipMemcpyHostToDevice)); CHECK_HIP_ERROR(hipMemcpy(d_beta, &h_beta, sizeof(T), hipMemcpyHostToDevice)); if(argus.unit_check) { // Testing using host pointer mode CHECK_HIPSPARSE_ERROR(hipsparseSetPointerMode(handle, HIPSPARSE_POINTER_MODE_HOST)); CHECK_HIPSPARSE_ERROR(hipsparseXbsrmm(handle, dirA, transA, transB, mb, n, kb, nnzb, &h_alpha, descr, dbsr_valA, dbsr_row_ptrA, dbsr_col_indA, block_dim, dB, ldb, &h_beta, dC_1, ldc)); // Testing using device pointer mode CHECK_HIPSPARSE_ERROR(hipsparseSetPointerMode(handle, HIPSPARSE_POINTER_MODE_DEVICE)); CHECK_HIPSPARSE_ERROR(hipsparseXbsrmm(handle, dirA, transA, transB, mb, n, kb, nnzb, d_alpha, descr, dbsr_valA, dbsr_row_ptrA, dbsr_col_indA, block_dim, dB, ldb, d_beta, dC_2, ldc)); // Copy output from device to host CHECK_HIP_ERROR(hipMemcpy(hC_1.data(), dC_1, sizeof(T) * nnz_C, hipMemcpyDeviceToHost)); CHECK_HIP_ERROR(hipMemcpy(hC_2.data(), dC_2, sizeof(T) * nnz_C, hipMemcpyDeviceToHost)); // Host bsrmm host_bsrmm(mb, n, kb, block_dim, dirA, transA, transB, h_alpha, hbsr_row_ptrA, hbsr_col_indA, hbsr_valA, hB, ldb, h_beta, hC_gold, ldc, idx_base); // Unit check unit_check_near(1, nnz_C, 1, hC_gold.data(), hC_1.data()); unit_check_near(1, nnz_C, 1, hC_gold.data(), hC_2.data()); } return HIPSPARSE_STATUS_SUCCESS; } #endif // TESTING_BSRMM_HPP