Skip to content

Commit 083c367

Browse files
yhmtsaiyhmtsai
committed
add rocsparse spgemm symbolic/numeric test
1 parent a1dc1ba commit 083c367

File tree

2 files changed

+326
-0
lines changed

2 files changed

+326
-0
lines changed

‎test/gtest/CMakeLists.txt‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -15,10 +15,12 @@ elseif(ENABLE_ROCSPARSE)
1515
# general one is handle in another pr
1616
set_source_files_properties(rocsparse/spmv_test.cpp PROPERTIES LANGUAGE HIP)
1717
set_source_files_properties(rocsparse/spgemm_test.cpp PROPERTIES LANGUAGE HIP)
18+
set_source_files_properties(rocsparse/spgemm_reuse_test.cpp PROPERTIES LANGUAGE HIP)
1819
set_source_files_properties(rocsparse/spgemm_4args_test.cpp PROPERTIES LANGUAGE HIP)
1920
add_executable(spblas-tests
2021
rocsparse/spmv_test.cpp
2122
rocsparse/spgemm_test.cpp
23+
rocsparse/spgemm_reuse_test.cpp
2224
rocsparse/spgemm_4args_test.cpp)
2325
endif()
2426

‎test/gtest/rocsparse/spgemm_reuse_test.cpp‎

Lines changed: 324 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,324 @@
1+
2+
#include "../util.hpp"
3+
#include <spblas/spblas.hpp>
4+
5+
#include <gtest/gtest.h>
6+
#include <thrust/device_vector.h>
7+
8+
using value_t = float;
9+
using index_t = spblas::index_t;
10+
using offset_t = spblas::offset_t;
11+
12+
TEST(CsrView, SpGEMMReuse) {
13+
for (auto&& [m, k, nnz] : util::dims) {
14+
for (auto&& n : {m, k}) {
15+
auto [a_values, a_rowptr, a_colind, a_shape, a_nnz] =
16+
spblas::generate_csr<value_t, index_t, offset_t>(m, k, nnz);
17+
thrust::device_vector<value_t> d_a_values(a_values);
18+
thrust::device_vector<offset_t> d_a_rowptr(a_rowptr);
19+
thrust::device_vector<index_t> d_a_colind(a_colind);
20+
spblas::csr_view<value_t, index_t, offset_t> d_a(
21+
d_a_values.data().get(), d_a_rowptr.data().get(),
22+
d_a_colind.data().get(), a_shape, a_nnz);
23+
spblas::csr_view<value_t, index_t, offset_t> a(a_values, a_rowptr,
24+
a_colind, a_shape, a_nnz);
25+
26+
auto [b_values, b_rowptr, b_colind, b_shape, b_nnz] =
27+
spblas::generate_csr<value_t, index_t, offset_t>(k, n, nnz);
28+
thrust::device_vector<value_t> d_b_values(b_values);
29+
thrust::device_vector<offset_t> d_b_rowptr(b_rowptr);
30+
thrust::device_vector<index_t> d_b_colind(b_colind);
31+
spblas::csr_view<value_t, index_t, offset_t> d_b(
32+
d_b_values.data().get(), d_b_rowptr.data().get(),
33+
d_b_colind.data().get(), b_shape, b_nnz);
34+
spblas::csr_view<value_t, index_t, offset_t> b(b_values, b_rowptr,
35+
b_colind, b_shape, b_nnz);
36+
37+
thrust::device_vector<offset_t> d_c_rowptr(m + 1);
38+
39+
spblas::csr_view<value_t, index_t, offset_t> d_c(
40+
nullptr, d_c_rowptr.data().get(), nullptr, {m, n}, 0);
41+
42+
spblas::spgemm_state_t state;
43+
spblas::multiply_symbolic_compute(state, d_a, d_b, d_c);
44+
auto nnz = state.result_nnz();
45+
thrust::device_vector<value_t> d_c_values(nnz);
46+
thrust::device_vector<index_t> d_c_colind(nnz);
47+
std::span<value_t> d_c_values_span(d_c_values.data().get(), nnz);
48+
std::span<offset_t> d_c_rowptr_span(d_c_rowptr.data().get(), m + 1);
49+
std::span<index_t> d_c_colind_span(d_c_colind.data().get(), nnz);
50+
d_c.update(d_c_values_span, d_c_rowptr_span, d_c_colind_span, {m, n},
51+
nnz);
52+
53+
spblas::multiply_symbolic_fill(state, d_a, d_b, d_c);
54+
std::mt19937 g(0);
55+
for (int i = 0; i < 3; i++) {
56+
// we can change the value of a and b but only need to call
57+
// multiply_numeric answer here.
58+
if (i != 0) {
59+
// regenerate value of a and b;
60+
std::uniform_real_distribution val_dist(0.0, 100.0);
61+
for (auto& v : a_values) {
62+
v = val_dist(g);
63+
}
64+
for (auto& v : b_values) {
65+
v = val_dist(g);
66+
}
67+
thrust::copy(a_values.begin(), a_values.end(), d_a_values.begin());
68+
thrust::copy(b_values.begin(), b_values.end(), d_b_values.begin());
69+
}
70+
spblas::multiply_numeric(state, d_a, d_b, d_c);
71+
std::vector<value_t> c_values(nnz);
72+
std::vector<offset_t> c_rowptr(m + 1);
73+
std::vector<index_t> c_colind(nnz);
74+
thrust::copy(d_c_values.begin(), d_c_values.end(), c_values.begin());
75+
thrust::copy(d_c_rowptr.begin(), d_c_rowptr.end(), c_rowptr.begin());
76+
thrust::copy(d_c_colind.begin(), d_c_colind.end(), c_colind.begin());
77+
spblas::csr_view<value_t, index_t, offset_t> c(c_values, c_rowptr,
78+
c_colind, {m, n}, nnz);
79+
80+
spblas::__backend::spa_accumulator<value_t, index_t> c_row_ref(
81+
spblas::__backend::shape(c)[1]);
82+
83+
spblas::__backend::spa_accumulator<value_t, index_t> c_row_acc(
84+
spblas::__backend::shape(c)[1]);
85+
86+
for (auto&& [i, a_row] : spblas::__backend::rows(a)) {
87+
c_row_ref.clear();
88+
for (auto&& [k, a_v] : a_row) {
89+
auto&& b_row = spblas::__backend::lookup_row(b, k);
90+
91+
for (auto&& [j, b_v] : b_row) {
92+
c_row_ref[j] += a_v * b_v;
93+
}
94+
}
95+
96+
auto&& c_row = spblas::__backend::lookup_row(c, i);
97+
98+
// Accumulate output into `c_row_acc` so that we can allow
99+
// duplicate column indices.
100+
c_row_acc.clear();
101+
for (auto&& [j, c_v] : c_row) {
102+
c_row_acc[j] += c_v;
103+
}
104+
105+
for (auto&& [j, c_v] : c_row) {
106+
EXPECT_EQ_(c_row_ref[j], c_row_acc[j]);
107+
}
108+
109+
EXPECT_EQ(c_row_ref.size(), c_row_acc.size());
110+
}
111+
}
112+
}
113+
}
114+
}
115+
116+
TEST(CsrView, SpGEMMReuse_AScaled) {
117+
value_t alpha = 2.0f;
118+
for (auto&& [m, k, nnz] : util::dims) {
119+
for (auto&& n : {m, k}) {
120+
auto [a_values, a_rowptr, a_colind, a_shape, a_nnz] =
121+
spblas::generate_csr<value_t, index_t, offset_t>(m, k, nnz);
122+
thrust::device_vector<value_t> d_a_values(a_values);
123+
thrust::device_vector<offset_t> d_a_rowptr(a_rowptr);
124+
thrust::device_vector<index_t> d_a_colind(a_colind);
125+
spblas::csr_view<value_t, index_t, offset_t> d_a(
126+
d_a_values.data().get(), d_a_rowptr.data().get(),
127+
d_a_colind.data().get(), a_shape, a_nnz);
128+
spblas::csr_view<value_t, index_t, offset_t> a(a_values, a_rowptr,
129+
a_colind, a_shape, a_nnz);
130+
131+
auto [b_values, b_rowptr, b_colind, b_shape, b_nnz] =
132+
spblas::generate_csr<value_t, index_t, offset_t>(k, n, nnz);
133+
thrust::device_vector<value_t> d_b_values(b_values);
134+
thrust::device_vector<offset_t> d_b_rowptr(b_rowptr);
135+
thrust::device_vector<index_t> d_b_colind(b_colind);
136+
spblas::csr_view<value_t, index_t, offset_t> d_b(
137+
d_b_values.data().get(), d_b_rowptr.data().get(),
138+
d_b_colind.data().get(), b_shape, b_nnz);
139+
spblas::csr_view<value_t, index_t, offset_t> b(b_values, b_rowptr,
140+
b_colind, b_shape, b_nnz);
141+
142+
thrust::device_vector<offset_t> d_c_rowptr(m + 1);
143+
144+
spblas::csr_view<value_t, index_t, offset_t> d_c(
145+
nullptr, d_c_rowptr.data().get(), nullptr, {m, n}, 0);
146+
147+
spblas::spgemm_state_t state;
148+
spblas::multiply_symbolic_compute(state, scaled(alpha, d_a), d_b, d_c);
149+
auto nnz = state.result_nnz();
150+
thrust::device_vector<value_t> d_c_values(nnz);
151+
thrust::device_vector<index_t> d_c_colind(nnz);
152+
std::span<value_t> d_c_values_span(d_c_values.data().get(), nnz);
153+
std::span<offset_t> d_c_rowptr_span(d_c_rowptr.data().get(), m + 1);
154+
std::span<index_t> d_c_colind_span(d_c_colind.data().get(), nnz);
155+
d_c.update(d_c_values_span, d_c_rowptr_span, d_c_colind_span, {m, n},
156+
nnz);
157+
158+
spblas::multiply_symbolic_fill(state, scaled(alpha, d_a), d_b, d_c);
159+
std::mt19937 g(0);
160+
for (int i = 0; i < 3; i++) {
161+
// we can change the value of a and b but only need to call
162+
// multiply_numeric answer here.
163+
if (i != 0) {
164+
// regenerate value of a and b;
165+
std::uniform_real_distribution val_dist(0.0, 100.0);
166+
for (auto& v : a_values) {
167+
v = val_dist(g);
168+
}
169+
for (auto& v : b_values) {
170+
v = val_dist(g);
171+
}
172+
thrust::copy(a_values.begin(), a_values.end(), d_a_values.begin());
173+
thrust::copy(b_values.begin(), b_values.end(), d_b_values.begin());
174+
}
175+
spblas::multiply_numeric(state, scaled(alpha, d_a), d_b, d_c);
176+
std::vector<value_t> c_values(nnz);
177+
std::vector<offset_t> c_rowptr(m + 1);
178+
std::vector<index_t> c_colind(nnz);
179+
thrust::copy(d_c_values.begin(), d_c_values.end(), c_values.begin());
180+
thrust::copy(d_c_rowptr.begin(), d_c_rowptr.end(), c_rowptr.begin());
181+
thrust::copy(d_c_colind.begin(), d_c_colind.end(), c_colind.begin());
182+
spblas::csr_view<value_t, index_t, offset_t> c(c_values, c_rowptr,
183+
c_colind, {m, n}, nnz);
184+
185+
spblas::__backend::spa_accumulator<value_t, index_t> c_row_ref(
186+
spblas::__backend::shape(c)[1]);
187+
188+
spblas::__backend::spa_accumulator<value_t, index_t> c_row_acc(
189+
spblas::__backend::shape(c)[1]);
190+
191+
for (auto&& [i, a_row] : spblas::__backend::rows(a)) {
192+
c_row_ref.clear();
193+
for (auto&& [k, a_v] : a_row) {
194+
auto&& b_row = spblas::__backend::lookup_row(b, k);
195+
196+
for (auto&& [j, b_v] : b_row) {
197+
c_row_ref[j] += alpha * a_v * b_v;
198+
}
199+
}
200+
201+
auto&& c_row = spblas::__backend::lookup_row(c, i);
202+
203+
// Accumulate output into `c_row_acc` so that we can allow
204+
// duplicate column indices.
205+
c_row_acc.clear();
206+
for (auto&& [j, c_v] : c_row) {
207+
c_row_acc[j] += c_v;
208+
}
209+
210+
for (auto&& [j, c_v] : c_row) {
211+
EXPECT_EQ_(c_row_ref[j], c_row_acc[j]);
212+
}
213+
214+
EXPECT_EQ(c_row_ref.size(), c_row_acc.size());
215+
}
216+
}
217+
}
218+
}
219+
}
220+
221+
TEST(CsrView, SpGEMMReuse_BScaled) {
222+
value_t alpha = 2.0f;
223+
for (auto&& [m, k, nnz] : util::dims) {
224+
for (auto&& n : {m, k}) {
225+
auto [a_values, a_rowptr, a_colind, a_shape, a_nnz] =
226+
spblas::generate_csr<value_t, index_t, offset_t>(m, k, nnz);
227+
thrust::device_vector<value_t> d_a_values(a_values);
228+
thrust::device_vector<offset_t> d_a_rowptr(a_rowptr);
229+
thrust::device_vector<index_t> d_a_colind(a_colind);
230+
spblas::csr_view<value_t, index_t, offset_t> d_a(
231+
d_a_values.data().get(), d_a_rowptr.data().get(),
232+
d_a_colind.data().get(), a_shape, a_nnz);
233+
spblas::csr_view<value_t, index_t, offset_t> a(a_values, a_rowptr,
234+
a_colind, a_shape, a_nnz);
235+
236+
auto [b_values, b_rowptr, b_colind, b_shape, b_nnz] =
237+
spblas::generate_csr<value_t, index_t, offset_t>(k, n, nnz);
238+
thrust::device_vector<value_t> d_b_values(b_values);
239+
thrust::device_vector<offset_t> d_b_rowptr(b_rowptr);
240+
thrust::device_vector<index_t> d_b_colind(b_colind);
241+
spblas::csr_view<value_t, index_t, offset_t> d_b(
242+
d_b_values.data().get(), d_b_rowptr.data().get(),
243+
d_b_colind.data().get(), b_shape, b_nnz);
244+
spblas::csr_view<value_t, index_t, offset_t> b(b_values, b_rowptr,
245+
b_colind, b_shape, b_nnz);
246+
247+
thrust::device_vector<offset_t> d_c_rowptr(m + 1);
248+
249+
spblas::csr_view<value_t, index_t, offset_t> d_c(
250+
nullptr, d_c_rowptr.data().get(), nullptr, {m, n}, 0);
251+
252+
spblas::spgemm_state_t state;
253+
spblas::multiply_symbolic_compute(state, d_a, scaled(alpha, d_b), d_c);
254+
auto nnz = state.result_nnz();
255+
thrust::device_vector<value_t> d_c_values(nnz);
256+
thrust::device_vector<index_t> d_c_colind(nnz);
257+
std::span<value_t> d_c_values_span(d_c_values.data().get(), nnz);
258+
std::span<offset_t> d_c_rowptr_span(d_c_rowptr.data().get(), m + 1);
259+
std::span<index_t> d_c_colind_span(d_c_colind.data().get(), nnz);
260+
d_c.update(d_c_values_span, d_c_rowptr_span, d_c_colind_span, {m, n},
261+
nnz);
262+
263+
spblas::multiply_symbolic_fill(state, d_a, scaled(alpha, d_b), d_c);
264+
std::mt19937 g(0);
265+
for (int i = 0; i < 3; i++) {
266+
// we can change the value of a and b but only need to call
267+
// multiply_numeric answer here.
268+
if (i != 0) {
269+
// regenerate value of a and b;
270+
std::uniform_real_distribution val_dist(0.0, 100.0);
271+
for (auto& v : a_values) {
272+
v = val_dist(g);
273+
}
274+
for (auto& v : b_values) {
275+
v = val_dist(g);
276+
}
277+
thrust::copy(a_values.begin(), a_values.end(), d_a_values.begin());
278+
thrust::copy(b_values.begin(), b_values.end(), d_b_values.begin());
279+
}
280+
spblas::multiply_numeric(state, d_a, scaled(alpha, d_b), d_c);
281+
std::vector<value_t> c_values(nnz);
282+
std::vector<offset_t> c_rowptr(m + 1);
283+
std::vector<index_t> c_colind(nnz);
284+
thrust::copy(d_c_values.begin(), d_c_values.end(), c_values.begin());
285+
thrust::copy(d_c_rowptr.begin(), d_c_rowptr.end(), c_rowptr.begin());
286+
thrust::copy(d_c_colind.begin(), d_c_colind.end(), c_colind.begin());
287+
spblas::csr_view<value_t, index_t, offset_t> c(c_values, c_rowptr,
288+
c_colind, {m, n}, nnz);
289+
290+
spblas::__backend::spa_accumulator<value_t, index_t> c_row_ref(
291+
spblas::__backend::shape(c)[1]);
292+
293+
spblas::__backend::spa_accumulator<value_t, index_t> c_row_acc(
294+
spblas::__backend::shape(c)[1]);
295+
296+
for (auto&& [i, a_row] : spblas::__backend::rows(a)) {
297+
c_row_ref.clear();
298+
for (auto&& [k, a_v] : a_row) {
299+
auto&& b_row = spblas::__backend::lookup_row(b, k);
300+
301+
for (auto&& [j, b_v] : b_row) {
302+
c_row_ref[j] += alpha * a_v * b_v;
303+
}
304+
}
305+
306+
auto&& c_row = spblas::__backend::lookup_row(c, i);
307+
308+
// Accumulate output into `c_row_acc` so that we can allow
309+
// duplicate column indices.
310+
c_row_acc.clear();
311+
for (auto&& [j, c_v] : c_row) {
312+
c_row_acc[j] += c_v;
313+
}
314+
315+
for (auto&& [j, c_v] : c_row) {
316+
EXPECT_EQ_(c_row_ref[j], c_row_acc[j]);
317+
}
318+
319+
EXPECT_EQ(c_row_ref.size(), c_row_acc.size());
320+
}
321+
}
322+
}
323+
}
324+
}

0 commit comments

Comments
 (0)