Lines Matching full:if
30 if (is_composite) { in CeedOperatorDestroy_Hip_gen()
35 if (impl->streams[i]) CeedCallHip(ceed, hipStreamDestroy(impl->streams[i])); in CeedOperatorDestroy_Hip_gen()
39 if (impl->module) CeedCallHip(ceed, hipModuleUnload(impl->module)); in CeedOperatorDestroy_Hip_gen()
40 if (impl->module_assemble_full) CeedCallHip(ceed, hipModuleUnload(impl->module_assemble_full)); in CeedOperatorDestroy_Hip_gen()
41 …if (impl->module_assemble_diagonal) CeedCallHip(ceed, hipModuleUnload(impl->module_assemble_diagon… in CeedOperatorDestroy_Hip_gen()
42 …if (impl->module_assemble_qfunction) CeedCallHip(ceed, hipModuleUnload(impl->module_assemble_qfunc… in CeedOperatorDestroy_Hip_gen()
43 if (impl->points.num_per_elem) CeedCallHip(ceed, hipFree((void **)impl->points.num_per_elem)); in CeedOperatorDestroy_Hip_gen()
66 if (!(*is_run_good)) return CEED_ERROR_SUCCESS; in CeedOperatorApplyAddCore_Hip_gen()
79 if (eval_mode == CEED_EVAL_WEIGHT) { // Skip in CeedOperatorApplyAddCore_Hip_gen()
88 if (is_active) data->fields.inputs[i] = input_arr; in CeedOperatorApplyAddCore_Hip_gen()
97 if (eval_mode == CEED_EVAL_WEIGHT) { // Skip in CeedOperatorApplyAddCore_Hip_gen()
106 if (is_active) data->fields.outputs[i] = output_arr; in CeedOperatorApplyAddCore_Hip_gen()
112 // Point coordinates, if needed in CeedOperatorApplyAddCore_Hip_gen()
114 if (is_at_points) { in CeedOperatorApplyAddCore_Hip_gen()
123 if (num_elem != data->points.num_elem) { in CeedOperatorApplyAddCore_Hip_gen()
137 if (data->points.num_per_elem) CeedCallHip(ceed, hipFree((void **)data->points.num_per_elem)); in CeedOperatorApplyAddCore_Hip_gen()
154 if (is_tensor) { in CeedOperatorApplyAddCore_Hip_gen()
162 if (data->dim == 1 || !is_tensor) { in CeedOperatorApplyAddCore_Hip_gen()
168 } else if (data->dim == 2) { in CeedOperatorApplyAddCore_Hip_gen()
174 } else if (data->dim == 3) { in CeedOperatorApplyAddCore_Hip_gen()
185 if (eval_mode == CEED_EVAL_WEIGHT) { // Skip in CeedOperatorApplyAddCore_Hip_gen()
192 if (!is_active) CeedCallBackend(CeedVectorRestoreArrayRead(vec, &data->fields.inputs[i])); in CeedOperatorApplyAddCore_Hip_gen()
200 if (eval_mode == CEED_EVAL_WEIGHT) { // Skip in CeedOperatorApplyAddCore_Hip_gen()
207 if (!is_active) CeedCallBackend(CeedVectorRestoreArray(vec, &data->fields.outputs[i])); in CeedOperatorApplyAddCore_Hip_gen()
212 // Restore point coordinates, if needed in CeedOperatorApplyAddCore_Hip_gen()
213 if (is_at_points) { in CeedOperatorApplyAddCore_Hip_gen()
227 if (!(*is_run_good)) data->use_fallback = true; in CeedOperatorApplyAddCore_Hip_gen()
237 …if (input_vec != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorGetArrayRead(input_vec, CEED_MEM_DEVI… in CeedOperatorApplyAdd_Hip_gen()
238 …if (output_vec != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorGetArray(output_vec, CEED_MEM_DEVICE… in CeedOperatorApplyAdd_Hip_gen()
240 …if (input_vec != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorRestoreArrayRead(input_vec, &input_ar… in CeedOperatorApplyAdd_Hip_gen()
241 …if (output_vec != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorRestoreArray(output_vec, &output_arr… in CeedOperatorApplyAdd_Hip_gen()
244 if (!is_run_good) { in CeedOperatorApplyAdd_Hip_gen()
268 …if (input_vec != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorGetArrayRead(input_vec, CEED_MEM_DEVI… in CeedOperatorApplyAddComposite_Hip_gen()
269 …if (output_vec != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorGetArray(output_vec, CEED_MEM_DEVICE… in CeedOperatorApplyAddComposite_Hip_gen()
275 if (num_elem > 0) { in CeedOperatorApplyAddComposite_Hip_gen()
276 …if (!impl->streams[stream_index]) CeedCallHip(ceed, hipStreamCreate(&impl->streams[stream_index])); in CeedOperatorApplyAddComposite_Hip_gen()
283 if (is_sequential) CeedCallHip(ceed, hipStreamSynchronize(impl->streams[0])); in CeedOperatorApplyAddComposite_Hip_gen()
286 if (impl->streams[i]) { in CeedOperatorApplyAddComposite_Hip_gen()
287 if (is_run_good[i]) CeedCallHip(ceed, hipStreamSynchronize(impl->streams[i])); in CeedOperatorApplyAddComposite_Hip_gen()
291 …if (input_vec != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorRestoreArrayRead(input_vec, &input_ar… in CeedOperatorApplyAddComposite_Hip_gen()
292 …if (output_vec != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorRestoreArray(output_vec, &output_arr… in CeedOperatorApplyAddComposite_Hip_gen()
297 if (!is_run_good[i]) { in CeedOperatorApplyAddComposite_Hip_gen()
321 if (!data->assemble_qfunction && !data->use_assembly_fallback) { in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
325 …if (is_build_good) CeedCallBackend(CeedOperatorBuildKernelLinearAssembleQFunction_Hip_gen(op, &is_… in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
326 if (!is_build_good) data->use_assembly_fallback = true; in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
330 if (!data->use_assembly_fallback) { in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
351 if (eval_mode == CEED_EVAL_WEIGHT) { // Skip in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
360 if (is_active) data->fields.inputs[i] = NULL; in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
369 // Build objects if needed in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
370 if (build_objects) { in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
381 // Check if active input in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
382 if (vec == CEED_VECTOR_ACTIVE) { in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
399 // Check if active output in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
400 if (vec == CEED_VECTOR_ACTIVE) { in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
432 if (is_tensor) { in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
440 if (data->dim == 1 || !is_tensor) { in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
446 } else if (data->dim == 2) { in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
452 } else if (data->dim == 3) { in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
463 if (eval_mode == CEED_EVAL_WEIGHT) { // Skip in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
470 if (!is_active) CeedCallBackend(CeedVectorRestoreArrayRead(vec, &data->fields.inputs[i])); in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
483 if (!is_run_good) { in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
485 if (build_objects) { in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
493 // Fallback, if needed in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
494 if (data->use_assembly_fallback) { in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
524 if (!data->assemble_diagonal && !data->use_assembly_fallback) { in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
532 if (num_active_bases_in == num_active_bases_out) { in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
534 …if (is_build_good) CeedCallBackend(CeedOperatorBuildKernelDiagonalAssemblyAtPoints_Hip_gen(op, &is… in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
536 if (!is_build_good) data->use_assembly_fallback = true; in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
540 if (!data->use_assembly_fallback) { in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
561 if (eval_mode == CEED_EVAL_WEIGHT) { // Skip in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
570 if (is_active) data->fields.inputs[i] = NULL; in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
585 if (num_elem != data->points.num_elem) { in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
599 … if (data->points.num_per_elem) CeedCallHip(ceed, hipFree((void **)data->points.num_per_elem)); in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
620 if (data->dim == 1) { in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
626 } else if (data->dim == 2) { in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
632 } else if (data->dim == 3) { in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
644 if (eval_mode == CEED_EVAL_WEIGHT) { // Skip in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
651 if (!is_active) CeedCallBackend(CeedVectorRestoreArrayRead(vec, &data->fields.inputs[i])); in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
673 if (!is_run_good) data->use_assembly_fallback = true; in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
677 // Fallback, if needed in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
678 if (data->use_assembly_fallback) { in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
700 if (!data->assemble_full && !data->use_assembly_fallback) { in CeedOperatorAssembleSingleAtPoints_Hip_gen()
708 if (num_active_bases_in == num_active_bases_out) { in CeedOperatorAssembleSingleAtPoints_Hip_gen()
710 …if (is_build_good) CeedCallBackend(CeedOperatorBuildKernelFullAssemblyAtPoints_Hip_gen(op, &is_bui… in CeedOperatorAssembleSingleAtPoints_Hip_gen()
712 if (!is_build_good) { in CeedOperatorAssembleSingleAtPoints_Hip_gen()
719 if (!data->use_assembly_fallback) { in CeedOperatorAssembleSingleAtPoints_Hip_gen()
741 if (eval_mode == CEED_EVAL_WEIGHT) { // Skip in CeedOperatorAssembleSingleAtPoints_Hip_gen()
750 if (is_active) data->fields.inputs[i] = NULL; in CeedOperatorAssembleSingleAtPoints_Hip_gen()
765 if (num_elem != data->points.num_elem) { in CeedOperatorAssembleSingleAtPoints_Hip_gen()
779 … if (data->points.num_per_elem) CeedCallHip(ceed, hipFree((void **)data->points.num_per_elem)); in CeedOperatorAssembleSingleAtPoints_Hip_gen()
802 if (data->dim == 1) { in CeedOperatorAssembleSingleAtPoints_Hip_gen()
808 } else if (data->dim == 2) { in CeedOperatorAssembleSingleAtPoints_Hip_gen()
814 } else if (data->dim == 3) { in CeedOperatorAssembleSingleAtPoints_Hip_gen()
826 if (eval_mode == CEED_EVAL_WEIGHT) { // Skip in CeedOperatorAssembleSingleAtPoints_Hip_gen()
833 if (!is_active) CeedCallBackend(CeedVectorRestoreArrayRead(vec, &data->fields.inputs[i])); in CeedOperatorAssembleSingleAtPoints_Hip_gen()
855 if (!is_run_good) { in CeedOperatorAssembleSingleAtPoints_Hip_gen()
862 // Fallback, if needed in CeedOperatorAssembleSingleAtPoints_Hip_gen()
863 if (data->use_assembly_fallback) { in CeedOperatorAssembleSingleAtPoints_Hip_gen()
886 if (is_composite) { in CeedOperatorCreate_Hip_gen()
892 if (is_at_points) { in CeedOperatorCreate_Hip_gen()
896 if (!is_at_points) { in CeedOperatorCreate_Hip_gen()