Line data Source code
1 : //* This file is part of the MOOSE framework 2 : //* https://mooseframework.inl.gov 3 : //* 4 : //* All rights reserved, see COPYRIGHT for full restrictions 5 : //* https://github.com/idaholab/moose/blob/master/COPYRIGHT 6 : //* 7 : //* Licensed under LGPL 2.1, please see LICENSE for details 8 : //* https://www.gnu.org/licenses/lgpl-2.1.html 9 : 10 : #include "KokkosExtraIDIntegralVectorPostprocessor.h" 11 : 12 : #include "MooseVariable.h" 13 : #include "MooseMesh.h" 14 : #include "MooseMeshUtils.h" 15 : 16 : #include "libmesh/mesh_base.h" 17 : 18 363195 : registerKokkosUserObject("MooseApp", KokkosExtraIDIntegralVectorPostprocessor); 19 : 20 : InputParameters 21 2334 : KokkosExtraIDIntegralVectorPostprocessor::validParams() 22 : { 23 2334 : InputParameters params = ElementVectorPostprocessor::validParams(); 24 9336 : params.addCoupledVar("variable", 25 : "The names of the variables that this VectorPostprocessor operates on"); 26 9336 : params.addParam<std::vector<MaterialPropertyName>>( 27 : "mat_prop", "The names of material properties that this VectorPostprocessor operates on"); 28 9336 : params.addRequiredParam<std::vector<ExtraElementIDName>>( 29 : "id_name", "List of extra element ID names by which to separate integral(s)."); 30 9336 : params.addParam<bool>("average", false, "Whether or not to compute volume average"); 31 9336 : MooseEnum mode("atomic reduction", "atomic"); 32 9336 : params.addParam<MooseEnum>("calculation_mode", 33 : mode, 34 : "Whether to use atomic addition or reduction for parallel summation"); 35 2334 : params.addClassDescription("Integrates or averages variables based on extra element IDs"); 36 4668 : return params; 37 2334 : } 38 : 39 118 : KokkosExtraIDIntegralVectorPostprocessor::KokkosExtraIDIntegralVectorPostprocessor( 40 838 : const InputParameters & parameters) 41 : : ElementVectorPostprocessor(parameters), 42 140 : _mesh(_fe_problem.mesh()), 43 140 : _average(getParam<bool>("average")), 44 140 : _mode(getParam<MooseEnum>("calculation_mode")), 45 280 : _nvar(isParamValid("variable") ? coupledComponents("variable") : 0), 46 196 : _nprop(isParamValid("mat_prop") ? getParam<std::vector<MaterialPropertyName>>("mat_prop").size() 47 : : 0), 48 196 : _prop_names(isParamValid("mat_prop") ? getParam<std::vector<MaterialPropertyName>>("mat_prop") 49 : : std::vector<MaterialPropertyName>()), 50 140 : _extra_id(getParam<std::vector<ExtraElementIDName>>("id_name")), 51 140 : _n_extra_id(_extra_id.size()) 52 : { 53 118 : if (!_nvar && !_nprop) 54 0 : mooseError("Neither 'variable' nor 'mat_prop' was specified."); 55 : 56 118 : if (_nvar + _nprop > _cache_size) 57 0 : mooseError("Cannot have more than ", 58 : _cache_size, 59 : " variables and material properties at the same time."); 60 : 61 : // create map of element ids to parsed vpp ids 62 48 : _unique_vpp_id_map = 63 70 : MooseMeshUtils::getExtraIDUniqueCombinationMap(_mesh.getMesh(), blockIDs(), _extra_id); 64 : 65 : // set up variable vector containing parsed extra ids 66 328 : for (unsigned int i = 0; i < _n_extra_id; ++i) 67 : { 68 210 : const auto id_type = _extra_id[i]; 69 210 : auto & p = declareVector("Level-" + std::to_string(i) + "-" + id_type); 70 : // collect local map of local vpp id to extra id 71 210 : std::map<dof_id_type, dof_id_type> extra_ids; 72 897234 : for (const auto elem : _mesh.getMesh().active_local_element_ptr_range()) 73 : { 74 897024 : if (!hasBlocks(elem->subdomain_id())) 75 0 : continue; 76 897024 : auto vpp_id = libmesh_map_find(_unique_vpp_id_map, elem->id()); 77 897024 : if (extra_ids.find(vpp_id) == extra_ids.end()) 78 8500 : extra_ids[vpp_id] = elem->get_extra_integer(getElementIDIndexByName(id_type)); 79 210 : } 80 : // create global map of local vpp id to extra id 81 210 : comm().set_union(extra_ids); 82 210 : p.resize(extra_ids.size()); 83 12090 : for (auto it : extra_ids) 84 11880 : p[it.first] = it.second; 85 210 : _extra_ids.push_back(&p); 86 210 : } 87 : 88 118 : _vector_size = _extra_ids[0]->size(); 89 : 90 118 : std::vector<std::string> vector_names; 91 : 92 118 : _integrals.create(_nvar + _prop_names.size()); 93 118 : _props.create(_prop_names.size()); 94 : 95 118 : unsigned int i = 0; 96 : 97 : // declare vectors containing variable integral values 98 302 : for (; i < _nvar; ++i) 99 : { 100 368 : vector_names.push_back(getVar("variable", i)->name()); 101 184 : auto & p = declareVector(vector_names.back()); 102 184 : p.resize(_vector_size); 103 184 : if (_mode == "atomic") 104 104 : _integrals[i].createDevice(p.size()); 105 : else 106 80 : _integrals[i].init(p.size()); 107 184 : _integrals[i].aliasHost(p.data()); 108 : } 109 : 110 118 : if (_nvar) 111 354 : _var_values = kokkosCoupledValues("variable"); 112 : 113 : // declare vectors containing material property integral values 114 198 : for (unsigned int iprop = 0; iprop < _nprop; ++iprop, ++i) 115 : { 116 80 : vector_names.push_back(_prop_names[iprop]); 117 80 : _props[iprop] = getKokkosMaterialPropertyByName<Real>(vector_names.back()); 118 80 : auto & p = declareVector(vector_names.back()); 119 80 : p.resize(_vector_size); 120 80 : if (_mode == "atomic") 121 52 : _integrals[i].createDevice(p.size()); 122 : else 123 28 : _integrals[i].init(p.size()); 124 80 : _integrals[i].aliasHost(p.data()); 125 : } 126 : 127 118 : _integrals.copyToDevice(); 128 118 : _props.copyToDevice(); 129 118 : } 130 : 131 : void 132 118 : KokkosExtraIDIntegralVectorPostprocessor::initialSetup() 133 : { 134 118 : _unique_vpp_ids.create(kokkosMesh().getNumSubdomains()); 135 : 136 354 : for (const auto block : blockIDs()) 137 : { 138 236 : const auto sid = kokkosMesh().getContiguousSubdomainID(block); 139 : 140 236 : _unique_vpp_ids[sid].create(kokkosMesh().getNumSubdomainLocalElements(block)); 141 236 : _unique_vpp_ids[sid].offset(*kokkosMesh().getSubdomainContiguousElementIDRange(block).begin()); 142 : 143 504044 : for (const auto elem : _mesh.getMesh().active_local_subdomain_elements_ptr_range(block)) 144 : { 145 503808 : const auto eid = kokkosMesh().getContiguousElementID(elem); 146 : 147 503808 : _unique_vpp_ids[sid](eid) = libmesh_map_find(_unique_vpp_id_map, elem->id()); 148 236 : } 149 : 150 236 : _unique_vpp_ids[sid].moveToDevice(); 151 : } 152 : 153 118 : _unique_vpp_ids.copyToDevice(); 154 118 : } 155 : 156 : void 157 118 : KokkosExtraIDIntegralVectorPostprocessor::initialize() 158 : { 159 118 : if (_mode == "atomic") 160 : { 161 221 : for (auto & integral : _integrals) 162 156 : integral = 0; 163 : 164 65 : if (_average) 165 : { 166 13 : _volumes.create(_vector_size); 167 13 : _volumes = 0; 168 : } 169 : } 170 : else 171 53 : allocateReductionBuffer((_integrals.size() + _average) * _vector_size); 172 118 : } 173 : 174 : void 175 118 : KokkosExtraIDIntegralVectorPostprocessor::compute() 176 : { 177 118 : if (_mode == "atomic") 178 65 : computeUserObject(); 179 : else 180 53 : computeReducer(); 181 118 : } 182 : 183 : void 184 118 : KokkosExtraIDIntegralVectorPostprocessor::finalize() 185 : { 186 118 : if (_mode == "atomic") 187 : { 188 221 : for (auto & integral : _integrals) 189 : { 190 156 : integral.copyToHost(); 191 : 192 156 : libmesh_call_mpi(MPI_Allreduce(MPI_IN_PLACE, 193 : integral.data(), 194 : integral.size(), 195 : libMesh::Parallel::StandardType<Real>(), 196 : libMesh::Parallel::OpFunction<Real>::sum(), 197 : comm().get())); 198 : } 199 : 200 65 : if (_average) 201 : { 202 13 : _volumes.copyToHost(); 203 : 204 13 : libmesh_call_mpi(MPI_Allreduce(MPI_IN_PLACE, 205 : _volumes.data(), 206 : _volumes.size(), 207 : libMesh::Parallel::StandardType<Real>(), 208 : libMesh::Parallel::OpFunction<Real>::sum(), 209 : comm().get())); 210 : 211 65 : for (auto & integral : _integrals) 212 3380 : for (unsigned int i = 0; i < integral.size(); ++i) 213 3328 : integral[i] /= _volumes[i]; 214 : } 215 : } 216 : else 217 : { 218 53 : libmesh_call_mpi(MPI_Allreduce(MPI_IN_PLACE, 219 : _reduction_buffer.data(), 220 : _reduction_buffer.size(), 221 : libMesh::Parallel::StandardType<Real>(), 222 : libMesh::Parallel::OpFunction<Real>::sum(), 223 : comm().get())); 224 : 225 161 : for (unsigned int i = 0; i < _integrals.size(); ++i) 226 6240 : for (unsigned int j = 0; j < _vector_size; ++j) 227 : { 228 12264 : _integrals[i][j] = _reduction_buffer[j + i * _vector_size]; 229 : 230 6132 : if (_average) 231 3584 : _integrals[i][j] /= _reduction_buffer[j + _integrals.size() * _vector_size]; 232 : } 233 : } 234 118 : }