| /libCEED/include/ceed/jit-source/hip/ |
| H A D | hip-ref-restriction-curl-oriented.h | 17 …for (CeedInt node = blockIdx.x * blockDim.x + threadIdx.x; node < RSTR_NUM_ELEM * RSTR_ELEM_SIZE; … in CurlOrientedNoTranspose() local 18 const CeedInt loc_node = node % RSTR_ELEM_SIZE; in CurlOrientedNoTranspose() 19 const CeedInt elem = node / RSTR_ELEM_SIZE; in CurlOrientedNoTranspose() 20 const CeedInt ind_dl = loc_node > 0 ? indices[node - 1] : 0; in CurlOrientedNoTranspose() 21 const CeedInt ind_d = indices[node]; in CurlOrientedNoTranspose() 22 const CeedInt ind_du = loc_node < (RSTR_ELEM_SIZE - 1) ? indices[node + 1] : 0; in CurlOrientedNoTranspose() 23 const CeedInt8 curl_orient_dl = curl_orients[3 * node + 0]; in CurlOrientedNoTranspose() 24 const CeedInt8 curl_orient_d = curl_orients[3 * node + 1]; in CurlOrientedNoTranspose() 25 const CeedInt8 curl_orient_du = curl_orients[3 * node + 2]; in CurlOrientedNoTranspose() 42 …for (CeedInt node = blockIdx.x * blockDim.x + threadIdx.x; node < RSTR_NUM_ELEM * RSTR_ELEM_SIZE; … in CurlOrientedUnsignedNoTranspose() local [all …]
|
| H A D | hip-ref-restriction-strided.h | 16 …for (CeedInt node = blockIdx.x * blockDim.x + threadIdx.x; node < RSTR_NUM_ELEM * RSTR_ELEM_SIZE; … in StridedNoTranspose() local 17 const CeedInt loc_node = node % RSTR_ELEM_SIZE; in StridedNoTranspose() 18 const CeedInt elem = node / RSTR_ELEM_SIZE; in StridedNoTranspose() 31 …for (CeedInt node = blockIdx.x * blockDim.x + threadIdx.x; node < RSTR_NUM_ELEM * RSTR_ELEM_SIZE; … in StridedTranspose() local 32 const CeedInt loc_node = node % RSTR_ELEM_SIZE; in StridedTranspose() 33 const CeedInt elem = node / RSTR_ELEM_SIZE; in StridedTranspose()
|
| H A D | hip-ref-restriction-oriented.h | 17 …for (CeedInt node = blockIdx.x * blockDim.x + threadIdx.x; node < RSTR_NUM_ELEM * RSTR_ELEM_SIZE; … in OrientedNoTranspose() local 18 const CeedInt ind = indices[node]; in OrientedNoTranspose() 19 const bool orient = orients[node]; in OrientedNoTranspose() 20 const CeedInt loc_node = node % RSTR_ELEM_SIZE; in OrientedNoTranspose() 21 const CeedInt elem = node / RSTR_ELEM_SIZE; in OrientedNoTranspose() 35 …for (CeedInt node = blockIdx.x * blockDim.x + threadIdx.x; node < RSTR_NUM_ELEM * RSTR_ELEM_SIZE; … in OrientedTranspose() local 36 const CeedInt ind = indices[node]; in OrientedTranspose() 37 const bool orient = orients[node]; in OrientedTranspose() 38 const CeedInt loc_node = node % RSTR_ELEM_SIZE; in OrientedTranspose() 39 const CeedInt elem = node / RSTR_ELEM_SIZE; in OrientedTranspose()
|
| H A D | hip-ref-restriction-offset.h | 16 …for (CeedInt node = blockIdx.x * blockDim.x + threadIdx.x; node < RSTR_NUM_ELEM * RSTR_ELEM_SIZE; … in OffsetNoTranspose() local 17 const CeedInt ind = indices[node]; in OffsetNoTranspose() 18 const CeedInt loc_node = node % RSTR_ELEM_SIZE; in OffsetNoTranspose() 19 const CeedInt elem = node / RSTR_ELEM_SIZE; in OffsetNoTranspose() 32 …for (CeedInt node = blockIdx.x * blockDim.x + threadIdx.x; node < RSTR_NUM_ELEM * RSTR_ELEM_SIZE; … in OffsetTranspose() local 33 const CeedInt ind = indices[node]; in OffsetTranspose() 34 const CeedInt loc_node = node % RSTR_ELEM_SIZE; in OffsetTranspose() 35 const CeedInt elem = node / RSTR_ELEM_SIZE; in OffsetTranspose()
|
| H A D | hip-shared-basis-read-write-templates.h | 31 const CeedInt node = data.t_id_x; in ReadElementStrided1d() local 32 const CeedInt ind = node * strides_node + elem * strides_elem; in ReadElementStrided1d() 47 const CeedInt node = data.t_id_x; in WriteElementStrided1d() local 48 const CeedInt ind = node * strides_node + elem * strides_elem; in WriteElementStrided1d() 60 const CeedInt node = data.t_id_x; in SumElementStrided1d() local 61 const CeedInt ind = node * strides_node + elem * strides_elem; in SumElementStrided1d() 80 const CeedInt node = data.t_id_x + data.t_id_y * P_1D; in ReadElementStrided2d() local 81 const CeedInt ind = node * strides_node + elem * strides_elem; in ReadElementStrided2d() 96 const CeedInt node = data.t_id_x + data.t_id_y * P_1D; in WriteElementStrided2d() local 97 const CeedInt ind = node * strides_node + elem * strides_elem; in WriteElementStrided2d() [all …]
|
| H A D | hip-ref-restriction-at-points.h | 18 …for (CeedInt node = blockIdx.x * blockDim.x + threadIdx.x; node < RSTR_NUM_ELEM * RSTR_ELEM_SIZE; … in AtPointsTranspose() local 19 const CeedInt ind = indices[node]; in AtPointsTranspose() 20 const CeedInt loc_node = node % RSTR_ELEM_SIZE; in AtPointsTranspose() 21 const CeedInt elem = node / RSTR_ELEM_SIZE; in AtPointsTranspose()
|
| H A D | hip-gen-templates.h | 76 const CeedInt node = data.t_id_x; in ReadLVecStandard1d() local 77 const CeedInt ind = indices[node + elem * P_1D]; in ReadLVecStandard1d() 89 const CeedInt node = data.t_id_x; in ReadLVecStrided1d() local 90 const CeedInt ind = node * STRIDES_NODE + elem * STRIDES_ELEM; in ReadLVecStrided1d() 103 const CeedInt node = data.t_id_x; in WriteLVecStandard1d() local 104 const CeedInt ind = indices[node + elem * P_1D]; in WriteLVecStandard1d() 165 const CeedInt node = data.t_id_x; in WriteLVecStrided1d() local 166 const CeedInt ind = node * STRIDES_NODE + elem * STRIDES_ELEM; in WriteLVecStrided1d() 197 const CeedInt node = data.t_id_x + data.t_id_y * P_1D; in ReadLVecStandard2d() local 198 const CeedInt ind = indices[node + elem * P_1D * P_1D]; in ReadLVecStandard2d() [all …]
|
| /libCEED/include/ceed/jit-source/cuda/ |
| H A D | cuda-ref-restriction-curl-oriented.h | 17 …for (CeedInt node = blockIdx.x * blockDim.x + threadIdx.x; node < RSTR_NUM_ELEM * RSTR_ELEM_SIZE; … in CurlOrientedNoTranspose() local 18 const CeedInt loc_node = node % RSTR_ELEM_SIZE; in CurlOrientedNoTranspose() 19 const CeedInt elem = node / RSTR_ELEM_SIZE; in CurlOrientedNoTranspose() 20 const CeedInt ind_dl = loc_node > 0 ? indices[node - 1] : 0; in CurlOrientedNoTranspose() 21 const CeedInt ind_d = indices[node]; in CurlOrientedNoTranspose() 22 const CeedInt ind_du = loc_node < (RSTR_ELEM_SIZE - 1) ? indices[node + 1] : 0; in CurlOrientedNoTranspose() 23 const CeedInt8 curl_orient_dl = curl_orients[3 * node + 0]; in CurlOrientedNoTranspose() 24 const CeedInt8 curl_orient_d = curl_orients[3 * node + 1]; in CurlOrientedNoTranspose() 25 const CeedInt8 curl_orient_du = curl_orients[3 * node + 2]; in CurlOrientedNoTranspose() 42 …for (CeedInt node = blockIdx.x * blockDim.x + threadIdx.x; node < RSTR_NUM_ELEM * RSTR_ELEM_SIZE; … in CurlOrientedUnsignedNoTranspose() local [all …]
|
| H A D | cuda-ref-restriction-strided.h | 16 …for (CeedInt node = blockIdx.x * blockDim.x + threadIdx.x; node < RSTR_NUM_ELEM * RSTR_ELEM_SIZE; … in StridedNoTranspose() local 17 const CeedInt loc_node = node % RSTR_ELEM_SIZE; in StridedNoTranspose() 18 const CeedInt elem = node / RSTR_ELEM_SIZE; in StridedNoTranspose() 31 …for (CeedInt node = blockIdx.x * blockDim.x + threadIdx.x; node < RSTR_NUM_ELEM * RSTR_ELEM_SIZE; … in StridedTranspose() local 32 const CeedInt loc_node = node % RSTR_ELEM_SIZE; in StridedTranspose() 33 const CeedInt elem = node / RSTR_ELEM_SIZE; in StridedTranspose()
|
| H A D | cuda-ref-restriction-oriented.h | 17 …for (CeedInt node = blockIdx.x * blockDim.x + threadIdx.x; node < RSTR_NUM_ELEM * RSTR_ELEM_SIZE; … in OrientedNoTranspose() local 18 const CeedInt ind = indices[node]; in OrientedNoTranspose() 19 const bool orient = orients[node]; in OrientedNoTranspose() 20 const CeedInt loc_node = node % RSTR_ELEM_SIZE; in OrientedNoTranspose() 21 const CeedInt elem = node / RSTR_ELEM_SIZE; in OrientedNoTranspose() 35 …for (CeedInt node = blockIdx.x * blockDim.x + threadIdx.x; node < RSTR_NUM_ELEM * RSTR_ELEM_SIZE; … in OrientedTranspose() local 36 const CeedInt ind = indices[node]; in OrientedTranspose() 37 const bool orient = orients[node]; in OrientedTranspose() 38 const CeedInt loc_node = node % RSTR_ELEM_SIZE; in OrientedTranspose() 39 const CeedInt elem = node / RSTR_ELEM_SIZE; in OrientedTranspose()
|
| H A D | cuda-ref-restriction-offset.h | 16 …for (CeedInt node = blockIdx.x * blockDim.x + threadIdx.x; node < RSTR_NUM_ELEM * RSTR_ELEM_SIZE; … in OffsetNoTranspose() local 17 const CeedInt ind = indices[node]; in OffsetNoTranspose() 18 const CeedInt loc_node = node % RSTR_ELEM_SIZE; in OffsetNoTranspose() 19 const CeedInt elem = node / RSTR_ELEM_SIZE; in OffsetNoTranspose() 32 …for (CeedInt node = blockIdx.x * blockDim.x + threadIdx.x; node < RSTR_NUM_ELEM * RSTR_ELEM_SIZE; … in OffsetTranspose() local 33 const CeedInt ind = indices[node]; in OffsetTranspose() 34 const CeedInt loc_node = node % RSTR_ELEM_SIZE; in OffsetTranspose() 35 const CeedInt elem = node / RSTR_ELEM_SIZE; in OffsetTranspose()
|
| H A D | cuda-shared-basis-read-write-templates.h | 31 const CeedInt node = data.t_id_x; in ReadElementStrided1d() local 32 const CeedInt ind = node * strides_node + elem * strides_elem; in ReadElementStrided1d() 47 const CeedInt node = data.t_id_x; in WriteElementStrided1d() local 48 const CeedInt ind = node * strides_node + elem * strides_elem; in WriteElementStrided1d() 60 const CeedInt node = data.t_id_x; in SumElementStrided1d() local 61 const CeedInt ind = node * strides_node + elem * strides_elem; in SumElementStrided1d() 80 const CeedInt node = data.t_id_x + data.t_id_y * P_1D; in ReadElementStrided2d() local 81 const CeedInt ind = node * strides_node + elem * strides_elem; in ReadElementStrided2d() 96 const CeedInt node = data.t_id_x + data.t_id_y * P_1D; in WriteElementStrided2d() local 97 const CeedInt ind = node * strides_node + elem * strides_elem; in WriteElementStrided2d() [all …]
|
| H A D | cuda-ref-restriction-at-points.h | 18 …for (CeedInt node = blockIdx.x * blockDim.x + threadIdx.x; node < RSTR_NUM_ELEM * RSTR_ELEM_SIZE; … in AtPointsTranspose() local 19 const CeedInt ind = indices[node]; in AtPointsTranspose() 20 const CeedInt loc_node = node % RSTR_ELEM_SIZE; in AtPointsTranspose() 21 const CeedInt elem = node / RSTR_ELEM_SIZE; in AtPointsTranspose()
|
| H A D | cuda-gen-templates.h | 76 const CeedInt node = data.t_id_x; in ReadLVecStandard1d() local 77 const CeedInt ind = indices[node + elem * P_1D]; in ReadLVecStandard1d() 90 const CeedInt node = data.t_id_x; in ReadLVecStrided1d() local 91 const CeedInt ind = node * STRIDES_NODE + elem * STRIDES_ELEM; in ReadLVecStrided1d() 104 const CeedInt node = data.t_id_x; in WriteLVecStandard1d() local 105 const CeedInt ind = indices[node + elem * P_1D]; in WriteLVecStandard1d() 166 const CeedInt node = data.t_id_x; in WriteLVecStrided1d() local 167 const CeedInt ind = node * STRIDES_NODE + elem * STRIDES_ELEM; in WriteLVecStrided1d() 198 const CeedInt node = data.t_id_x + data.t_id_y * P_1D; in ReadLVecStandard2d() local 199 const CeedInt ind = indices[node + elem * P_1D * P_1D]; in ReadLVecStandard2d() [all …]
|
| /libCEED/doc/img/tex/ |
| H A D | libCEEDBackends.tex | 20 \node at (0.8,6.1) {\large Application}; 27 node[pos=.5,align=center,color=black] {Ratel}; 36 node[pos=.5,align=center,color=black] {PETSc}; 45 node[pos=.5,align=center,color=black] {Nek5000}; 54 node[pos=.5,align=center,color=black] {MFEM}; 60 \node at (0.8,6.1) {\large Library}; 66 node[pos=.5,align=center,color=black] {libCEED}; 80 \node at (0.95,6.1) {\large Backends}; 88 node[pos=.5,align=center,color=black] {Pure C}; 97 node[pos=.5,align=center,color=black] {AVX}; [all …]
|
| /libCEED/doc/img/ |
| H A D | libCEEDBackends.tex | 22 \node at (1.0,6.1) {\large Application}; 30 node[pos=.5,align=center,color=black] {PETSc}; 39 node[pos=.5,align=center,color=black] {Ratel}; 48 node[pos=.5,align=center,color=black] {HONEE}; 57 node[pos=.5,align=center,color=black] {MFEM}; 63 \node at (0.8,6.1) {\large Library}; 69 node[pos=.5,align=center,color=black] {libCEED}; 83 \node at (0.95,6.1) {\large Backends}; 91 node[pos=.5,align=center,color=black] {Pure C}; 100 node[pos=.5,align=center,color=black] {AVX}; [all …]
|
| /libCEED/include/ceed/jit-source/sycl/ |
| H A D | sycl-shared-basis-read-write-templates.h | 35 const CeedInt node = item_id_x; in ReadElementStrided1d() local 36 const CeedInt ind = node * strides_node + elem * strides_elem; in ReadElementStrided1d() 53 const CeedInt node = item_id_x; in WriteElementStrided1d() local 54 const CeedInt ind = node * strides_node + elem * strides_elem; in WriteElementStrided1d() 76 const CeedInt node = item_id_x + item_id_y * P_1D; in ReadElementStrided2d() local 77 const CeedInt ind = node * strides_node + elem * strides_elem; in ReadElementStrided2d() 95 const CeedInt node = item_id_x + item_id_y * P_1D; in WriteElementStrided2d() local 96 const CeedInt ind = node * strides_node + elem * strides_elem; in WriteElementStrided2d() 119 const CeedInt node = item_id_x + item_id_y * P_1D + z * P_1D * P_1D; in ReadElementStrided3d() local 120 const CeedInt ind = node * strides_node + elem * strides_elem; in ReadElementStrided3d() [all …]
|
| H A D | sycl-gen-templates.h | 39 const CeedInt node = item_id_x; in readDofsOffset1d() local 40 const CeedInt ind = indices[node + elem * P_1D]; in readDofsOffset1d() 57 const CeedInt node = item_id_x; in readDofsStrided1d() local 58 const CeedInt ind = node * strides_node + elem * strides_elem; in readDofsStrided1d() 74 const CeedInt node = item_id_x; in writeDofsOffset1d() local 75 const CeedInt ind = indices[node + elem * P_1D]; in writeDofsOffset1d() 91 const CeedInt node = item_id_x; in writeDofsStrided1d() local 92 const CeedInt ind = node * strides_node + elem * strides_elem; in writeDofsStrided1d() 113 const CeedInt node = item_id_x + item_id_y * P_1D; in readDofsOffset2d() local 114 const CeedInt ind = indices[node + elem * P_1D * P_1D]; in readDofsOffset2d() [all …]
|
| /libCEED/backends/sycl-ref/ |
| H A D | ceed-sycl-restriction.sycl.cpp | 40 sycl_queue.parallel_for<CeedElemRestrSyclStridedNT>(kernel_range, e, [=](sycl::id<1> node) { in CeedElemRestrictionStridedNoTranspose_Sycl() argument 41 const CeedInt loc_node = node % elem_size; in CeedElemRestrictionStridedNoTranspose_Sycl() 42 const CeedInt elem = node / elem_size; in CeedElemRestrictionStridedNoTranspose_Sycl() 67 sycl_queue.parallel_for<CeedElemRestrSyclOffsetNT>(kernel_range, e, [=](sycl::id<1> node) { in CeedElemRestrictionOffsetNoTranspose_Sycl() argument 68 const CeedInt ind = indices[node]; in CeedElemRestrictionOffsetNoTranspose_Sycl() 69 const CeedInt loc_node = node % elem_size; in CeedElemRestrictionOffsetNoTranspose_Sycl() 70 const CeedInt elem = node / elem_size; in CeedElemRestrictionOffsetNoTranspose_Sycl() 96 sycl_queue.parallel_for<CeedElemRestrSyclStridedT>(kernel_range, e, [=](sycl::id<1> node) { in CeedElemRestrictionStridedTranspose_Sycl() argument 97 const CeedInt loc_node = node % elem_size; in CeedElemRestrictionStridedTranspose_Sycl() 98 const CeedInt elem = node / elem_size; in CeedElemRestrictionStridedTranspose_Sycl() [all …]
|
| /libCEED/julia/LibCEED.jl/src/ |
| H A D | ElemRestriction.jl | 57 - `ncomp`: Number of field components per interpolation node (1 for scalar fields) 58 - `compstride`: Stride between components for the same L-vector "node". Data for node $i$, 121 - `ncomp`: Number of field components per interpolation node (1 for scalar fields) 122 - `compstride`: Stride between components for the same L-vector "node". Data for node $i$, 189 - `ncomp`: Number of field components per interpolation node (1 for scalar fields) 190 - `compstride`: Stride between components for the same L-vector "node". Data for node $i$, 248 - `ncomp`: Number of field components per interpolation node (1 for scalar fields) 251 - `strides`: Array for strides between [nodes, components, elements]. Data for node $i$,
|
| /libCEED/benchmarks/ |
| H A D | README.md | 15 of processors per node. 30 bound of the problem sizes, per compute node; the default value is 3*2^20.
|
| H A D | benchmark.sh | 170 -p|--proc-node)
|
| /libCEED/examples/fluids/src/ |
| H A D | setupts.c | 86 for (PetscInt node = 0; node < dof / num_comp; node++) { in Surface_Forces_NS() local 88 reaction_force[w * dim + j] -= r[node].momentum[j]; in Surface_Forces_NS()
|
| /libCEED/interface/ |
| H A D | ceed-basis.c | 2343 for (CeedInt node = 0; node < basis->P; node++) { in CeedBasisGetInterp() local 2344 CeedInt p = (node / CeedIntPow(basis->P_1d, d)) % basis->P_1d; in CeedBasisGetInterp() 2347 basis->interp[qpt * (basis->P) + node] *= basis->interp_1d[q * basis->P_1d + p]; in CeedBasisGetInterp() 2397 for (CeedInt node = 0; node < basis->P; node++) { in CeedBasisGetGrad() local 2398 CeedInt p = (node / CeedIntPow(basis->P_1d, d)) % basis->P_1d; in CeedBasisGetGrad() 2401 …if (i == d) basis->grad[(i * basis->Q + qpt) * (basis->P) + node] *= basis->grad_1d[q * basis->P_1… in CeedBasisGetGrad() 2402 …else basis->grad[(i * basis->Q + qpt) * (basis->P) + node] *= basis->interp_1d[q * basis->P_1d + p… in CeedBasisGetGrad()
|
| /libCEED/doc/sphinx/source/ |
| H A D | libCEEDdev.md | 97 …Data for node `i`, component `j`, element `k` can be found in the L-vector at index `offsets[i + k… 99 …Data for node `i`, component `j`, element `k` can be found in the L-vector at index `i*strides[0] … 104 …If the backend uses a strided E-vector layout, then the data for node `i`, component `j`, element …
|