1 // Copyright (c) 2017-2022, Lawrence Livermore National Security, LLC and other CEED contributors. 2 // All Rights Reserved. See the top-level LICENSE and NOTICE files for details. 3 // 4 // SPDX-License-Identifier: BSD-2-Clause 5 // 6 // This file is part of CEED: http://github.com/ceed 7 8 #include <ceed.h> 9 #include <ceed/backend.h> 10 #include <ceed/jit-tools.h> 11 #include <assert.h> 12 #include <stdbool.h> 13 #include <string.h> 14 #include <hip/hip_runtime.h> 15 16 #include "../hip/ceed-hip-common.h" 17 #include "../hip/ceed-hip-compile.h" 18 #include "ceed-hip-ref.h" 19 20 //------------------------------------------------------------------------------ 21 // Destroy operator 22 //------------------------------------------------------------------------------ 23 static int CeedOperatorDestroy_Hip(CeedOperator op) { 24 CeedOperator_Hip *impl; 25 26 CeedCallBackend(CeedOperatorGetData(op, &impl)); 27 28 // Apply data 29 for (CeedInt i = 0; i < impl->num_inputs + impl->num_outputs; i++) { 30 CeedCallBackend(CeedVectorDestroy(&impl->e_vecs[i])); 31 } 32 CeedCallBackend(CeedFree(&impl->e_vecs)); 33 34 for (CeedInt i = 0; i < impl->num_inputs; i++) { 35 CeedCallBackend(CeedVectorDestroy(&impl->q_vecs_in[i])); 36 } 37 CeedCallBackend(CeedFree(&impl->q_vecs_in)); 38 39 for (CeedInt i = 0; i < impl->num_outputs; i++) { 40 CeedCallBackend(CeedVectorDestroy(&impl->q_vecs_out[i])); 41 } 42 CeedCallBackend(CeedFree(&impl->q_vecs_out)); 43 44 // QFunction assembly data 45 for (CeedInt i = 0; i < impl->num_active_in; i++) { 46 CeedCallBackend(CeedVectorDestroy(&impl->qf_active_in[i])); 47 } 48 CeedCallBackend(CeedFree(&impl->qf_active_in)); 49 50 // Diag data 51 if (impl->diag) { 52 Ceed ceed; 53 54 CeedCallBackend(CeedOperatorGetCeed(op, &ceed)); 55 if (impl->diag->module) { 56 CeedCallHip(ceed, hipModuleUnload(impl->diag->module)); 57 } 58 if (impl->diag->module_point_block) { 59 CeedCallHip(ceed, hipModuleUnload(impl->diag->module_point_block)); 60 } 61 CeedCallHip(ceed, hipFree(impl->diag->d_eval_modes_in)); 62 CeedCallHip(ceed, hipFree(impl->diag->d_eval_modes_out)); 63 CeedCallHip(ceed, hipFree(impl->diag->d_identity)); 64 CeedCallHip(ceed, hipFree(impl->diag->d_interp_in)); 65 CeedCallHip(ceed, hipFree(impl->diag->d_interp_out)); 66 CeedCallHip(ceed, hipFree(impl->diag->d_grad_in)); 67 CeedCallHip(ceed, hipFree(impl->diag->d_grad_out)); 68 CeedCallHip(ceed, hipFree(impl->diag->d_div_in)); 69 CeedCallHip(ceed, hipFree(impl->diag->d_div_out)); 70 CeedCallHip(ceed, hipFree(impl->diag->d_curl_in)); 71 CeedCallHip(ceed, hipFree(impl->diag->d_curl_out)); 72 CeedCallBackend(CeedElemRestrictionDestroy(&impl->diag->diag_rstr)); 73 CeedCallBackend(CeedElemRestrictionDestroy(&impl->diag->point_block_diag_rstr)); 74 CeedCallBackend(CeedVectorDestroy(&impl->diag->elem_diag)); 75 CeedCallBackend(CeedVectorDestroy(&impl->diag->point_block_elem_diag)); 76 } 77 CeedCallBackend(CeedFree(&impl->diag)); 78 79 if (impl->asmb) { 80 Ceed ceed; 81 82 CeedCallBackend(CeedOperatorGetCeed(op, &ceed)); 83 CeedCallHip(ceed, hipModuleUnload(impl->asmb->module)); 84 CeedCallHip(ceed, hipFree(impl->asmb->d_B_in)); 85 CeedCallHip(ceed, hipFree(impl->asmb->d_B_out)); 86 } 87 CeedCallBackend(CeedFree(&impl->asmb)); 88 89 CeedCallBackend(CeedFree(&impl)); 90 return CEED_ERROR_SUCCESS; 91 } 92 93 //------------------------------------------------------------------------------ 94 // Setup infields or outfields 95 //------------------------------------------------------------------------------ 96 static int CeedOperatorSetupFields_Hip(CeedQFunction qf, CeedOperator op, bool is_input, CeedVector *e_vecs, CeedVector *q_vecs, CeedInt start_e, 97 CeedInt num_fields, CeedInt Q, CeedInt num_elem) { 98 Ceed ceed; 99 CeedQFunctionField *qf_fields; 100 CeedOperatorField *op_fields; 101 102 CeedCallBackend(CeedOperatorGetCeed(op, &ceed)); 103 if (is_input) { 104 CeedCallBackend(CeedOperatorGetFields(op, NULL, &op_fields, NULL, NULL)); 105 CeedCallBackend(CeedQFunctionGetFields(qf, NULL, &qf_fields, NULL, NULL)); 106 } else { 107 CeedCallBackend(CeedOperatorGetFields(op, NULL, NULL, NULL, &op_fields)); 108 CeedCallBackend(CeedQFunctionGetFields(qf, NULL, NULL, NULL, &qf_fields)); 109 } 110 111 // Loop over fields 112 for (CeedInt i = 0; i < num_fields; i++) { 113 bool is_strided = false, skip_restriction = false; 114 CeedSize q_size; 115 CeedInt size; 116 CeedEvalMode eval_mode; 117 CeedBasis basis; 118 119 CeedCallBackend(CeedQFunctionFieldGetEvalMode(qf_fields[i], &eval_mode)); 120 if (eval_mode != CEED_EVAL_WEIGHT) { 121 CeedElemRestriction elem_rstr; 122 123 // Check whether this field can skip the element restriction: 124 // Must be passive input, with eval_mode NONE, and have a strided restriction with CEED_STRIDES_BACKEND. 125 CeedCallBackend(CeedOperatorFieldGetElemRestriction(op_fields[i], &elem_rstr)); 126 127 // First, check whether the field is input or output: 128 if (is_input) { 129 CeedVector vec; 130 131 // Check for passive input 132 CeedCallBackend(CeedOperatorFieldGetVector(op_fields[i], &vec)); 133 if (vec != CEED_VECTOR_ACTIVE) { 134 // Check eval_mode 135 if (eval_mode == CEED_EVAL_NONE) { 136 // Check for strided restriction 137 CeedCallBackend(CeedElemRestrictionIsStrided(elem_rstr, &is_strided)); 138 if (is_strided) { 139 // Check if vector is already in preferred backend ordering 140 CeedCallBackend(CeedElemRestrictionHasBackendStrides(elem_rstr, &skip_restriction)); 141 } 142 } 143 } 144 } 145 if (skip_restriction) { 146 // We do not need an E-Vector, but will use the input field vector's data directly in the operator application. 147 e_vecs[i + start_e] = NULL; 148 } else { 149 CeedCallBackend(CeedElemRestrictionCreateVector(elem_rstr, NULL, &e_vecs[i + start_e])); 150 } 151 } 152 153 switch (eval_mode) { 154 case CEED_EVAL_NONE: 155 CeedCallBackend(CeedQFunctionFieldGetSize(qf_fields[i], &size)); 156 q_size = (CeedSize)num_elem * Q * size; 157 CeedCallBackend(CeedVectorCreate(ceed, q_size, &q_vecs[i])); 158 break; 159 case CEED_EVAL_INTERP: 160 case CEED_EVAL_GRAD: 161 case CEED_EVAL_DIV: 162 case CEED_EVAL_CURL: 163 CeedCallBackend(CeedQFunctionFieldGetSize(qf_fields[i], &size)); 164 q_size = (CeedSize)num_elem * Q * size; 165 CeedCallBackend(CeedVectorCreate(ceed, q_size, &q_vecs[i])); 166 break; 167 case CEED_EVAL_WEIGHT: // Only on input fields 168 CeedCallBackend(CeedOperatorFieldGetBasis(op_fields[i], &basis)); 169 q_size = (CeedSize)num_elem * Q; 170 CeedCallBackend(CeedVectorCreate(ceed, q_size, &q_vecs[i])); 171 CeedCallBackend(CeedBasisApply(basis, num_elem, CEED_NOTRANSPOSE, CEED_EVAL_WEIGHT, CEED_VECTOR_NONE, q_vecs[i])); 172 break; 173 } 174 } 175 return CEED_ERROR_SUCCESS; 176 } 177 178 //------------------------------------------------------------------------------ 179 // CeedOperator needs to connect all the named fields (be they active or passive) to the named inputs and outputs of its CeedQFunction. 180 //------------------------------------------------------------------------------ 181 static int CeedOperatorSetup_Hip(CeedOperator op) { 182 Ceed ceed; 183 bool is_setup_done; 184 CeedInt Q, num_elem, num_input_fields, num_output_fields; 185 CeedQFunctionField *qf_input_fields, *qf_output_fields; 186 CeedQFunction qf; 187 CeedOperatorField *op_input_fields, *op_output_fields; 188 CeedOperator_Hip *impl; 189 190 CeedCallBackend(CeedOperatorIsSetupDone(op, &is_setup_done)); 191 if (is_setup_done) return CEED_ERROR_SUCCESS; 192 193 CeedCallBackend(CeedOperatorGetCeed(op, &ceed)); 194 CeedCallBackend(CeedOperatorGetData(op, &impl)); 195 CeedCallBackend(CeedOperatorGetQFunction(op, &qf)); 196 CeedCallBackend(CeedOperatorGetNumQuadraturePoints(op, &Q)); 197 CeedCallBackend(CeedOperatorGetNumElements(op, &num_elem)); 198 CeedCallBackend(CeedOperatorGetFields(op, &num_input_fields, &op_input_fields, &num_output_fields, &op_output_fields)); 199 CeedCallBackend(CeedQFunctionGetFields(qf, NULL, &qf_input_fields, NULL, &qf_output_fields)); 200 201 // Allocate 202 CeedCallBackend(CeedCalloc(num_input_fields + num_output_fields, &impl->e_vecs)); 203 CeedCallBackend(CeedCalloc(CEED_FIELD_MAX, &impl->q_vecs_in)); 204 CeedCallBackend(CeedCalloc(CEED_FIELD_MAX, &impl->q_vecs_out)); 205 impl->num_inputs = num_input_fields; 206 impl->num_outputs = num_output_fields; 207 208 // Set up infield and outfield e_vecs and q_vecs 209 // Infields 210 CeedCallBackend(CeedOperatorSetupFields_Hip(qf, op, true, impl->e_vecs, impl->q_vecs_in, 0, num_input_fields, Q, num_elem)); 211 // Outfields 212 CeedCallBackend(CeedOperatorSetupFields_Hip(qf, op, false, impl->e_vecs, impl->q_vecs_out, num_input_fields, num_output_fields, Q, num_elem)); 213 214 CeedCallBackend(CeedOperatorSetSetupDone(op)); 215 return CEED_ERROR_SUCCESS; 216 } 217 218 //------------------------------------------------------------------------------ 219 // Setup Operator Inputs 220 //------------------------------------------------------------------------------ 221 static inline int CeedOperatorSetupInputs_Hip(CeedInt num_input_fields, CeedQFunctionField *qf_input_fields, CeedOperatorField *op_input_fields, 222 CeedVector in_vec, const bool skip_active, CeedScalar *e_data[2 * CEED_FIELD_MAX], 223 CeedOperator_Hip *impl, CeedRequest *request) { 224 for (CeedInt i = 0; i < num_input_fields; i++) { 225 CeedEvalMode eval_mode; 226 CeedVector vec; 227 CeedElemRestriction elem_rstr; 228 229 // Get input vector 230 CeedCallBackend(CeedOperatorFieldGetVector(op_input_fields[i], &vec)); 231 if (vec == CEED_VECTOR_ACTIVE) { 232 if (skip_active) continue; 233 else vec = in_vec; 234 } 235 236 CeedCallBackend(CeedQFunctionFieldGetEvalMode(qf_input_fields[i], &eval_mode)); 237 if (eval_mode == CEED_EVAL_WEIGHT) { // Skip 238 } else { 239 // Get input vector 240 CeedCallBackend(CeedOperatorFieldGetVector(op_input_fields[i], &vec)); 241 // Get input element restriction 242 CeedCallBackend(CeedOperatorFieldGetElemRestriction(op_input_fields[i], &elem_rstr)); 243 if (vec == CEED_VECTOR_ACTIVE) vec = in_vec; 244 // Restrict, if necessary 245 if (!impl->e_vecs[i]) { 246 // No restriction for this field; read data directly from vec. 247 CeedCallBackend(CeedVectorGetArrayRead(vec, CEED_MEM_DEVICE, (const CeedScalar **)&e_data[i])); 248 } else { 249 CeedCallBackend(CeedElemRestrictionApply(elem_rstr, CEED_NOTRANSPOSE, vec, impl->e_vecs[i], request)); 250 // Get evec 251 CeedCallBackend(CeedVectorGetArrayRead(impl->e_vecs[i], CEED_MEM_DEVICE, (const CeedScalar **)&e_data[i])); 252 } 253 } 254 } 255 return CEED_ERROR_SUCCESS; 256 } 257 258 //------------------------------------------------------------------------------ 259 // Input Basis Action 260 //------------------------------------------------------------------------------ 261 static inline int CeedOperatorInputBasis_Hip(CeedInt num_elem, CeedQFunctionField *qf_input_fields, CeedOperatorField *op_input_fields, 262 CeedInt num_input_fields, const bool skip_active, CeedScalar *e_data[2 * CEED_FIELD_MAX], 263 CeedOperator_Hip *impl) { 264 for (CeedInt i = 0; i < num_input_fields; i++) { 265 CeedInt elem_size, size; 266 CeedEvalMode eval_mode; 267 CeedElemRestriction elem_rstr; 268 CeedBasis basis; 269 270 // Skip active input 271 if (skip_active) { 272 CeedVector vec; 273 274 CeedCallBackend(CeedOperatorFieldGetVector(op_input_fields[i], &vec)); 275 if (vec == CEED_VECTOR_ACTIVE) continue; 276 } 277 // Get elem_size, eval_mode, size 278 CeedCallBackend(CeedOperatorFieldGetElemRestriction(op_input_fields[i], &elem_rstr)); 279 CeedCallBackend(CeedElemRestrictionGetElementSize(elem_rstr, &elem_size)); 280 CeedCallBackend(CeedQFunctionFieldGetEvalMode(qf_input_fields[i], &eval_mode)); 281 CeedCallBackend(CeedQFunctionFieldGetSize(qf_input_fields[i], &size)); 282 // Basis action 283 switch (eval_mode) { 284 case CEED_EVAL_NONE: 285 CeedCallBackend(CeedVectorSetArray(impl->q_vecs_in[i], CEED_MEM_DEVICE, CEED_USE_POINTER, e_data[i])); 286 break; 287 case CEED_EVAL_INTERP: 288 case CEED_EVAL_GRAD: 289 case CEED_EVAL_DIV: 290 case CEED_EVAL_CURL: 291 CeedCallBackend(CeedOperatorFieldGetBasis(op_input_fields[i], &basis)); 292 CeedCallBackend(CeedBasisApply(basis, num_elem, CEED_NOTRANSPOSE, eval_mode, impl->e_vecs[i], impl->q_vecs_in[i])); 293 break; 294 case CEED_EVAL_WEIGHT: 295 break; // No action 296 } 297 } 298 return CEED_ERROR_SUCCESS; 299 } 300 301 //------------------------------------------------------------------------------ 302 // Restore Input Vectors 303 //------------------------------------------------------------------------------ 304 static inline int CeedOperatorRestoreInputs_Hip(CeedInt num_input_fields, CeedQFunctionField *qf_input_fields, CeedOperatorField *op_input_fields, 305 const bool skip_active, CeedScalar *e_data[2 * CEED_FIELD_MAX], CeedOperator_Hip *impl) { 306 for (CeedInt i = 0; i < num_input_fields; i++) { 307 CeedEvalMode eval_mode; 308 CeedVector vec; 309 310 // Skip active input 311 if (skip_active) { 312 CeedCallBackend(CeedOperatorFieldGetVector(op_input_fields[i], &vec)); 313 if (vec == CEED_VECTOR_ACTIVE) continue; 314 } 315 CeedCallBackend(CeedQFunctionFieldGetEvalMode(qf_input_fields[i], &eval_mode)); 316 if (eval_mode == CEED_EVAL_WEIGHT) { // Skip 317 } else { 318 if (!impl->e_vecs[i]) { // This was a skip_restriction case 319 CeedCallBackend(CeedOperatorFieldGetVector(op_input_fields[i], &vec)); 320 CeedCallBackend(CeedVectorRestoreArrayRead(vec, (const CeedScalar **)&e_data[i])); 321 } else { 322 CeedCallBackend(CeedVectorRestoreArrayRead(impl->e_vecs[i], (const CeedScalar **)&e_data[i])); 323 } 324 } 325 } 326 return CEED_ERROR_SUCCESS; 327 } 328 329 //------------------------------------------------------------------------------ 330 // Apply and add to output 331 //------------------------------------------------------------------------------ 332 static int CeedOperatorApplyAdd_Hip(CeedOperator op, CeedVector in_vec, CeedVector out_vec, CeedRequest *request) { 333 CeedInt Q, num_elem, elem_size, num_input_fields, num_output_fields, size; 334 CeedScalar *e_data[2 * CEED_FIELD_MAX] = {NULL}; 335 CeedQFunctionField *qf_input_fields, *qf_output_fields; 336 CeedQFunction qf; 337 CeedOperatorField *op_input_fields, *op_output_fields; 338 CeedOperator_Hip *impl; 339 340 CeedCallBackend(CeedOperatorGetData(op, &impl)); 341 CeedCallBackend(CeedOperatorGetQFunction(op, &qf)); 342 CeedCallBackend(CeedOperatorGetNumQuadraturePoints(op, &Q)); 343 CeedCallBackend(CeedOperatorGetNumElements(op, &num_elem)); 344 CeedCallBackend(CeedOperatorGetFields(op, &num_input_fields, &op_input_fields, &num_output_fields, &op_output_fields)); 345 CeedCallBackend(CeedQFunctionGetFields(qf, NULL, &qf_input_fields, NULL, &qf_output_fields)); 346 347 // Setup 348 CeedCallBackend(CeedOperatorSetup_Hip(op)); 349 350 // Input Evecs and Restriction 351 CeedCallBackend(CeedOperatorSetupInputs_Hip(num_input_fields, qf_input_fields, op_input_fields, in_vec, false, e_data, impl, request)); 352 353 // Input basis apply if needed 354 CeedCallBackend(CeedOperatorInputBasis_Hip(num_elem, qf_input_fields, op_input_fields, num_input_fields, false, e_data, impl)); 355 356 // Output pointers, as necessary 357 for (CeedInt i = 0; i < num_output_fields; i++) { 358 CeedEvalMode eval_mode; 359 360 CeedCallBackend(CeedQFunctionFieldGetEvalMode(qf_output_fields[i], &eval_mode)); 361 if (eval_mode == CEED_EVAL_NONE) { 362 // Set the output Q-Vector to use the E-Vector data directly. 363 CeedCallBackend(CeedVectorGetArrayWrite(impl->e_vecs[i + impl->num_inputs], CEED_MEM_DEVICE, &e_data[i + num_input_fields])); 364 CeedCallBackend(CeedVectorSetArray(impl->q_vecs_out[i], CEED_MEM_DEVICE, CEED_USE_POINTER, e_data[i + num_input_fields])); 365 } 366 } 367 368 // Q function 369 CeedCallBackend(CeedQFunctionApply(qf, num_elem * Q, impl->q_vecs_in, impl->q_vecs_out)); 370 371 // Output basis apply if needed 372 for (CeedInt i = 0; i < num_output_fields; i++) { 373 CeedEvalMode eval_mode; 374 CeedElemRestriction elem_rstr; 375 CeedBasis basis; 376 377 // Get elem_size, eval_mode, size 378 CeedCallBackend(CeedOperatorFieldGetElemRestriction(op_output_fields[i], &elem_rstr)); 379 CeedCallBackend(CeedElemRestrictionGetElementSize(elem_rstr, &elem_size)); 380 CeedCallBackend(CeedQFunctionFieldGetEvalMode(qf_output_fields[i], &eval_mode)); 381 CeedCallBackend(CeedQFunctionFieldGetSize(qf_output_fields[i], &size)); 382 // Basis action 383 switch (eval_mode) { 384 case CEED_EVAL_NONE: 385 break; // No action 386 case CEED_EVAL_INTERP: 387 case CEED_EVAL_GRAD: 388 case CEED_EVAL_DIV: 389 case CEED_EVAL_CURL: 390 CeedCallBackend(CeedOperatorFieldGetBasis(op_output_fields[i], &basis)); 391 CeedCallBackend(CeedBasisApply(basis, num_elem, CEED_TRANSPOSE, eval_mode, impl->q_vecs_out[i], impl->e_vecs[i + impl->num_inputs])); 392 break; 393 // LCOV_EXCL_START 394 case CEED_EVAL_WEIGHT: { 395 Ceed ceed; 396 397 CeedCallBackend(CeedOperatorGetCeed(op, &ceed)); 398 return CeedError(ceed, CEED_ERROR_BACKEND, "CEED_EVAL_WEIGHT cannot be an output evaluation mode"); 399 // LCOV_EXCL_STOP 400 } 401 } 402 } 403 404 // Output restriction 405 for (CeedInt i = 0; i < num_output_fields; i++) { 406 CeedEvalMode eval_mode; 407 CeedVector vec; 408 CeedElemRestriction elem_rstr; 409 410 // Restore evec 411 CeedCallBackend(CeedQFunctionFieldGetEvalMode(qf_output_fields[i], &eval_mode)); 412 if (eval_mode == CEED_EVAL_NONE) { 413 CeedCallBackend(CeedVectorRestoreArray(impl->e_vecs[i + impl->num_inputs], &e_data[i + num_input_fields])); 414 } 415 // Get output vector 416 CeedCallBackend(CeedOperatorFieldGetVector(op_output_fields[i], &vec)); 417 // Restrict 418 CeedCallBackend(CeedOperatorFieldGetElemRestriction(op_output_fields[i], &elem_rstr)); 419 // Active 420 if (vec == CEED_VECTOR_ACTIVE) vec = out_vec; 421 422 CeedCallBackend(CeedElemRestrictionApply(elem_rstr, CEED_TRANSPOSE, impl->e_vecs[i + impl->num_inputs], vec, request)); 423 } 424 425 // Restore input arrays 426 CeedCallBackend(CeedOperatorRestoreInputs_Hip(num_input_fields, qf_input_fields, op_input_fields, false, e_data, impl)); 427 return CEED_ERROR_SUCCESS; 428 } 429 430 //------------------------------------------------------------------------------ 431 // Linear QFunction Assembly Core 432 //------------------------------------------------------------------------------ 433 static inline int CeedOperatorLinearAssembleQFunctionCore_Hip(CeedOperator op, bool build_objects, CeedVector *assembled, CeedElemRestriction *rstr, 434 CeedRequest *request) { 435 Ceed ceed, ceed_parent; 436 CeedInt num_active_in, num_active_out, Q, num_elem, num_input_fields, num_output_fields, size; 437 CeedScalar *assembled_array, *e_data[2 * CEED_FIELD_MAX] = {NULL}; 438 CeedVector *active_inputs; 439 CeedQFunctionField *qf_input_fields, *qf_output_fields; 440 CeedQFunction qf; 441 CeedOperatorField *op_input_fields, *op_output_fields; 442 CeedOperator_Hip *impl; 443 444 CeedCallBackend(CeedOperatorGetCeed(op, &ceed)); 445 CeedCallBackend(CeedOperatorGetFallbackParentCeed(op, &ceed_parent)); 446 CeedCallBackend(CeedOperatorGetData(op, &impl)); 447 CeedCallBackend(CeedOperatorGetNumQuadraturePoints(op, &Q)); 448 CeedCallBackend(CeedOperatorGetNumElements(op, &num_elem)); 449 CeedCallBackend(CeedOperatorGetQFunction(op, &qf)); 450 CeedCallBackend(CeedQFunctionGetFields(qf, NULL, &qf_input_fields, NULL, &qf_output_fields)); 451 CeedCallBackend(CeedOperatorGetFields(op, &num_input_fields, &op_input_fields, &num_output_fields, &op_output_fields)); 452 active_inputs = impl->qf_active_in; 453 num_active_in = impl->num_active_in, num_active_out = impl->num_active_out; 454 455 // Setup 456 CeedCallBackend(CeedOperatorSetup_Hip(op)); 457 458 // Input Evecs and Restriction 459 CeedCallBackend(CeedOperatorSetupInputs_Hip(num_input_fields, qf_input_fields, op_input_fields, NULL, true, e_data, impl, request)); 460 461 // Count number of active input fields 462 if (!num_active_in) { 463 for (CeedInt i = 0; i < num_input_fields; i++) { 464 CeedScalar *q_vec_array; 465 CeedVector vec; 466 467 // Get input vector 468 CeedCallBackend(CeedOperatorFieldGetVector(op_input_fields[i], &vec)); 469 // Check if active input 470 if (vec == CEED_VECTOR_ACTIVE) { 471 CeedCallBackend(CeedQFunctionFieldGetSize(qf_input_fields[i], &size)); 472 CeedCallBackend(CeedVectorSetValue(impl->q_vecs_in[i], 0.0)); 473 CeedCallBackend(CeedVectorGetArray(impl->q_vecs_in[i], CEED_MEM_DEVICE, &q_vec_array)); 474 CeedCallBackend(CeedRealloc(num_active_in + size, &active_inputs)); 475 for (CeedInt field = 0; field < size; field++) { 476 CeedSize q_size = (CeedSize)Q * num_elem; 477 478 CeedCallBackend(CeedVectorCreate(ceed, q_size, &active_inputs[num_active_in + field])); 479 CeedCallBackend( 480 CeedVectorSetArray(active_inputs[num_active_in + field], CEED_MEM_DEVICE, CEED_USE_POINTER, &q_vec_array[field * Q * num_elem])); 481 } 482 num_active_in += size; 483 CeedCallBackend(CeedVectorRestoreArray(impl->q_vecs_in[i], &q_vec_array)); 484 } 485 } 486 impl->num_active_in = num_active_in; 487 impl->qf_active_in = active_inputs; 488 } 489 490 // Count number of active output fields 491 if (!num_active_out) { 492 for (CeedInt i = 0; i < num_output_fields; i++) { 493 CeedVector vec; 494 495 // Get output vector 496 CeedCallBackend(CeedOperatorFieldGetVector(op_output_fields[i], &vec)); 497 // Check if active output 498 if (vec == CEED_VECTOR_ACTIVE) { 499 CeedCallBackend(CeedQFunctionFieldGetSize(qf_output_fields[i], &size)); 500 num_active_out += size; 501 } 502 } 503 impl->num_active_out = num_active_out; 504 } 505 506 // Check sizes 507 CeedCheck(num_active_in > 0 && num_active_out > 0, ceed, CEED_ERROR_BACKEND, "Cannot assemble QFunction without active inputs and outputs"); 508 509 // Build objects if needed 510 if (build_objects) { 511 CeedSize l_size = (CeedSize)num_elem * Q * num_active_in * num_active_out; 512 CeedInt strides[3] = {1, num_elem * Q, Q}; /* *NOPAD* */ 513 514 // Create output restriction 515 CeedCallBackend(CeedElemRestrictionCreateStrided(ceed_parent, num_elem, Q, num_active_in * num_active_out, 516 num_active_in * num_active_out * num_elem * Q, strides, rstr)); 517 // Create assembled vector 518 CeedCallBackend(CeedVectorCreate(ceed_parent, l_size, assembled)); 519 } 520 CeedCallBackend(CeedVectorSetValue(*assembled, 0.0)); 521 CeedCallBackend(CeedVectorGetArray(*assembled, CEED_MEM_DEVICE, &assembled_array)); 522 523 // Input basis apply 524 CeedCallBackend(CeedOperatorInputBasis_Hip(num_elem, qf_input_fields, op_input_fields, num_input_fields, true, e_data, impl)); 525 526 // Assemble QFunction 527 for (CeedInt in = 0; in < num_active_in; in++) { 528 // Set Inputs 529 CeedCallBackend(CeedVectorSetValue(active_inputs[in], 1.0)); 530 if (num_active_in > 1) { 531 CeedCallBackend(CeedVectorSetValue(active_inputs[(in + num_active_in - 1) % num_active_in], 0.0)); 532 } 533 // Set Outputs 534 for (CeedInt out = 0; out < num_output_fields; out++) { 535 CeedVector vec; 536 537 // Get output vector 538 CeedCallBackend(CeedOperatorFieldGetVector(op_output_fields[out], &vec)); 539 // Check if active output 540 if (vec == CEED_VECTOR_ACTIVE) { 541 CeedCallBackend(CeedVectorSetArray(impl->q_vecs_out[out], CEED_MEM_DEVICE, CEED_USE_POINTER, assembled_array)); 542 CeedCallBackend(CeedQFunctionFieldGetSize(qf_output_fields[out], &size)); 543 assembled_array += size * Q * num_elem; // Advance the pointer by the size of the output 544 } 545 } 546 // Apply QFunction 547 CeedCallBackend(CeedQFunctionApply(qf, Q * num_elem, impl->q_vecs_in, impl->q_vecs_out)); 548 } 549 550 // Un-set output q_vecs to prevent accidental overwrite of Assembled 551 for (CeedInt out = 0; out < num_output_fields; out++) { 552 CeedVector vec; 553 554 // Get output vector 555 CeedCallBackend(CeedOperatorFieldGetVector(op_output_fields[out], &vec)); 556 // Check if active output 557 if (vec == CEED_VECTOR_ACTIVE) { 558 CeedCallBackend(CeedVectorTakeArray(impl->q_vecs_out[out], CEED_MEM_DEVICE, NULL)); 559 } 560 } 561 562 // Restore input arrays 563 CeedCallBackend(CeedOperatorRestoreInputs_Hip(num_input_fields, qf_input_fields, op_input_fields, true, e_data, impl)); 564 565 // Restore output 566 CeedCallBackend(CeedVectorRestoreArray(*assembled, &assembled_array)); 567 return CEED_ERROR_SUCCESS; 568 } 569 570 //------------------------------------------------------------------------------ 571 // Assemble Linear QFunction 572 //------------------------------------------------------------------------------ 573 static int CeedOperatorLinearAssembleQFunction_Hip(CeedOperator op, CeedVector *assembled, CeedElemRestriction *rstr, CeedRequest *request) { 574 return CeedOperatorLinearAssembleQFunctionCore_Hip(op, true, assembled, rstr, request); 575 } 576 577 //------------------------------------------------------------------------------ 578 // Update Assembled Linear QFunction 579 //------------------------------------------------------------------------------ 580 static int CeedOperatorLinearAssembleQFunctionUpdate_Hip(CeedOperator op, CeedVector assembled, CeedElemRestriction rstr, CeedRequest *request) { 581 return CeedOperatorLinearAssembleQFunctionCore_Hip(op, false, &assembled, &rstr, request); 582 } 583 584 //------------------------------------------------------------------------------ 585 // Assemble Diagonal Setup 586 //------------------------------------------------------------------------------ 587 static inline int CeedOperatorAssembleDiagonalSetup_Hip(CeedOperator op) { 588 Ceed ceed; 589 CeedInt num_input_fields, num_output_fields, num_eval_modes_in = 0, num_eval_modes_out = 0; 590 CeedInt q_comp, num_nodes, num_qpts; 591 CeedEvalMode *eval_modes_in = NULL, *eval_modes_out = NULL; 592 CeedBasis basis_in = NULL, basis_out = NULL; 593 CeedQFunctionField *qf_fields; 594 CeedQFunction qf; 595 CeedOperatorField *op_fields; 596 CeedOperator_Hip *impl; 597 598 CeedCallBackend(CeedOperatorGetCeed(op, &ceed)); 599 CeedCallBackend(CeedOperatorGetQFunction(op, &qf)); 600 CeedCallBackend(CeedQFunctionGetNumArgs(qf, &num_input_fields, &num_output_fields)); 601 602 // Determine active input basis 603 CeedCallBackend(CeedOperatorGetFields(op, NULL, &op_fields, NULL, NULL)); 604 CeedCallBackend(CeedQFunctionGetFields(qf, NULL, &qf_fields, NULL, NULL)); 605 for (CeedInt i = 0; i < num_input_fields; i++) { 606 CeedVector vec; 607 608 CeedCallBackend(CeedOperatorFieldGetVector(op_fields[i], &vec)); 609 if (vec == CEED_VECTOR_ACTIVE) { 610 CeedBasis basis; 611 CeedEvalMode eval_mode; 612 613 CeedCallBackend(CeedOperatorFieldGetBasis(op_fields[i], &basis)); 614 CeedCheck(!basis_in || basis_in == basis, ceed, CEED_ERROR_BACKEND, 615 "Backend does not implement operator diagonal assembly with multiple active bases"); 616 basis_in = basis; 617 CeedCallBackend(CeedQFunctionFieldGetEvalMode(qf_fields[i], &eval_mode)); 618 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis_in, eval_mode, &q_comp)); 619 if (eval_mode != CEED_EVAL_WEIGHT) { 620 // q_comp = 1 if CEED_EVAL_NONE, CEED_EVAL_WEIGHT caught by QF assembly 621 CeedCallBackend(CeedRealloc(num_eval_modes_in + q_comp, &eval_modes_in)); 622 for (CeedInt d = 0; d < q_comp; d++) eval_modes_in[num_eval_modes_in + d] = eval_mode; 623 num_eval_modes_in += q_comp; 624 } 625 } 626 } 627 628 // Determine active output basis 629 CeedCallBackend(CeedOperatorGetFields(op, NULL, NULL, NULL, &op_fields)); 630 CeedCallBackend(CeedQFunctionGetFields(qf, NULL, NULL, NULL, &qf_fields)); 631 for (CeedInt i = 0; i < num_output_fields; i++) { 632 CeedVector vec; 633 634 CeedCallBackend(CeedOperatorFieldGetVector(op_fields[i], &vec)); 635 if (vec == CEED_VECTOR_ACTIVE) { 636 CeedBasis basis; 637 CeedEvalMode eval_mode; 638 639 CeedCallBackend(CeedOperatorFieldGetBasis(op_fields[i], &basis)); 640 CeedCheck(!basis_out || basis_out == basis, ceed, CEED_ERROR_BACKEND, 641 "Backend does not implement operator diagonal assembly with multiple active bases"); 642 basis_out = basis; 643 CeedCallBackend(CeedQFunctionFieldGetEvalMode(qf_fields[i], &eval_mode)); 644 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis_out, eval_mode, &q_comp)); 645 if (eval_mode != CEED_EVAL_WEIGHT) { 646 // q_comp = 1 if CEED_EVAL_NONE, CEED_EVAL_WEIGHT caught by QF assembly 647 CeedCallBackend(CeedRealloc(num_eval_modes_out + q_comp, &eval_modes_out)); 648 for (CeedInt d = 0; d < q_comp; d++) eval_modes_out[num_eval_modes_out + d] = eval_mode; 649 num_eval_modes_out += q_comp; 650 } 651 } 652 } 653 654 // Operator data struct 655 CeedCallBackend(CeedOperatorGetData(op, &impl)); 656 CeedCallBackend(CeedCalloc(1, &impl->diag)); 657 CeedOperatorDiag_Hip *diag = impl->diag; 658 659 // Basis matrices 660 CeedCallBackend(CeedBasisGetNumNodes(basis_in, &num_nodes)); 661 if (basis_in == CEED_BASIS_NONE) num_qpts = num_nodes; 662 else CeedCallBackend(CeedBasisGetNumQuadraturePoints(basis_in, &num_qpts)); 663 const CeedInt interp_bytes = num_nodes * num_qpts * sizeof(CeedScalar); 664 const CeedInt eval_modes_bytes = sizeof(CeedEvalMode); 665 bool has_eval_none = false; 666 667 // CEED_EVAL_NONE 668 for (CeedInt i = 0; i < num_eval_modes_in; i++) has_eval_none = has_eval_none || (eval_modes_in[i] == CEED_EVAL_NONE); 669 for (CeedInt i = 0; i < num_eval_modes_out; i++) has_eval_none = has_eval_none || (eval_modes_out[i] == CEED_EVAL_NONE); 670 if (has_eval_none) { 671 CeedScalar *identity = NULL; 672 673 CeedCallBackend(CeedCalloc(num_nodes * num_qpts, &identity)); 674 for (CeedInt i = 0; i < (num_nodes < num_qpts ? num_nodes : num_qpts); i++) identity[i * num_nodes + i] = 1.0; 675 CeedCallHip(ceed, hipMalloc((void **)&diag->d_identity, interp_bytes)); 676 CeedCallHip(ceed, hipMemcpy(diag->d_identity, identity, interp_bytes, hipMemcpyHostToDevice)); 677 CeedCallBackend(CeedFree(&identity)); 678 } 679 680 // CEED_EVAL_INTERP, CEED_EVAL_GRAD, CEED_EVAL_DIV, and CEED_EVAL_CURL 681 for (CeedInt in = 0; in < 2; in++) { 682 CeedFESpace fespace; 683 CeedBasis basis = in ? basis_in : basis_out; 684 685 CeedCallBackend(CeedBasisGetFESpace(basis, &fespace)); 686 switch (fespace) { 687 case CEED_FE_SPACE_H1: { 688 CeedInt q_comp_interp, q_comp_grad; 689 const CeedScalar *interp, *grad; 690 CeedScalar *d_interp, *d_grad; 691 692 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis, CEED_EVAL_INTERP, &q_comp_interp)); 693 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis, CEED_EVAL_GRAD, &q_comp_grad)); 694 695 CeedCallBackend(CeedBasisGetInterp(basis, &interp)); 696 CeedCallHip(ceed, hipMalloc((void **)&d_interp, interp_bytes * q_comp_interp)); 697 CeedCallHip(ceed, hipMemcpy(d_interp, interp, interp_bytes * q_comp_interp, hipMemcpyHostToDevice)); 698 CeedCallBackend(CeedBasisGetGrad(basis, &grad)); 699 CeedCallHip(ceed, hipMalloc((void **)&d_grad, interp_bytes * q_comp_grad)); 700 CeedCallHip(ceed, hipMemcpy(d_grad, grad, interp_bytes * q_comp_grad, hipMemcpyHostToDevice)); 701 if (in) { 702 diag->d_interp_in = d_interp; 703 diag->d_grad_in = d_grad; 704 } else { 705 diag->d_interp_out = d_interp; 706 diag->d_grad_out = d_grad; 707 } 708 } break; 709 case CEED_FE_SPACE_HDIV: { 710 CeedInt q_comp_interp, q_comp_div; 711 const CeedScalar *interp, *div; 712 CeedScalar *d_interp, *d_div; 713 714 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis, CEED_EVAL_INTERP, &q_comp_interp)); 715 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis, CEED_EVAL_DIV, &q_comp_div)); 716 717 CeedCallBackend(CeedBasisGetInterp(basis, &interp)); 718 CeedCallHip(ceed, hipMalloc((void **)&d_interp, interp_bytes * q_comp_interp)); 719 CeedCallHip(ceed, hipMemcpy(d_interp, interp, interp_bytes * q_comp_interp, hipMemcpyHostToDevice)); 720 CeedCallBackend(CeedBasisGetDiv(basis, &div)); 721 CeedCallHip(ceed, hipMalloc((void **)&d_div, interp_bytes * q_comp_div)); 722 CeedCallHip(ceed, hipMemcpy(d_div, div, interp_bytes * q_comp_div, hipMemcpyHostToDevice)); 723 if (in) { 724 diag->d_interp_in = d_interp; 725 diag->d_div_in = d_div; 726 } else { 727 diag->d_interp_out = d_interp; 728 diag->d_div_out = d_div; 729 } 730 } break; 731 case CEED_FE_SPACE_HCURL: { 732 CeedInt q_comp_interp, q_comp_curl; 733 const CeedScalar *interp, *curl; 734 CeedScalar *d_interp, *d_curl; 735 736 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis, CEED_EVAL_INTERP, &q_comp_interp)); 737 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis, CEED_EVAL_CURL, &q_comp_curl)); 738 739 CeedCallBackend(CeedBasisGetInterp(basis, &interp)); 740 CeedCallHip(ceed, hipMalloc((void **)&d_interp, interp_bytes * q_comp_interp)); 741 CeedCallHip(ceed, hipMemcpy(d_interp, interp, interp_bytes * q_comp_interp, hipMemcpyHostToDevice)); 742 CeedCallBackend(CeedBasisGetCurl(basis, &curl)); 743 CeedCallHip(ceed, hipMalloc((void **)&d_curl, interp_bytes * q_comp_curl)); 744 CeedCallHip(ceed, hipMemcpy(d_curl, curl, interp_bytes * q_comp_curl, hipMemcpyHostToDevice)); 745 if (in) { 746 diag->d_interp_in = d_interp; 747 diag->d_curl_in = d_curl; 748 } else { 749 diag->d_interp_out = d_interp; 750 diag->d_curl_out = d_curl; 751 } 752 } break; 753 } 754 } 755 756 // Arrays of eval_modes 757 CeedCallHip(ceed, hipMalloc((void **)&diag->d_eval_modes_in, num_eval_modes_in * eval_modes_bytes)); 758 CeedCallHip(ceed, hipMemcpy(diag->d_eval_modes_in, eval_modes_in, num_eval_modes_in * eval_modes_bytes, hipMemcpyHostToDevice)); 759 CeedCallHip(ceed, hipMalloc((void **)&diag->d_eval_modes_out, num_eval_modes_out * eval_modes_bytes)); 760 CeedCallHip(ceed, hipMemcpy(diag->d_eval_modes_out, eval_modes_out, num_eval_modes_out * eval_modes_bytes, hipMemcpyHostToDevice)); 761 CeedCallBackend(CeedFree(&eval_modes_in)); 762 CeedCallBackend(CeedFree(&eval_modes_out)); 763 return CEED_ERROR_SUCCESS; 764 } 765 766 //------------------------------------------------------------------------------ 767 // Assemble Diagonal Setup (Compilation) 768 //------------------------------------------------------------------------------ 769 static inline int CeedOperatorAssembleDiagonalSetupCompile_Hip(CeedOperator op, CeedInt use_ceedsize_idx, const bool is_point_block) { 770 Ceed ceed; 771 char *diagonal_kernel_path, *diagonal_kernel_source; 772 CeedInt num_input_fields, num_output_fields, num_eval_modes_in = 0, num_eval_modes_out = 0; 773 CeedInt num_comp, q_comp, num_nodes, num_qpts; 774 CeedBasis basis_in = NULL, basis_out = NULL; 775 CeedQFunctionField *qf_fields; 776 CeedQFunction qf; 777 CeedOperatorField *op_fields; 778 CeedOperator_Hip *impl; 779 780 CeedCallBackend(CeedOperatorGetCeed(op, &ceed)); 781 CeedCallBackend(CeedOperatorGetQFunction(op, &qf)); 782 CeedCallBackend(CeedQFunctionGetNumArgs(qf, &num_input_fields, &num_output_fields)); 783 784 // Determine active input basis 785 CeedCallBackend(CeedOperatorGetFields(op, NULL, &op_fields, NULL, NULL)); 786 CeedCallBackend(CeedQFunctionGetFields(qf, NULL, &qf_fields, NULL, NULL)); 787 for (CeedInt i = 0; i < num_input_fields; i++) { 788 CeedVector vec; 789 790 CeedCallBackend(CeedOperatorFieldGetVector(op_fields[i], &vec)); 791 if (vec == CEED_VECTOR_ACTIVE) { 792 CeedEvalMode eval_mode; 793 794 CeedCallBackend(CeedOperatorFieldGetBasis(op_fields[i], &basis_in)); 795 CeedCallBackend(CeedQFunctionFieldGetEvalMode(qf_fields[i], &eval_mode)); 796 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis_in, eval_mode, &q_comp)); 797 if (eval_mode != CEED_EVAL_WEIGHT) { 798 num_eval_modes_in += q_comp; 799 } 800 } 801 } 802 803 // Determine active output basis 804 CeedCallBackend(CeedOperatorGetFields(op, NULL, NULL, NULL, &op_fields)); 805 CeedCallBackend(CeedQFunctionGetFields(qf, NULL, NULL, NULL, &qf_fields)); 806 for (CeedInt i = 0; i < num_output_fields; i++) { 807 CeedVector vec; 808 809 CeedCallBackend(CeedOperatorFieldGetVector(op_fields[i], &vec)); 810 if (vec == CEED_VECTOR_ACTIVE) { 811 CeedEvalMode eval_mode; 812 813 CeedCallBackend(CeedOperatorFieldGetBasis(op_fields[i], &basis_out)); 814 CeedCallBackend(CeedQFunctionFieldGetEvalMode(qf_fields[i], &eval_mode)); 815 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis_out, eval_mode, &q_comp)); 816 if (eval_mode != CEED_EVAL_WEIGHT) { 817 num_eval_modes_out += q_comp; 818 } 819 } 820 } 821 822 // Operator data struct 823 CeedCallBackend(CeedOperatorGetData(op, &impl)); 824 CeedOperatorDiag_Hip *diag = impl->diag; 825 826 // Assemble kernel 827 hipModule_t *module = is_point_block ? &diag->module_point_block : &diag->module; 828 CeedInt elems_per_block = 1; 829 CeedCallBackend(CeedBasisGetNumNodes(basis_in, &num_nodes)); 830 CeedCallBackend(CeedBasisGetNumComponents(basis_in, &num_comp)); 831 if (basis_in == CEED_BASIS_NONE) num_qpts = num_nodes; 832 else CeedCallBackend(CeedBasisGetNumQuadraturePoints(basis_in, &num_qpts)); 833 CeedCallBackend(CeedGetJitAbsolutePath(ceed, "ceed/jit-source/hip/hip-ref-operator-assemble-diagonal.h", &diagonal_kernel_path)); 834 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Diagonal Assembly Kernel Source -----\n"); 835 CeedCallBackend(CeedLoadSourceToBuffer(ceed, diagonal_kernel_path, &diagonal_kernel_source)); 836 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Diagonal Assembly Source Complete! -----\n"); 837 CeedCallHip(ceed, CeedCompile_Hip(ceed, diagonal_kernel_source, module, 8, "NUM_EVAL_MODES_IN", num_eval_modes_in, "NUM_EVAL_MODES_OUT", 838 num_eval_modes_out, "NUM_COMP", num_comp, "NUM_NODES", num_nodes, "NUM_QPTS", num_qpts, "USE_CEEDSIZE", 839 use_ceedsize_idx, "USE_POINT_BLOCK", is_point_block ? 1 : 0, "BLOCK_SIZE", num_nodes * elems_per_block)); 840 CeedCallHip(ceed, CeedGetKernel_Hip(ceed, *module, "LinearDiagonal", is_point_block ? &diag->LinearPointBlock : &diag->LinearDiagonal)); 841 CeedCallBackend(CeedFree(&diagonal_kernel_path)); 842 CeedCallBackend(CeedFree(&diagonal_kernel_source)); 843 return CEED_ERROR_SUCCESS; 844 } 845 846 //------------------------------------------------------------------------------ 847 // Assemble Diagonal Core 848 //------------------------------------------------------------------------------ 849 static inline int CeedOperatorAssembleDiagonalCore_Hip(CeedOperator op, CeedVector assembled, CeedRequest *request, const bool is_point_block) { 850 Ceed ceed; 851 CeedInt num_elem, num_nodes; 852 CeedScalar *elem_diag_array; 853 const CeedScalar *assembled_qf_array; 854 CeedVector assembled_qf = NULL, elem_diag; 855 CeedElemRestriction assembled_rstr = NULL, rstr_in, rstr_out, diag_rstr; 856 CeedOperator_Hip *impl; 857 858 CeedCallBackend(CeedOperatorGetCeed(op, &ceed)); 859 CeedCallBackend(CeedOperatorGetData(op, &impl)); 860 861 // Assemble QFunction 862 CeedCallBackend(CeedOperatorLinearAssembleQFunctionBuildOrUpdate(op, &assembled_qf, &assembled_rstr, request)); 863 CeedCallBackend(CeedElemRestrictionDestroy(&assembled_rstr)); 864 CeedCallBackend(CeedVectorGetArrayRead(assembled_qf, CEED_MEM_DEVICE, &assembled_qf_array)); 865 866 // Setup 867 if (!impl->diag) CeedCallBackend(CeedOperatorAssembleDiagonalSetup_Hip(op)); 868 CeedOperatorDiag_Hip *diag = impl->diag; 869 870 assert(diag != NULL); 871 872 // Assemble kernel if needed 873 if ((!is_point_block && !diag->LinearDiagonal) || (is_point_block && !diag->LinearPointBlock)) { 874 CeedSize assembled_length, assembled_qf_length; 875 CeedInt use_ceedsize_idx = 0; 876 CeedCallBackend(CeedVectorGetLength(assembled, &assembled_length)); 877 CeedCallBackend(CeedVectorGetLength(assembled_qf, &assembled_qf_length)); 878 if ((assembled_length > INT_MAX) || (assembled_qf_length > INT_MAX)) use_ceedsize_idx = 1; 879 880 CeedCallBackend(CeedOperatorAssembleDiagonalSetupCompile_Hip(op, use_ceedsize_idx, is_point_block)); 881 } 882 883 // Restriction and diagonal vector 884 CeedCallBackend(CeedOperatorGetActiveElemRestrictions(op, &rstr_in, &rstr_out)); 885 CeedCheck(rstr_in == rstr_out, ceed, CEED_ERROR_BACKEND, 886 "Cannot assemble operator diagonal with different input and output active element restrictions"); 887 if (!is_point_block && !diag->diag_rstr) { 888 CeedCallBackend(CeedElemRestrictionCreateUnsignedCopy(rstr_out, &diag->diag_rstr)); 889 CeedCallBackend(CeedElemRestrictionCreateVector(diag->diag_rstr, NULL, &diag->elem_diag)); 890 } else if (is_point_block && !diag->point_block_diag_rstr) { 891 CeedCallBackend(CeedOperatorCreateActivePointBlockRestriction(rstr_out, &diag->point_block_diag_rstr)); 892 CeedCallBackend(CeedElemRestrictionCreateVector(diag->point_block_diag_rstr, NULL, &diag->point_block_elem_diag)); 893 } 894 diag_rstr = is_point_block ? diag->point_block_diag_rstr : diag->diag_rstr; 895 elem_diag = is_point_block ? diag->point_block_elem_diag : diag->elem_diag; 896 CeedCallBackend(CeedVectorSetValue(elem_diag, 0.0)); 897 898 // Only assemble diagonal if the basis has nodes, otherwise inputs are null pointers 899 CeedCallBackend(CeedElemRestrictionGetElementSize(diag_rstr, &num_nodes)); 900 if (num_nodes > 0) { 901 // Assemble element operator diagonals 902 CeedCallBackend(CeedVectorGetArray(elem_diag, CEED_MEM_DEVICE, &elem_diag_array)); 903 CeedCallBackend(CeedElemRestrictionGetNumElements(diag_rstr, &num_elem)); 904 905 // Compute the diagonal of B^T D B 906 CeedInt elems_per_block = 1; 907 CeedInt grid = CeedDivUpInt(num_elem, elems_per_block); 908 void *args[] = {(void *)&num_elem, &diag->d_identity, &diag->d_interp_in, &diag->d_grad_in, &diag->d_div_in, 909 &diag->d_curl_in, &diag->d_interp_out, &diag->d_grad_out, &diag->d_div_out, &diag->d_curl_out, 910 &diag->d_eval_modes_in, &diag->d_eval_modes_out, &assembled_qf_array, &elem_diag_array}; 911 912 if (is_point_block) { 913 CeedCallBackend(CeedRunKernelDim_Hip(ceed, diag->LinearPointBlock, grid, num_nodes, 1, elems_per_block, args)); 914 } else { 915 CeedCallBackend(CeedRunKernelDim_Hip(ceed, diag->LinearDiagonal, grid, num_nodes, 1, elems_per_block, args)); 916 } 917 918 // Restore arrays 919 CeedCallBackend(CeedVectorRestoreArray(elem_diag, &elem_diag_array)); 920 CeedCallBackend(CeedVectorRestoreArrayRead(assembled_qf, &assembled_qf_array)); 921 } 922 923 // Assemble local operator diagonal 924 CeedCallBackend(CeedElemRestrictionApply(diag_rstr, CEED_TRANSPOSE, elem_diag, assembled, request)); 925 926 // Cleanup 927 CeedCallBackend(CeedVectorDestroy(&assembled_qf)); 928 return CEED_ERROR_SUCCESS; 929 } 930 931 //------------------------------------------------------------------------------ 932 // Assemble Linear Diagonal 933 //------------------------------------------------------------------------------ 934 static int CeedOperatorLinearAssembleAddDiagonal_Hip(CeedOperator op, CeedVector assembled, CeedRequest *request) { 935 CeedCallBackend(CeedOperatorAssembleDiagonalCore_Hip(op, assembled, request, false)); 936 return CEED_ERROR_SUCCESS; 937 } 938 939 //------------------------------------------------------------------------------ 940 // Assemble Linear Point Block Diagonal 941 //------------------------------------------------------------------------------ 942 static int CeedOperatorLinearAssembleAddPointBlockDiagonal_Hip(CeedOperator op, CeedVector assembled, CeedRequest *request) { 943 CeedCallBackend(CeedOperatorAssembleDiagonalCore_Hip(op, assembled, request, true)); 944 return CEED_ERROR_SUCCESS; 945 } 946 947 //------------------------------------------------------------------------------ 948 // Single Operator Assembly Setup 949 //------------------------------------------------------------------------------ 950 static int CeedSingleOperatorAssembleSetup_Hip(CeedOperator op, CeedInt use_ceedsize_idx) { 951 Ceed ceed; 952 char *assembly_kernel_path, *assembly_kernel_source; 953 CeedInt num_input_fields, num_output_fields, num_eval_modes_in = 0, num_eval_modes_out = 0; 954 CeedInt elem_size_in, num_qpts_in, num_comp_in, elem_size_out, num_qpts_out, num_comp_out, q_comp; 955 CeedEvalMode *eval_modes_in = NULL, *eval_modes_out = NULL; 956 CeedElemRestriction rstr_in = NULL, rstr_out = NULL; 957 CeedBasis basis_in = NULL, basis_out = NULL; 958 CeedQFunctionField *qf_fields; 959 CeedQFunction qf; 960 CeedOperatorField *input_fields, *output_fields; 961 CeedOperator_Hip *impl; 962 963 CeedCallBackend(CeedOperatorGetCeed(op, &ceed)); 964 CeedCallBackend(CeedOperatorGetData(op, &impl)); 965 966 // Get intput and output fields 967 CeedCallBackend(CeedOperatorGetFields(op, &num_input_fields, &input_fields, &num_output_fields, &output_fields)); 968 969 // Determine active input basis eval mode 970 CeedCallBackend(CeedOperatorGetQFunction(op, &qf)); 971 CeedCallBackend(CeedQFunctionGetFields(qf, NULL, &qf_fields, NULL, NULL)); 972 for (CeedInt i = 0; i < num_input_fields; i++) { 973 CeedVector vec; 974 975 CeedCallBackend(CeedOperatorFieldGetVector(input_fields[i], &vec)); 976 if (vec == CEED_VECTOR_ACTIVE) { 977 CeedBasis basis; 978 CeedEvalMode eval_mode; 979 980 CeedCallBackend(CeedOperatorFieldGetBasis(input_fields[i], &basis)); 981 CeedCheck(!basis_in || basis_in == basis, ceed, CEED_ERROR_BACKEND, "Backend does not implement operator assembly with multiple active bases"); 982 basis_in = basis; 983 CeedCallBackend(CeedOperatorFieldGetElemRestriction(input_fields[i], &rstr_in)); 984 CeedCallBackend(CeedElemRestrictionGetElementSize(rstr_in, &elem_size_in)); 985 if (basis_in == CEED_BASIS_NONE) num_qpts_in = elem_size_in; 986 else CeedCallBackend(CeedBasisGetNumQuadraturePoints(basis_in, &num_qpts_in)); 987 CeedCallBackend(CeedQFunctionFieldGetEvalMode(qf_fields[i], &eval_mode)); 988 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis_in, eval_mode, &q_comp)); 989 if (eval_mode != CEED_EVAL_WEIGHT) { 990 // q_comp = 1 if CEED_EVAL_NONE, CEED_EVAL_WEIGHT caught by QF Assembly 991 CeedCallBackend(CeedRealloc(num_eval_modes_in + q_comp, &eval_modes_in)); 992 for (CeedInt d = 0; d < q_comp; d++) { 993 eval_modes_in[num_eval_modes_in + d] = eval_mode; 994 } 995 num_eval_modes_in += q_comp; 996 } 997 } 998 } 999 1000 // Determine active output basis; basis_out and rstr_out only used if same as input, TODO 1001 CeedCallBackend(CeedQFunctionGetFields(qf, NULL, NULL, NULL, &qf_fields)); 1002 for (CeedInt i = 0; i < num_output_fields; i++) { 1003 CeedVector vec; 1004 1005 CeedCallBackend(CeedOperatorFieldGetVector(output_fields[i], &vec)); 1006 if (vec == CEED_VECTOR_ACTIVE) { 1007 CeedBasis basis; 1008 CeedEvalMode eval_mode; 1009 1010 CeedCallBackend(CeedOperatorFieldGetBasis(output_fields[i], &basis)); 1011 CeedCheck(!basis_out || basis_out == basis, ceed, CEED_ERROR_BACKEND, 1012 "Backend does not implement operator assembly with multiple active bases"); 1013 basis_out = basis; 1014 CeedCallBackend(CeedOperatorFieldGetElemRestriction(output_fields[i], &rstr_out)); 1015 CeedCallBackend(CeedElemRestrictionGetElementSize(rstr_out, &elem_size_out)); 1016 if (basis_out == CEED_BASIS_NONE) num_qpts_out = elem_size_out; 1017 else CeedCallBackend(CeedBasisGetNumQuadraturePoints(basis_out, &num_qpts_out)); 1018 CeedCheck(num_qpts_in == num_qpts_out, ceed, CEED_ERROR_UNSUPPORTED, 1019 "Active input and output bases must have the same number of quadrature points"); 1020 CeedCallBackend(CeedQFunctionFieldGetEvalMode(qf_fields[i], &eval_mode)); 1021 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis_out, eval_mode, &q_comp)); 1022 if (eval_mode != CEED_EVAL_WEIGHT) { 1023 // q_comp = 1 if CEED_EVAL_NONE, CEED_EVAL_WEIGHT caught by QF Assembly 1024 CeedCallBackend(CeedRealloc(num_eval_modes_out + q_comp, &eval_modes_out)); 1025 for (CeedInt d = 0; d < q_comp; d++) { 1026 eval_modes_out[num_eval_modes_out + d] = eval_mode; 1027 } 1028 num_eval_modes_out += q_comp; 1029 } 1030 } 1031 } 1032 CeedCheck(num_eval_modes_in > 0 && num_eval_modes_out > 0, ceed, CEED_ERROR_UNSUPPORTED, "Cannot assemble operator without inputs/outputs"); 1033 1034 CeedCallBackend(CeedCalloc(1, &impl->asmb)); 1035 CeedOperatorAssemble_Hip *asmb = impl->asmb; 1036 asmb->elems_per_block = 1; 1037 asmb->block_size_x = elem_size_in; 1038 asmb->block_size_y = elem_size_out; 1039 1040 bool fallback = asmb->block_size_x * asmb->block_size_y * asmb->elems_per_block > 1024; 1041 1042 if (fallback) { 1043 // Use fallback kernel with 1D threadblock 1044 asmb->block_size_y = 1; 1045 } 1046 1047 // Compile kernels 1048 CeedCallBackend(CeedElemRestrictionGetNumComponents(rstr_in, &num_comp_in)); 1049 CeedCallBackend(CeedElemRestrictionGetNumComponents(rstr_out, &num_comp_out)); 1050 CeedCallBackend(CeedGetJitAbsolutePath(ceed, "ceed/jit-source/hip/hip-ref-operator-assemble.h", &assembly_kernel_path)); 1051 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Assembly Kernel Source -----\n"); 1052 CeedCallBackend(CeedLoadSourceToBuffer(ceed, assembly_kernel_path, &assembly_kernel_source)); 1053 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Assembly Source Complete! -----\n"); 1054 CeedCallBackend(CeedCompile_Hip(ceed, assembly_kernel_source, &asmb->module, 10, "NUM_EVAL_MODES_IN", num_eval_modes_in, "NUM_EVAL_MODES_OUT", 1055 num_eval_modes_out, "NUM_COMP_IN", num_comp_in, "NUM_COMP_OUT", num_comp_out, "NUM_NODES_IN", elem_size_in, 1056 "NUM_NODES_OUT", elem_size_out, "NUM_QPTS", num_qpts_in, "BLOCK_SIZE", 1057 asmb->block_size_x * asmb->block_size_y * asmb->elems_per_block, "BLOCK_SIZE_Y", asmb->block_size_y, "USE_CEEDSIZE", 1058 use_ceedsize_idx)); 1059 CeedCallBackend(CeedGetKernel_Hip(ceed, asmb->module, "LinearAssemble", &asmb->LinearAssemble)); 1060 CeedCallBackend(CeedFree(&assembly_kernel_path)); 1061 CeedCallBackend(CeedFree(&assembly_kernel_source)); 1062 1063 // Load into B_in, in order that they will be used in eval_modes_in 1064 { 1065 const CeedInt in_bytes = elem_size_in * num_qpts_in * num_eval_modes_in * sizeof(CeedScalar); 1066 CeedInt d_in = 0; 1067 CeedEvalMode eval_modes_in_prev = CEED_EVAL_NONE; 1068 bool has_eval_none = false; 1069 CeedScalar *identity = NULL; 1070 1071 for (CeedInt i = 0; i < num_eval_modes_in; i++) { 1072 has_eval_none = has_eval_none || (eval_modes_in[i] == CEED_EVAL_NONE); 1073 } 1074 if (has_eval_none) { 1075 CeedCallBackend(CeedCalloc(elem_size_in * num_qpts_in, &identity)); 1076 for (CeedInt i = 0; i < (elem_size_in < num_qpts_in ? elem_size_in : num_qpts_in); i++) identity[i * elem_size_in + i] = 1.0; 1077 } 1078 1079 CeedCallHip(ceed, hipMalloc((void **)&asmb->d_B_in, in_bytes)); 1080 for (CeedInt i = 0; i < num_eval_modes_in; i++) { 1081 const CeedScalar *h_B_in; 1082 1083 CeedCallBackend(CeedOperatorGetBasisPointer(basis_in, eval_modes_in[i], identity, &h_B_in)); 1084 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis_in, eval_modes_in[i], &q_comp)); 1085 if (q_comp > 1) { 1086 if (i == 0 || eval_modes_in[i] != eval_modes_in_prev) d_in = 0; 1087 else h_B_in = &h_B_in[(++d_in) * elem_size_in * num_qpts_in]; 1088 } 1089 eval_modes_in_prev = eval_modes_in[i]; 1090 1091 CeedCallHip(ceed, hipMemcpy(&asmb->d_B_in[i * elem_size_in * num_qpts_in], h_B_in, elem_size_in * num_qpts_in * sizeof(CeedScalar), 1092 hipMemcpyHostToDevice)); 1093 } 1094 1095 if (identity) { 1096 CeedCallBackend(CeedFree(&identity)); 1097 } 1098 } 1099 1100 // Load into B_out, in order that they will be used in eval_modes_out 1101 { 1102 const CeedInt out_bytes = elem_size_out * num_qpts_out * num_eval_modes_out * sizeof(CeedScalar); 1103 CeedInt d_out = 0; 1104 CeedEvalMode eval_modes_out_prev = CEED_EVAL_NONE; 1105 bool has_eval_none = false; 1106 CeedScalar *identity = NULL; 1107 1108 for (CeedInt i = 0; i < num_eval_modes_out; i++) { 1109 has_eval_none = has_eval_none || (eval_modes_out[i] == CEED_EVAL_NONE); 1110 } 1111 if (has_eval_none) { 1112 CeedCallBackend(CeedCalloc(elem_size_out * num_qpts_out, &identity)); 1113 for (CeedInt i = 0; i < (elem_size_out < num_qpts_out ? elem_size_out : num_qpts_out); i++) identity[i * elem_size_out + i] = 1.0; 1114 } 1115 1116 CeedCallHip(ceed, hipMalloc((void **)&asmb->d_B_out, out_bytes)); 1117 for (CeedInt i = 0; i < num_eval_modes_out; i++) { 1118 const CeedScalar *h_B_out; 1119 1120 CeedCallBackend(CeedOperatorGetBasisPointer(basis_out, eval_modes_out[i], identity, &h_B_out)); 1121 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis_out, eval_modes_out[i], &q_comp)); 1122 if (q_comp > 1) { 1123 if (i == 0 || eval_modes_out[i] != eval_modes_out_prev) d_out = 0; 1124 else h_B_out = &h_B_out[(++d_out) * elem_size_out * num_qpts_out]; 1125 } 1126 eval_modes_out_prev = eval_modes_out[i]; 1127 1128 CeedCallHip(ceed, hipMemcpy(&asmb->d_B_out[i * elem_size_out * num_qpts_out], h_B_out, elem_size_out * num_qpts_out * sizeof(CeedScalar), 1129 hipMemcpyHostToDevice)); 1130 } 1131 1132 if (identity) { 1133 CeedCallBackend(CeedFree(&identity)); 1134 } 1135 } 1136 return CEED_ERROR_SUCCESS; 1137 } 1138 1139 //------------------------------------------------------------------------------ 1140 // Assemble matrix data for COO matrix of assembled operator. 1141 // The sparsity pattern is set by CeedOperatorLinearAssembleSymbolic. 1142 // 1143 // Note that this (and other assembly routines) currently assume only one active input restriction/basis per operator (could have multiple basis eval 1144 // modes). 1145 // TODO: allow multiple active input restrictions/basis objects 1146 //------------------------------------------------------------------------------ 1147 static int CeedSingleOperatorAssemble_Hip(CeedOperator op, CeedInt offset, CeedVector values) { 1148 Ceed ceed; 1149 CeedSize values_length = 0, assembled_qf_length = 0; 1150 CeedInt use_ceedsize_idx = 0, num_elem_in, num_elem_out, elem_size_in, elem_size_out; 1151 CeedScalar *values_array; 1152 const CeedScalar *assembled_qf_array; 1153 CeedVector assembled_qf = NULL; 1154 CeedElemRestriction assembled_rstr = NULL, rstr_in, rstr_out; 1155 CeedRestrictionType rstr_type_in, rstr_type_out; 1156 const bool *orients_in = NULL, *orients_out = NULL; 1157 const CeedInt8 *curl_orients_in = NULL, *curl_orients_out = NULL; 1158 CeedOperator_Hip *impl; 1159 1160 CeedCallBackend(CeedOperatorGetCeed(op, &ceed)); 1161 CeedCallBackend(CeedOperatorGetData(op, &impl)); 1162 1163 // Assemble QFunction 1164 CeedCallBackend(CeedOperatorLinearAssembleQFunctionBuildOrUpdate(op, &assembled_qf, &assembled_rstr, CEED_REQUEST_IMMEDIATE)); 1165 CeedCallBackend(CeedElemRestrictionDestroy(&assembled_rstr)); 1166 CeedCallBackend(CeedVectorGetArrayRead(assembled_qf, CEED_MEM_DEVICE, &assembled_qf_array)); 1167 1168 CeedCallBackend(CeedVectorGetLength(values, &values_length)); 1169 CeedCallBackend(CeedVectorGetLength(assembled_qf, &assembled_qf_length)); 1170 if ((values_length > INT_MAX) || (assembled_qf_length > INT_MAX)) use_ceedsize_idx = 1; 1171 1172 // Setup 1173 if (!impl->asmb) CeedCallBackend(CeedSingleOperatorAssembleSetup_Hip(op, use_ceedsize_idx)); 1174 CeedOperatorAssemble_Hip *asmb = impl->asmb; 1175 1176 assert(asmb != NULL); 1177 1178 // Assemble element operator 1179 CeedCallBackend(CeedVectorGetArray(values, CEED_MEM_DEVICE, &values_array)); 1180 values_array += offset; 1181 1182 CeedCallBackend(CeedOperatorGetActiveElemRestrictions(op, &rstr_in, &rstr_out)); 1183 CeedCallBackend(CeedElemRestrictionGetNumElements(rstr_in, &num_elem_in)); 1184 CeedCallBackend(CeedElemRestrictionGetElementSize(rstr_in, &elem_size_in)); 1185 1186 CeedCallBackend(CeedElemRestrictionGetType(rstr_in, &rstr_type_in)); 1187 if (rstr_type_in == CEED_RESTRICTION_ORIENTED) { 1188 CeedCallBackend(CeedElemRestrictionGetOrientations(rstr_in, CEED_MEM_DEVICE, &orients_in)); 1189 } else if (rstr_type_in == CEED_RESTRICTION_CURL_ORIENTED) { 1190 CeedCallBackend(CeedElemRestrictionGetCurlOrientations(rstr_in, CEED_MEM_DEVICE, &curl_orients_in)); 1191 } 1192 1193 if (rstr_in != rstr_out) { 1194 CeedCallBackend(CeedElemRestrictionGetNumElements(rstr_out, &num_elem_out)); 1195 CeedCheck(num_elem_in == num_elem_out, ceed, CEED_ERROR_UNSUPPORTED, 1196 "Active input and output operator restrictions must have the same number of elements"); 1197 CeedCallBackend(CeedElemRestrictionGetElementSize(rstr_out, &elem_size_out)); 1198 1199 CeedCallBackend(CeedElemRestrictionGetType(rstr_out, &rstr_type_out)); 1200 if (rstr_type_out == CEED_RESTRICTION_ORIENTED) { 1201 CeedCallBackend(CeedElemRestrictionGetOrientations(rstr_out, CEED_MEM_DEVICE, &orients_out)); 1202 } else if (rstr_type_out == CEED_RESTRICTION_CURL_ORIENTED) { 1203 CeedCallBackend(CeedElemRestrictionGetCurlOrientations(rstr_out, CEED_MEM_DEVICE, &curl_orients_out)); 1204 } 1205 } else { 1206 elem_size_out = elem_size_in; 1207 orients_out = orients_in; 1208 curl_orients_out = curl_orients_in; 1209 } 1210 1211 // Compute B^T D B 1212 CeedInt shared_mem = 1213 ((curl_orients_in || curl_orients_out ? elem_size_in * elem_size_out : 0) + (curl_orients_in ? elem_size_in * asmb->block_size_y : 0)) * 1214 sizeof(CeedScalar); 1215 CeedInt grid = CeedDivUpInt(num_elem_in, asmb->elems_per_block); 1216 void *args[] = {(void *)&num_elem_in, &asmb->d_B_in, &asmb->d_B_out, &orients_in, &curl_orients_in, 1217 &orients_out, &curl_orients_out, &assembled_qf_array, &values_array}; 1218 1219 CeedCallBackend( 1220 CeedRunKernelDimShared_Hip(ceed, asmb->LinearAssemble, grid, asmb->block_size_x, asmb->block_size_y, asmb->elems_per_block, shared_mem, args)); 1221 1222 // Restore arrays 1223 CeedCallBackend(CeedVectorRestoreArray(values, &values_array)); 1224 CeedCallBackend(CeedVectorRestoreArrayRead(assembled_qf, &assembled_qf_array)); 1225 1226 // Cleanup 1227 CeedCallBackend(CeedVectorDestroy(&assembled_qf)); 1228 if (rstr_type_in == CEED_RESTRICTION_ORIENTED) { 1229 CeedCallBackend(CeedElemRestrictionRestoreOrientations(rstr_in, &orients_in)); 1230 } else if (rstr_type_in == CEED_RESTRICTION_CURL_ORIENTED) { 1231 CeedCallBackend(CeedElemRestrictionRestoreCurlOrientations(rstr_in, &curl_orients_in)); 1232 } 1233 if (rstr_in != rstr_out) { 1234 if (rstr_type_out == CEED_RESTRICTION_ORIENTED) { 1235 CeedCallBackend(CeedElemRestrictionRestoreOrientations(rstr_out, &orients_out)); 1236 } else if (rstr_type_out == CEED_RESTRICTION_CURL_ORIENTED) { 1237 CeedCallBackend(CeedElemRestrictionRestoreCurlOrientations(rstr_out, &curl_orients_out)); 1238 } 1239 } 1240 return CEED_ERROR_SUCCESS; 1241 } 1242 1243 //------------------------------------------------------------------------------ 1244 // Create operator 1245 //------------------------------------------------------------------------------ 1246 int CeedOperatorCreate_Hip(CeedOperator op) { 1247 Ceed ceed; 1248 CeedOperator_Hip *impl; 1249 1250 CeedCallBackend(CeedOperatorGetCeed(op, &ceed)); 1251 CeedCallBackend(CeedCalloc(1, &impl)); 1252 CeedCallBackend(CeedOperatorSetData(op, impl)); 1253 CeedCallBackend(CeedSetBackendFunction(ceed, "Operator", op, "LinearAssembleQFunction", CeedOperatorLinearAssembleQFunction_Hip)); 1254 CeedCallBackend(CeedSetBackendFunction(ceed, "Operator", op, "LinearAssembleQFunctionUpdate", CeedOperatorLinearAssembleQFunctionUpdate_Hip)); 1255 CeedCallBackend(CeedSetBackendFunction(ceed, "Operator", op, "LinearAssembleAddDiagonal", CeedOperatorLinearAssembleAddDiagonal_Hip)); 1256 CeedCallBackend( 1257 CeedSetBackendFunction(ceed, "Operator", op, "LinearAssembleAddPointBlockDiagonal", CeedOperatorLinearAssembleAddPointBlockDiagonal_Hip)); 1258 CeedCallBackend(CeedSetBackendFunction(ceed, "Operator", op, "LinearAssembleSingle", CeedSingleOperatorAssemble_Hip)); 1259 CeedCallBackend(CeedSetBackendFunction(ceed, "Operator", op, "ApplyAdd", CeedOperatorApplyAdd_Hip)); 1260 CeedCallBackend(CeedSetBackendFunction(ceed, "Operator", op, "Destroy", CeedOperatorDestroy_Hip)); 1261 return CEED_ERROR_SUCCESS; 1262 } 1263 1264 //------------------------------------------------------------------------------ 1265