59   if (impl->diag) {  in CeedOperatorDestroy_Hip()
63     if (impl->diag->module) {  in CeedOperatorDestroy_Hip()
64       CeedCallHip(ceed, hipModuleUnload(impl->diag->module));  in CeedOperatorDestroy_Hip()
66     if (impl->diag->module_point_block) {  in CeedOperatorDestroy_Hip()
67       CeedCallHip(ceed, hipModuleUnload(impl->diag->module_point_block));  in CeedOperatorDestroy_Hip()
69     CeedCallHip(ceed, hipFree(impl->diag->d_eval_modes_in));  in CeedOperatorDestroy_Hip()
70     CeedCallHip(ceed, hipFree(impl->diag->d_eval_modes_out));  in CeedOperatorDestroy_Hip()
71     CeedCallHip(ceed, hipFree(impl->diag->d_identity));  in CeedOperatorDestroy_Hip()
72     CeedCallHip(ceed, hipFree(impl->diag->d_interp_in));  in CeedOperatorDestroy_Hip()
73     CeedCallHip(ceed, hipFree(impl->diag->d_interp_out));  in CeedOperatorDestroy_Hip()
74     CeedCallHip(ceed, hipFree(impl->diag->d_grad_in));  in CeedOperatorDestroy_Hip()
75     CeedCallHip(ceed, hipFree(impl->diag->d_grad_out));  in CeedOperatorDestroy_Hip()
76     CeedCallHip(ceed, hipFree(impl->diag->d_div_in));  in CeedOperatorDestroy_Hip()
77     CeedCallHip(ceed, hipFree(impl->diag->d_div_out));  in CeedOperatorDestroy_Hip()
78     CeedCallHip(ceed, hipFree(impl->diag->d_curl_in));  in CeedOperatorDestroy_Hip()
79     CeedCallHip(ceed, hipFree(impl->diag->d_curl_out));  in CeedOperatorDestroy_Hip()
81     CeedCallBackend(CeedVectorDestroy(&impl->diag->elem_diag));  in CeedOperatorDestroy_Hip()
82     CeedCallBackend(CeedVectorDestroy(&impl->diag->point_block_elem_diag));  in CeedOperatorDestroy_Hip()
83     CeedCallBackend(CeedElemRestrictionDestroy(&impl->diag->diag_rstr));  in CeedOperatorDestroy_Hip()
84     CeedCallBackend(CeedElemRestrictionDestroy(&impl->diag->point_block_diag_rstr));  in CeedOperatorDestroy_Hip()
86   CeedCallBackend(CeedFree(&impl->diag));  in CeedOperatorDestroy_Hip()
1199   CeedCallBackend(CeedCalloc(1, &impl->diag));  in CeedOperatorAssembleDiagonalSetup_Hip()
1200   CeedOperatorDiag_Hip *diag = impl->diag;  in CeedOperatorAssembleDiagonalSetup_Hip()  local
1218     CeedCallHip(ceed, hipMalloc((void **)&diag->d_identity, interp_bytes));  in CeedOperatorAssembleDiagonalSetup_Hip()
1219     CeedCallHip(ceed, hipMemcpy(diag->d_identity, identity, interp_bytes, hipMemcpyHostToDevice));  in CeedOperatorAssembleDiagonalSetup_Hip()
1245           diag->d_interp_in = d_interp;  in CeedOperatorAssembleDiagonalSetup_Hip()
1246           diag->d_grad_in   = d_grad;  in CeedOperatorAssembleDiagonalSetup_Hip()
1248           diag->d_interp_out = d_interp;  in CeedOperatorAssembleDiagonalSetup_Hip()
1249           diag->d_grad_out   = d_grad;  in CeedOperatorAssembleDiagonalSetup_Hip()
1267           diag->d_interp_in = d_interp;  in CeedOperatorAssembleDiagonalSetup_Hip()
1268           diag->d_div_in    = d_div;  in CeedOperatorAssembleDiagonalSetup_Hip()
1270           diag->d_interp_out = d_interp;  in CeedOperatorAssembleDiagonalSetup_Hip()
1271           diag->d_div_out    = d_div;  in CeedOperatorAssembleDiagonalSetup_Hip()
1289           diag->d_interp_in = d_interp;  in CeedOperatorAssembleDiagonalSetup_Hip()
1290           diag->d_curl_in   = d_curl;  in CeedOperatorAssembleDiagonalSetup_Hip()
1292           diag->d_interp_out = d_interp;  in CeedOperatorAssembleDiagonalSetup_Hip()
1293           diag->d_curl_out   = d_curl;  in CeedOperatorAssembleDiagonalSetup_Hip()
1300 …CeedCallHip(ceed, hipMalloc((void **)&diag->d_eval_modes_in, num_eval_modes_in * eval_modes_bytes)…  in CeedOperatorAssembleDiagonalSetup_Hip()
1301 …CeedCallHip(ceed, hipMemcpy(diag->d_eval_modes_in, eval_modes_in, num_eval_modes_in * eval_modes_b…  in CeedOperatorAssembleDiagonalSetup_Hip()
1302 …CeedCallHip(ceed, hipMalloc((void **)&diag->d_eval_modes_out, num_eval_modes_out * eval_modes_byte…  in CeedOperatorAssembleDiagonalSetup_Hip()
1303 …CeedCallHip(ceed, hipMemcpy(diag->d_eval_modes_out, eval_modes_out, num_eval_modes_out * eval_mode…  in CeedOperatorAssembleDiagonalSetup_Hip()
1378   CeedOperatorDiag_Hip *diag = impl->diag;  in CeedOperatorAssembleDiagonalSetupCompile_Hip()  local
1382 …hipModule_t *module                   = is_point_block ? &diag->module_point_block : &diag->module;  in CeedOperatorAssembleDiagonalSetupCompile_Hip()
1392 …ernel_Hip(ceed, *module, "LinearDiagonal", is_point_block ? &diag->LinearPointBlock : &diag->Linea…  in CeedOperatorAssembleDiagonalSetupCompile_Hip()
1421   if (!impl->diag) CeedCallBackend(CeedOperatorAssembleDiagonalSetup_Hip(op));  in CeedOperatorAssembleDiagonalCore_Hip()
1422   CeedOperatorDiag_Hip *diag = impl->diag;  in CeedOperatorAssembleDiagonalCore_Hip()  local
1424   assert(diag != NULL);  in CeedOperatorAssembleDiagonalCore_Hip()
1427   if ((!is_point_block && !diag->LinearDiagonal) || (is_point_block && !diag->LinearPointBlock)) {  in CeedOperatorAssembleDiagonalCore_Hip()
1441   if (!is_point_block && !diag->diag_rstr) {  in CeedOperatorAssembleDiagonalCore_Hip()
1442     CeedCallBackend(CeedElemRestrictionCreateUnsignedCopy(rstr_out, &diag->diag_rstr));  in CeedOperatorAssembleDiagonalCore_Hip()
1443     CeedCallBackend(CeedElemRestrictionCreateVector(diag->diag_rstr, NULL, &diag->elem_diag));  in CeedOperatorAssembleDiagonalCore_Hip()
1444   } else if (is_point_block && !diag->point_block_diag_rstr) {  in CeedOperatorAssembleDiagonalCore_Hip()
1445 …CeedCallBackend(CeedOperatorCreateActivePointBlockRestriction(rstr_out, &diag->point_block_diag_rs…  in CeedOperatorAssembleDiagonalCore_Hip()
1446 …CeedCallBackend(CeedElemRestrictionCreateVector(diag->point_block_diag_rstr, NULL, &diag->point_bl…  in CeedOperatorAssembleDiagonalCore_Hip()
1450   diag_rstr = is_point_block ? diag->point_block_diag_rstr : diag->diag_rstr;  in CeedOperatorAssembleDiagonalCore_Hip()
1451   elem_diag = is_point_block ? diag->point_block_elem_diag : diag->elem_diag;  in CeedOperatorAssembleDiagonalCore_Hip()
1464 …gs[]          = {(void *)&num_elem,      &diag->d_identity,       &diag->d_interp_in,  &diag->d_gr…  in CeedOperatorAssembleDiagonalCore_Hip()
1465 …&diag->d_curl_in,       &diag->d_interp_out,     &diag->d_grad_out,   &diag->d_div_out, &diag->d_c…  in CeedOperatorAssembleDiagonalCore_Hip()
1466 …           &diag->d_eval_modes_in, &diag->d_eval_modes_out, &assembled_qf_array, &elem_diag_array};  in CeedOperatorAssembleDiagonalCore_Hip()
1469 …CeedCallBackend(CeedRunKernelDim_Hip(ceed, diag->LinearPointBlock, grid, num_nodes, 1, elems_per_b…  in CeedOperatorAssembleDiagonalCore_Hip()
1471 …CeedCallBackend(CeedRunKernelDim_Hip(ceed, diag->LinearDiagonal, grid, num_nodes, 1, elems_per_blo…  in CeedOperatorAssembleDiagonalCore_Hip()