Home
last modified time | relevance | path

Searched refs:slice (Results 1 – 24 of 24) sorted by relevance

/libCEED/include/ceed/jit-source/cuda/
H A Dcuda-shared-basis-nontensor.h19 extern __shared__ CeedScalar slice[]; in Interp()
26 data.slice = slice + data.t_id_z * BASIS_T_1D; in Interp()
46 extern __shared__ CeedScalar slice[]; in InterpTranspose()
53 data.slice = slice + data.t_id_z * BASIS_T_1D; in InterpTranspose()
73 extern __shared__ CeedScalar slice[]; in InterpTransposeAdd()
80 data.slice = slice + data.t_id_z * BASIS_T_1D; in InterpTransposeAdd()
102 extern __shared__ CeedScalar slice[]; in Grad()
109 data.slice = slice + data.t_id_z * BASIS_T_1D; in Grad()
129 extern __shared__ CeedScalar slice[]; in GradTranspose()
136 data.slice = slice + data.t_id_z * BASIS_T_1D; in GradTranspose()
[all …]
H A Dcuda-shared-basis-tensor-at-points-templates.h52 if (data.t_id_x < Q_1D) data.slice[data.t_id_x] = r_C[comp]; in InterpAtPoints1d()
56 r_V[comp] += chebyshev_x[i] * data.slice[i]; in InterpAtPoints1d()
72 if (data.t_id_x < Q_1D) data.slice[data.t_id_x] = 0.0; in InterpTransposeAtPoints1d()
77 …atomicAdd_block(&data.slice[comp * Q_1D + (i + data.t_id_x) % Q_1D], chebyshev_x[(i + data.t_id_x)… in InterpTransposeAtPoints1d()
82 if (data.t_id_x < Q_1D) r_C[comp] += data.slice[data.t_id_x]; in InterpTransposeAtPoints1d()
99 if (data.t_id_x < Q_1D) data.slice[data.t_id_x] = r_C[comp]; in GradAtPoints1d()
103 r_V[comp] += chebyshev_x[i] * data.slice[i]; in GradAtPoints1d()
119 if (data.t_id_x < Q_1D) data.slice[data.t_id_x] = 0.0; in GradTransposeAtPoints1d()
124 …atomicAdd_block(&data.slice[comp * Q_1D + (i + data.t_id_x) % Q_1D], chebyshev_x[(i + data.t_id_x)… in GradTransposeAtPoints1d()
129 if (data.t_id_x < Q_1D) r_C[comp] += data.slice[data.t_id_x]; in GradTransposeAtPoints1d()
[all …]
H A Dcuda-shared-basis-tensor.h19 extern __shared__ CeedScalar slice[]; in Interp()
26 data.slice = slice + data.t_id_z * BASIS_T_1D * (BASIS_DIM > 1 ? BASIS_T_1D : 1); in Interp()
58 extern __shared__ CeedScalar slice[]; in InterpCollocated()
65 data.slice = slice + data.t_id_z * BASIS_T_1D * (BASIS_DIM > 1 ? BASIS_T_1D : 1); in InterpCollocated()
88 extern __shared__ CeedScalar slice[]; in InterpTranspose()
95 data.slice = slice + data.t_id_z * BASIS_T_1D * (BASIS_DIM > 1 ? BASIS_T_1D : 1); in InterpTranspose()
127 extern __shared__ CeedScalar slice[]; in InterpCollocatedTranspose()
134 data.slice = slice + data.t_id_z * BASIS_T_1D * (BASIS_DIM > 1 ? BASIS_T_1D : 1); in InterpCollocatedTranspose()
157 extern __shared__ CeedScalar slice[]; in InterpTransposeAdd()
164 data.slice = slice + data.t_id_z * BASIS_T_1D * (BASIS_DIM > 1 ? BASIS_T_1D : 1); in InterpTransposeAdd()
[all …]
H A Dcuda-shared-basis-tensor-flattened-templates.h23 data.slice[t_id_x + t_id_y * T_1D] = *U; in ContractX2dFlattened()
28 *V += B[i + t_id_x * P_1D] * data.slice[i + t_id_y * T_1D]; // Contract x direction in ContractX2dFlattened()
40 data.slice[t_id_x + t_id_y * T_1D] = *U; in ContractY2dFlattened()
45 *V += B[i + t_id_y * P_1D] * data.slice[t_id_x + i * T_1D]; // Contract y direction in ContractY2dFlattened()
57 data.slice[t_id_x + t_id_y * T_1D] = *U; in ContractTransposeY2dFlattened()
62 *V += B[t_id_y + i * P_1D] * data.slice[t_id_x + i * T_1D]; // Contract y direction in ContractTransposeY2dFlattened()
74 data.slice[t_id_x + t_id_y * T_1D] = *U; in ContractTransposeX2dFlattened()
79 *V += B[t_id_x + i * P_1D] * data.slice[i + t_id_y * T_1D]; // Contract x direction in ContractTransposeX2dFlattened()
91 data.slice[t_id_x + t_id_y * T_1D] = *U; in ContractTransposeAddX2dFlattened()
95 *V += B[t_id_x + i * P_1D] * data.slice[i + t_id_y * T_1D]; // Contract x direction in ContractTransposeAddX2dFlattened()
[all …]
H A Dcuda-shared-basis-tensor-templates.h22 data.slice[data.t_id_x] = *U; in ContractX1d()
27 *V += B[i + data.t_id_x * P_1D] * data.slice[i]; // Contract x direction in ContractX1d()
38 data.slice[data.t_id_x] = *U; in ContractTransposeX1d()
43 *V += B[data.t_id_x + i * P_1D] * data.slice[i]; // Contract x direction in ContractTransposeX1d()
131 data.slice[data.t_id_x + data.t_id_y * T_1D] = *U; in ContractX2d()
136 *V += B[i + data.t_id_x * P_1D] * data.slice[i + data.t_id_y * T_1D]; // Contract x direction in ContractX2d()
147 data.slice[data.t_id_x + data.t_id_y * T_1D] = *U; in ContractY2d()
152 *V += B[i + data.t_id_y * P_1D] * data.slice[data.t_id_x + i * T_1D]; // Contract y direction in ContractY2d()
163 data.slice[data.t_id_x + data.t_id_y * T_1D] = *U; in ContractTransposeY2d()
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 Dcuda-shared-basis-tensor-at-points.h25 extern __shared__ CeedScalar slice[]; in InterpAtPoints()
32 data.slice = slice + data.t_id_z * BASIS_T_1D * (BASIS_DIM > 1 ? BASIS_T_1D : 1); in InterpAtPoints()
81 extern __shared__ CeedScalar slice[]; in InterpTransposeAtPoints()
88 data.slice = slice + data.t_id_z * BASIS_T_1D * (BASIS_DIM > 1 ? BASIS_T_1D : 1); in InterpTransposeAtPoints()
151 extern __shared__ CeedScalar slice[]; in InterpTransposeAddAtPoints()
158 data.slice = slice + data.t_id_z * BASIS_T_1D * (BASIS_DIM > 1 ? BASIS_T_1D : 1); in InterpTransposeAddAtPoints()
212 extern __shared__ CeedScalar slice[]; in GradAtPoints()
219 data.slice = slice + data.t_id_z * BASIS_T_1D * (BASIS_DIM > 1 ? BASIS_T_1D : 1); in GradAtPoints()
268 extern __shared__ CeedScalar slice[]; in GradTransposeAtPoints()
275 data.slice = slice + data.t_id_z * BASIS_T_1D * (BASIS_DIM > 1 ? BASIS_T_1D : 1); in GradTransposeAtPoints()
[all …]
H A Dcuda-shared-basis-nontensor-templates.h17 data.slice[data.t_id_x] = *U; in Contract1d()
22 *V += B[i + data.t_id_x * P_1D] * data.slice[i]; // Contract x direction in Contract1d()
33 data.slice[data.t_id_x] = *U; in ContractTranspose1d()
37 *V += B[data.t_id_x + i * P_1D] * data.slice[i]; // Contract x direction in ContractTranspose1d()
H A Dcuda-types.h38 CeedScalar *slice; member
H A Dcuda-gen-templates.h484 data.slice[data.t_id_x + data.t_id_y * T_1D] = r_U[q + comp * Q_1D]; in GradColloSlice3d()
489 … r_V[comp + 0 * NUM_COMP] += c_G[i + data.t_id_x * Q_1D] * data.slice[i + data.t_id_y * T_1D]; in GradColloSlice3d()
494 … r_V[comp + 1 * NUM_COMP] += c_G[i + data.t_id_y * Q_1D] * data.slice[data.t_id_x + i * T_1D]; in GradColloSlice3d()
514 data.slice[data.t_id_x + data.t_id_y * T_1D] = r_U[comp + 0 * NUM_COMP]; in GradColloSliceTranspose3d()
518 r_V[q + comp * Q_1D] += c_G[data.t_id_x + i * Q_1D] * data.slice[i + data.t_id_y * T_1D]; in GradColloSliceTranspose3d()
522 data.slice[data.t_id_x + data.t_id_y * T_1D] = r_U[comp + 1 * NUM_COMP]; in GradColloSliceTranspose3d()
525 r_V[q + comp * Q_1D] += c_G[data.t_id_y + i * Q_1D] * data.slice[data.t_id_x + i * T_1D]; in GradColloSliceTranspose3d()
/libCEED/include/ceed/jit-source/hip/
H A Dhip-shared-basis-nontensor.h20 extern __shared__ CeedScalar slice[]; in __launch_bounds__()
27 data.slice = slice + data.t_id_z * BASIS_T_1D; in __launch_bounds__()
47 extern __shared__ CeedScalar slice[]; in __launch_bounds__()
54 data.slice = slice + data.t_id_z * BASIS_T_1D; in __launch_bounds__()
74 extern __shared__ CeedScalar slice[]; in __launch_bounds__()
81 data.slice = slice + data.t_id_z * BASIS_T_1D; in __launch_bounds__()
104 extern __shared__ CeedScalar slice[]; in __launch_bounds__()
111 data.slice = slice + data.t_id_z * BASIS_T_1D; in __launch_bounds__()
131 extern __shared__ CeedScalar slice[]; in __launch_bounds__()
138 data.slice = slice + data.t_id_z * BASIS_T_1D; in __launch_bounds__()
[all …]
H A Dhip-shared-basis-tensor-at-points-templates.h53 if (data.t_id_x < Q_1D) data.slice[data.t_id_x] = r_C[comp]; in InterpAtPoints1d()
57 r_V[comp] += chebyshev_x[i] * data.slice[i]; in InterpAtPoints1d()
73 if (data.t_id_x < Q_1D) data.slice[data.t_id_x] = 0.0; in InterpTransposeAtPoints1d()
78 …atomicAdd(&data.slice[comp * Q_1D + (i + data.t_id_x) % Q_1D], chebyshev_x[(i + data.t_id_x) % Q_1… in InterpTransposeAtPoints1d()
83 if (data.t_id_x < Q_1D) r_C[comp] += data.slice[data.t_id_x]; in InterpTransposeAtPoints1d()
100 if (data.t_id_x < Q_1D) data.slice[data.t_id_x] = r_C[comp]; in GradAtPoints1d()
104 r_V[comp] += chebyshev_x[i] * data.slice[i]; in GradAtPoints1d()
120 if (data.t_id_x < Q_1D) data.slice[data.t_id_x] = 0.0; in GradTransposeAtPoints1d()
125 …atomicAdd(&data.slice[comp * Q_1D + (i + data.t_id_x) % Q_1D], chebyshev_x[(i + data.t_id_x) % Q_1… in GradTransposeAtPoints1d()
130 if (data.t_id_x < Q_1D) r_C[comp] += data.slice[data.t_id_x]; in GradTransposeAtPoints1d()
[all …]
H A Dhip-shared-basis-tensor.h20 extern __shared__ CeedScalar slice[]; in __launch_bounds__()
27 data.slice = slice + data.t_id_z * BASIS_T_1D * (BASIS_DIM > 1 ? BASIS_T_1D : 1); in __launch_bounds__()
59 extern __shared__ CeedScalar slice[]; in __launch_bounds__()
66 data.slice = slice + data.t_id_z * BASIS_T_1D * (BASIS_DIM > 1 ? BASIS_T_1D : 1); in __launch_bounds__()
89 extern __shared__ CeedScalar slice[]; in __launch_bounds__()
96 data.slice = slice + data.t_id_z * BASIS_T_1D * (BASIS_DIM > 1 ? BASIS_T_1D : 1); in __launch_bounds__()
128 extern __shared__ CeedScalar slice[]; in __launch_bounds__()
135 data.slice = slice + data.t_id_z * BASIS_T_1D * (BASIS_DIM > 1 ? BASIS_T_1D : 1); in __launch_bounds__()
158 extern __shared__ CeedScalar slice[]; in __launch_bounds__()
165 data.slice = slice + data.t_id_z * BASIS_T_1D * (BASIS_DIM > 1 ? BASIS_T_1D : 1); in __launch_bounds__()
[all …]
H A Dhip-shared-basis-tensor-flattened-templates.h23 data.slice[t_id_x + t_id_y * T_1D] = *U; in ContractX2dFlattened()
28 *V += B[i + t_id_x * P_1D] * data.slice[i + t_id_y * T_1D]; // Contract x direction in ContractX2dFlattened()
40 data.slice[t_id_x + t_id_y * T_1D] = *U; in ContractY2dFlattened()
45 *V += B[i + t_id_y * P_1D] * data.slice[t_id_x + i * T_1D]; // Contract y direction in ContractY2dFlattened()
57 data.slice[t_id_x + t_id_y * T_1D] = *U; in ContractTransposeY2dFlattened()
62 *V += B[t_id_y + i * P_1D] * data.slice[t_id_x + i * T_1D]; // Contract y direction in ContractTransposeY2dFlattened()
74 data.slice[t_id_x + t_id_y * T_1D] = *U; in ContractTransposeX2dFlattened()
79 *V += B[t_id_x + i * P_1D] * data.slice[i + t_id_y * T_1D]; // Contract x direction in ContractTransposeX2dFlattened()
91 data.slice[t_id_x + t_id_y * T_1D] = *U; in ContractTransposeAddX2dFlattened()
95 *V += B[t_id_x + i * P_1D] * data.slice[i + t_id_y * T_1D]; // Contract x direction in ContractTransposeAddX2dFlattened()
[all …]
H A Dhip-shared-basis-tensor-templates.h22 data.slice[data.t_id_x] = *U; in ContractX1d()
27 *V += B[i + data.t_id_x * P_1D] * data.slice[i]; // Contract x direction in ContractX1d()
38 data.slice[data.t_id_x] = *U; in ContractTransposeX1d()
43 *V += B[data.t_id_x + i * P_1D] * data.slice[i]; // Contract x direction in ContractTransposeX1d()
131 data.slice[data.t_id_x + data.t_id_y * T_1D] = *U; in ContractX2d()
136 *V += B[i + data.t_id_x * P_1D] * data.slice[i + data.t_id_y * T_1D]; // Contract x direction in ContractX2d()
147 data.slice[data.t_id_x + data.t_id_y * T_1D] = *U; in ContractY2d()
152 *V += B[i + data.t_id_y * P_1D] * data.slice[data.t_id_x + i * T_1D]; // Contract y direction in ContractY2d()
163 data.slice[data.t_id_x + data.t_id_y * T_1D] = *U; in ContractTransposeY2d()
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 Dhip-shared-basis-tensor-at-points.h26 extern __shared__ CeedScalar slice[]; in __launch_bounds__()
33 data.slice = slice + data.t_id_z * BASIS_T_1D * (BASIS_DIM > 1 ? BASIS_T_1D : 1); in __launch_bounds__()
82 extern __shared__ CeedScalar slice[]; in __launch_bounds__()
89 data.slice = slice + data.t_id_z * BASIS_T_1D * (BASIS_DIM > 1 ? BASIS_T_1D : 1); in __launch_bounds__()
152 extern __shared__ CeedScalar slice[]; in __launch_bounds__()
159 data.slice = slice + data.t_id_z * BASIS_T_1D * (BASIS_DIM > 1 ? BASIS_T_1D : 1); in __launch_bounds__()
214 extern __shared__ CeedScalar slice[]; in __launch_bounds__()
221 data.slice = slice + data.t_id_z * BASIS_T_1D * (BASIS_DIM > 1 ? BASIS_T_1D : 1); in __launch_bounds__()
270 extern __shared__ CeedScalar slice[]; in __launch_bounds__()
277 data.slice = slice + data.t_id_z * BASIS_T_1D * (BASIS_DIM > 1 ? BASIS_T_1D : 1); in __launch_bounds__()
[all …]
H A Dhip-shared-basis-nontensor-templates.h17 data.slice[data.t_id_x] = *U; in Contract1d()
22 *V += B[i + data.t_id_x * P_1D] * data.slice[i]; // Contract x direction in Contract1d()
33 data.slice[data.t_id_x] = *U; in ContractTranspose1d()
37 *V += B[data.t_id_x + i * P_1D] * data.slice[i]; // Contract x direction in ContractTranspose1d()
H A Dhip-types.h38 CeedScalar *slice; member
H A Dhip-gen-templates.h481 data.slice[data.t_id_x + data.t_id_y * T_1D] = r_U[q + comp * Q_1D]; in GradColloSlice3d()
486 … r_V[comp + 0 * NUM_COMP] += c_G[i + data.t_id_x * Q_1D] * data.slice[i + data.t_id_y * T_1D]; in GradColloSlice3d()
491 … r_V[comp + 1 * NUM_COMP] += c_G[i + data.t_id_y * Q_1D] * data.slice[data.t_id_x + i * T_1D]; in GradColloSlice3d()
512 data.slice[data.t_id_x + data.t_id_y * T_1D] = r_U[comp + 0 * NUM_COMP]; in GradColloSliceTranspose3d()
515 r_V[q + comp * Q_1D] += c_G[data.t_id_x + i * Q_1D] * data.slice[i + data.t_id_y * T_1D]; in GradColloSliceTranspose3d()
519 data.slice[data.t_id_x + data.t_id_y * T_1D] = r_U[comp + 1 * NUM_COMP]; in GradColloSliceTranspose3d()
522 r_V[q + comp * Q_1D] += c_G[data.t_id_y + i * Q_1D] * data.slice[data.t_id_x + i * T_1D]; in GradColloSliceTranspose3d()
/libCEED/examples/rust-qfunctions/ex1-volume-rs/src/
H A Dlib.rs51 let in_slice = core::slice::from_raw_parts(in_, 2); in build_mass_rs()
60 let w = core::slice::from_raw_parts(w_ptr, q as usize); in build_mass_rs()
62 let out_slice = core::slice::from_raw_parts_mut(out, 1); in build_mass_rs()
63 let q_data = core::slice::from_raw_parts_mut(out_slice[0], q as usize); in build_mass_rs()
106 let in_slice = core::slice::from_raw_parts(in_, 2); in apply_mass_rs()
111 let u = core::slice::from_raw_parts(u_ptr, q as usize); in apply_mass_rs()
112 let q_data = core::slice::from_raw_parts(q_data_ptr, q as usize); in apply_mass_rs()
114 let out_slice = core::slice::from_raw_parts_mut(out, 1); in apply_mass_rs()
117 let v = core::slice::from_raw_parts_mut(v_ptr, q as usize); in apply_mass_rs()
/libCEED/rust/libceed/src/
H A Dvector.rs153 slice: &'a mut [crate::Scalar], in from_vector_and_slice_mut()
155 assert_eq!(vec.length(), slice.len()); in from_vector_and_slice_mut()
165 slice.as_ptr() as *mut crate::Scalar, in from_vector_and_slice_mut()
170 _slice: slice, in from_vector_and_slice_mut()
450 pub fn set_slice(&mut self, slice: &[crate::Scalar]) -> crate::Result<i32> { in set_slice()
451 assert_eq!(self.length(), slice.len()); in set_slice()
461 slice.as_ptr() as *mut crate::Scalar, in set_slice()
513 slice: &'b mut [crate::Scalar], in wrap_slice_mut()
515 VectorSliceWrapper::from_vector_and_slice_mut(self, slice) in wrap_slice_mut()
824 unsafe { std::slice::from_raw_parts(self.array, self.vec.len()) } in deref()
[all …]
H A Dqfunction.rs485 std::slice::from_raw_parts(inputs_ptr as *const QFunctionField, num_inputs as usize) in inputs()
505 std::slice::from_raw_parts(outputs_ptr as *const QFunctionField, num_outputs as usize) in outputs()
536 std::slice::from_raw_parts(inputs, MAX_QFUNCTION_FIELDS); in trampoline()
543 std::slice::from_raw_parts(x, trampoline_data.input_sizes[i] * q as usize) in trampoline()
551 std::slice::from_raw_parts(outputs, MAX_QFUNCTION_FIELDS); in trampoline()
559 std::slice::from_raw_parts_mut(x, trampoline_data.output_sizes[i] * q as usize) in trampoline()
H A Doperator.rs874 std::slice::from_raw_parts( in inputs()
944 std::slice::from_raw_parts( in outputs()
H A Dlib.rs401 pub fn vector_from_slice<'a>(&self, slice: &[crate::Scalar]) -> Result<Vector<'a>> { in vector_from_slice()
402 Vector::from_slice(self, slice) in vector_from_slice()
/libCEED/julia/LibCEED.jl/src/
H A DUserQFunction.jl49 slice = Expr(:ref, arr_name_gen, idx, (:(:) for i = 1:ndims)...)
52 array_views[i] = :($arr_name = $slice)
55 array_views[i] = :($arr_name = LibCEED.SArray{$S}(@view $slice))
58 array_views[i] = :($arr_name = @view $slice)