| /libCEED/backends/xsmm/ |
| H A D | ceed-xsmm-tensor.c | 17 static int CeedTensorContractApply_Xsmm(CeedTensorContract contract, CeedInt A, CeedInt B, CeedInt … in CeedTensorContractApply_Xsmm() argument 25 …? libxsmm_create_gemm_shape(J, A, B, !t_mode ? B : J, B, J, LIBXSMM_DATATYPE_F64, LIBXSMM_DATATYPE… in CeedTensorContractApply_Xsmm() 27 …: libxsmm_create_gemm_shape(J, A, B, !t_mode ? B : J, B, J, LIBXSMM_DATATYPE_F32, LIBXSMM_DATATYPE… in CeedTensorContractApply_Xsmm() 45 …? libxsmm_create_gemm_shape(C, J, B, C, !t_mode ? B : J, C, LIBXSMM_DATATYPE_F64, LIBXSMM_DATATYPE… in CeedTensorContractApply_Xsmm() 47 …: libxsmm_create_gemm_shape(C, J, B, C, !t_mode ? B : J, C, LIBXSMM_DATATYPE_F32, LIBXSMM_DATATYPE… in CeedTensorContractApply_Xsmm() 57 gemm_param.a.primary = (CeedScalar *)&u[a * B * C]; in CeedTensorContractApply_Xsmm()
|
| /libCEED/backends/avx/ |
| H A D | ceed-avx-tensor.c | 42 static inline int CeedTensorContract_Avx_Blocked(CeedTensorContract contract, CeedInt A, CeedInt B,… in CeedTensorContract_Avx_Blocked() argument 45 CeedInt t_stride_0 = B, t_stride_1 = 1; in CeedTensorContract_Avx_Blocked() 60 for (CeedInt b = 0; b < B; b++) { in CeedTensorContract_Avx_Blocked() 64 fmadd(vv[jj][cc], tqv, loadu(&u[(a * B + b) * C + c + cc * 4])); in CeedTensorContract_Avx_Blocked() 83 for (CeedInt b = 0; b < B; b++) { in CeedTensorContract_Avx_Blocked() 88 fmadd(vv[jj][cc], tqv, loadu(&u[(a * B + b) * C + c + cc * 4])); in CeedTensorContract_Avx_Blocked() 104 …ensorContract_Avx_Remainder(CeedTensorContract contract, CeedInt A, CeedInt B, CeedInt C, CeedInt … in CeedTensorContract_Avx_Remainder() argument 107 CeedInt t_stride_0 = B, t_stride_1 = 1; in CeedTensorContract_Avx_Remainder() 124 for (CeedInt b = 0; b < B; b++) { in CeedTensorContract_Avx_Remainder() 127 if (C - c == 1) tqu = set(0.0, 0.0, 0.0, u[(a * B + b) * C + c + 0]); in CeedTensorContract_Avx_Remainder() [all …]
|
| /libCEED/backends/opt/ |
| H A D | ceed-opt-tensor.c | 16 …ensorContractApply_Core_Opt(CeedTensorContract contract, CeedInt A, CeedInt B, CeedInt C, CeedInt … in CeedTensorContractApply_Core_Opt() argument 19 CeedInt t_stride_0 = B, t_stride_1 = 1; in CeedTensorContractApply_Core_Opt() 27 for (CeedInt b = 0; b < B; b++) { in CeedTensorContractApply_Core_Opt() 30 for (CeedInt c = 0; c < C; c++) v[(a * J + j) * C + c] += tq * u[(a * B + b) * C + c]; in CeedTensorContractApply_Core_Opt() 40 static int CeedTensorContractApply_Opt(CeedTensorContract contract, CeedInt A, CeedInt B, CeedInt C… in CeedTensorContractApply_Opt() argument 46 if (C == 1) return CeedTensorContractApply_Core_Opt(contract, A, B, 1, J, t, t_mode, add, u, v); in CeedTensorContractApply_Opt() 47 else return CeedTensorContractApply_Core_Opt(contract, A, B, C, J, t, t_mode, add, u, v); in CeedTensorContractApply_Opt()
|
| /libCEED/interface/ |
| H A D | ceed-tensor.c | 94 int CeedTensorContractApply(CeedTensorContract contract, CeedInt A, CeedInt B, CeedInt C, CeedInt J… in CeedTensorContractApply() argument 96 CeedCall(contract->Apply(contract, A, B, C, J, t, t_mode, add, u, v)); in CeedTensorContractApply() 124 int CeedTensorContractStridedApply(CeedTensorContract contract, CeedInt A, CeedInt B, CeedInt C, Ce… in CeedTensorContractStridedApply() argument 128 … CeedCall(contract->Apply(contract, A, J, C, B, t + d * B * J, t_mode, add, u + d * A * J * C, v)); in CeedTensorContractStridedApply() 132 … CeedCall(contract->Apply(contract, A, B, C, J, t + d * B * J, t_mode, add, u, v + d * A * J * C)); in CeedTensorContractStridedApply()
|
| /libCEED/include/ceed/jit-source/sycl/ |
| H A D | sycl-shared-basis-tensor-templates.h | 19 …P_1D, const CeedInt Q_1D, private const CeedScalar *restrict U, local const CeedScalar *restrict B, in ContractX1d() argument 29 *V += B[i + item_id_x * P_1D] * scratch[i]; // Contract x direction in ContractX1d() 38 …P_1D, const CeedInt Q_1D, private const CeedScalar *restrict U, local const CeedScalar *restrict B, in ContractTransposeX1d() argument 48 *V += B[item_id_x + i * P_1D] * scratch[i]; // Contract x direction in ContractTransposeX1d() 109 …P_1D, const CeedInt Q_1D, private const CeedScalar *restrict U, local const CeedScalar *restrict B, in ContractX2d() argument 120 *V += B[i + item_id_x * P_1D] * scratch[i + item_id_y * T_1D]; // Contract x direction in ContractX2d() 129 …P_1D, const CeedInt Q_1D, private const CeedScalar *restrict U, local const CeedScalar *restrict B, in ContractY2d() argument 140 *V += B[i + item_id_y * P_1D] * scratch[item_id_x + i * T_1D]; // Contract y direction in ContractY2d() 149 …P_1D, const CeedInt Q_1D, private const CeedScalar *restrict U, local const CeedScalar *restrict B, in ContractTransposeY2d() argument 160 *V += B[item_id_y + i * P_1D] * scratch[item_id_x + i * T_1D]; // Contract y direction in ContractTransposeY2d() [all …]
|
| H A D | sycl-shared-basis-read-write-templates.h | 15 inline void loadMatrix(const CeedInt N, const CeedScalar *restrict d_B, CeedScalar *restrict B) { in loadMatrix() argument 18 for (CeedInt i = item_id; i < N; i += group_size) B[i] = d_B[i]; in loadMatrix()
|
| /libCEED/examples/fluids/qfunctions/ |
| H A D | utils.h | 44 …id CopyMat3(const CeedScalar A[3][3], CeedScalar B[3][3]) { CopyN((const CeedScalar *)A, (CeedScal… in CopyMat3() 91 CeedScalar *B) { in MatDiagNM() argument 94 …CeedInt i = 0; i < N; i++) { CeedPragmaSIMD for (CeedInt j = 0; j < M; j++) B[i * M + j] += D[i] *… in MatDiagNM() 97 …CeedInt i = 0; i < M; i++) { CeedPragmaSIMD for (CeedInt j = 0; j < N; j++) B[i * N + j] += D[i] *… in MatDiagNM() 104 …edScalar A[3][3], const CeedScalar D[3], const CeedTransposeMode transpose_A, CeedScalar B[3][3]) { in MatDiag3() 105 MatDiagNM((const CeedScalar *)A, (const CeedScalar *)D, 3, 3, transpose_A, (CeedScalar *)B); in MatDiag3() 108 CEED_QFUNCTION_HELPER void MatMatN(const CeedScalar *A, const CeedScalar *B, const CeedInt N, const… in MatMatN() argument 116 … CeedPragmaSIMD for (CeedInt k = 0; k < N; k++) C[i * N + j] += A[i * N + k] * B[k * N + j]; in MatMatN() 123 … CeedPragmaSIMD for (CeedInt k = 0; k < N; k++) C[i * N + j] += A[i * N + k] * B[j * N + k]; in MatMatN() 134 … CeedPragmaSIMD for (CeedInt k = 0; k < N; k++) C[i * N + j] += A[k * N + i] * B[k * N + j]; in MatMatN() [all …]
|
| /libCEED/backends/ref/ |
| H A D | ceed-ref-tensor.c | 16 static int CeedTensorContractApply_Ref(CeedTensorContract contract, CeedInt A, CeedInt B, CeedInt C… in CeedTensorContractApply_Ref() argument 18 CeedInt t_stride_0 = B, t_stride_1 = 1; in CeedTensorContractApply_Ref() 30 for (CeedInt b = 0; b < B; b++) { in CeedTensorContractApply_Ref() 33 for (CeedInt c = 0; c < C; c++) v[(a * J + j) * C + c] += tq * u[(a * B + b) * C + c]; in CeedTensorContractApply_Ref()
|
| /libCEED/include/ceed/jit-source/cuda/ |
| H A D | cuda-shared-basis-tensor-templates.h | 20 inline __device__ void ContractX1d(SharedData_Cuda &data, const CeedScalar *U, const CeedScalar *B,… in ContractX1d() argument 27 *V += B[i + data.t_id_x * P_1D] * data.slice[i]; // Contract x direction in ContractX1d() 36 …tractTransposeX1d(SharedData_Cuda &data, const CeedScalar *U, const CeedScalar *B, CeedScalar *V) { in ContractTransposeX1d() argument 43 *V += B[data.t_id_x + i * P_1D] * data.slice[i]; // Contract x direction in ContractTransposeX1d() 129 inline __device__ void ContractX2d(SharedData_Cuda &data, const CeedScalar *U, const CeedScalar *B,… in ContractX2d() argument 136 *V += B[i + data.t_id_x * P_1D] * data.slice[i + data.t_id_y * T_1D]; // Contract x direction in ContractX2d() 145 inline __device__ void ContractY2d(SharedData_Cuda &data, const CeedScalar *U, const CeedScalar *B,… in ContractY2d() argument 152 *V += B[i + data.t_id_y * P_1D] * data.slice[data.t_id_x + i * T_1D]; // Contract y direction in ContractY2d() 161 …tractTransposeY2d(SharedData_Cuda &data, const CeedScalar *U, const CeedScalar *B, CeedScalar *V) { in ContractTransposeY2d() argument 168 *V += B[data.t_id_y + i * P_1D] * data.slice[data.t_id_x + i * T_1D]; // Contract y direction in ContractTransposeY2d() [all …]
|
| H A D | cuda-shared-basis-nontensor-templates.h | 16 inline __device__ void Contract1d(SharedData_Cuda &data, const CeedScalar *U, const CeedScalar *B, … in Contract1d() argument 22 *V += B[i + data.t_id_x * P_1D] * data.slice[i]; // Contract x direction in Contract1d() 32 …ntractTranspose1d(SharedData_Cuda &data, const CeedScalar *U, const CeedScalar *B, CeedScalar *V) { in ContractTranspose1d() argument 37 *V += B[data.t_id_x + i * P_1D] * data.slice[i]; // Contract x direction in ContractTranspose1d()
|
| H A D | cuda-shared-basis-tensor-flattened-templates.h | 20 …haredData_Cuda &data, const int t_id_x, const int t_id_y, const CeedScalar *U, const CeedScalar *B, in ContractX2dFlattened() argument 28 *V += B[i + t_id_x * P_1D] * data.slice[i + t_id_y * T_1D]; // Contract x direction in ContractX2dFlattened() 37 …haredData_Cuda &data, const int t_id_x, const int t_id_y, const CeedScalar *U, const CeedScalar *B, in ContractY2dFlattened() argument 45 *V += B[i + t_id_y * P_1D] * data.slice[t_id_x + i * T_1D]; // Contract y direction in ContractY2dFlattened() 55 const CeedScalar *B, CeedScalar *V) { in ContractTransposeY2dFlattened() argument 62 *V += B[t_id_y + i * P_1D] * data.slice[t_id_x + i * T_1D]; // Contract y direction in ContractTransposeY2dFlattened() 72 const CeedScalar *B, CeedScalar *V) { in ContractTransposeX2dFlattened() argument 79 *V += B[t_id_x + i * P_1D] * data.slice[i + t_id_y * T_1D]; // Contract x direction in ContractTransposeX2dFlattened() 89 const CeedScalar *B, CeedScalar *V) { in ContractTransposeAddX2dFlattened() argument 95 *V += B[t_id_x + i * P_1D] * data.slice[i + t_id_y * T_1D]; // Contract x direction in ContractTransposeAddX2dFlattened() [all …]
|
| /libCEED/include/ceed/jit-source/hip/ |
| H A D | hip-shared-basis-tensor-templates.h | 20 inline __device__ void ContractX1d(SharedData_Hip &data, const CeedScalar *U, const CeedScalar *B, … in ContractX1d() argument 27 *V += B[i + data.t_id_x * P_1D] * data.slice[i]; // Contract x direction in ContractX1d() 36 …ntractTransposeX1d(SharedData_Hip &data, const CeedScalar *U, const CeedScalar *B, CeedScalar *V) { in ContractTransposeX1d() argument 43 *V += B[data.t_id_x + i * P_1D] * data.slice[i]; // Contract x direction in ContractTransposeX1d() 129 inline __device__ void ContractX2d(SharedData_Hip &data, const CeedScalar *U, const CeedScalar *B, … in ContractX2d() argument 136 *V += B[i + data.t_id_x * P_1D] * data.slice[i + data.t_id_y * T_1D]; // Contract x direction in ContractX2d() 145 inline __device__ void ContractY2d(SharedData_Hip &data, const CeedScalar *U, const CeedScalar *B, … in ContractY2d() argument 152 *V += B[i + data.t_id_y * P_1D] * data.slice[data.t_id_x + i * T_1D]; // Contract y direction in ContractY2d() 161 …ntractTransposeY2d(SharedData_Hip &data, const CeedScalar *U, const CeedScalar *B, CeedScalar *V) { in ContractTransposeY2d() argument 168 *V += B[data.t_id_y + i * P_1D] * data.slice[data.t_id_x + i * T_1D]; // Contract y direction in ContractTransposeY2d() [all …]
|
| H A D | hip-shared-basis-nontensor-templates.h | 16 inline __device__ void Contract1d(SharedData_Hip &data, const CeedScalar *U, const CeedScalar *B, C… in Contract1d() argument 22 *V += B[i + data.t_id_x * P_1D] * data.slice[i]; // Contract x direction in Contract1d() 32 …ontractTranspose1d(SharedData_Hip &data, const CeedScalar *U, const CeedScalar *B, CeedScalar *V) { in ContractTranspose1d() argument 37 *V += B[data.t_id_x + i * P_1D] * data.slice[i]; // Contract x direction in ContractTranspose1d()
|
| H A D | hip-shared-basis-tensor-flattened-templates.h | 20 …SharedData_Hip &data, const int t_id_x, const int t_id_y, const CeedScalar *U, const CeedScalar *B, in ContractX2dFlattened() argument 28 *V += B[i + t_id_x * P_1D] * data.slice[i + t_id_y * T_1D]; // Contract x direction in ContractX2dFlattened() 37 …SharedData_Hip &data, const int t_id_x, const int t_id_y, const CeedScalar *U, const CeedScalar *B, in ContractY2dFlattened() argument 45 *V += B[i + t_id_y * P_1D] * data.slice[t_id_x + i * T_1D]; // Contract y direction in ContractY2dFlattened() 55 const CeedScalar *B, CeedScalar *V) { in ContractTransposeY2dFlattened() argument 62 *V += B[t_id_y + i * P_1D] * data.slice[t_id_x + i * T_1D]; // Contract y direction in ContractTransposeY2dFlattened() 72 const CeedScalar *B, CeedScalar *V) { in ContractTransposeX2dFlattened() argument 79 *V += B[t_id_x + i * P_1D] * data.slice[i + t_id_y * T_1D]; // Contract x direction in ContractTransposeX2dFlattened() 89 const CeedScalar *B, CeedScalar *V) { in ContractTransposeAddX2dFlattened() argument 95 *V += B[t_id_x + i * P_1D] * data.slice[i + t_id_y * T_1D]; // Contract x direction in ContractTransposeAddX2dFlattened() [all …]
|
| /libCEED/include/ceed/jit-source/magma/ |
| H A D | magma-common-tensor.h | 231 template <int B, int J> 233 if (tx < B) { in read_T_notrans_gm2sm() 235 sT[i * B + tx] = dT[i * B + tx]; in read_T_notrans_gm2sm() 245 template <int B, int J> 248 for (int i = 0; i < B; i++) { in read_T_trans_gm2sm() 249 sT[tx * B + i] = dT[i * J + tx]; in read_T_trans_gm2sm()
|
| /libCEED/tests/ |
| H A D | t580-operator.h | 14 …aMatTransposeMatMult2x2(const CeedScalar alpha, const CeedScalar A[2][2], const CeedScalar B[2][2], in AlphaMatTransposeMatMult2x2() 20 C[j][k] += alpha * A[m][j] * B[m][k]; in AlphaMatTransposeMatMult2x2()
|
| /libCEED/doc/sphinx/source/api/ |
| H A D | index.rst | 25 O --> B(CeedBasis) 27 B --> V(CeedVector)
|
| /libCEED/examples/mfem/ |
| H A D | bp3.cpp | 152 mfem::Vector X, B; in main() local 153 diff.FormLinearSystem(ess_tdof_list, sol, b, D, X, B); in main() 166 cg.Mult(B, X); in main()
|
| /libCEED/backends/sycl-gen/ |
| H A D | ceed-sycl-gen.hpp | 24 Fields_Sycl *B; member
|
| /libCEED/backends/cuda-gen/ |
| H A D | ceed-cuda-gen.h | 24 Fields_Cuda B; member
|
| /libCEED/backends/hip-gen/ |
| H A D | ceed-hip-gen.h | 25 Fields_Hip B; member
|
| /libCEED/doc/bib/ |
| H A D | references.bib | 47 author = {{V}aleria {B}arra and {J}ed {B}rown and {J}eremy {T}hompson and {Y}ohann {D}udouit},
|
| /libCEED/examples/fluids/problems/ |
| H A D | channel.c | 84 channel_ctx->B = body_force_scale * 2 * umax * newtonian_ig_ctx->mu / (H * H); in NS_CHANNEL() 91 CeedScalar g[] = {channel_ctx->B / rho, 0., 0.}; in NS_CHANNEL()
|
| /libCEED/doc/sphinx/source/ |
| H A D | libCEEDapi.md | 57 - Basis (Dofs-to-Qpts) evaluator $\bm{B}$ 60 …t and trial space differ, they get their own versions of $\bm{P}$, $\bm{\mathcal{E}}$ and $\bm{B}$. 158 Since the global operator $\bm{A}$ is just a series of variational restrictions with $\bm{B}$, $\bm… 163 …r portions of it) and evaluate the actions of $\bm{P}$, $\bm{\mathcal{E}}$ and $\bm{B}$ on-the-fly. 164 …uadrature points on *quad* and *hex* elements to perform the action of $\bm{B}$ without storing it… 173 …m{\mathcal{E}}$ and $\bm{B}$, the operator evaluation is decoupled on their ranges, so $\bm{P}$, … 177 …B}$ and $\bm{D}$ clearly separate the MPI parallelism in the operator ($\bm{P}$) from the unstruct… 178 …athcal{E}}$, dense/structured linear algebra (tensor contractions) for $\bm{B}$ and parallel point… 191 …lly describes in the *frontend* the operators $\bm{\bm{\mathcal{E}}}$, $\bm{B}$, and $\bm{D}$ and … 197 …bCEED Operators, through backend implementations of $\bm{\bm{\mathcal{E}}}$, $\bm{B}$, and $\bm{D}$ [all …]
|
| /libCEED/julia/LibCEED.jl/docs/src/ |
| H A D | UserQFunctions.md | 19 The mass operator on each element can be written as $B^\intercal D B$, where $B$ 56 When adding the inputs and outputs, `CEED_EVAL_INTERP` indicates that the $B$ 106 action of the diffusion operator. When written in the form $B^\intercal D B$, in 107 this case $B$ represents the basis gradient matrix, and $D$ represents
|