Skip to content

Commit 195147f

Browse files
yhmtsaiyhmtsai
committed
use cuda2hip.sh to get hip files
1 parent b2863d5 commit 195147f

4 files changed

Lines changed: 334 additions & 8 deletions

File tree

‎hip/factorization/par_ilu_kernels.hip.cpp‎

Lines changed: 94 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -33,10 +33,19 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
3333
#include "core/factorization/par_ilu_kernels.hpp"
3434

3535

36-
#include <ginkgo/core/base/exception_helpers.hpp>
36+
#include <hip/hip_runtime.h>
37+
38+
39+
#include <ginkgo/core/base/array.hpp>
40+
#include <ginkgo/core/base/std_extensions.hpp>
3741
#include <ginkgo/core/matrix/coo.hpp>
3842

3943

44+
#include "hip/base/math.hip.hpp"
45+
#include "hip/base/types.hip.hpp"
46+
#include "hip/components/prefix_sum.hip.hpp"
47+
48+
4049
namespace gko {
4150
namespace kernels {
4251
namespace hip {
@@ -48,11 +57,50 @@ namespace hip {
4857
namespace par_ilu_factorization {
4958

5059

60+
constexpr int default_block_size{512};
61+
62+
63+
#include "common/factorization/par_ilu_kernels.hpp.inc"
64+
65+
5166
template <typename ValueType, typename IndexType>
5267
void initialize_row_ptrs_l_u(
5368
std::shared_ptr<const HipExecutor> exec,
5469
const matrix::Csr<ValueType, IndexType> *system_matrix,
55-
IndexType *l_row_ptrs, IndexType *u_row_ptrs) GKO_NOT_IMPLEMENTED;
70+
IndexType *l_row_ptrs, IndexType *u_row_ptrs)
71+
{
72+
const size_type num_rows{system_matrix->get_size()[0]};
73+
const size_type num_row_ptrs{num_rows + 1};
74+
75+
const dim3 block_size{default_block_size, 1, 1};
76+
const uint32 number_blocks =
77+
ceildiv(num_rows, static_cast<size_type>(block_size.x));
78+
const dim3 grid_dim{number_blocks, 1, 1};
79+
80+
hipLaunchKernelGGL(kernel::count_nnz_per_l_u_row, dim3(grid_dim),
81+
dim3(block_size), 0, 0, num_rows,
82+
as_hip_type(system_matrix->get_const_row_ptrs()),
83+
as_hip_type(system_matrix->get_const_col_idxs()),
84+
as_hip_type(system_matrix->get_const_values()),
85+
as_hip_type(l_row_ptrs), as_hip_type(u_row_ptrs));
86+
87+
Array<IndexType> block_sum(exec, grid_dim.x);
88+
auto block_sum_ptr = block_sum.get_data();
89+
90+
hipLaunchKernelGGL(HIP_KERNEL_NAME(start_prefix_sum<default_block_size>),
91+
dim3(grid_dim), dim3(block_size), 0, 0, num_row_ptrs,
92+
as_hip_type(l_row_ptrs), as_hip_type(block_sum_ptr));
93+
hipLaunchKernelGGL(HIP_KERNEL_NAME(finalize_prefix_sum<default_block_size>),
94+
dim3(grid_dim), dim3(block_size), 0, 0, num_row_ptrs,
95+
as_hip_type(l_row_ptrs), as_hip_type(block_sum_ptr));
96+
97+
hipLaunchKernelGGL(HIP_KERNEL_NAME(start_prefix_sum<default_block_size>),
98+
dim3(grid_dim), dim3(block_size), 0, 0, num_row_ptrs,
99+
as_hip_type(u_row_ptrs), as_hip_type(block_sum_ptr));
100+
hipLaunchKernelGGL(HIP_KERNEL_NAME(finalize_prefix_sum<default_block_size>),
101+
dim3(grid_dim), dim3(block_size), 0, 0, num_row_ptrs,
102+
as_hip_type(u_row_ptrs), as_hip_type(block_sum_ptr));
103+
}
56104

57105
GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(
58106
GKO_DECLARE_PAR_ILU_INITIALIZE_ROW_PTRS_L_U_KERNEL);
@@ -63,18 +111,56 @@ void initialize_l_u(std::shared_ptr<const HipExecutor> exec,
63111
const matrix::Csr<ValueType, IndexType> *system_matrix,
64112
matrix::Csr<ValueType, IndexType> *csr_l,
65113
matrix::Csr<ValueType, IndexType> *csr_u)
66-
GKO_NOT_IMPLEMENTED;
114+
{
115+
const size_type num_rows{system_matrix->get_size()[0]};
116+
const dim3 block_size{default_block_size, 1, 1};
117+
const dim3 grid_dim{static_cast<uint32>(ceildiv(
118+
num_rows, static_cast<size_type>(block_size.x))),
119+
1, 1};
120+
121+
hipLaunchKernelGGL(
122+
kernel::initialize_l_u, dim3(grid_dim), dim3(block_size), 0, 0,
123+
num_rows, as_hip_type(system_matrix->get_const_row_ptrs()),
124+
as_hip_type(system_matrix->get_const_col_idxs()),
125+
as_hip_type(system_matrix->get_const_values()),
126+
as_hip_type(csr_l->get_const_row_ptrs()),
127+
as_hip_type(csr_l->get_col_idxs()), as_hip_type(csr_l->get_values()),
128+
as_hip_type(csr_u->get_const_row_ptrs()),
129+
as_hip_type(csr_u->get_col_idxs()), as_hip_type(csr_u->get_values()));
130+
}
67131

68132
GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(
69133
GKO_DECLARE_PAR_ILU_INITIALIZE_L_U_KERNEL);
70134

71135

72136
template <typename ValueType, typename IndexType>
73-
void compute_l_u_factors(
74-
std::shared_ptr<const HipExecutor> exec, size_type iterations,
75-
const matrix::Coo<ValueType, IndexType> *system_matrix,
76-
matrix::Csr<ValueType, IndexType> *l_factor,
77-
matrix::Csr<ValueType, IndexType> *u_factor) GKO_NOT_IMPLEMENTED;
137+
void compute_l_u_factors(std::shared_ptr<const HipExecutor> exec,
138+
size_type iterations,
139+
const matrix::Coo<ValueType, IndexType> *system_matrix,
140+
matrix::Csr<ValueType, IndexType> *l_factor,
141+
matrix::Csr<ValueType, IndexType> *u_factor)
142+
{
143+
iterations = (iterations == 0) ? 10 : iterations;
144+
const auto num_elements = system_matrix->get_num_stored_elements();
145+
const dim3 block_size{default_block_size, 1, 1};
146+
const dim3 grid_dim{
147+
static_cast<uint32>(
148+
ceildiv(num_elements, static_cast<size_type>(block_size.x))),
149+
1, 1};
150+
for (size_type i = 0; i < iterations; ++i) {
151+
hipLaunchKernelGGL(kernel::compute_l_u_factors, dim3(grid_dim),
152+
dim3(block_size), 0, 0, num_elements,
153+
as_hip_type(system_matrix->get_const_row_idxs()),
154+
as_hip_type(system_matrix->get_const_col_idxs()),
155+
as_hip_type(system_matrix->get_const_values()),
156+
as_hip_type(l_factor->get_const_row_ptrs()),
157+
as_hip_type(l_factor->get_const_col_idxs()),
158+
as_hip_type(l_factor->get_values()),
159+
as_hip_type(u_factor->get_const_row_ptrs()),
160+
as_hip_type(u_factor->get_const_col_idxs()),
161+
as_hip_type(u_factor->get_values()));
162+
}
163+
}
78164

79165
GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(
80166
GKO_DECLARE_PAR_ILU_COMPUTE_L_U_FACTORS_KERNEL);

‎hip/test/CMakeLists.txt‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,7 @@
11
include(${CMAKE_SOURCE_DIR}/cmake/create_test.cmake)
22

33
add_subdirectory(base)
4+
add_subdirectory(factorization)
45
add_subdirectory(matrix)
56
add_subdirectory(stop)
67
add_subdirectory(utils)

‎hip/test/factorization/CMakeLists.txt‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1 @@
1+
ginkgo_create_hip_test(par_ilu_kernels)

‎hip/test/factorization/par_ilu_kernels.hip.cpp‎

Lines changed: 238 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,238 @@
1+
/*******************************<GINKGO LICENSE>******************************
2+
Copyright (c) 2017-2019, the Ginkgo authors
3+
All rights reserved.
4+
5+
Redistribution and use in source and binary forms, with or without
6+
modification, are permitted provided that the following conditions
7+
are met:
8+
9+
1. Redistributions of source code must retain the above copyright
10+
notice, this list of conditions and the following disclaimer.
11+
12+
2. Redistributions in binary form must reproduce the above copyright
13+
notice, this list of conditions and the following disclaimer in the
14+
documentation and/or other materials provided with the distribution.
15+
16+
3. Neither the name of the copyright holder nor the names of its
17+
contributors may be used to endorse or promote products derived from
18+
this software without specific prior written permission.
19+
20+
THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS
21+
IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED
22+
TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A
23+
PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
24+
HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
25+
SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT
26+
LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
27+
DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY
28+
THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
29+
(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
30+
OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
31+
******************************<GINKGO LICENSE>*******************************/
32+
33+
#include "core/factorization/par_ilu_kernels.hpp"
34+
35+
36+
#include <algorithm>
37+
#include <fstream>
38+
#include <memory>
39+
#include <string>
40+
41+
42+
#include <gtest/gtest.h>
43+
44+
45+
#include <ginkgo/core/base/array.hpp>
46+
#include <ginkgo/core/base/executor.hpp>
47+
#include <ginkgo/core/matrix/coo.hpp>
48+
#include <ginkgo/core/matrix/csr.hpp>
49+
#include <ginkgo/core/matrix/dense.hpp>
50+
51+
52+
#include "core/test/utils.hpp"
53+
#include "matrices/config.hpp"
54+
55+
56+
namespace {
57+
58+
59+
class ParIlu : public ::testing::Test {
60+
protected:
61+
using value_type = gko::default_precision;
62+
using index_type = gko::int32;
63+
using Dense = gko::matrix::Dense<value_type>;
64+
using Coo = gko::matrix::Coo<value_type, index_type>;
65+
using Csr = gko::matrix::Csr<value_type, index_type>;
66+
67+
ParIlu()
68+
: ref(gko::ReferenceExecutor::create()),
69+
hip(gko::HipExecutor::create(0, ref)),
70+
csr_ref(nullptr),
71+
csr_hip(nullptr)
72+
{}
73+
74+
void SetUp() override
75+
{
76+
std::string file_name(gko::matrices::location_ani4_mtx);
77+
auto input_file = std::ifstream(file_name, std::ios::in);
78+
if (!input_file) {
79+
FAIL() << "Could not find the file \"" << file_name
80+
<< "\", which is required for this test.\n";
81+
}
82+
csr_ref = gko::read<Csr>(input_file, ref);
83+
auto csr_hip_temp = Csr::create(hip);
84+
csr_hip_temp->copy_from(gko::lend(csr_ref));
85+
csr_hip = gko::give(csr_hip_temp);
86+
}
87+
88+
std::shared_ptr<gko::ReferenceExecutor> ref;
89+
std::shared_ptr<gko::HipExecutor> hip;
90+
std::shared_ptr<const Csr> csr_ref;
91+
std::shared_ptr<const Csr> csr_hip;
92+
93+
void initialize_row_ptrs(index_type *l_row_ptrs_ref,
94+
index_type *u_row_ptrs_ref,
95+
index_type *l_row_ptrs_hip,
96+
index_type *u_row_ptrs_hip)
97+
{
98+
gko::kernels::reference::par_ilu_factorization::initialize_row_ptrs_l_u(
99+
ref, gko::lend(csr_ref), l_row_ptrs_ref, u_row_ptrs_ref);
100+
gko::kernels::hip::par_ilu_factorization::initialize_row_ptrs_l_u(
101+
hip, gko::lend(csr_hip), l_row_ptrs_hip, u_row_ptrs_hip);
102+
}
103+
104+
void initialize_lu(std::unique_ptr<Csr> *l_ref, std::unique_ptr<Csr> *u_ref,
105+
std::unique_ptr<Csr> *l_hip, std::unique_ptr<Csr> *u_hip)
106+
{
107+
auto num_row_ptrs = csr_ref->get_size()[0] + 1;
108+
gko::Array<index_type> l_row_ptrs_ref{ref, num_row_ptrs};
109+
gko::Array<index_type> u_row_ptrs_ref{ref, num_row_ptrs};
110+
gko::Array<index_type> l_row_ptrs_hip{hip, num_row_ptrs};
111+
gko::Array<index_type> u_row_ptrs_hip{hip, num_row_ptrs};
112+
113+
initialize_row_ptrs(
114+
l_row_ptrs_ref.get_data(), u_row_ptrs_ref.get_data(),
115+
l_row_ptrs_hip.get_data(), u_row_ptrs_hip.get_data());
116+
// Since `initialize_row_ptrs` was already tested, it is expected that
117+
// `*_ref` and `*_hip` contain identical values
118+
auto l_nnz = l_row_ptrs_ref.get_const_data()[num_row_ptrs - 1];
119+
auto u_nnz = u_row_ptrs_ref.get_const_data()[num_row_ptrs - 1];
120+
121+
*l_ref = Csr::create(ref, csr_ref->get_size(), l_nnz);
122+
*u_ref = Csr::create(ref, csr_ref->get_size(), u_nnz);
123+
*l_hip = Csr::create(hip, csr_hip->get_size(), l_nnz);
124+
*u_hip = Csr::create(hip, csr_hip->get_size(), u_nnz);
125+
// Copy the already initialized `row_ptrs` to the new matrices
126+
ref->copy_from(gko::lend(ref), num_row_ptrs, l_row_ptrs_ref.get_data(),
127+
(*l_ref)->get_row_ptrs());
128+
ref->copy_from(gko::lend(ref), num_row_ptrs, u_row_ptrs_ref.get_data(),
129+
(*u_ref)->get_row_ptrs());
130+
hip->copy_from(gko::lend(hip), num_row_ptrs, l_row_ptrs_hip.get_data(),
131+
(*l_hip)->get_row_ptrs());
132+
hip->copy_from(gko::lend(hip), num_row_ptrs, u_row_ptrs_hip.get_data(),
133+
(*u_hip)->get_row_ptrs());
134+
135+
gko::kernels::reference::par_ilu_factorization::initialize_l_u(
136+
ref, gko::lend(csr_ref), gko::lend(*l_ref), gko::lend(*u_ref));
137+
gko::kernels::hip::par_ilu_factorization::initialize_l_u(
138+
hip, gko::lend(csr_hip), gko::lend(*l_hip), gko::lend(*u_hip));
139+
}
140+
141+
template <typename ToType, typename FromType>
142+
static std::unique_ptr<ToType> static_unique_ptr_cast(
143+
std::unique_ptr<FromType> &&from)
144+
{
145+
return std::unique_ptr<ToType>{static_cast<ToType *>(from.release())};
146+
}
147+
148+
void compute_lu(std::unique_ptr<Csr> *l_ref, std::unique_ptr<Csr> *u_ref,
149+
std::unique_ptr<Csr> *l_hip, std::unique_ptr<Csr> *u_hip,
150+
gko::size_type iterations = 0)
151+
{
152+
auto coo_ref = Coo::create(ref);
153+
csr_ref->convert_to(gko::lend(coo_ref));
154+
auto coo_hip = Coo::create(hip);
155+
csr_hip->convert_to(gko::lend(coo_hip));
156+
initialize_lu(l_ref, u_ref, l_hip, u_hip);
157+
auto u_transpose_lin_op_ref = (*u_ref)->transpose();
158+
auto u_transpose_csr_ref =
159+
static_unique_ptr_cast<Csr>(std::move(u_transpose_lin_op_ref));
160+
auto u_transpose_lin_op_hip = (*u_hip)->transpose();
161+
auto u_transpose_csr_hip =
162+
static_unique_ptr_cast<Csr>(std::move(u_transpose_lin_op_hip));
163+
164+
gko::kernels::reference::par_ilu_factorization::compute_l_u_factors(
165+
ref, iterations, gko::lend(coo_ref), gko::lend(*l_ref),
166+
gko::lend(u_transpose_csr_ref));
167+
gko::kernels::hip::par_ilu_factorization::compute_l_u_factors(
168+
hip, iterations, gko::lend(coo_hip), gko::lend(*l_hip),
169+
gko::lend(u_transpose_csr_hip));
170+
auto u_lin_op_ref = u_transpose_csr_ref->transpose();
171+
*u_ref = static_unique_ptr_cast<Csr>(std::move(u_lin_op_ref));
172+
auto u_lin_op_hip = u_transpose_csr_hip->transpose();
173+
*u_hip = static_unique_ptr_cast<Csr>(std::move(u_lin_op_hip));
174+
}
175+
};
176+
177+
178+
TEST_F(ParIlu, KernelInitializeRowPtrsLUEquivalentToRef)
179+
{
180+
auto num_row_ptrs = csr_ref->get_size()[0] + 1;
181+
gko::Array<index_type> l_row_ptrs_array_ref(ref, num_row_ptrs);
182+
gko::Array<index_type> u_row_ptrs_array_ref(ref, num_row_ptrs);
183+
gko::Array<index_type> l_row_ptrs_array_hip(hip, num_row_ptrs);
184+
gko::Array<index_type> u_row_ptrs_array_hip(hip, num_row_ptrs);
185+
186+
initialize_row_ptrs(
187+
l_row_ptrs_array_ref.get_data(), u_row_ptrs_array_ref.get_data(),
188+
l_row_ptrs_array_hip.get_data(), u_row_ptrs_array_hip.get_data());
189+
190+
GKO_ASSERT_ARRAY_EQ(&l_row_ptrs_array_ref, &l_row_ptrs_array_hip);
191+
GKO_ASSERT_ARRAY_EQ(&u_row_ptrs_array_ref, &u_row_ptrs_array_hip);
192+
}
193+
194+
195+
TEST_F(ParIlu, KernelInitializeParILUIsEquivalentToRef)
196+
{
197+
std::unique_ptr<Csr> l_ref{};
198+
std::unique_ptr<Csr> u_ref{};
199+
std::unique_ptr<Csr> l_hip{};
200+
std::unique_ptr<Csr> u_hip{};
201+
202+
initialize_lu(&l_ref, &u_ref, &l_hip, &u_hip);
203+
204+
GKO_ASSERT_MTX_NEAR(l_ref, l_hip, 1e-14);
205+
GKO_ASSERT_MTX_NEAR(u_ref, u_hip, 1e-14);
206+
}
207+
208+
209+
TEST_F(ParIlu, KernelComputeParILUIsEquivalentToRef)
210+
{
211+
std::unique_ptr<Csr> l_ref{};
212+
std::unique_ptr<Csr> u_ref{};
213+
std::unique_ptr<Csr> l_hip{};
214+
std::unique_ptr<Csr> u_hip{};
215+
216+
compute_lu(&l_ref, &u_ref, &l_hip, &u_hip);
217+
218+
GKO_ASSERT_MTX_NEAR(l_ref, l_hip, 5e-2);
219+
GKO_ASSERT_MTX_NEAR(u_ref, u_hip, 5e-2);
220+
}
221+
222+
223+
TEST_F(ParIlu, KernelComputeParILUWithMoreIterationsIsEquivalentToRef)
224+
{
225+
std::unique_ptr<Csr> l_ref{};
226+
std::unique_ptr<Csr> u_ref{};
227+
std::unique_ptr<Csr> l_hip{};
228+
std::unique_ptr<Csr> u_hip{};
229+
gko::size_type iterations{200};
230+
231+
compute_lu(&l_ref, &u_ref, &l_hip, &u_hip, iterations);
232+
233+
GKO_ASSERT_MTX_NEAR(l_ref, l_hip, 1e-14);
234+
GKO_ASSERT_MTX_NEAR(u_ref, u_hip, 1e-14);
235+
}
236+
237+
238+
} // namespace

0 commit comments

Comments
 (0)