Lines Matching refs:diag
60 if (impl->diag) { in CeedOperatorDestroy_Cuda()
64 if (impl->diag->module) { in CeedOperatorDestroy_Cuda()
65 CeedCallCuda(ceed, cuModuleUnload(impl->diag->module)); in CeedOperatorDestroy_Cuda()
67 if (impl->diag->module_point_block) { in CeedOperatorDestroy_Cuda()
68 CeedCallCuda(ceed, cuModuleUnload(impl->diag->module_point_block)); in CeedOperatorDestroy_Cuda()
70 CeedCallCuda(ceed, cudaFree(impl->diag->d_eval_modes_in)); in CeedOperatorDestroy_Cuda()
71 CeedCallCuda(ceed, cudaFree(impl->diag->d_eval_modes_out)); in CeedOperatorDestroy_Cuda()
72 CeedCallCuda(ceed, cudaFree(impl->diag->d_identity)); in CeedOperatorDestroy_Cuda()
73 CeedCallCuda(ceed, cudaFree(impl->diag->d_interp_in)); in CeedOperatorDestroy_Cuda()
74 CeedCallCuda(ceed, cudaFree(impl->diag->d_interp_out)); in CeedOperatorDestroy_Cuda()
75 CeedCallCuda(ceed, cudaFree(impl->diag->d_grad_in)); in CeedOperatorDestroy_Cuda()
76 CeedCallCuda(ceed, cudaFree(impl->diag->d_grad_out)); in CeedOperatorDestroy_Cuda()
77 CeedCallCuda(ceed, cudaFree(impl->diag->d_div_in)); in CeedOperatorDestroy_Cuda()
78 CeedCallCuda(ceed, cudaFree(impl->diag->d_div_out)); in CeedOperatorDestroy_Cuda()
79 CeedCallCuda(ceed, cudaFree(impl->diag->d_curl_in)); in CeedOperatorDestroy_Cuda()
80 CeedCallCuda(ceed, cudaFree(impl->diag->d_curl_out)); in CeedOperatorDestroy_Cuda()
82 CeedCallBackend(CeedVectorDestroy(&impl->diag->elem_diag)); in CeedOperatorDestroy_Cuda()
83 CeedCallBackend(CeedVectorDestroy(&impl->diag->point_block_elem_diag)); in CeedOperatorDestroy_Cuda()
84 CeedCallBackend(CeedElemRestrictionDestroy(&impl->diag->diag_rstr)); in CeedOperatorDestroy_Cuda()
85 CeedCallBackend(CeedElemRestrictionDestroy(&impl->diag->point_block_diag_rstr)); in CeedOperatorDestroy_Cuda()
87 CeedCallBackend(CeedFree(&impl->diag)); in CeedOperatorDestroy_Cuda()
1202 CeedCallBackend(CeedCalloc(1, &impl->diag)); in CeedOperatorAssembleDiagonalSetup_Cuda()
1203 CeedOperatorDiag_Cuda *diag = impl->diag; in CeedOperatorAssembleDiagonalSetup_Cuda() local
1221 CeedCallCuda(ceed, cudaMalloc((void **)&diag->d_identity, interp_bytes)); in CeedOperatorAssembleDiagonalSetup_Cuda()
1222 … CeedCallCuda(ceed, cudaMemcpy(diag->d_identity, identity, interp_bytes, cudaMemcpyHostToDevice)); in CeedOperatorAssembleDiagonalSetup_Cuda()
1248 diag->d_interp_in = d_interp; in CeedOperatorAssembleDiagonalSetup_Cuda()
1249 diag->d_grad_in = d_grad; in CeedOperatorAssembleDiagonalSetup_Cuda()
1251 diag->d_interp_out = d_interp; in CeedOperatorAssembleDiagonalSetup_Cuda()
1252 diag->d_grad_out = d_grad; in CeedOperatorAssembleDiagonalSetup_Cuda()
1270 diag->d_interp_in = d_interp; in CeedOperatorAssembleDiagonalSetup_Cuda()
1271 diag->d_div_in = d_div; in CeedOperatorAssembleDiagonalSetup_Cuda()
1273 diag->d_interp_out = d_interp; in CeedOperatorAssembleDiagonalSetup_Cuda()
1274 diag->d_div_out = d_div; in CeedOperatorAssembleDiagonalSetup_Cuda()
1292 diag->d_interp_in = d_interp; in CeedOperatorAssembleDiagonalSetup_Cuda()
1293 diag->d_curl_in = d_curl; in CeedOperatorAssembleDiagonalSetup_Cuda()
1295 diag->d_interp_out = d_interp; in CeedOperatorAssembleDiagonalSetup_Cuda()
1296 diag->d_curl_out = d_curl; in CeedOperatorAssembleDiagonalSetup_Cuda()
1303 …CeedCallCuda(ceed, cudaMalloc((void **)&diag->d_eval_modes_in, num_eval_modes_in * eval_modes_byte… in CeedOperatorAssembleDiagonalSetup_Cuda()
1304 …CeedCallCuda(ceed, cudaMemcpy(diag->d_eval_modes_in, eval_modes_in, num_eval_modes_in * eval_modes… in CeedOperatorAssembleDiagonalSetup_Cuda()
1305 …CeedCallCuda(ceed, cudaMalloc((void **)&diag->d_eval_modes_out, num_eval_modes_out * eval_modes_by… in CeedOperatorAssembleDiagonalSetup_Cuda()
1306 …CeedCallCuda(ceed, cudaMemcpy(diag->d_eval_modes_out, eval_modes_out, num_eval_modes_out * eval_mo… in CeedOperatorAssembleDiagonalSetup_Cuda()
1381 CeedOperatorDiag_Cuda *diag = impl->diag; in CeedOperatorAssembleDiagonalSetupCompile_Cuda() local
1385 CUmodule *module = is_point_block ? &diag->module_point_block : &diag->module; in CeedOperatorAssembleDiagonalSetupCompile_Cuda()
1395 …rnel_Cuda(ceed, *module, "LinearDiagonal", is_point_block ? &diag->LinearPointBlock : &diag->Linea… in CeedOperatorAssembleDiagonalSetupCompile_Cuda()
1424 if (!impl->diag) CeedCallBackend(CeedOperatorAssembleDiagonalSetup_Cuda(op)); in CeedOperatorAssembleDiagonalCore_Cuda()
1425 CeedOperatorDiag_Cuda *diag = impl->diag; in CeedOperatorAssembleDiagonalCore_Cuda() local
1427 assert(diag != NULL); in CeedOperatorAssembleDiagonalCore_Cuda()
1430 if ((!is_point_block && !diag->LinearDiagonal) || (is_point_block && !diag->LinearPointBlock)) { in CeedOperatorAssembleDiagonalCore_Cuda()
1444 if (!is_point_block && !diag->diag_rstr) { in CeedOperatorAssembleDiagonalCore_Cuda()
1445 CeedCallBackend(CeedElemRestrictionCreateUnsignedCopy(rstr_out, &diag->diag_rstr)); in CeedOperatorAssembleDiagonalCore_Cuda()
1446 CeedCallBackend(CeedElemRestrictionCreateVector(diag->diag_rstr, NULL, &diag->elem_diag)); in CeedOperatorAssembleDiagonalCore_Cuda()
1447 } else if (is_point_block && !diag->point_block_diag_rstr) { in CeedOperatorAssembleDiagonalCore_Cuda()
1448 …CeedCallBackend(CeedOperatorCreateActivePointBlockRestriction(rstr_out, &diag->point_block_diag_rs… in CeedOperatorAssembleDiagonalCore_Cuda()
1449 …CeedCallBackend(CeedElemRestrictionCreateVector(diag->point_block_diag_rstr, NULL, &diag->point_bl… in CeedOperatorAssembleDiagonalCore_Cuda()
1453 diag_rstr = is_point_block ? diag->point_block_diag_rstr : diag->diag_rstr; in CeedOperatorAssembleDiagonalCore_Cuda()
1454 elem_diag = is_point_block ? diag->point_block_elem_diag : diag->elem_diag; in CeedOperatorAssembleDiagonalCore_Cuda()
1467 …gs[] = {(void *)&num_elem, &diag->d_identity, &diag->d_interp_in, &diag->d_gr… in CeedOperatorAssembleDiagonalCore_Cuda()
1468 …&diag->d_curl_in, &diag->d_interp_out, &diag->d_grad_out, &diag->d_div_out, &diag->d_c… in CeedOperatorAssembleDiagonalCore_Cuda()
1469 … &diag->d_eval_modes_in, &diag->d_eval_modes_out, &assembled_qf_array, &elem_diag_array}; in CeedOperatorAssembleDiagonalCore_Cuda()
1472 …CeedCallBackend(CeedRunKernelDim_Cuda(ceed, diag->LinearPointBlock, grid, num_nodes, 1, elems_per_… in CeedOperatorAssembleDiagonalCore_Cuda()
1474 …CeedCallBackend(CeedRunKernelDim_Cuda(ceed, diag->LinearDiagonal, grid, num_nodes, 1, elems_per_bl… in CeedOperatorAssembleDiagonalCore_Cuda()