Lines Matching refs:data
62 CeedOperator_Hip_gen *data; in CeedOperatorApplyAddCore_Hip_gen() local
69 CeedCallBackend(CeedOperatorGetData(op, &data)); in CeedOperatorApplyAddCore_Hip_gen()
80 data->fields.inputs[i] = NULL; in CeedOperatorApplyAddCore_Hip_gen()
88 if (is_active) data->fields.inputs[i] = input_arr; in CeedOperatorApplyAddCore_Hip_gen()
89 else CeedCallBackend(CeedVectorGetArrayRead(vec, CEED_MEM_DEVICE, &data->fields.inputs[i])); in CeedOperatorApplyAddCore_Hip_gen()
98 data->fields.outputs[i] = NULL; in CeedOperatorApplyAddCore_Hip_gen()
106 if (is_active) data->fields.outputs[i] = output_arr; in CeedOperatorApplyAddCore_Hip_gen()
107 else CeedCallBackend(CeedVectorGetArray(vec, CEED_MEM_DEVICE, &data->fields.outputs[i])); in CeedOperatorApplyAddCore_Hip_gen()
119 CeedCallBackend(CeedVectorGetArrayRead(vec, CEED_MEM_DEVICE, &data->points.coords)); in CeedOperatorApplyAddCore_Hip_gen()
123 if (num_elem != data->points.num_elem) { in CeedOperatorApplyAddCore_Hip_gen()
128 data->points.num_elem = 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()
138 CeedCallHip(ceed, hipMalloc((void **)&data->points.num_per_elem, num_bytes)); in CeedOperatorApplyAddCore_Hip_gen()
139 …CeedCallHip(ceed, hipMemcpy((void *)data->points.num_per_elem, points_per_elem, num_bytes, hipMemc… in CeedOperatorApplyAddCore_Hip_gen()
149 …rgs[] = {(void *)&num_elem, &qf_data->d_c, &data->indices, &data->fields, &data->B, &data->G, &dat… in CeedOperatorApplyAddCore_Hip_gen()
152 …CeedInt block_sizes[3] = {data->thread_1d, ((!is_tensor || data->dim == 1) ? 1 : data->thread_1d),… in CeedOperatorApplyAddCore_Hip_gen()
155 …CeedCallBackend(BlockGridCalculate_Hip_gen(data->dim, num_elem, data->max_P_1d, data->Q_1d, block_… in CeedOperatorApplyAddCore_Hip_gen()
157 CeedInt elems_per_block = 64 * data->thread_1d > 256 ? 256 / data->thread_1d : 64; in CeedOperatorApplyAddCore_Hip_gen()
162 if (data->dim == 1 || !is_tensor) { in CeedOperatorApplyAddCore_Hip_gen()
164 CeedInt sharedMem = block_sizes[2] * data->thread_1d * sizeof(CeedScalar); in CeedOperatorApplyAddCore_Hip_gen()
166 …CeedCallBackend(CeedTryRunKernelDimShared_Hip(ceed, data->op, stream, grid, block_sizes[0], block_… in CeedOperatorApplyAddCore_Hip_gen()
168 } else if (data->dim == 2) { in CeedOperatorApplyAddCore_Hip_gen()
170 CeedInt sharedMem = block_sizes[2] * data->thread_1d * data->thread_1d * sizeof(CeedScalar); in CeedOperatorApplyAddCore_Hip_gen()
172 …CeedCallBackend(CeedTryRunKernelDimShared_Hip(ceed, data->op, stream, grid, block_sizes[0], block_… in CeedOperatorApplyAddCore_Hip_gen()
174 } else if (data->dim == 3) { in CeedOperatorApplyAddCore_Hip_gen()
176 CeedInt sharedMem = block_sizes[2] * data->thread_1d * data->thread_1d * sizeof(CeedScalar); in CeedOperatorApplyAddCore_Hip_gen()
178 …CeedCallBackend(CeedTryRunKernelDimShared_Hip(ceed, data->op, stream, grid, block_sizes[0], block_… in CeedOperatorApplyAddCore_Hip_gen()
192 if (!is_active) CeedCallBackend(CeedVectorRestoreArrayRead(vec, &data->fields.inputs[i])); in CeedOperatorApplyAddCore_Hip_gen()
207 if (!is_active) CeedCallBackend(CeedVectorRestoreArray(vec, &data->fields.outputs[i])); in CeedOperatorApplyAddCore_Hip_gen()
217 CeedCallBackend(CeedVectorRestoreArrayRead(vec, &data->points.coords)); in CeedOperatorApplyAddCore_Hip_gen()
227 if (!(*is_run_good)) data->use_fallback = true; in CeedOperatorApplyAddCore_Hip_gen()
315 CeedOperator_Hip_gen *data; in CeedOperatorLinearAssembleQFunctionCore_Hip_gen() local
318 CeedCallBackend(CeedOperatorGetData(op, &data)); in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
321 if (!data->assemble_qfunction && !data->use_assembly_fallback) { 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()
352 data->fields.inputs[i] = NULL; in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
360 if (is_active) data->fields.inputs[i] = NULL; in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
361 else CeedCallBackend(CeedVectorGetArrayRead(vec, CEED_MEM_DEVICE, &data->fields.inputs[i])); in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
427 …rgs[] = {(void *)&num_elem, &qf_data->d_c, &data->indices, &data->fields, &data->B, &data->G, &dat… in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
430 …CeedInt block_sizes[3] = {data->thread_1d, ((!is_tensor || data->dim == 1) ? 1 : data->thread_1d),… in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
433 …CeedCallBackend(BlockGridCalculate_Hip_gen(data->dim, num_elem, data->max_P_1d, data->Q_1d, block_… in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
435 CeedInt elems_per_block = 64 * data->thread_1d > 256 ? 256 / data->thread_1d : 64; in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
440 if (data->dim == 1 || !is_tensor) { in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
442 CeedInt sharedMem = block_sizes[2] * data->thread_1d * sizeof(CeedScalar); in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
444 …CeedCallBackend(CeedTryRunKernelDimShared_Hip(ceed, data->assemble_qfunction, NULL, grid, block_si… in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
446 } else if (data->dim == 2) { in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
448 CeedInt sharedMem = block_sizes[2] * data->thread_1d * data->thread_1d * sizeof(CeedScalar); in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
450 …CeedCallBackend(CeedTryRunKernelDimShared_Hip(ceed, data->assemble_qfunction, NULL, grid, block_si… in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
452 } else if (data->dim == 3) { in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
454 CeedInt sharedMem = block_sizes[2] * data->thread_1d * data->thread_1d * sizeof(CeedScalar); in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
456 …CeedCallBackend(CeedTryRunKernelDimShared_Hip(ceed, data->assemble_qfunction, NULL, grid, block_si… in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
470 if (!is_active) CeedCallBackend(CeedVectorRestoreArrayRead(vec, &data->fields.inputs[i])); in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
484 data->use_assembly_fallback = true; in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
494 if (data->use_assembly_fallback) { in CeedOperatorLinearAssembleQFunctionCore_Hip_gen()
518 CeedOperator_Hip_gen *data; in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen() local
521 CeedCallBackend(CeedOperatorGetData(op, &data)); in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
524 if (!data->assemble_diagonal && !data->use_assembly_fallback) { 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()
562 data->fields.inputs[i] = NULL; in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
570 if (is_active) data->fields.inputs[i] = NULL; in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
571 else CeedCallBackend(CeedVectorGetArrayRead(vec, CEED_MEM_DEVICE, &data->fields.inputs[i])); in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
581 CeedCallBackend(CeedVectorGetArrayRead(vec, CEED_MEM_DEVICE, &data->points.coords)); in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
585 if (num_elem != data->points.num_elem) { in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
590 data->points.num_elem = 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()
600 CeedCallHip(ceed, hipMalloc((void **)&data->points.num_per_elem, num_bytes)); in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
601 …CeedCallHip(ceed, hipMemcpy((void *)data->points.num_per_elem, points_per_elem, num_bytes, hipMemc… in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
614 …rgs[] = {(void *)&num_elem, &qf_data->d_c, &data->indices, &data->fields, &data->B, &data->G, &dat… in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
616 CeedInt block_sizes[3] = {data->thread_1d, (data->dim == 1 ? 1 : data->thread_1d), -1}; in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
618 …CeedCallBackend(BlockGridCalculate_Hip_gen(data->dim, num_elem, data->max_P_1d, data->Q_1d, block_… in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
620 if (data->dim == 1) { in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
622 CeedInt sharedMem = block_sizes[2] * data->thread_1d * sizeof(CeedScalar); in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
624 …CeedCallBackend(CeedTryRunKernelDimShared_Hip(ceed, data->assemble_diagonal, NULL, grid, block_siz… in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
626 } else if (data->dim == 2) { in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
628 CeedInt sharedMem = block_sizes[2] * data->thread_1d * data->thread_1d * sizeof(CeedScalar); in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
630 …CeedCallBackend(CeedTryRunKernelDimShared_Hip(ceed, data->assemble_diagonal, NULL, grid, block_siz… in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
632 } else if (data->dim == 3) { in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
634 CeedInt sharedMem = block_sizes[2] * data->thread_1d * data->thread_1d * sizeof(CeedScalar); in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
636 …CeedCallBackend(CeedTryRunKernelDimShared_Hip(ceed, data->assemble_diagonal, NULL, grid, block_siz… in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
651 if (!is_active) CeedCallBackend(CeedVectorRestoreArrayRead(vec, &data->fields.inputs[i])); in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
661 CeedCallBackend(CeedVectorRestoreArrayRead(vec, &data->points.coords)); in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
673 if (!is_run_good) data->use_assembly_fallback = true; in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
678 if (data->use_assembly_fallback) { in CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen()
694 CeedOperator_Hip_gen *data; in CeedOperatorAssembleSingleAtPoints_Hip_gen() local
697 CeedCallBackend(CeedOperatorGetData(op, &data)); in CeedOperatorAssembleSingleAtPoints_Hip_gen()
700 if (!data->assemble_full && !data->use_assembly_fallback) { in CeedOperatorAssembleSingleAtPoints_Hip_gen()
714 data->use_assembly_fallback = true; in CeedOperatorAssembleSingleAtPoints_Hip_gen()
719 if (!data->use_assembly_fallback) { in CeedOperatorAssembleSingleAtPoints_Hip_gen()
742 data->fields.inputs[i] = NULL; in CeedOperatorAssembleSingleAtPoints_Hip_gen()
750 if (is_active) data->fields.inputs[i] = NULL; in CeedOperatorAssembleSingleAtPoints_Hip_gen()
751 else CeedCallBackend(CeedVectorGetArrayRead(vec, CEED_MEM_DEVICE, &data->fields.inputs[i])); in CeedOperatorAssembleSingleAtPoints_Hip_gen()
761 CeedCallBackend(CeedVectorGetArrayRead(vec, CEED_MEM_DEVICE, &data->points.coords)); in CeedOperatorAssembleSingleAtPoints_Hip_gen()
765 if (num_elem != data->points.num_elem) { in CeedOperatorAssembleSingleAtPoints_Hip_gen()
770 data->points.num_elem = 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()
780 CeedCallHip(ceed, hipMalloc((void **)&data->points.num_per_elem, num_bytes)); in CeedOperatorAssembleSingleAtPoints_Hip_gen()
781 …CeedCallHip(ceed, hipMemcpy((void *)data->points.num_per_elem, points_per_elem, num_bytes, hipMemc… in CeedOperatorAssembleSingleAtPoints_Hip_gen()
795 …void *opargs[] = {(void *)&num_elem, &qf_data->d_c, &data->indices, &data->fields, &data-… in CeedOperatorAssembleSingleAtPoints_Hip_gen()
796 &data->G, &data->W, &data->points, &assembled_offset_array}; in CeedOperatorAssembleSingleAtPoints_Hip_gen()
798 CeedInt block_sizes[3] = {data->thread_1d, (data->dim == 1 ? 1 : data->thread_1d), -1}; in CeedOperatorAssembleSingleAtPoints_Hip_gen()
800 …CeedCallBackend(BlockGridCalculate_Hip_gen(data->dim, num_elem, data->max_P_1d, data->Q_1d, block_… in CeedOperatorAssembleSingleAtPoints_Hip_gen()
802 if (data->dim == 1) { in CeedOperatorAssembleSingleAtPoints_Hip_gen()
804 CeedInt sharedMem = block_sizes[2] * data->thread_1d * sizeof(CeedScalar); in CeedOperatorAssembleSingleAtPoints_Hip_gen()
806 …CeedCallBackend(CeedTryRunKernelDimShared_Hip(ceed, data->assemble_full, NULL, grid, block_sizes[0… in CeedOperatorAssembleSingleAtPoints_Hip_gen()
808 } else if (data->dim == 2) { in CeedOperatorAssembleSingleAtPoints_Hip_gen()
810 CeedInt sharedMem = block_sizes[2] * data->thread_1d * data->thread_1d * sizeof(CeedScalar); in CeedOperatorAssembleSingleAtPoints_Hip_gen()
812 …CeedCallBackend(CeedTryRunKernelDimShared_Hip(ceed, data->assemble_full, NULL, grid, block_sizes[0… in CeedOperatorAssembleSingleAtPoints_Hip_gen()
814 } else if (data->dim == 3) { in CeedOperatorAssembleSingleAtPoints_Hip_gen()
816 CeedInt sharedMem = block_sizes[2] * data->thread_1d * data->thread_1d * sizeof(CeedScalar); in CeedOperatorAssembleSingleAtPoints_Hip_gen()
818 …CeedCallBackend(CeedTryRunKernelDimShared_Hip(ceed, data->assemble_full, NULL, grid, block_sizes[0… in CeedOperatorAssembleSingleAtPoints_Hip_gen()
833 if (!is_active) CeedCallBackend(CeedVectorRestoreArrayRead(vec, &data->fields.inputs[i])); in CeedOperatorAssembleSingleAtPoints_Hip_gen()
843 CeedCallBackend(CeedVectorRestoreArrayRead(vec, &data->points.coords)); in CeedOperatorAssembleSingleAtPoints_Hip_gen()
857 data->use_assembly_fallback = true; in CeedOperatorAssembleSingleAtPoints_Hip_gen()
863 if (data->use_assembly_fallback) { in CeedOperatorAssembleSingleAtPoints_Hip_gen()