| /libCEED/backends/memcheck/ |
| H A D | ceed-memcheck-qfunctioncontext.c | 20 CeedQFunctionContext_Memcheck *impl; in CeedQFunctionContextHasValidData_Memcheck() local
22 CeedCallBackend(CeedQFunctionContextGetBackendData(ctx, &impl)); in CeedQFunctionContextHasValidData_Memcheck()
23 *has_valid_data = !!impl->data_allocated; in CeedQFunctionContextHasValidData_Memcheck()
31 CeedQFunctionContext_Memcheck *impl; in CeedQFunctionContextHasBorrowedDataOfType_Memcheck() local
35 CeedCallBackend(CeedQFunctionContextGetBackendData(ctx, &impl)); in CeedQFunctionContextHasBorrowedDataOfType_Memcheck()
36 *has_borrowed_data_of_type = !!impl->data_borrowed; in CeedQFunctionContextHasBorrowedDataOfType_Memcheck()
45 CeedQFunctionContext_Memcheck *impl; in CeedQFunctionContextSetData_Memcheck() local
49 CeedCallBackend(CeedQFunctionContextGetBackendData(ctx, &impl)); in CeedQFunctionContextSetData_Memcheck()
53 if (impl->data_allocated) { in CeedQFunctionContextSetData_Memcheck()
54 memset(impl->data_allocated, -42, ctx_size); in CeedQFunctionContextSetData_Memcheck()
[all …]
|
| H A D | ceed-memcheck-vector.c | 22 CeedVector_Memcheck *impl; in CeedVectorHasValidArray_Memcheck() local
24 CeedCallBackend(CeedVectorGetData(vec, &impl)); in CeedVectorHasValidArray_Memcheck()
25 *has_valid_array = !!impl->array_allocated; in CeedVectorHasValidArray_Memcheck()
33 CeedVector_Memcheck *impl; in CeedVectorHasBorrowedArrayOfType_Memcheck() local
37 CeedCallBackend(CeedVectorGetData(vec, &impl)); in CeedVectorHasBorrowedArrayOfType_Memcheck()
38 *has_borrowed_array_of_type = !!impl->array_borrowed; in CeedVectorHasBorrowedArrayOfType_Memcheck()
47 CeedVector_Memcheck *impl; in CeedVectorSetArray_Memcheck() local
51 CeedCallBackend(CeedVectorGetData(vec, &impl)); in CeedVectorSetArray_Memcheck()
55 if (impl->array_allocated) { in CeedVectorSetArray_Memcheck()
56 for (CeedSize i = 0; i < length; i++) impl->array_allocated[i] = NAN; in CeedVectorSetArray_Memcheck()
[all …]
|
| H A D | ceed-memcheck-restriction.c | 77 CeedElemRestriction_Memcheck *impl; in CeedElemRestrictionApplyOffsetNoTranspose_Memcheck_Core() local
79 CeedCallBackend(CeedElemRestrictionGetData(rstr, &impl)); in CeedElemRestrictionApplyOffsetNoTranspose_Memcheck_Core()
83 …vv[elem_size * (k * block_size + e * num_comp) + i - v_offset] = uu[impl->offsets[i + e * elem_siz… in CeedElemRestrictionApplyOffsetNoTranspose_Memcheck_Core()
95 CeedElemRestriction_Memcheck *impl; in CeedElemRestrictionApplyOrientedNoTranspose_Memcheck_Core() local
97 CeedCallBackend(CeedElemRestrictionGetData(rstr, &impl)); in CeedElemRestrictionApplyOrientedNoTranspose_Memcheck_Core()
102 …uu[impl->offsets[i + e * elem_size] + k * comp_stride] * (impl->orients[i + e * elem_size] ? -1.0 … in CeedElemRestrictionApplyOrientedNoTranspose_Memcheck_Core()
114 CeedElemRestriction_Memcheck *impl; in CeedElemRestrictionApplyCurlOrientedNoTranspose_Memcheck_Core() local
116 CeedCallBackend(CeedElemRestrictionGetData(rstr, &impl)); in CeedElemRestrictionApplyCurlOrientedNoTranspose_Memcheck_Core()
123 uu[impl->offsets[j + n * block_size + e * elem_size] + k * comp_stride] * in CeedElemRestrictionApplyCurlOrientedNoTranspose_Memcheck_Core()
124 impl->curl_orients[j + (3 * n + 1) * block_size + e * 3 * elem_size] + in CeedElemRestrictionApplyCurlOrientedNoTranspose_Memcheck_Core()
[all …]
|
| H A D | ceed-memcheck-qfunction.c | 25 CeedQFunction_Memcheck *impl; in CeedQFunctionApply_Memcheck() local
27 CeedCallBackend(CeedQFunctionGetData(qf, &impl)); in CeedQFunctionApply_Memcheck()
37 CeedCallBackend(CeedVectorGetArrayRead(U[i], CEED_MEM_HOST, &impl->inputs[i])); in CeedQFunctionApply_Memcheck()
42 input_block_ids[i] = VALGRIND_CREATE_BLOCK(impl->inputs[i], len, name); in CeedQFunctionApply_Memcheck()
50 CeedCallBackend(CeedVectorGetArrayWrite(V[i], CEED_MEM_HOST, &impl->outputs[i])); in CeedQFunctionApply_Memcheck()
53 VALGRIND_MAKE_MEM_UNDEFINED(impl->outputs[i], len); in CeedQFunctionApply_Memcheck()
56 output_block_ids[i] = VALGRIND_CREATE_BLOCK(impl->outputs[i], len, name); in CeedQFunctionApply_Memcheck()
60 CeedCallBackend(f(ctx_data, Q, impl->inputs, impl->outputs)); in CeedQFunctionApply_Memcheck()
64 CeedCallBackend(CeedVectorRestoreArrayRead(U[i], &impl->inputs[i])); in CeedQFunctionApply_Memcheck()
83 CeedCheck(!isnan(impl->outputs[i][j]), CeedQFunctionReturnCeed(qf), CEED_ERROR_BACKEND, in CeedQFunctionApply_Memcheck()
[all …]
|
| /libCEED/backends/hip-ref/ |
| H A D | ceed-hip-ref-qfunctioncontext.c | 23 CeedQFunctionContext_Hip *impl; in CeedQFunctionContextSyncH2D_Hip() local
26 CeedCallBackend(CeedQFunctionContextGetBackendData(ctx, &impl)); in CeedQFunctionContextSyncH2D_Hip()
28 CeedCheck(impl->h_data, ceed, CEED_ERROR_BACKEND, "No valid host data to sync to device"); in CeedQFunctionContextSyncH2D_Hip()
31 if (impl->d_data_borrowed) { in CeedQFunctionContextSyncH2D_Hip()
32 impl->d_data = impl->d_data_borrowed; in CeedQFunctionContextSyncH2D_Hip()
33 } else if (impl->d_data_owned) { in CeedQFunctionContextSyncH2D_Hip()
34 impl->d_data = impl->d_data_owned; in CeedQFunctionContextSyncH2D_Hip()
36 CeedCallHip(ceed, hipMalloc((void **)&impl->d_data_owned, ctx_size)); in CeedQFunctionContextSyncH2D_Hip()
37 impl->d_data = impl->d_data_owned; in CeedQFunctionContextSyncH2D_Hip()
39 CeedCallHip(ceed, hipMemcpy(impl->d_data, impl->h_data, ctx_size, hipMemcpyHostToDevice)); in CeedQFunctionContextSyncH2D_Hip()
[all …]
|
| H A D | ceed-hip-ref-restriction.c | 28 CeedElemRestriction_Hip *impl; in CeedElemRestrictionSetupCompile_Hip() local
30 CeedCallBackend(CeedElemRestrictionGetData(rstr, &impl)); in CeedElemRestrictionSetupCompile_Hip()
41 is_deterministic = impl->d_l_vec_indices != NULL; in CeedElemRestrictionSetupCompile_Hip()
54 …CeedCallBackend(CeedCompile_Hip(ceed, restriction_kernel_source, &impl->module, 6, "RSTR_ELEM_SIZE… in CeedElemRestrictionSetupCompile_Hip()
57 …CeedCallBackend(CeedGetKernel_Hip(ceed, impl->module, "StridedNoTranspose", &impl->ApplyNoTranspos… in CeedElemRestrictionSetupCompile_Hip()
58 … CeedCallBackend(CeedGetKernel_Hip(ceed, impl->module, "StridedTranspose", &impl->ApplyTranspose)); in CeedElemRestrictionSetupCompile_Hip()
63 …CeedCallBackend(CeedCompile_Hip(ceed, restriction_kernel_source, &impl->module, 6, "RSTR_ELEM_SIZE… in CeedElemRestrictionSetupCompile_Hip()
64 … "RSTR_NUM_COMP", num_comp, "RSTR_NUM_NODES", impl->num_nodes, "RSTR_COMP_STRIDE", comp_stride, in CeedElemRestrictionSetupCompile_Hip()
66 …CeedCallBackend(CeedGetKernel_Hip(ceed, impl->module, "OffsetNoTranspose", &impl->ApplyNoTranspose… in CeedElemRestrictionSetupCompile_Hip()
67 … CeedCallBackend(CeedGetKernel_Hip(ceed, impl->module, "OffsetTranspose", &impl->ApplyTranspose)); in CeedElemRestrictionSetupCompile_Hip()
[all …]
|
| H A D | ceed-hip-ref-vector.c | 22 CeedVector_Hip *impl; in CeedVectorNeedSync_Hip() local
25 CeedCallBackend(CeedVectorGetData(vec, &impl)); in CeedVectorNeedSync_Hip()
29 *need_sync = has_valid_array && !impl->h_array; in CeedVectorNeedSync_Hip()
32 *need_sync = has_valid_array && !impl->d_array; in CeedVectorNeedSync_Hip()
44 CeedVector_Hip *impl; in CeedVectorSyncH2D_Hip() local
46 CeedCallBackend(CeedVectorGetData(vec, &impl)); in CeedVectorSyncH2D_Hip()
48 …CeedCheck(impl->h_array, CeedVectorReturnCeed(vec), CEED_ERROR_BACKEND, "No valid host data to syn… in CeedVectorSyncH2D_Hip()
52 if (impl->d_array_borrowed) { in CeedVectorSyncH2D_Hip()
53 impl->d_array = impl->d_array_borrowed; in CeedVectorSyncH2D_Hip()
54 } else if (impl->d_array_owned) { in CeedVectorSyncH2D_Hip()
[all …]
|
| H A D | ceed-hip-ref-operator.c | 24 CeedOperator_Hip *impl; in CeedOperatorDestroy_Hip() local
26 CeedCallBackend(CeedOperatorGetData(op, &impl)); in CeedOperatorDestroy_Hip()
29 CeedCallBackend(CeedFree(&impl->num_points)); in CeedOperatorDestroy_Hip()
30 CeedCallBackend(CeedFree(&impl->skip_rstr_in)); in CeedOperatorDestroy_Hip()
31 CeedCallBackend(CeedFree(&impl->skip_rstr_out)); in CeedOperatorDestroy_Hip()
32 CeedCallBackend(CeedFree(&impl->apply_add_basis_out)); in CeedOperatorDestroy_Hip()
33 CeedCallBackend(CeedFree(&impl->input_field_order)); in CeedOperatorDestroy_Hip()
34 CeedCallBackend(CeedFree(&impl->output_field_order)); in CeedOperatorDestroy_Hip()
35 CeedCallBackend(CeedFree(&impl->input_states)); in CeedOperatorDestroy_Hip()
37 for (CeedInt i = 0; i < impl->num_inputs; i++) { in CeedOperatorDestroy_Hip()
[all …]
|
| /libCEED/backends/cuda-ref/ |
| H A D | ceed-cuda-ref-qfunctioncontext.c | 23 CeedQFunctionContext_Cuda *impl; in CeedQFunctionContextSyncH2D_Cuda() local
26 CeedCallBackend(CeedQFunctionContextGetBackendData(ctx, &impl)); in CeedQFunctionContextSyncH2D_Cuda()
28 CeedCheck(impl->h_data, ceed, CEED_ERROR_BACKEND, "No valid host data to sync to device"); in CeedQFunctionContextSyncH2D_Cuda()
31 if (impl->d_data_borrowed) { in CeedQFunctionContextSyncH2D_Cuda()
32 impl->d_data = impl->d_data_borrowed; in CeedQFunctionContextSyncH2D_Cuda()
33 } else if (impl->d_data_owned) { in CeedQFunctionContextSyncH2D_Cuda()
34 impl->d_data = impl->d_data_owned; in CeedQFunctionContextSyncH2D_Cuda()
36 CeedCallCuda(ceed, cudaMalloc((void **)&impl->d_data_owned, ctx_size)); in CeedQFunctionContextSyncH2D_Cuda()
37 impl->d_data = impl->d_data_owned; in CeedQFunctionContextSyncH2D_Cuda()
39 CeedCallCuda(ceed, cudaMemcpy(impl->d_data, impl->h_data, ctx_size, cudaMemcpyHostToDevice)); in CeedQFunctionContextSyncH2D_Cuda()
[all …]
|
| H A D | ceed-cuda-ref-restriction.c | 29 CeedElemRestriction_Cuda *impl; in CeedElemRestrictionSetupCompile_Cuda() local
31 CeedCallBackend(CeedElemRestrictionGetData(rstr, &impl)); in CeedElemRestrictionSetupCompile_Cuda()
42 is_deterministic = impl->d_l_vec_indices != NULL; in CeedElemRestrictionSetupCompile_Cuda()
55 …CeedCallBackend(CeedCompile_Cuda(ceed, restriction_kernel_source, &impl->module, 6, "RSTR_ELEM_SIZ… in CeedElemRestrictionSetupCompile_Cuda()
58 …CeedCallBackend(CeedGetKernel_Cuda(ceed, impl->module, "StridedNoTranspose", &impl->ApplyNoTranspo… in CeedElemRestrictionSetupCompile_Cuda()
59 …CeedCallBackend(CeedGetKernel_Cuda(ceed, impl->module, "StridedTranspose", &impl->ApplyTranspose)); in CeedElemRestrictionSetupCompile_Cuda()
64 …CeedCallBackend(CeedCompile_Cuda(ceed, restriction_kernel_source, &impl->module, 6, "RSTR_ELEM_SIZ… in CeedElemRestrictionSetupCompile_Cuda()
65 … "RSTR_NUM_COMP", num_comp, "RSTR_NUM_NODES", impl->num_nodes, "RSTR_COMP_STRIDE", comp_stride, in CeedElemRestrictionSetupCompile_Cuda()
67 …CeedCallBackend(CeedGetKernel_Cuda(ceed, impl->module, "OffsetNoTranspose", &impl->ApplyNoTranspos… in CeedElemRestrictionSetupCompile_Cuda()
68 … CeedCallBackend(CeedGetKernel_Cuda(ceed, impl->module, "OffsetTranspose", &impl->ApplyTranspose)); in CeedElemRestrictionSetupCompile_Cuda()
[all …]
|
| H A D | ceed-cuda-ref-vector.c | 23 CeedVector_Cuda *impl; in CeedVectorNeedSync_Cuda() local
25 CeedCallBackend(CeedVectorGetData(vec, &impl)); in CeedVectorNeedSync_Cuda()
29 *need_sync = has_valid_array && !impl->h_array; in CeedVectorNeedSync_Cuda()
32 *need_sync = has_valid_array && !impl->d_array; in CeedVectorNeedSync_Cuda()
44 CeedVector_Cuda *impl; in CeedVectorSyncH2D_Cuda() local
46 CeedCallBackend(CeedVectorGetData(vec, &impl)); in CeedVectorSyncH2D_Cuda()
48 …CeedCheck(impl->h_array, CeedVectorReturnCeed(vec), CEED_ERROR_BACKEND, "No valid host data to syn… in CeedVectorSyncH2D_Cuda()
52 if (impl->d_array_borrowed) { in CeedVectorSyncH2D_Cuda()
53 impl->d_array = impl->d_array_borrowed; in CeedVectorSyncH2D_Cuda()
54 } else if (impl->d_array_owned) { in CeedVectorSyncH2D_Cuda()
[all …]
|
| H A D | ceed-cuda-ref-operator.c | 25 CeedOperator_Cuda *impl; in CeedOperatorDestroy_Cuda() local
27 CeedCallBackend(CeedOperatorGetData(op, &impl)); in CeedOperatorDestroy_Cuda()
30 CeedCallBackend(CeedFree(&impl->num_points)); in CeedOperatorDestroy_Cuda()
31 CeedCallBackend(CeedFree(&impl->skip_rstr_in)); in CeedOperatorDestroy_Cuda()
32 CeedCallBackend(CeedFree(&impl->skip_rstr_out)); in CeedOperatorDestroy_Cuda()
33 CeedCallBackend(CeedFree(&impl->apply_add_basis_out)); in CeedOperatorDestroy_Cuda()
34 CeedCallBackend(CeedFree(&impl->input_field_order)); in CeedOperatorDestroy_Cuda()
35 CeedCallBackend(CeedFree(&impl->output_field_order)); in CeedOperatorDestroy_Cuda()
36 CeedCallBackend(CeedFree(&impl->input_states)); in CeedOperatorDestroy_Cuda()
38 for (CeedInt i = 0; i < impl->num_inputs; i++) { in CeedOperatorDestroy_Cuda()
[all …]
|
| /libCEED/backends/sycl-ref/ |
| H A D | ceed-sycl-ref-qfunctioncontext.sycl.cpp | 23 CeedQFunctionContext_Sycl *impl; in CeedQFunctionContextSyncH2D_Sycl() local
25 CeedCallBackend(CeedQFunctionContextGetBackendData(ctx, &impl)); in CeedQFunctionContextSyncH2D_Sycl()
28 CeedCheck(impl->h_data, ceed, CEED_ERROR_BACKEND, "No valid host data to sync to device"); in CeedQFunctionContextSyncH2D_Sycl()
32 if (impl->d_data_borrowed) { in CeedQFunctionContextSyncH2D_Sycl()
33 impl->d_data = impl->d_data_borrowed; in CeedQFunctionContextSyncH2D_Sycl()
34 } else if (impl->d_data_owned) { in CeedQFunctionContextSyncH2D_Sycl()
35 impl->d_data = impl->d_data_owned; in CeedQFunctionContextSyncH2D_Sycl()
37 …CeedCallSycl(ceed, impl->d_data_owned = sycl::malloc_device(ctx_size, sycl_data->sycl_device, sycl… in CeedQFunctionContextSyncH2D_Sycl()
38 impl->d_data = impl->d_data_owned; in CeedQFunctionContextSyncH2D_Sycl()
43 sycl::event copy_event = sycl_data->sycl_queue.memcpy(impl->d_data, impl->h_data, ctx_size, e); in CeedQFunctionContextSyncH2D_Sycl()
[all …]
|
| H A D | ceed-sycl-vector.sycl.cpp | 22 CeedVector_Sycl *impl; in CeedVectorNeedSync_Sycl() local
24 CeedCallBackend(CeedVectorGetData(vec, &impl)); in CeedVectorNeedSync_Sycl()
28 *need_sync = has_valid_array && !impl->h_array; in CeedVectorNeedSync_Sycl()
31 *need_sync = has_valid_array && !impl->d_array; in CeedVectorNeedSync_Sycl()
44 CeedVector_Sycl *impl; in CeedVectorSyncH2D_Sycl() local
48 CeedCallBackend(CeedVectorGetData(vec, &impl)); in CeedVectorSyncH2D_Sycl()
50 CeedCheck(impl->h_array, ceed, CEED_ERROR_BACKEND, "No valid host data to sync to device"); in CeedVectorSyncH2D_Sycl()
53 if (impl->d_array_borrowed) { in CeedVectorSyncH2D_Sycl()
54 impl->d_array = impl->d_array_borrowed; in CeedVectorSyncH2D_Sycl()
55 } else if (impl->d_array_owned) { in CeedVectorSyncH2D_Sycl()
[all …]
|
| H A D | ceed-sycl-restriction.sycl.cpp | 27 …NoTranspose_Sycl(sycl::queue &sycl_queue, const CeedElemRestriction_Sycl *impl, const CeedScalar *… in CeedElemRestrictionStridedNoTranspose_Sycl() argument
29 const CeedInt elem_size = impl->elem_size; in CeedElemRestrictionStridedNoTranspose_Sycl()
30 const CeedInt num_elem = impl->num_elem; in CeedElemRestrictionStridedNoTranspose_Sycl()
31 const CeedInt num_comp = impl->num_comp; in CeedElemRestrictionStridedNoTranspose_Sycl()
32 const CeedInt stride_nodes = impl->strides[0]; in CeedElemRestrictionStridedNoTranspose_Sycl()
33 const CeedInt stride_comp = impl->strides[1]; in CeedElemRestrictionStridedNoTranspose_Sycl()
34 const CeedInt stride_elem = impl->strides[2]; in CeedElemRestrictionStridedNoTranspose_Sycl()
54 …NoTranspose_Sycl(sycl::queue &sycl_queue, const CeedElemRestriction_Sycl *impl, const CeedScalar *… in CeedElemRestrictionOffsetNoTranspose_Sycl() argument
56 const CeedInt elem_size = impl->elem_size; in CeedElemRestrictionOffsetNoTranspose_Sycl()
57 const CeedInt num_elem = impl->num_elem; in CeedElemRestrictionOffsetNoTranspose_Sycl()
[all …]
|
| H A D | ceed-sycl-ref-basis.sycl.cpp | 39 …::queue &sycl_queue, const SyclModule_t &sycl_module, CeedInt num_elem, const CeedBasis_Sycl *impl, in CeedBasisApplyInterp_Sycl() argument
41 const CeedInt buf_len = impl->buf_len; in CeedBasisApplyInterp_Sycl()
42 const CeedInt op_len = impl->op_len; in CeedBasisApplyInterp_Sycl()
43 const CeedScalar *interp_1d = impl->d_interp_1d; in CeedBasisApplyInterp_Sycl()
47 const CeedInt work_group_size = CeedIntMin(impl->num_qpts, max_work_group_size); in CeedBasisApplyInterp_Sycl()
142 …::queue &sycl_queue, const SyclModule_t &sycl_module, CeedInt num_elem, const CeedBasis_Sycl *impl, in CeedBasisApplyGrad_Sycl() argument
144 const CeedInt buf_len = impl->buf_len; in CeedBasisApplyGrad_Sycl()
145 const CeedInt op_len = impl->op_len; in CeedBasisApplyGrad_Sycl()
146 const CeedScalar *interp_1d = impl->d_interp_1d; in CeedBasisApplyGrad_Sycl()
147 const CeedScalar *grad_1d = impl->d_grad_1d; in CeedBasisApplyGrad_Sycl()
[all …]
|
| /libCEED/backends/ref/ |
| H A D | ceed-ref-qfunctioncontext.c | 19 CeedQFunctionContext_Ref *impl; in CeedQFunctionContextHasValidData_Ref() local
21 CeedCallBackend(CeedQFunctionContextGetBackendData(ctx, &impl)); in CeedQFunctionContextHasValidData_Ref()
22 *has_valid_data = impl->data; in CeedQFunctionContextHasValidData_Ref()
30 CeedQFunctionContext_Ref *impl; in CeedQFunctionContextHasBorrowedDataOfType_Ref() local
32 CeedCallBackend(CeedQFunctionContextGetBackendData(ctx, &impl)); in CeedQFunctionContextHasBorrowedDataOfType_Ref()
34 *has_borrowed_data_of_type = impl->data_borrowed; in CeedQFunctionContextHasBorrowedDataOfType_Ref()
43 CeedQFunctionContext_Ref *impl; in CeedQFunctionContextSetData_Ref() local
45 CeedCallBackend(CeedQFunctionContextGetBackendData(ctx, &impl)); in CeedQFunctionContextSetData_Ref()
50 CeedCallBackend(CeedFree(&impl->data_owned)); in CeedQFunctionContextSetData_Ref()
53 CeedCallBackend(CeedMallocArray(1, ctx_size, &impl->data_owned)); in CeedQFunctionContextSetData_Ref()
[all …]
|
| H A D | ceed-ref-vector.c | 19 CeedVector_Ref *impl; in CeedVectorHasValidArray_Ref() local
21 CeedCallBackend(CeedVectorGetData(vec, &impl)); in CeedVectorHasValidArray_Ref()
23 *has_valid_array = impl->array; in CeedVectorHasValidArray_Ref()
31 CeedVector_Ref *impl; in CeedVectorHasBorrowedArrayOfType_Ref() local
33 CeedCallBackend(CeedVectorGetData(vec, &impl)); in CeedVectorHasBorrowedArrayOfType_Ref()
35 *has_borrowed_array_of_type = impl->array_borrowed; in CeedVectorHasBorrowedArrayOfType_Ref()
44 CeedVector_Ref *impl; in CeedVectorSetArray_Ref() local
46 CeedCallBackend(CeedVectorGetData(vec, &impl)); in CeedVectorSetArray_Ref()
51 …CeedCallBackend(CeedSetHostCeedScalarArray(array, copy_mode, length, (const CeedScalar **)&impl->a… in CeedVectorSetArray_Ref()
52 … (const CeedScalar **)&impl->array_borrowed, (const CeedScalar **)&impl->array)); in CeedVectorSetArray_Ref()
[all …]
|
| H A D | ceed-ref-operator.c | 151 CeedOperator_Ref *impl; in CeedOperatorSetup_Ref() local
156 CeedCallBackend(CeedOperatorGetData(op, &impl)); in CeedOperatorSetup_Ref()
159 CeedCallBackend(CeedQFunctionIsIdentity(qf, &impl->is_identity_qf)); in CeedOperatorSetup_Ref()
164 CeedCallBackend(CeedCalloc(num_input_fields + num_output_fields, &impl->e_vecs_full)); in CeedOperatorSetup_Ref()
166 CeedCallBackend(CeedCalloc(CEED_FIELD_MAX, &impl->skip_rstr_in)); in CeedOperatorSetup_Ref()
167 CeedCallBackend(CeedCalloc(CEED_FIELD_MAX, &impl->skip_rstr_out)); in CeedOperatorSetup_Ref()
168 CeedCallBackend(CeedCalloc(CEED_FIELD_MAX, &impl->e_data_out_indices)); in CeedOperatorSetup_Ref()
169 CeedCallBackend(CeedCalloc(CEED_FIELD_MAX, &impl->apply_add_basis_out)); in CeedOperatorSetup_Ref()
170 CeedCallBackend(CeedCalloc(CEED_FIELD_MAX, &impl->input_states)); in CeedOperatorSetup_Ref()
171 CeedCallBackend(CeedCalloc(CEED_FIELD_MAX, &impl->e_vecs_in)); in CeedOperatorSetup_Ref()
[all …]
|
| H A D | ceed-ref-qfunction.c | 21 CeedQFunction_Ref *impl; in CeedQFunctionApply_Ref() local
23 CeedCallBackend(CeedQFunctionGetData(qf, &impl)); in CeedQFunctionApply_Ref()
29 CeedCallBackend(CeedVectorGetArrayRead(U[i], CEED_MEM_HOST, &impl->inputs[i])); in CeedQFunctionApply_Ref()
32 CeedCallBackend(CeedVectorGetArrayWrite(V[i], CEED_MEM_HOST, &impl->outputs[i])); in CeedQFunctionApply_Ref()
35 CeedCallBackend(f(ctx_data, Q, impl->inputs, impl->outputs)); in CeedQFunctionApply_Ref()
38 CeedCallBackend(CeedVectorRestoreArrayRead(U[i], &impl->inputs[i])); in CeedQFunctionApply_Ref()
41 CeedCallBackend(CeedVectorRestoreArray(V[i], &impl->outputs[i])); in CeedQFunctionApply_Ref()
51 CeedQFunction_Ref *impl; in CeedQFunctionDestroy_Ref() local
53 CeedCallBackend(CeedQFunctionGetData(qf, &impl)); in CeedQFunctionDestroy_Ref()
54 CeedCallBackend(CeedFree(&impl->inputs)); in CeedQFunctionDestroy_Ref()
[all …]
|
| /libCEED/backends/magma/ |
| H A D | ceed-magma-basis.c | 36 CeedBasis_Magma *impl; in CeedBasisApplyCore_Magma() local
40 CeedCallBackend(CeedBasisGetData(basis, &impl)); in CeedBasisApplyCore_Magma()
117 …void *args[] = {&impl->d_interp_1d, &d_u, &u_elem_stride, &u_comp_stride, &d_v, &v_elem_stride, … in CeedBasisApplyCore_Magma()
120 …eedCallBackend(CeedRunKernelDimSharedMagma(ceed, apply_add ? impl->InterpTransposeAdd : impl->Inte… in CeedBasisApplyCore_Magma()
123 …CeedCallBackend(CeedRunKernelDimSharedMagma(ceed, impl->Interp, NULL, grid, num_threads, num_t_col… in CeedBasisApplyCore_Magma()
194 …void *args[] = {&impl->d_interp_1d, &impl->d_grad_1d, &d_u, &u_elem_stride, &u_comp_str… in CeedBasisApplyCore_Magma()
198 …CeedCallBackend(CeedRunKernelDimSharedMagma(ceed, apply_add ? impl->GradTransposeAdd : impl->GradT… in CeedBasisApplyCore_Magma()
201 …CeedCallBackend(CeedRunKernelDimSharedMagma(ceed, impl->Grad, NULL, grid, num_threads, num_t_col, … in CeedBasisApplyCore_Magma()
206 …CeedCheck(impl->d_q_weight_1d, ceed, CEED_ERROR_BACKEND, "%s not supported; q_weight_1d not set", … in CeedBasisApplyCore_Magma()
231 void *args[] = {&impl->d_q_weight_1d, &d_v, &elem_dofs_size, &num_elem}; in CeedBasisApplyCore_Magma()
[all …]
|
| /libCEED/backends/sycl-shared/ |
| H A D | ceed-sycl-shared-basis.sycl.cpp | 42 CeedBasis_Sycl_shared *impl; in CeedBasisApplyTensor_Sycl_shared() local
46 CeedCallBackend(CeedBasisGetData(basis, &impl)); in CeedBasisApplyTensor_Sycl_shared()
56 CeedInt *lrange = impl->interp_local_range; in CeedBasisApplyTensor_Sycl_shared()
64 …ycl::kernel *interp_kernel = (t_mode == CEED_TRANSPOSE) ? impl->interp_transpose_kernel : impl->in… in CeedBasisApplyTensor_Sycl_shared()
71 cgh.set_args(num_elem, impl->d_interp_1d, d_u, d_v); in CeedBasisApplyTensor_Sycl_shared()
77 CeedInt *lrange = impl->grad_local_range; in CeedBasisApplyTensor_Sycl_shared()
85 …cl::kernel *grad_kernel = (t_mode == CEED_TRANSPOSE) ? impl->grad_transpose_kernel : impl->gra… in CeedBasisApplyTensor_Sycl_shared()
86 … const CeedScalar *d_grad_1d = (impl->d_collo_grad_1d) ? impl->d_collo_grad_1d : impl->d_grad_1d; in CeedBasisApplyTensor_Sycl_shared()
94 cgh.set_args(num_elem, impl->d_interp_1d, d_grad_1d, d_u, d_v); in CeedBasisApplyTensor_Sycl_shared()
99 CeedInt *lrange = impl->weight_local_range; in CeedBasisApplyTensor_Sycl_shared()
[all …]
|
| /libCEED/backends/blocked/ |
| H A D | ceed-blocked-operator.c | 209 CeedOperator_Blocked *impl; in CeedOperatorSetup_Blocked() local
214 CeedCallBackend(CeedOperatorGetData(op, &impl)); in CeedOperatorSetup_Blocked()
217 CeedCallBackend(CeedQFunctionIsIdentity(qf, &impl->is_identity_qf)); in CeedOperatorSetup_Blocked()
222 CeedCallBackend(CeedCalloc(num_input_fields + num_output_fields, &impl->block_rstr)); in CeedOperatorSetup_Blocked()
223 CeedCallBackend(CeedCalloc(num_input_fields + num_output_fields, &impl->e_vecs_full)); in CeedOperatorSetup_Blocked()
225 CeedCallBackend(CeedCalloc(CEED_FIELD_MAX, &impl->skip_rstr_in)); in CeedOperatorSetup_Blocked()
226 CeedCallBackend(CeedCalloc(CEED_FIELD_MAX, &impl->skip_rstr_out)); in CeedOperatorSetup_Blocked()
227 CeedCallBackend(CeedCalloc(CEED_FIELD_MAX, &impl->e_data_out_indices)); in CeedOperatorSetup_Blocked()
228 CeedCallBackend(CeedCalloc(CEED_FIELD_MAX, &impl->apply_add_basis_out)); in CeedOperatorSetup_Blocked()
229 CeedCallBackend(CeedCalloc(CEED_FIELD_MAX, &impl->input_states)); in CeedOperatorSetup_Blocked()
[all …]
|
| /libCEED/backends/opt/ |
| H A D | ceed-opt-operator.c | 213 CeedOperator_Opt *impl; in CeedOperatorSetup_Opt() local
221 CeedCallBackend(CeedOperatorGetData(op, &impl)); in CeedOperatorSetup_Opt()
224 CeedCallBackend(CeedQFunctionIsIdentity(qf, &impl->is_identity_qf)); in CeedOperatorSetup_Opt()
230 CeedCallBackend(CeedCalloc(num_input_fields + num_output_fields, &impl->block_rstr)); in CeedOperatorSetup_Opt()
231 CeedCallBackend(CeedCalloc(num_input_fields + num_output_fields, &impl->e_vecs_full)); in CeedOperatorSetup_Opt()
233 CeedCallBackend(CeedCalloc(CEED_FIELD_MAX, &impl->skip_rstr_in)); in CeedOperatorSetup_Opt()
234 CeedCallBackend(CeedCalloc(CEED_FIELD_MAX, &impl->skip_rstr_out)); in CeedOperatorSetup_Opt()
235 CeedCallBackend(CeedCalloc(CEED_FIELD_MAX, &impl->apply_add_basis_out)); in CeedOperatorSetup_Opt()
236 CeedCallBackend(CeedCalloc(CEED_FIELD_MAX, &impl->input_states)); in CeedOperatorSetup_Opt()
237 CeedCallBackend(CeedCalloc(CEED_FIELD_MAX, &impl->e_vecs_in)); in CeedOperatorSetup_Opt()
[all …]
|
| /libCEED/backends/sycl-gen/ |
| H A D | ceed-sycl-gen-operator.sycl.cpp | 20 CeedOperator_Sycl_gen *impl; in CeedOperatorDestroy_Sycl_gen() local
22 CeedCallBackend(CeedOperatorGetData(op, &impl)); in CeedOperatorDestroy_Sycl_gen()
23 CeedCallBackend(CeedFree(&impl)); in CeedOperatorDestroy_Sycl_gen()
40 CeedOperator_Sycl_gen *impl; in CeedOperatorApplyAdd_Sycl_gen() local
60 CeedCallBackend(CeedOperatorGetData(op, &impl)); in CeedOperatorApplyAdd_Sycl_gen()
74 impl->fields->inputs[i] = NULL; in CeedOperatorApplyAdd_Sycl_gen()
83 CeedCallBackend(CeedVectorGetArrayRead(vec, CEED_MEM_DEVICE, &impl->fields->inputs[i])); in CeedOperatorApplyAdd_Sycl_gen()
92 impl->fields->outputs[i] = NULL; in CeedOperatorApplyAdd_Sycl_gen()
111 CeedCallBackend(CeedVectorGetArray(vec, CEED_MEM_DEVICE, &impl->fields->outputs[i])); in CeedOperatorApplyAdd_Sycl_gen()
113 impl->fields->outputs[i] = impl->fields->outputs[index]; in CeedOperatorApplyAdd_Sycl_gen()
[all …]
|