Lines Matching full:if
30 if (impl->module) CeedCallCuda(ceed, cuModuleUnload(impl->module)); in CeedOperatorDestroy_Cuda_gen()
31 if (impl->module_assemble_full) CeedCallCuda(ceed, cuModuleUnload(impl->module_assemble_full)); in CeedOperatorDestroy_Cuda_gen()
32 …if (impl->module_assemble_diagonal) CeedCallCuda(ceed, cuModuleUnload(impl->module_assemble_diagon… in CeedOperatorDestroy_Cuda_gen()
33 …if (impl->module_assemble_qfunction) CeedCallCuda(ceed, cuModuleUnload(impl->module_assemble_qfunc… in CeedOperatorDestroy_Cuda_gen()
34 if (impl->points.num_per_elem) CeedCallCuda(ceed, cudaFree((void **)impl->points.num_per_elem)); in CeedOperatorDestroy_Cuda_gen()
72 // If instead, we had packed 3 elements, we'd have 3*49=147 useful threads occupying 160 slots, and…
87 if (i_waste < waste || (i_waste == waste && threads_per_elem * i <= 128)) { in BlockGridCalculate()
120 if (!(*is_run_good)) return CEED_ERROR_SUCCESS; in CeedOperatorApplyAddCore_Cuda_gen()
134 if (eval_mode == CEED_EVAL_WEIGHT) { // Skip in CeedOperatorApplyAddCore_Cuda_gen()
143 if (is_active) data->fields.inputs[i] = input_arr; in CeedOperatorApplyAddCore_Cuda_gen()
152 if (eval_mode == CEED_EVAL_WEIGHT) { // Skip in CeedOperatorApplyAddCore_Cuda_gen()
161 if (is_active) data->fields.outputs[i] = output_arr; in CeedOperatorApplyAddCore_Cuda_gen()
167 // Point coordinates, if needed in CeedOperatorApplyAddCore_Cuda_gen()
169 if (is_at_points) { in CeedOperatorApplyAddCore_Cuda_gen()
178 if (num_elem != data->points.num_elem) { in CeedOperatorApplyAddCore_Cuda_gen()
192 … if (data->points.num_per_elem) CeedCallCuda(ceed, cudaFree((void **)data->points.num_per_elem)); in CeedOperatorApplyAddCore_Cuda_gen()
211 if (is_tensor) { in CeedOperatorApplyAddCore_Cuda_gen()
227 if (eval_mode == CEED_EVAL_WEIGHT) { // Skip in CeedOperatorApplyAddCore_Cuda_gen()
234 if (!is_active) CeedCallBackend(CeedVectorRestoreArrayRead(vec, &data->fields.inputs[i])); in CeedOperatorApplyAddCore_Cuda_gen()
242 if (eval_mode == CEED_EVAL_WEIGHT) { // Skip in CeedOperatorApplyAddCore_Cuda_gen()
249 if (!is_active) CeedCallBackend(CeedVectorRestoreArray(vec, &data->fields.outputs[i])); in CeedOperatorApplyAddCore_Cuda_gen()
254 // Restore point coordinates, if needed in CeedOperatorApplyAddCore_Cuda_gen()
255 if (is_at_points) { in CeedOperatorApplyAddCore_Cuda_gen()
269 if (!(*is_run_good)) data->use_fallback = true; in CeedOperatorApplyAddCore_Cuda_gen()
279 …if (input_vec != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorGetArrayRead(input_vec, CEED_MEM_DEVI… in CeedOperatorApplyAdd_Cuda_gen()
280 …if (output_vec != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorGetArray(output_vec, CEED_MEM_DEVICE… in CeedOperatorApplyAdd_Cuda_gen()
282 …if (input_vec != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorRestoreArrayRead(input_vec, &input_ar… in CeedOperatorApplyAdd_Cuda_gen()
283 …if (output_vec != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorRestoreArray(output_vec, &output_arr… in CeedOperatorApplyAdd_Cuda_gen()
286 if (!is_run_good) { in CeedOperatorApplyAdd_Cuda_gen()
309 …if (input_vec != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorGetArrayRead(input_vec, CEED_MEM_DEVI… in CeedOperatorApplyAddComposite_Cuda_gen()
310 …if (output_vec != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorGetArray(output_vec, CEED_MEM_DEVICE… in CeedOperatorApplyAddComposite_Cuda_gen()
311 if (is_sequential) CeedCallCuda(ceed, cudaStreamCreate(&stream)); in CeedOperatorApplyAddComposite_Cuda_gen()
316 if (num_elem > 0) { in CeedOperatorApplyAddComposite_Cuda_gen()
317 if (!is_sequential) CeedCallCuda(ceed, cudaStreamCreate(&stream)); in CeedOperatorApplyAddComposite_Cuda_gen()
319 if (!is_sequential) CeedCallCuda(ceed, cudaStreamDestroy(stream)); in CeedOperatorApplyAddComposite_Cuda_gen()
322 if (is_sequential) CeedCallCuda(ceed, cudaStreamDestroy(stream)); in CeedOperatorApplyAddComposite_Cuda_gen()
323 …if (input_vec != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorRestoreArrayRead(input_vec, &input_ar… in CeedOperatorApplyAddComposite_Cuda_gen()
324 …if (output_vec != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorRestoreArray(output_vec, &output_arr… in CeedOperatorApplyAddComposite_Cuda_gen()
329 if (!is_run_good[i]) { in CeedOperatorApplyAddComposite_Cuda_gen()
353 if (!data->assemble_qfunction && !data->use_assembly_fallback) { in CeedOperatorLinearAssembleQFunctionCore_Cuda_gen()
357 …if (is_build_good) CeedCallBackend(CeedOperatorBuildKernelLinearAssembleQFunction_Cuda_gen(op, &is… in CeedOperatorLinearAssembleQFunctionCore_Cuda_gen()
358 if (!is_build_good) data->use_assembly_fallback = true; in CeedOperatorLinearAssembleQFunctionCore_Cuda_gen()
362 if (!data->use_assembly_fallback) { in CeedOperatorLinearAssembleQFunctionCore_Cuda_gen()
383 if (eval_mode == CEED_EVAL_WEIGHT) { // Skip in CeedOperatorLinearAssembleQFunctionCore_Cuda_gen()
392 if (is_active) data->fields.inputs[i] = NULL; in CeedOperatorLinearAssembleQFunctionCore_Cuda_gen()
401 // Build objects if needed in CeedOperatorLinearAssembleQFunctionCore_Cuda_gen()
402 if (build_objects) { in CeedOperatorLinearAssembleQFunctionCore_Cuda_gen()
413 // Check if active input in CeedOperatorLinearAssembleQFunctionCore_Cuda_gen()
414 if (vec == CEED_VECTOR_ACTIVE) { in CeedOperatorLinearAssembleQFunctionCore_Cuda_gen()
431 // Check if active output in CeedOperatorLinearAssembleQFunctionCore_Cuda_gen()
432 if (vec == CEED_VECTOR_ACTIVE) { in CeedOperatorLinearAssembleQFunctionCore_Cuda_gen()
466 if (is_tensor) { in CeedOperatorLinearAssembleQFunctionCore_Cuda_gen()
484 if (eval_mode == CEED_EVAL_WEIGHT) { // Skip in CeedOperatorLinearAssembleQFunctionCore_Cuda_gen()
491 if (!is_active) CeedCallBackend(CeedVectorRestoreArrayRead(vec, &data->fields.inputs[i])); in CeedOperatorLinearAssembleQFunctionCore_Cuda_gen()
504 if (!is_run_good) { in CeedOperatorLinearAssembleQFunctionCore_Cuda_gen()
506 if (build_objects) { in CeedOperatorLinearAssembleQFunctionCore_Cuda_gen()
514 // Fallback, if needed in CeedOperatorLinearAssembleQFunctionCore_Cuda_gen()
515 if (data->use_assembly_fallback) { in CeedOperatorLinearAssembleQFunctionCore_Cuda_gen()
545 if (!data->assemble_diagonal && !data->use_assembly_fallback) { in CeedOperatorLinearAssembleAddDiagonalAtPoints_Cuda_gen()
553 if (num_active_bases_in == num_active_bases_out) { in CeedOperatorLinearAssembleAddDiagonalAtPoints_Cuda_gen()
555 …if (is_build_good) CeedCallBackend(CeedOperatorBuildKernelDiagonalAssemblyAtPoints_Cuda_gen(op, &i… in CeedOperatorLinearAssembleAddDiagonalAtPoints_Cuda_gen()
557 if (!is_build_good) data->use_assembly_fallback = true; in CeedOperatorLinearAssembleAddDiagonalAtPoints_Cuda_gen()
561 if (!data->use_assembly_fallback) { in CeedOperatorLinearAssembleAddDiagonalAtPoints_Cuda_gen()
582 if (eval_mode == CEED_EVAL_WEIGHT) { // Skip in CeedOperatorLinearAssembleAddDiagonalAtPoints_Cuda_gen()
591 if (is_active) data->fields.inputs[i] = NULL; in CeedOperatorLinearAssembleAddDiagonalAtPoints_Cuda_gen()
606 if (num_elem != data->points.num_elem) { in CeedOperatorLinearAssembleAddDiagonalAtPoints_Cuda_gen()
620 … if (data->points.num_per_elem) CeedCallCuda(ceed, cudaFree((void **)data->points.num_per_elem)); in CeedOperatorLinearAssembleAddDiagonalAtPoints_Cuda_gen()
652 if (eval_mode == CEED_EVAL_WEIGHT) { // Skip in CeedOperatorLinearAssembleAddDiagonalAtPoints_Cuda_gen()
659 if (!is_active) CeedCallBackend(CeedVectorRestoreArrayRead(vec, &data->fields.inputs[i])); in CeedOperatorLinearAssembleAddDiagonalAtPoints_Cuda_gen()
681 if (!is_run_good) data->use_assembly_fallback = true; in CeedOperatorLinearAssembleAddDiagonalAtPoints_Cuda_gen()
685 // Fallback, if needed in CeedOperatorLinearAssembleAddDiagonalAtPoints_Cuda_gen()
686 if (data->use_assembly_fallback) { in CeedOperatorLinearAssembleAddDiagonalAtPoints_Cuda_gen()
708 if (!data->assemble_full && !data->use_assembly_fallback) { in CeedOperatorAssembleSingleAtPoints_Cuda_gen()
716 if (num_active_bases_in == num_active_bases_out) { in CeedOperatorAssembleSingleAtPoints_Cuda_gen()
718 …if (is_build_good) CeedCallBackend(CeedOperatorBuildKernelFullAssemblyAtPoints_Cuda_gen(op, &is_bu… in CeedOperatorAssembleSingleAtPoints_Cuda_gen()
720 if (!is_build_good) data->use_assembly_fallback = true; in CeedOperatorAssembleSingleAtPoints_Cuda_gen()
724 if (!data->use_assembly_fallback) { in CeedOperatorAssembleSingleAtPoints_Cuda_gen()
745 if (eval_mode == CEED_EVAL_WEIGHT) { // Skip in CeedOperatorAssembleSingleAtPoints_Cuda_gen()
754 if (is_active) data->fields.inputs[i] = NULL; in CeedOperatorAssembleSingleAtPoints_Cuda_gen()
769 if (num_elem != data->points.num_elem) { in CeedOperatorAssembleSingleAtPoints_Cuda_gen()
783 … if (data->points.num_per_elem) CeedCallCuda(ceed, cudaFree((void **)data->points.num_per_elem)); in CeedOperatorAssembleSingleAtPoints_Cuda_gen()
817 if (eval_mode == CEED_EVAL_WEIGHT) { // Skip in CeedOperatorAssembleSingleAtPoints_Cuda_gen()
824 if (!is_active) CeedCallBackend(CeedVectorRestoreArrayRead(vec, &data->fields.inputs[i])); in CeedOperatorAssembleSingleAtPoints_Cuda_gen()
846 if (!is_run_good) data->use_assembly_fallback = true; in CeedOperatorAssembleSingleAtPoints_Cuda_gen()
850 // Fallback, if needed in CeedOperatorAssembleSingleAtPoints_Cuda_gen()
851 if (data->use_assembly_fallback) { in CeedOperatorAssembleSingleAtPoints_Cuda_gen()
874 if (is_composite) { in CeedOperatorCreate_Cuda_gen()
880 if (is_at_points) { in CeedOperatorCreate_Cuda_gen()
885 if (!is_at_points) { in CeedOperatorCreate_Cuda_gen()