Home
last modified time | relevance | path

Searched refs:elem_diag_array (Results 1 – 6 of 6) sorted by relevance

/libCEED/include/ceed/jit-source/cuda/
H A Dcuda-ref-operator-assemble-diagonal.h51 … const CeedScalar *__restrict__ assembled_qf_array, CeedScalar *__restrict__ elem_diag_array) { in __launch_bounds__()
96elem_diag_array[((comp_out * NUM_COMP + comp_in) * num_elem + e) * NUM_NODES + tid] += e_value; in __launch_bounds__()
110 elem_diag_array[(comp_out * num_elem + e) * NUM_NODES + tid] += e_value; in __launch_bounds__()
/libCEED/include/ceed/jit-source/hip/
H A Dhip-ref-operator-assemble-diagonal.h51 … const CeedScalar *__restrict__ assembled_qf_array, CeedScalar *__restrict__ elem_diag_array) { in __launch_bounds__()
96elem_diag_array[((comp_out * NUM_COMP + comp_in) * num_elem + e) * NUM_NODES + tid] += e_value; in __launch_bounds__()
110 elem_diag_array[(comp_out * num_elem + e) * NUM_NODES + tid] += e_value; in __launch_bounds__()
/libCEED/backends/sycl-ref/
H A Dceed-sycl-ref-operator.sycl.cpp833 …t CeedOperatorDiag_Sycl *diag, const CeedScalar *assembled_qf_array, CeedScalar *elem_diag_array) { in CeedOperatorLinearDiagonal_Sycl() argument
888elem_diag_array[((comp_out * num_comp + comp_in) * num_elem + e) * num_nodes + tid] += e_value; in CeedOperatorLinearDiagonal_Sycl()
900 elem_diag_array[(comp_out * num_elem + e) * num_nodes + tid] += e_value; in CeedOperatorLinearDiagonal_Sycl()
916 CeedScalar *elem_diag_array; in CeedOperatorAssembleDiagonalCore_Sycl() local
959 CeedCallBackend(CeedVectorGetArray(elem_diag, CEED_MEM_DEVICE, &elem_diag_array)); in CeedOperatorAssembleDiagonalCore_Sycl()
964 …_Sycl(sycl_data->sycl_queue, is_point_block, num_elem, diag, assembled_qf_array, elem_diag_array)); in CeedOperatorAssembleDiagonalCore_Sycl()
970 CeedCallBackend(CeedVectorRestoreArray(elem_diag, &elem_diag_array)); in CeedOperatorAssembleDiagonalCore_Sycl()
/libCEED/interface/
H A Dceed-preconditioning.c253 CeedScalar *elem_diag_array, *identity = NULL; in CeedOperatorLinearAssembleAddDiagonalSingle_Mesh() local
293 CeedCall(CeedVectorGetArray(elem_diag, CEED_MEM_HOST, &elem_diag_array)); in CeedOperatorLinearAssembleAddDiagonalSingle_Mesh()
355 elem_diag_array[((e * num_comp + c_out) * num_comp + c_in) * num_nodes + n] += in CeedOperatorLinearAssembleAddDiagonalSingle_Mesh()
366elem_diag_array[(e * num_comp + c_out) * num_nodes + n] += B_t[q * num_nodes + n] * qf_value * B[q… in CeedOperatorLinearAssembleAddDiagonalSingle_Mesh()
374 CeedCall(CeedVectorRestoreArray(elem_diag, &elem_diag_array)); in CeedOperatorLinearAssembleAddDiagonalSingle_Mesh()
/libCEED/backends/hip-ref/
H A Dceed-hip-ref-operator.c1406 CeedScalar *elem_diag_array; in CeedOperatorAssembleDiagonalCore_Hip() local
1459 CeedCallBackend(CeedVectorGetArray(elem_diag, CEED_MEM_DEVICE, &elem_diag_array)); in CeedOperatorAssembleDiagonalCore_Hip()
1466 … &diag->d_eval_modes_in, &diag->d_eval_modes_out, &assembled_qf_array, &elem_diag_array}; in CeedOperatorAssembleDiagonalCore_Hip()
1475 CeedCallBackend(CeedVectorRestoreArray(elem_diag, &elem_diag_array)); in CeedOperatorAssembleDiagonalCore_Hip()
/libCEED/backends/cuda-ref/
H A Dceed-cuda-ref-operator.c1409 CeedScalar *elem_diag_array; in CeedOperatorAssembleDiagonalCore_Cuda() local
1462 CeedCallBackend(CeedVectorGetArray(elem_diag, CEED_MEM_DEVICE, &elem_diag_array)); in CeedOperatorAssembleDiagonalCore_Cuda()
1469 … &diag->d_eval_modes_in, &diag->d_eval_modes_out, &assembled_qf_array, &elem_diag_array}; in CeedOperatorAssembleDiagonalCore_Cuda()
1478 CeedCallBackend(CeedVectorRestoreArray(elem_diag, &elem_diag_array)); in CeedOperatorAssembleDiagonalCore_Cuda()