Skip to content

Commit cd430f2

Browse files
fritzgoebelfritzgoebel
committed
some work on deluxe scaling
1 parent 0bc4a4d commit cd430f2

7 files changed

Lines changed: 360 additions & 233 deletions

File tree

‎common/cuda_hip/distributed/preconditioner/bddc_kernels.cpp‎

Lines changed: 12 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -69,21 +69,23 @@ __global__ __launch_bounds__(default_block_size) void fill_coarse_data(
6969
} // namespace kernel
7070

7171

72-
template <typename ValueType, typename IndexType>
72+
template <typename ValueType, typename LocalIndexType, typename GlobalIndexType>
7373
void classify_dofs(
7474
std::shared_ptr<const DefaultExecutor> exec,
75-
matrix::Dense<ValueType>* labels, const array<IndexType>& tags,
76-
comm_index_type local_part,
75+
matrix::Dense<ValueType>* labels, const array<GlobalIndexType>& tags,
76+
comm_index_type local_part, const LocalIndexType* row_ptrs,
77+
const LocalIndexType* col_idxs,
7778
array<experimental::distributed::preconditioner::dof_type>& dof_types,
78-
array<IndexType>& permutation_array, array<IndexType>& interface_sizes,
79-
array<ValueType>& unique_labels, array<IndexType>& unique_tags,
80-
array<ValueType>& owning_labels, array<IndexType>& owning_tags,
81-
size_type& n_inner_idxs, size_type& n_face_idxs, size_type& n_edge_idxs,
82-
size_type& n_vertices, size_type& n_faces, size_type& n_edges,
83-
size_type& n_constraints, int& n_owning_interfaces, bool use_faces,
79+
array<LocalIndexType>& permutation_array,
80+
array<LocalIndexType>& interface_sizes, array<ValueType>& unique_labels,
81+
array<GlobalIndexType>& unique_tags, array<ValueType>& owning_labels,
82+
array<GlobalIndexType>& owning_tags, size_type& n_inner_idxs,
83+
size_type& n_face_idxs, size_type& n_edge_idxs, size_type& n_vertices,
84+
size_type& n_faces, size_type& n_edges, size_type& n_constraints,
85+
int& n_owning_interfaces, bool use_faces,
8486
bool use_edges) GKO_NOT_IMPLEMENTED;
8587

86-
GKO_INSTANTIATE_FOR_EACH_NON_COMPLEX_VALUE_AND_INDEX_TYPE_BASE(
88+
GKO_INSTANTIATE_FOR_EACH_NON_COMPLEX_VALUE_AND_LOCAL_GLOBAL_INDEX_TYPE_BASE(
8789
GKO_DECLARE_CLASSIFY_DOFS);
8890

8991

‎core/device_hooks/common_kernels.inc.cpp‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -405,7 +405,8 @@ GKO_STUB_VALUE_AND_LOCAL_GLOBAL_INDEX_TYPE_BASE(
405405
namespace bddc {
406406

407407

408-
GKO_STUB_NON_COMPLEX_VALUE_AND_INDEX_TYPE_BASE(GKO_DECLARE_CLASSIFY_DOFS);
408+
GKO_STUB_NON_COMPLEX_VALUE_AND_LOCAL_GLOBAL_INDEX_TYPE(
409+
GKO_DECLARE_CLASSIFY_DOFS);
409410
GKO_STUB_NON_COMPLEX_VALUE_AND_INDEX_TYPE_BASE(
410411
GKO_DECLARE_GENERATE_CONSTRAINTS);
411412
GKO_STUB_VALUE_TYPE_BASE(GKO_DECLARE_FILL_COARSE_DATA);

‎core/distributed/preconditioner/bddc.cpp‎

Lines changed: 198 additions & 122 deletions
Large diffs are not rendered by default.

‎core/distributed/preconditioner/bddc_kernels.hpp‎

Lines changed: 21 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -22,18 +22,20 @@ namespace gko {
2222
namespace kernels {
2323

2424

25-
#define GKO_DECLARE_CLASSIFY_DOFS(ValueType, IndexType) \
25+
#define GKO_DECLARE_CLASSIFY_DOFS(ValueType, LocalIndexType, GlobalIndexType) \
2626
void classify_dofs( \
2727
std::shared_ptr<const DefaultExecutor> exec, \
28-
matrix::Dense<ValueType>* labels, const array<IndexType>& tags, \
29-
comm_index_type local_part, \
28+
matrix::Dense<ValueType>* labels, const array<GlobalIndexType>& tags, \
29+
comm_index_type local_part, const LocalIndexType* row_ptrs, \
30+
const LocalIndexType* col_idxs, \
3031
array<experimental::distributed::preconditioner::dof_type>& dof_types, \
31-
array<IndexType>& permutation_array, \
32-
array<IndexType>& interface_sizes, array<ValueType>& unique_labels, \
33-
array<IndexType>& unique_tags, array<ValueType>& owning_labels, \
34-
array<IndexType>& owning_tags, size_type& n_inner_idxs, \
35-
size_type& n_face_idxs, size_type& n_edge_idxs, size_type& n_vertices, \
36-
size_type& n_faces, size_type& n_edges, size_type& n_constraints, \
32+
array<LocalIndexType>& permutation_array, \
33+
array<LocalIndexType>& interface_sizes, \
34+
array<ValueType>& unique_labels, array<GlobalIndexType>& unique_tags, \
35+
array<ValueType>& owning_labels, array<GlobalIndexType>& owning_tags, \
36+
size_type& n_inner_idxs, size_type& n_face_idxs, \
37+
size_type& n_edge_idxs, size_type& n_vertices, size_type& n_faces, \
38+
size_type& n_edges, size_type& n_constraints, \
3739
int& n_owning_interfaces, bool use_faces, bool use_edges)
3840

3941

@@ -65,15 +67,16 @@ namespace kernels {
6567
array<IndexType>& permutation_array)
6668

6769

68-
#define GKO_DECLARE_ALL_AS_TEMPLATES \
69-
using comm_index_type = experimental::distributed::comm_index_type; \
70-
template <typename ValueType, typename IndexType> \
71-
GKO_DECLARE_CLASSIFY_DOFS(ValueType, IndexType); \
72-
template <typename ValueType, typename IndexType> \
73-
GKO_DECLARE_GENERATE_CONSTRAINTS(ValueType, IndexType); \
74-
template <typename ValueType> \
75-
GKO_DECLARE_FILL_COARSE_DATA(ValueType); \
76-
template <typename ValueType, typename IndexType> \
70+
#define GKO_DECLARE_ALL_AS_TEMPLATES \
71+
using comm_index_type = experimental::distributed::comm_index_type; \
72+
template <typename ValueType, typename LocalIndexType, \
73+
typename GlobalIndexType> \
74+
GKO_DECLARE_CLASSIFY_DOFS(ValueType, LocalIndexType, GlobalIndexType); \
75+
template <typename ValueType, typename IndexType> \
76+
GKO_DECLARE_GENERATE_CONSTRAINTS(ValueType, IndexType); \
77+
template <typename ValueType> \
78+
GKO_DECLARE_FILL_COARSE_DATA(ValueType); \
79+
template <typename ValueType, typename IndexType> \
7780
GKO_DECLARE_BUILD_COARSE_CONTRIBUTION(ValueType, IndexType)
7881

7982

‎dpcpp/distributed/preconditioner/bddc_kernels.dp.cpp‎

Lines changed: 12 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -19,21 +19,23 @@ namespace dpcpp {
1919
namespace bddc {
2020

2121

22-
template <typename ValueType, typename IndexType>
22+
template <typename ValueType, typename LocalIndexType, typename GlobalIndexType>
2323
void classify_dofs(
2424
std::shared_ptr<const DefaultExecutor> exec,
25-
matrix::Dense<ValueType>* labels, const array<IndexType>& tags,
26-
comm_index_type local_part,
25+
matrix::Dense<ValueType>* labels, const array<GlobalIndexType>& tags,
26+
comm_index_type local_part, const LocalIndexType* row_ptrs,
27+
const LocalIndexType* col_idxs,
2728
array<experimental::distributed::preconditioner::dof_type>& dof_types,
28-
array<IndexType>& permutation_array, array<IndexType>& interface_sizes,
29-
array<ValueType>& unique_labels, array<IndexType>& unique_tags,
30-
array<ValueType>& owning_labels, array<IndexType>& owning_tags,
31-
size_type& n_inner_idxs, size_type& n_face_idxs, size_type& n_edge_idxs,
32-
size_type& n_vertices, size_type& n_faces, size_type& n_edges,
33-
size_type& n_constraints, int& n_owning_interfaces, bool use_faces,
29+
array<LocalIndexType>& permutation_array,
30+
array<LocalIndexType>& interface_sizes, array<ValueType>& unique_labels,
31+
array<GlobalIndexType>& unique_tags, array<ValueType>& owning_labels,
32+
array<GlobalIndexType>& owning_tags, size_type& n_inner_idxs,
33+
size_type& n_face_idxs, size_type& n_edge_idxs, size_type& n_vertices,
34+
size_type& n_faces, size_type& n_edges, size_type& n_constraints,
35+
int& n_owning_interfaces, bool use_faces,
3436
bool use_edges) GKO_NOT_IMPLEMENTED;
3537

36-
GKO_INSTANTIATE_FOR_EACH_NON_COMPLEX_VALUE_AND_INDEX_TYPE_BASE(
38+
GKO_INSTANTIATE_FOR_EACH_NON_COMPLEX_VALUE_AND_LOCAL_GLOBAL_INDEX_TYPE_BASE(
3739
GKO_DECLARE_CLASSIFY_DOFS);
3840

3941

0 commit comments

Comments
 (0)