Skip to content

Commit

Permalink
Merge pull request #1928 from brian-kelley/FixSpaddPerftest
Browse files Browse the repository at this point in the history 
Fix SpAdd perf test when offset/ordinal is not int
  • Loading branch information
@lucbv
lucbv authored Aug 2, 2023
2 parents efac9bf + 8edc246 commit fe9a5ac
Show file tree
Hide file tree
Showing 2 changed files with 77 additions and 35 deletions.
4 changes: 3 additions & 1 deletion common/src/KokkosKernels_default_types.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -25,7 +25,8 @@ using default_lno_t = int;
#elif defined(KOKKOSKERNELS_INST_ORDINAL_INT64_T)
using default_lno_t = int64_t;
#else
using default_lno_t = int;
// Non-ETI build: default to int
using default_lno_t = int;
#endif
// Prefer int as the default offset type, because cuSPARSE doesn't support
// size_t for rowptrs.
Expand All @@ -34,6 +35,7 @@ using default_size_type = int;
#elif defined(KOKKOSKERNELS_INST_OFFSET_SIZE_T)
using default_size_type = size_t;
#else
// Non-ETI build: default to int
using default_size_type = int;
#endif

Expand Down
108 changes: 74 additions & 34 deletions perf_test/sparse/KokkosSparse_spadd.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -185,6 +185,31 @@ void run_experiment(int argc, char** argv, CommonInputParams) {
"If running MKL, can't output the result to file");
}

// Check that offset/ordinal types are compatible with any requested TPLs
#ifdef KOKKOSKERNELS_ENABLE_TPL_MKL
if (params.use_mkl) {
if constexpr (!std::is_same_v<int, MKL_INT>) {
throw std::runtime_error(
"MKL configured with long long int not supported in Kokkos Kernels");
}
if constexpr (!std::is_same_v<MKL_INT, lno_t> ||
!std::is_same_v<MKL_INT, size_type>) {
throw std::runtime_error(
"Must enable int as both ordinal and offset type in KokkosKernels to "
"call MKL SpAdd");
}
}
#endif

if (params.use_cusparse) {
if constexpr (!std::is_same_v<int, lno_t> ||
!std::is_same_v<int, size_type>) {
throw std::runtime_error(
"Must enable int as both ordinal and offset type in KokkosKernels to "
"call cuSPARSE SpAdd");
}
}

std::cout << "************************************* \n";
crsMat_t A;
crsMat_t B;
Expand Down Expand Up @@ -319,9 +344,11 @@ void run_experiment(int argc, char** argv, CommonInputParams) {
}
#endif
#ifdef KOKKOSKERNELS_ENABLE_TPL_MKL
sparse_matrix_t Amkl, Bmkl, Cmkl;
sparse_matrix_t Amkl = sparse_matrix_t(), Bmkl = sparse_matrix_t(),
Cmkl = sparse_matrix_t();
if (params.use_mkl) {
if constexpr (std::is_same_v<int, MKL_INT>) {
if constexpr (std::is_same_v<lno_t, MKL_INT> &&
std::is_same_v<size_type, MKL_INT>) {
KOKKOSKERNELS_MKL_SAFE_CALL(mkl_sparse_d_create_csr(
&Amkl, SPARSE_INDEX_BASE_ZERO, m, n, (int*)A.graph.row_map.data(),
(int*)A.graph.row_map.data() + 1, A.graph.entries.data(),
Expand All @@ -330,9 +357,6 @@ void run_experiment(int argc, char** argv, CommonInputParams) {
&Bmkl, SPARSE_INDEX_BASE_ZERO, m, n, (int*)B.graph.row_map.data(),
(int*)B.graph.row_map.data() + 1, B.graph.entries.data(),
B.values.data()));
} else {
throw std::runtime_error(
"MKL configured with long long int not supported in Kokkos Kernels");
}
}
#endif
Expand All @@ -347,22 +371,30 @@ void run_experiment(int argc, char** argv, CommonInputParams) {
c_nnz = addHandle->get_c_nnz();
} else if (params.use_cusparse) {
#ifdef KOKKOSKERNELS_ENABLE_TPL_CUSPARSE
// Symbolic phase: compute buffer size, then compute nnz
size_t bufferSize;
KOKKOS_CUSPARSE_SAFE_CALL(cusparseDcsrgeam2_bufferSizeExt(
cusparseHandle, A.numRows(), A.numCols(), &alphabeta, A_cusparse,
A.nnz(), A.values.data(), A.graph.row_map.data(),
A.graph.entries.data(), &alphabeta, B_cusparse, B.nnz(),
B.values.data(), B.graph.row_map.data(), B.graph.entries.data(),
C_cusparse, NULL, row_mapC.data(), NULL, &bufferSize));
// Allocate work buffer
KOKKOS_IMPL_CUDA_SAFE_CALL(
cudaMalloc((void**)&cusparseBuffer, bufferSize));
KOKKOS_CUSPARSE_SAFE_CALL(cusparseXcsrgeam2Nnz(
cusparseHandle, m, n, A_cusparse, A.nnz(), A.graph.row_map.data(),
A.graph.entries.data(), B_cusparse, B.nnz(), B.graph.row_map.data(),
B.graph.entries.data(), C_cusparse, row_mapC.data(), &c_nnz,
cusparseBuffer));
if constexpr (std::is_same_v<lno_t, int> &&
std::is_same_v<size_type, int>) {
// Symbolic phase: compute buffer size, then compute nnz
size_t bufferSize;
KOKKOS_CUSPARSE_SAFE_CALL(cusparseDcsrgeam2_bufferSizeExt(
cusparseHandle, A.numRows(), A.numCols(), &alphabeta, A_cusparse,
A.nnz(), A.values.data(), A.graph.row_map.data(),
A.graph.entries.data(), &alphabeta, B_cusparse, B.nnz(),
B.values.data(), B.graph.row_map.data(), B.graph.entries.data(),
C_cusparse, NULL, row_mapC.data(), NULL, &bufferSize));
// Allocate work buffer
KOKKOS_IMPL_CUDA_SAFE_CALL(
cudaMalloc((void**)&cusparseBuffer, bufferSize));
KOKKOS_CUSPARSE_SAFE_CALL(cusparseXcsrgeam2Nnz(
cusparseHandle, m, n, A_cusparse, A.nnz(), A.graph.row_map.data(),
A.graph.entries.data(), B_cusparse, B.nnz(), B.graph.row_map.data(),
B.graph.entries.data(), C_cusparse, row_mapC.data(), &c_nnz,
cusparseBuffer));
} else {
throw std::runtime_error(
"Must enable int as both ordinal and offset type in KokkosKernels "
"to "
"call cuSPARSE");
}
#endif
}
if (!params.use_mkl) {
Expand All @@ -381,24 +413,32 @@ void run_experiment(int argc, char** argv, CommonInputParams) {
for (int numericRep = 0; numericRep < params.numericRepeat; numericRep++) {
if (params.use_cusparse) {
#ifdef KOKKOSKERNELS_ENABLE_TPL_CUSPARSE
KOKKOS_CUSPARSE_SAFE_CALL(cusparseDcsrgeam2(
cusparseHandle, m, n, &alphabeta, A_cusparse, A.nnz(),
A.values.data(), A.graph.row_map.data(), A.graph.entries.data(),
&alphabeta, B_cusparse, B.nnz(), B.values.data(),
B.graph.row_map.data(), B.graph.entries.data(), C_cusparse,
valuesC.data(), row_mapC.data(), entriesC.data(), cusparseBuffer));
if constexpr (std::is_same_v<lno_t, int> &&
std::is_same_v<size_type, int>) {
KOKKOS_CUSPARSE_SAFE_CALL(cusparseDcsrgeam2(
cusparseHandle, m, n, &alphabeta, A_cusparse, A.nnz(),
A.values.data(), A.graph.row_map.data(), A.graph.entries.data(),
&alphabeta, B_cusparse, B.nnz(), B.values.data(),
B.graph.row_map.data(), B.graph.entries.data(), C_cusparse,
valuesC.data(), row_mapC.data(), entriesC.data(),
cusparseBuffer));
}
#endif
} else if (params.use_mkl) {
#ifdef KOKKOSKERNELS_ENABLE_TPL_MKL
KOKKOSKERNELS_MKL_SAFE_CALL(mkl_sparse_d_add(
SPARSE_OPERATION_NON_TRANSPOSE, Amkl, 1.0, Bmkl, &Cmkl));
KOKKOSKERNELS_MKL_SAFE_CALL(mkl_sparse_destroy(Cmkl));
if constexpr (std::is_same_v<lno_t, int> &&
std::is_same_v<size_type, int>) {
KOKKOSKERNELS_MKL_SAFE_CALL(mkl_sparse_d_add(
SPARSE_OPERATION_NON_TRANSPOSE, Amkl, 1.0, Bmkl, &Cmkl));
KOKKOSKERNELS_MKL_SAFE_CALL(mkl_sparse_destroy(Cmkl));
}
#endif
} else {
spadd_numeric(
&kh, A.graph.row_map, A.graph.entries, A.values, 1.0, // A, alpha
B.graph.row_map, B.graph.entries, B.values, 1.0, // B, beta
row_mapC, entriesC, valuesC); // C
spadd_numeric(&kh, A.graph.row_map, A.graph.entries, A.values,
1.0, // A, alpha
B.graph.row_map, B.graph.entries, B.values,
1.0, // B, beta
row_mapC, entriesC, valuesC); // C
}
}
numericTime += timer.seconds();
Expand Down

0 comments on commit fe9a5ac

Please sign in to comment.