Skip to content

[hipsparse] Add HIP graph testing to conversion routines #3818

New issue

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

Sign up for GitHub

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

Already on GitHub? Sign in to your account

Merged
jsandham merged 7 commits into from
Jan 27, 2026
Merged

[hipsparse] Add HIP graph testing to conversion routines #3818

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
c664546
Add HIP graph testing to conversion routines
jsandham Jan 13, 2026
adb9451
Clang formatting
jsandham Jan 13, 2026
16682a6
Merge branch 'develop' of https://github.com/ROCm/rocm-libraries into…
jsandham Jan 14, 2026
bc34087
Clang formatting
jsandham Jan 14, 2026
189ab9d
Fix cuda compilation
jsandham Jan 14, 2026
e4954e5
Merge branch 'develop' into test_hipgraph_conversion
jsandham Jan 14, 2026
8d4bcd2
Merge branch 'develop' into test_hipgraph_conversion
jsandham Jan 23, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
66 changes: 66 additions & 0 deletions projects/hipsparse/clients/include/hipsparse_graph.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -120,4 +120,70 @@ namespace testing
TESTING_TEMPLATE(SDDMM_bufferSize)
TESTING_TEMPLATE(SDDMM_preprocess)
#endif

/*
* ===========================================================================
* conversion SPARSE
* ===========================================================================
*/

// coo2csr and csr2coo are not templated on T
TESTING_TEMPLATE(Xcoo2csr)
TESTING_TEMPLATE(Xcsr2coo)

#if(!defined(CUDART_VERSION) || CUDART_VERSION >= 10010)
// csr2csc_ex2 available from CUDA 10.1
TESTING_TEMPLATE(Csr2cscEx2_bufferSize)
TESTING_TEMPLATE(Csr2cscEx2)
#endif

// bsr2csr
TESTING_COMPUTE_TEMPLATE(bsr2csr)

// csr2gebsr and gebsr2csr
TESTING_TEMPLATE(Xcsr2gebsrNnz)
TESTING_COMPUTE_TEMPLATE(csr2gebsr)
TESTING_COMPUTE_TEMPLATE(gebsr2csr)

// gebsr2gebsr and gebsr2gebsc
TESTING_TEMPLATE(Xgebsr2gebsrNnz)
TESTING_COMPUTE_TEMPLATE(gebsr2gebsr)
TESTING_COMPUTE_TEMPLATE(gebsr2gebsc)

// csr2csr_compress
#if(!defined(CUDART_VERSION) || CUDART_VERSION < 13000)
TESTING_COMPUTE_TEMPLATE(nnz_compress)
#endif
TESTING_COMPUTE_TEMPLATE(csr2csr_compress)

// dense to sparse
#if(!defined(CUDART_VERSION) || CUDART_VERSION < 12000)
TESTING_COMPUTE_TEMPLATE(dense2csr)
TESTING_COMPUTE_TEMPLATE(dense2csc)
TESTING_COMPUTE_TEMPLATE(csr2dense)
TESTING_COMPUTE_TEMPLATE(csc2dense)
#endif

// prune routines
#if(!defined(CUDART_VERSION) || CUDART_VERSION < 13000)
TESTING_COMPUTE_TEMPLATE(pruneDense2csr_bufferSize)
TESTING_COMPUTE_TEMPLATE(pruneDense2csrNnz)
TESTING_COMPUTE_TEMPLATE(pruneDense2csr)
TESTING_COMPUTE_TEMPLATE(pruneCsr2csr_bufferSize)
TESTING_COMPUTE_TEMPLATE(pruneCsr2csrNnz)
TESTING_COMPUTE_TEMPLATE(pruneCsr2csr)
TESTING_COMPUTE_TEMPLATE(pruneDense2csr_bufferSizeExt)
TESTING_COMPUTE_TEMPLATE(pruneCsr2csr_bufferSizeExt)
TESTING_COMPUTE_TEMPLATE(pruneDense2csrByPercentage_bufferSize)
TESTING_COMPUTE_TEMPLATE(pruneDense2csrByPercentage_bufferSizeExt)
TESTING_COMPUTE_TEMPLATE(pruneDense2csrNnzByPercentage)
TESTING_COMPUTE_TEMPLATE(pruneDense2csrByPercentage)
TESTING_COMPUTE_TEMPLATE(pruneCsr2csrByPercentage_bufferSize)
TESTING_COMPUTE_TEMPLATE(pruneCsr2csrByPercentage_bufferSizeExt)
TESTING_COMPUTE_TEMPLATE(pruneCsr2csrNnzByPercentage)
TESTING_COMPUTE_TEMPLATE(pruneCsr2csrByPercentage)
#endif

// nnz routines
TESTING_COMPUTE_TEMPLATE(nnz)
}
101 changes: 50 additions & 51 deletions projects/hipsparse/clients/include/testing_bsr2csr.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,6 +30,7 @@
#include "gbyte.hpp"
#include "hipsparse.hpp"
#include "hipsparse_arguments.hpp"
#include "hipsparse_graph.hpp"
#include "hipsparse_test_unique_ptr.hpp"
#include "unit.hpp"
#include "utility.hpp"
Expand All @@ -53,12 +54,11 @@ void testing_bsr2csr_bad_arg(const Arguments& argus)
hipsparseIndexBase_t bsr_idx_base = HIPSPARSE_INDEX_BASE_ZERO;
hipsparseDirection_t dir = HIPSPARSE_DIRECTION_ROW;

std::unique_ptr<handle_struct> unique_ptr_handle(new handle_struct);
hipsparseHandle_t handle = unique_ptr_handle->handle;
std::unique_ptr<descr_struct> unique_ptr_csr_descr(new descr_struct);
hipsparseMatDescr_t csr_descr = unique_ptr_csr_descr->descr;
std::unique_ptr<descr_struct> unique_ptr_bsr_descr(new descr_struct);
hipsparseMatDescr_t bsr_descr = unique_ptr_bsr_descr->descr;
hipsparseLocalHandle_t handle;
std::unique_ptr<descr_struct> unique_ptr_csr_descr(new descr_struct);
hipsparseMatDescr_t csr_descr = unique_ptr_csr_descr->descr;
std::unique_ptr<descr_struct> unique_ptr_bsr_descr(new descr_struct);
hipsparseMatDescr_t bsr_descr = unique_ptr_bsr_descr->descr;

hipsparseSetMatIndexBase(csr_descr, csr_idx_base);
hipsparseSetMatIndexBase(bsr_descr, bsr_idx_base);
Expand Down Expand Up @@ -267,12 +267,11 @@ void testing_bsr2csr(Arguments argus)
hipsparseDirection_t dir = argus.dirA;
std::string filename = argus.filename;

std::unique_ptr<handle_struct> unique_ptr_handle(new handle_struct);
hipsparseHandle_t handle = unique_ptr_handle->handle;
std::unique_ptr<descr_struct> unique_ptr_csr_descr(new descr_struct);
hipsparseMatDescr_t csr_descr = unique_ptr_csr_descr->descr;
std::unique_ptr<descr_struct> unique_ptr_bsr_descr(new descr_struct);
hipsparseMatDescr_t bsr_descr = unique_ptr_bsr_descr->descr;
hipsparseLocalHandle_t handle(argus);
std::unique_ptr<descr_struct> unique_ptr_csr_descr(new descr_struct);
hipsparseMatDescr_t csr_descr = unique_ptr_csr_descr->descr;
std::unique_ptr<descr_struct> unique_ptr_bsr_descr(new descr_struct);
hipsparseMatDescr_t bsr_descr = unique_ptr_bsr_descr->descr;

hipsparseSetMatIndexBase(csr_descr, csr_idx_base);
hipsparseSetMatIndexBase(bsr_descr, bsr_idx_base);
Expand Down Expand Up @@ -356,19 +355,19 @@ void testing_bsr2csr(Arguments argus)

if(argus.unit_check)
{
CHECK_HIPSPARSE_ERROR(hipsparseXbsr2csr(handle,
dir,
mb,
nb,
bsr_descr,
dbsr_val,
dbsr_row_ptr,
dbsr_col_ind,
block_dim,
csr_descr,
dcsr_val,
dcsr_row_ptr,
dcsr_col_ind));
CHECK_HIPSPARSE_ERROR(testing::hipsparseXbsr2csr<T>(handle,
dir,
mb,
nb,
bsr_descr,
dbsr_val,
dbsr_row_ptr,
dbsr_col_ind,
block_dim,
csr_descr,
dcsr_val,
dcsr_row_ptr,
dcsr_col_ind));

// Copy output from device to host
CHECK_HIP_ERROR(hipMemcpy(
Expand Down Expand Up @@ -417,39 +416,39 @@ void testing_bsr2csr(Arguments argus)
// Warm up
for(int iter = 0; iter < number_cold_calls; ++iter)
{
CHECK_HIPSPARSE_ERROR(hipsparseXbsr2csr(handle,
dir,
mb,
nb,
bsr_descr,
dbsr_val,
dbsr_row_ptr,
dbsr_col_ind,
block_dim,
csr_descr,
dcsr_val,
dcsr_row_ptr,
dcsr_col_ind));
CHECK_HIPSPARSE_ERROR(testing::hipsparseXbsr2csr<T>(handle,
dir,
mb,
nb,
bsr_descr,
dbsr_val,
dbsr_row_ptr,
dbsr_col_ind,
block_dim,
csr_descr,
dcsr_val,
dcsr_row_ptr,
dcsr_col_ind));
}

double gpu_time_used = get_time_us();

// Performance run
for(int iter = 0; iter < number_hot_calls; ++iter)
{
CHECK_HIPSPARSE_ERROR(hipsparseXbsr2csr(handle,
dir,
mb,
nb,
bsr_descr,
dbsr_val,
dbsr_row_ptr,
dbsr_col_ind,
block_dim,
csr_descr,
dcsr_val,
dcsr_row_ptr,
dcsr_col_ind));
CHECK_HIPSPARSE_ERROR(testing::hipsparseXbsr2csr<T>(handle,
dir,
mb,
nb,
bsr_descr,
dbsr_val,
dbsr_row_ptr,
dbsr_col_ind,
block_dim,
csr_descr,
dcsr_val,
dcsr_row_ptr,
dcsr_col_ind));
}

gpu_time_used = (get_time_us() - gpu_time_used) / number_hot_calls;
Expand Down
13 changes: 6 additions & 7 deletions projects/hipsparse/clients/include/testing_coo2csr.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,6 +30,7 @@
#include "gbyte.hpp"
#include "hipsparse.hpp"
#include "hipsparse_arguments.hpp"
#include "hipsparse_graph.hpp"
#include "hipsparse_test_unique_ptr.hpp"
#include "unit.hpp"
#include "utility.hpp"
Expand All @@ -50,8 +51,7 @@ void testing_coo2csr_bad_arg(const Arguments& argus)
int safe_size = 100;
hipsparseIndexBase_t idx_base = HIPSPARSE_INDEX_BASE_ZERO;

std::unique_ptr<handle_struct> unique_ptr_handle(new handle_struct);
hipsparseHandle_t handle = unique_ptr_handle->handle;
hipsparseLocalHandle_t handle;

auto coo_row_ind_managed
= hipsparse_unique_ptr{device_malloc(sizeof(int) * safe_size), device_free};
Expand Down Expand Up @@ -80,8 +80,7 @@ void testing_coo2csr(Arguments argus)
hipsparseIndexBase_t idx_base = argus.baseA;
std::string filename = argus.filename;

std::unique_ptr<handle_struct> unique_ptr_handle(new handle_struct);
hipsparseHandle_t handle = unique_ptr_handle->handle;
hipsparseLocalHandle_t handle(argus);

srand(12345ULL);

Expand Down Expand Up @@ -116,7 +115,7 @@ void testing_coo2csr(Arguments argus)
if(argus.unit_check)
{
CHECK_HIPSPARSE_ERROR(
hipsparseXcoo2csr(handle, dcoo_row_ind, nnz, m, dcsr_row_ptr, idx_base));
testing::hipsparseXcoo2csr(handle, dcoo_row_ind, nnz, m, dcsr_row_ptr, idx_base));

// Copy output from device to host
CHECK_HIP_ERROR(hipMemcpy(
Expand Down Expand Up @@ -148,7 +147,7 @@ void testing_coo2csr(Arguments argus)
for(int iter = 0; iter < number_cold_calls; ++iter)
{
CHECK_HIPSPARSE_ERROR(
hipsparseXcoo2csr(handle, dcoo_row_ind, nnz, m, dcsr_row_ptr, idx_base));
testing::hipsparseXcoo2csr(handle, dcoo_row_ind, nnz, m, dcsr_row_ptr, idx_base));
}

double gpu_time_used = get_time_us();
Expand All @@ -157,7 +156,7 @@ void testing_coo2csr(Arguments argus)
for(int iter = 0; iter < number_hot_calls; ++iter)
{
CHECK_HIPSPARSE_ERROR(
hipsparseXcoo2csr(handle, dcoo_row_ind, nnz, m, dcsr_row_ptr, idx_base));
testing::hipsparseXcoo2csr(handle, dcoo_row_ind, nnz, m, dcsr_row_ptr, idx_base));
}

gpu_time_used = (get_time_us() - gpu_time_used) / number_hot_calls;
Expand Down
13 changes: 6 additions & 7 deletions projects/hipsparse/clients/include/testing_csr2coo.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,6 +30,7 @@
#include "gbyte.hpp"
#include "hipsparse.hpp"
#include "hipsparse_arguments.hpp"
#include "hipsparse_graph.hpp"
#include "hipsparse_test_unique_ptr.hpp"
#include "unit.hpp"
#include "utility.hpp"
Expand All @@ -50,8 +51,7 @@ void testing_csr2coo_bad_arg(const Arguments& argus)
int safe_size = 100;
hipsparseIndexBase_t idx_base = HIPSPARSE_INDEX_BASE_ZERO;

std::unique_ptr<handle_struct> unique_ptr_handle(new handle_struct);
hipsparseHandle_t handle = unique_ptr_handle->handle;
hipsparseLocalHandle_t handle;

auto csr_row_ptr_managed
= hipsparse_unique_ptr{device_malloc(sizeof(int) * safe_size), device_free};
Expand Down Expand Up @@ -80,8 +80,7 @@ void testing_csr2coo(Arguments argus)
hipsparseIndexBase_t idx_base = argus.baseA;
std::string filename = argus.filename;

std::unique_ptr<handle_struct> unique_ptr_handle(new handle_struct);
hipsparseHandle_t handle = unique_ptr_handle->handle;
hipsparseLocalHandle_t handle(argus);

srand(12345ULL);

Expand Down Expand Up @@ -110,7 +109,7 @@ void testing_csr2coo(Arguments argus)
if(argus.unit_check)
{
CHECK_HIPSPARSE_ERROR(
hipsparseXcsr2coo(handle, dcsr_row_ptr, nnz, m, dcoo_row_ind, idx_base));
testing::hipsparseXcsr2coo(handle, dcsr_row_ptr, nnz, m, dcoo_row_ind, idx_base));

// Copy output from device to host
std::vector<int> hcoo_row_ind(nnz);
Expand Down Expand Up @@ -143,7 +142,7 @@ void testing_csr2coo(Arguments argus)
for(int iter = 0; iter < number_cold_calls; ++iter)
{
CHECK_HIPSPARSE_ERROR(
hipsparseXcsr2coo(handle, dcsr_row_ptr, nnz, m, dcoo_row_ind, idx_base));
testing::hipsparseXcsr2coo(handle, dcsr_row_ptr, nnz, m, dcoo_row_ind, idx_base));
}

double gpu_time_used = get_time_us();
Expand All @@ -152,7 +151,7 @@ void testing_csr2coo(Arguments argus)
for(int iter = 0; iter < number_hot_calls; ++iter)
{
CHECK_HIPSPARSE_ERROR(
hipsparseXcsr2coo(handle, dcsr_row_ptr, nnz, m, dcoo_row_ind, idx_base));
testing::hipsparseXcsr2coo(handle, dcsr_row_ptr, nnz, m, dcoo_row_ind, idx_base));
}

gpu_time_used = (get_time_us() - gpu_time_used) / number_hot_calls;
Expand Down
Loading
Loading