Skip to content

Commit b8f2251

Browse files
wdeconinckwdeconinck
committed
Add HIP backend for LinearAlgebraSparse
1 parent 3d7f4fe commit b8f2251

File tree

6 files changed

+361
-0
lines changed

6 files changed

+361
-0
lines changed

‎CMakeLists.txt

+7
Original file line numberDiff line numberDiff line change
@@ -263,6 +263,13 @@ ecbuild_add_option( FEATURE CUDA
263263
DESCRIPTION "CUDA GPU linear algebra operations"
264264
REQUIRED_PACKAGES CUDAToolkit )
265265

266+
#### HIP
267+
268+
ecbuild_add_option( FEATURE HIP
269+
DEFAULT OFF
270+
DESCRIPTION "HIP GPU linear algebra operations"
271+
REQUIRED_PACKAGES hip hipsparse )
272+
266273
### ViennaCL
267274

268275
ecbuild_add_option( FEATURE VIENNACL

‎src/eckit/linalg/CMakeLists.txt

+8
Original file line numberDiff line numberDiff line change
@@ -50,6 +50,14 @@ if( eckit_HAVE_CUDA )
5050
list( APPEND eckit_la_plibs CUDA::cudart CUDA::cusparse CUDA::cublas )
5151
endif()
5252

53+
if( eckit_HAVE_HIP )
54+
list( APPEND eckit_la_srcs
55+
sparse/LinearAlgebraHIP.cc
56+
sparse/LinearAlgebraHIP.h )
57+
list( APPEND eckit_la_plibs hip::host roc::hipsparse )
58+
endif()
59+
60+
5361
if( eckit_HAVE_EIGEN )
5462
list( APPEND eckit_la_srcs
5563
dense/LinearAlgebraEigen.cc

‎src/eckit/linalg/detail/HIP.h

+38
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,38 @@
1+
/*
2+
* (C) Copyright 2025- ECMWF.
3+
*
4+
* This software is licensed under the terms of the Apache Licence Version 2.0
5+
* which can be obtained at http://www.apache.org/licenses/LICENSE-2.0.
6+
* In applying this licence, ECMWF does not waive the privileges and immunities
7+
* granted to it by virtue of its status as an intergovernmental organisation
8+
* nor does it submit to any jurisdiction.
9+
*/
10+
11+
12+
// There is a name clash because hip_runtime defines DEPRECATED, and eckit as well
13+
#ifdef DEPRECATED
14+
#undef DEPRECATED
15+
#endif
16+
17+
#include <hip/hip_runtime.h>
18+
#include <hip/library_types.h>
19+
#include <hipsparse/hipsparse.h>
20+
21+
// There is a name clash because hip_runtime defines DEPRECATED, and eckit as well
22+
#ifdef DEPRECATED
23+
#undef DEPRECATED
24+
#endif
25+
26+
#define CALL_HIP(e) \
27+
{ \
28+
hipError_t error; \
29+
if ((error = e) != hipSuccess) \
30+
printf("%s failed with error code %d @ %s +%d\n", #e, error, __FILE__, __LINE__), exit(EXIT_FAILURE); \
31+
}
32+
33+
#define CALL_HIPSPARSE(e) \
34+
{ \
35+
hipsparseStatus_t error; \
36+
if ((error = e) != HIPSPARSE_STATUS_SUCCESS) \
37+
printf("%s failed with error code %d @ %s +%d\n", #e, error, __FILE__, __LINE__), exit(EXIT_FAILURE); \
38+
}

‎src/eckit/linalg/sparse/LinearAlgebraHIP.cc

+267
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,267 @@
1+
/*
2+
* (C) Copyright 2025- ECMWF.
3+
*
4+
* This software is licensed under the terms of the Apache Licence Version 2.0
5+
* which can be obtained at http://www.apache.org/licenses/LICENSE-2.0.
6+
* In applying this licence, ECMWF does not waive the privileges and immunities
7+
* granted to it by virtue of its status as an intergovernmental organisation
8+
* nor does it submit to any jurisdiction.
9+
*/
10+
11+
12+
#include "eckit/linalg/sparse/LinearAlgebraHIP.h"
13+
14+
#include <ostream>
15+
16+
#include "eckit/exception/Exceptions.h"
17+
#include "eckit/linalg/Matrix.h"
18+
#include "eckit/linalg/SparseMatrix.h"
19+
#include "eckit/linalg/Vector.h"
20+
#include "eckit/linalg/detail/HIP.h"
21+
#include "eckit/linalg/sparse/LinearAlgebraGeneric.h"
22+
23+
24+
namespace eckit {
25+
namespace linalg {
26+
namespace sparse {
27+
28+
29+
static const LinearAlgebraHIP __la("hip");
30+
31+
32+
void LinearAlgebraHIP::print(std::ostream& out) const {
33+
out << "LinearAlgebraHIP[]";
34+
}
35+
36+
37+
void LinearAlgebraHIP::spmv(const SparseMatrix& A, const Vector& x, Vector& y) const {
38+
ASSERT(x.size() == A.cols() && y.size() == A.rows());
39+
// We expect indices to be 0-based
40+
ASSERT(A.outer()[0] == 0);
41+
const Size sizeArowptr = (A.rows() + 1) * sizeof(Index);
42+
const Size sizeAcolidx = A.nonZeros() * sizeof(Index);
43+
const Size sizeAvalues = A.nonZeros() * sizeof(Scalar);
44+
const Size sizex = A.cols() * sizeof(Scalar);
45+
const Size sizey = A.rows() * sizeof(Scalar);
46+
47+
Index* d_A_rowptr; ///< device memory matrix A row pointers
48+
Index* d_A_colidx; ///< device memory matrix A col indices
49+
Scalar* d_A_values; ///< device memory matrix A values
50+
Scalar* d_x; ///< device memory vector x
51+
Scalar* d_y; ///< device memory vector y
52+
53+
CALL_HIP(hipMalloc((void**)&d_A_rowptr, sizeArowptr));
54+
CALL_HIP(hipMalloc((void**)&d_A_colidx, sizeAcolidx));
55+
CALL_HIP(hipMalloc((void**)&d_A_values, sizeAvalues));
56+
CALL_HIP(hipMalloc((void**)&d_x, sizex));
57+
CALL_HIP(hipMalloc((void**)&d_y, sizey));
58+
59+
CALL_HIP(hipMemcpy(d_A_rowptr, A.outer(), sizeArowptr, hipMemcpyHostToDevice));
60+
CALL_HIP(hipMemcpy(d_A_colidx, A.inner(), sizeAcolidx, hipMemcpyHostToDevice));
61+
CALL_HIP(hipMemcpy(d_A_values, A.data(), sizeAvalues, hipMemcpyHostToDevice));
62+
CALL_HIP(hipMemcpy(d_x, x.data(), sizex, hipMemcpyHostToDevice));
63+
64+
hipsparseHandle_t handle;
65+
CALL_HIPSPARSE(hipsparseCreate(&handle));
66+
67+
hipsparseSpMatDescr_t matA;
68+
CALL_HIPSPARSE( hipsparseCreateCsr(
69+
&matA,
70+
A.rows(), A.cols(), A.nonZeros(),
71+
d_A_rowptr,
72+
d_A_colidx,
73+
d_A_values,
74+
HIPSPARSE_INDEX_32I,
75+
HIPSPARSE_INDEX_32I,
76+
HIPSPARSE_INDEX_BASE_ZERO,
77+
HIP_R_64F) );
78+
79+
hipsparseDnVecDescr_t vecX;
80+
CALL_HIPSPARSE( hipsparseCreateDnVec(
81+
&vecX,
82+
x.size(),
83+
d_x,
84+
HIP_R_64F) );
85+
86+
hipsparseDnVecDescr_t vecY;
87+
CALL_HIPSPARSE( hipsparseCreateDnVec(
88+
&vecY,
89+
y.size(),
90+
d_y,
91+
HIP_R_64F) );
92+
93+
const Scalar alpha = 1.0;
94+
const Scalar beta = 0.0;
95+
96+
// Determine buffer size
97+
size_t bufferSize = 0;
98+
CALL_HIPSPARSE( hipsparseSpMV_bufferSize(
99+
handle,
100+
HIPSPARSE_OPERATION_NON_TRANSPOSE,
101+
&alpha,
102+
matA,
103+
vecX,
104+
&beta,
105+
vecY,
106+
HIP_R_64F,
107+
HIPSPARSE_SPMV_ALG_DEFAULT,
108+
&bufferSize) );
109+
110+
// Allocate buffer
111+
char* buffer;
112+
CALL_HIP( hipMalloc(&buffer, bufferSize) );
113+
114+
// Perform SpMV
115+
// y = alpha * A * x + beta * y
116+
CALL_HIPSPARSE( hipsparseSpMV(
117+
handle,
118+
HIPSPARSE_OPERATION_NON_TRANSPOSE,
119+
&alpha,
120+
matA,
121+
vecX,
122+
&beta,
123+
vecY,
124+
HIP_R_64F,
125+
HIPSPARSE_SPMV_ALG_DEFAULT,
126+
buffer) );
127+
128+
// Copy result back to host
129+
CALL_HIP(hipMemcpy(y.data(), d_y, sizey, hipMemcpyDeviceToHost));
130+
131+
CALL_HIPSPARSE( hipsparseDestroyDnVec(vecY) );
132+
CALL_HIPSPARSE( hipsparseDestroyDnVec(vecX) );
133+
CALL_HIPSPARSE( hipsparseDestroySpMat(matA) );
134+
CALL_HIPSPARSE( hipsparseDestroy(handle) );
135+
136+
137+
CALL_HIP(hipFree(d_A_rowptr));
138+
CALL_HIP(hipFree(d_A_colidx));
139+
CALL_HIP(hipFree(d_A_values));
140+
CALL_HIP(hipFree(d_x));
141+
CALL_HIP(hipFree(d_y));
142+
}
143+
144+
145+
void LinearAlgebraHIP::spmm(const SparseMatrix& A, const Matrix& B, Matrix& C) const {
146+
ASSERT(A.cols() == B.rows() && A.rows() == C.rows() && B.cols() == C.cols());
147+
// We expect indices to be 0-based
148+
ASSERT(A.outer()[0] == 0);
149+
const Size sizeArowptr = (A.rows() + 1) * sizeof(Index);
150+
const Size sizeAcolidx = A.nonZeros() * sizeof(Index);
151+
const Size sizeAvalues = A.nonZeros() * sizeof(Scalar);
152+
const Size sizeB = B.rows() * B.cols() * sizeof(Scalar);
153+
const Size sizeC = A.rows() * B.cols() * sizeof(Scalar);
154+
155+
Index* d_A_rowptr; ///< device memory matrix A row pointers
156+
Index* d_A_colidx; ///< device memory matrix A col indices
157+
Scalar* d_A_values; ///< device memory matrix A values
158+
Scalar* d_B; ///< device memory matrix B
159+
Scalar* d_C; ///< device memory matrix C
160+
161+
CALL_HIP(hipMalloc((void**)&d_A_rowptr, sizeArowptr));
162+
CALL_HIP(hipMalloc((void**)&d_A_colidx, sizeAcolidx));
163+
CALL_HIP(hipMalloc((void**)&d_A_values, sizeAvalues));
164+
CALL_HIP(hipMalloc((void**)&d_B, sizeB));
165+
CALL_HIP(hipMalloc((void**)&d_C, sizeC));
166+
167+
CALL_HIP(hipMemcpy(d_A_rowptr, A.outer(), sizeArowptr, hipMemcpyHostToDevice));
168+
CALL_HIP(hipMemcpy(d_A_colidx, A.inner(), sizeAcolidx, hipMemcpyHostToDevice));
169+
CALL_HIP(hipMemcpy(d_A_values, A.data(), sizeAvalues, hipMemcpyHostToDevice));
170+
CALL_HIP(hipMemcpy(d_B, B.data(), sizeB, hipMemcpyHostToDevice));
171+
172+
hipsparseHandle_t handle;
173+
CALL_HIPSPARSE(hipsparseCreate(&handle));
174+
175+
hipsparseSpMatDescr_t matA;
176+
CALL_HIPSPARSE( hipsparseCreateCsr(
177+
&matA,
178+
A.rows(), A.cols(), A.nonZeros(),
179+
d_A_rowptr,
180+
d_A_colidx,
181+
d_A_values,
182+
HIPSPARSE_INDEX_32I,
183+
HIPSPARSE_INDEX_32I,
184+
HIPSPARSE_INDEX_BASE_ZERO,
185+
HIP_R_64F) );
186+
187+
// Create dense matrix descriptors
188+
hipsparseDnMatDescr_t matB;
189+
CALL_HIPSPARSE(hipsparseCreateDnMat(
190+
&matB,
191+
B.rows(), // rows
192+
B.cols(), // cols
193+
B.rows(), // leading dimension
194+
d_B,
195+
HIP_R_64F,
196+
HIPSPARSE_ORDER_COL) );
197+
198+
hipsparseDnMatDescr_t matC;
199+
CALL_HIPSPARSE(hipsparseCreateDnMat(
200+
&matC,
201+
C.rows(), // rows
202+
C.cols(), // cols
203+
C.rows(), // leading dimension
204+
d_C,
205+
HIP_R_64F,
206+
HIPSPARSE_ORDER_COL) );
207+
208+
const Scalar alpha = 1.0;
209+
const Scalar beta = 0.0;
210+
211+
size_t bufferSize = 0;
212+
CALL_HIPSPARSE(hipsparseSpMM_bufferSize(
213+
handle,
214+
HIPSPARSE_OPERATION_NON_TRANSPOSE,
215+
HIPSPARSE_OPERATION_NON_TRANSPOSE,
216+
&alpha,
217+
matA,
218+
matB,
219+
&beta,
220+
matC,
221+
HIP_R_64F,
222+
HIPSPARSE_SPMM_ALG_DEFAULT,
223+
&bufferSize));
224+
225+
// Allocate buffer
226+
char* buffer;
227+
CALL_HIP(hipMalloc(&buffer, bufferSize));
228+
229+
// Perform SpMM
230+
CALL_HIPSPARSE(hipsparseSpMM(
231+
handle,
232+
HIPSPARSE_OPERATION_NON_TRANSPOSE,
233+
HIPSPARSE_OPERATION_NON_TRANSPOSE,
234+
&alpha,
235+
matA,
236+
matB,
237+
&beta,
238+
matC,
239+
HIP_R_64F,
240+
HIPSPARSE_SPMM_ALG_DEFAULT,
241+
buffer));
242+
243+
CALL_HIP(hipMemcpy(C.data(), d_C, sizeC, hipMemcpyDeviceToHost));
244+
245+
CALL_HIPSPARSE(hipsparseDestroy(handle));
246+
CALL_HIPSPARSE(hipsparseDestroyDnMat(matC));
247+
CALL_HIPSPARSE(hipsparseDestroyDnMat(matB));
248+
CALL_HIPSPARSE(hipsparseDestroySpMat(matA));
249+
250+
CALL_HIP(hipFree(buffer));
251+
CALL_HIP(hipFree(d_A_rowptr));
252+
CALL_HIP(hipFree(d_A_colidx));
253+
CALL_HIP(hipFree(d_A_values));
254+
CALL_HIP(hipFree(d_B));
255+
CALL_HIP(hipFree(d_C));
256+
}
257+
258+
259+
void LinearAlgebraHIP::dsptd(const Vector& x, const SparseMatrix& A, const Vector& y, SparseMatrix& B) const {
260+
static const sparse::LinearAlgebraGeneric generic;
261+
generic.dsptd(x, A, y, B);
262+
}
263+
264+
265+
} // namespace sparse
266+
} // namespace linalg
267+
} // namespace eckit

‎src/eckit/linalg/sparse/LinearAlgebraHIP.h

+36
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,36 @@
1+
/*
2+
* (C) Copyright 2025- ECMWF.
3+
*
4+
* This software is licensed under the terms of the Apache Licence Version 2.0
5+
* which can be obtained at http://www.apache.org/licenses/LICENSE-2.0.
6+
* In applying this licence, ECMWF does not waive the privileges and immunities
7+
* granted to it by virtue of its status as an intergovernmental organisation
8+
* nor does it submit to any jurisdiction.
9+
*/
10+
11+
12+
#pragma once
13+
14+
#include "eckit/linalg/LinearAlgebraSparse.h"
15+
16+
17+
namespace eckit {
18+
namespace linalg {
19+
namespace sparse {
20+
21+
22+
struct LinearAlgebraHIP final : public LinearAlgebraSparse {
23+
LinearAlgebraHIP() {}
24+
LinearAlgebraHIP(const std::string& name) :
25+
LinearAlgebraSparse(name) {}
26+
27+
void spmv(const SparseMatrix&, const Vector&, Vector&) const override;
28+
void spmm(const SparseMatrix&, const Matrix&, Matrix&) const override;
29+
void dsptd(const Vector&, const SparseMatrix&, const Vector&, SparseMatrix&) const override;
30+
void print(std::ostream&) const override;
31+
};
32+
33+
34+
} // namespace sparse
35+
} // namespace linalg
36+
} // namespace eckit

‎tests/linalg/CMakeLists.txt

+5
Original file line numberDiff line numberDiff line change
@@ -134,6 +134,11 @@ ecbuild_add_test( TARGET eckit_test_linalg_sparse_backend_cuda
134134
CONDITION eckit_HAVE_CUDA
135135
ARGS --log_level=message -linearAlgebraSparseBackend cuda --catch_system_errors=no )
136136

137+
ecbuild_add_test( TARGET eckit_test_linalg_sparse_backend_hip
138+
COMMAND eckit_test_linalg_sparse_backend
139+
CONDITION eckit_HAVE_HIP
140+
ARGS --log_level=message -linearAlgebraSparseBackend hip )
141+
137142
ecbuild_add_test( TARGET eckit_test_linalg_sparse_backend_eigen
138143
COMMAND eckit_test_linalg_sparse_backend
139144
CONDITION eckit_HAVE_EIGEN

0 commit comments

Comments
 (0)