| /libCEED/benchmarks/ |
| H A D | postprocess_base.py | 32 data = data_default.copy()
38 data = data_default.copy()
39 data['num_procs'] = int(
47 data['num_procs_node'] = int(
56 data = data.copy()
57 runs.append(data)
58 data['file'] = fileinput.filename()
59 data['test'] = line.split()[-2] + " " + line.split('-- ')[1]
60 data['bp'] = data['test'].rsplit()[-1]
61 data['case'] = 'scalar' if (('Problem 1' in line) or ('Problem 3' in line)
[all …]
|
| /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
22 data.slice[data.t_id_x] = *U; in ContractX1d()
25 if (data.t_id_x < Q_1D) { in ContractX1d()
27 *V += B[i + data.t_id_x * P_1D] * data.slice[i]; // Contract x direction in ContractX1d()
36 inline __device__ void ContractTransposeX1d(SharedData_Cuda &data, const CeedScalar *U, const CeedS… in ContractTransposeX1d() argument
38 data.slice[data.t_id_x] = *U; in ContractTransposeX1d()
41 if (data.t_id_x < P_1D) { in ContractTransposeX1d()
43 *V += B[data.t_id_x + i * P_1D] * data.slice[i]; // Contract x direction in ContractTransposeX1d()
52 inline __device__ void Interp1d(SharedData_Cuda &data, const CeedScalar *__restrict__ r_U, const Ce… in Interp1d() argument
54 ContractX1d<NUM_COMP, P_1D, Q_1D>(data, &r_U[comp], c_B, &r_V[comp]); in Interp1d()
[all …]
|
| H A D | cuda-shared-basis-nontensor.h | 21 SharedData_Cuda data; in Interp() local
22 data.t_id_x = threadIdx.x; in Interp()
23 data.t_id_y = threadIdx.y; in Interp()
24 data.t_id_z = threadIdx.z; in Interp()
25 data.t_id = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x; in Interp()
26 data.slice = slice + data.t_id_z * BASIS_T_1D; in Interp()
33 LoadMatrix<BASIS_P, BASIS_Q>(data, c_B, s_B); in Interp()
38 …ReadElementStrided1d<BASIS_NUM_COMP, BASIS_P>(data, elem, 1, BASIS_P * num_elem, BASIS_P, d_U, r_U… in Interp()
39 InterpNonTensor<BASIS_NUM_COMP, BASIS_P, BASIS_Q, BASIS_T_1D>(data, r_U, s_B, r_V); in Interp()
40 …WriteElementStrided1d<BASIS_NUM_COMP, BASIS_Q>(data, elem, 1, BASIS_Q * num_elem, BASIS_Q, r_V, d_… in Interp()
[all …]
|
| H A D | cuda-shared-basis-tensor.h | 21 SharedData_Cuda data; in Interp() local
22 data.t_id_x = threadIdx.x; in Interp()
23 data.t_id_y = threadIdx.y; in Interp()
24 data.t_id_z = threadIdx.z; in Interp()
25 data.t_id = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x; in Interp()
26 data.slice = slice + data.t_id_z * BASIS_T_1D * (BASIS_DIM > 1 ? BASIS_T_1D : 1); in Interp()
33 LoadMatrix<BASIS_P_1D, BASIS_Q_1D>(data, c_B, s_B); in Interp()
39 …ReadElementStrided1d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * num_elem, BASIS_P_1D,… in Interp()
40 Interp1d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, s_B, r_V); in Interp()
41 …WriteElementStrided1d<BASIS_NUM_COMP, BASIS_Q_1D>(data, elem, 1, BASIS_Q_1D * num_elem, BASIS_Q_1D… in Interp()
[all …]
|
| H A D | cuda-shared-basis-tensor-at-points-templates.h | 44 inline __device__ void InterpAtPoints1d(SharedData_Cuda &data, const CeedInt p, const CeedScalar *_… in InterpAtPoints1d() argument
52 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()
65 inline __device__ void InterpTransposeAtPoints1d(SharedData_Cuda &data, const CeedInt p, const Ceed… in InterpTransposeAtPoints1d() argument
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()
90 inline __device__ void GradAtPoints1d(SharedData_Cuda &data, const CeedInt p, const CeedScalar *__r… in GradAtPoints1d() argument
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()
[all …]
|
| H A D | cuda-shared-basis-tensor-flattened-templates.h | 20 inline __device__ void ContractX2dFlattened(SharedData_Cuda &data, const int t_id_x, const int t_id… in ContractX2dFlattened() argument
23 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()
37 inline __device__ void ContractY2dFlattened(SharedData_Cuda &data, const int t_id_x, const int t_id… in ContractY2dFlattened() argument
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()
54 inline __device__ void ContractTransposeY2dFlattened(SharedData_Cuda &data, const int t_id_x, const… in ContractTransposeY2dFlattened() argument
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()
71 inline __device__ void ContractTransposeX2dFlattened(SharedData_Cuda &data, const int t_id_x, const… in ContractTransposeX2dFlattened() argument
[all …]
|
| H A D | cuda-gen-templates.h | 16 inline __device__ void LoadMatrix(SharedData_Cuda &data, const CeedScalar *__restrict__ d_B, CeedSc… in LoadMatrix() argument
17 for (CeedInt i = data.t_id; i < P * Q; i += blockDim.x * blockDim.y * blockDim.z) B[i] = d_B[i]; in LoadMatrix()
28 inline __device__ void ReadPoint(SharedData_Cuda &data, const CeedInt elem, const CeedInt p, const … in ReadPoint() argument
41 inline __device__ void WritePoint(SharedData_Cuda &data, const CeedInt elem, const CeedInt p, const… in WritePoint() argument
60 inline __device__ void SetEVecStandard1d_Single(SharedData_Cuda &data, const CeedInt n, const CeedS… in SetEVecStandard1d_Single() argument
64 if (data.t_id_x == target_node) { in SetEVecStandard1d_Single()
73 inline __device__ void ReadLVecStandard1d(SharedData_Cuda &data, const CeedInt num_nodes, const Cee… in ReadLVecStandard1d() argument
75 if (data.t_id_x < P_1D) { in ReadLVecStandard1d()
76 const CeedInt node = data.t_id_x; in ReadLVecStandard1d()
87 inline __device__ void ReadLVecStrided1d(SharedData_Cuda &data, const CeedInt elem, const CeedScala… in ReadLVecStrided1d() argument
[all …]
|
| H A D | cuda-shared-basis-tensor-at-points.h | 27 SharedData_Cuda data; in InterpAtPoints() local
28 data.t_id_x = threadIdx.x; in InterpAtPoints()
29 data.t_id_y = threadIdx.y; in InterpAtPoints()
30 data.t_id_z = threadIdx.z; in InterpAtPoints()
31 data.t_id = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x; in InterpAtPoints()
32 data.slice = slice + data.t_id_z * BASIS_T_1D * (BASIS_DIM > 1 ? BASIS_T_1D : 1); in InterpAtPoints()
41 LoadMatrix<BASIS_P_1D, BASIS_Q_1D>(data, c_B, s_B); in InterpAtPoints()
48 …ReadElementStrided1d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * num_elem, BASIS_P_1D,… in InterpAtPoints()
49 Interp1d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, s_B, r_C); in InterpAtPoints()
51 …ReadElementStrided2d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * BASIS_P_1D * num_elem… in InterpAtPoints()
[all …]
|
| H A D | cuda-shared-basis-read-write-templates.h | 16 inline __device__ void LoadMatrix(SharedData_Cuda &data, const CeedScalar *__restrict__ d_B, CeedSc… in LoadMatrix() argument
17 for (CeedInt i = data.t_id; i < P * Q; i += blockDim.x * blockDim.y * blockDim.z) B[i] = d_B[i]; in LoadMatrix()
28 inline __device__ void ReadElementStrided1d(SharedData_Cuda &data, const CeedInt elem, const CeedIn… in ReadElementStrided1d() argument
30 if (data.t_id_x < P_1D) { in ReadElementStrided1d()
31 const CeedInt node = data.t_id_x; in ReadElementStrided1d()
44 inline __device__ void WriteElementStrided1d(SharedData_Cuda &data, const CeedInt elem, const CeedI… in WriteElementStrided1d() argument
46 if (data.t_id_x < P_1D) { in WriteElementStrided1d()
47 const CeedInt node = data.t_id_x; in WriteElementStrided1d()
57 inline __device__ void SumElementStrided1d(SharedData_Cuda &data, const CeedInt elem, const CeedInt… in SumElementStrided1d() argument
59 if (data.t_id_x < P_1D) { in SumElementStrided1d()
[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
22 data.slice[data.t_id_x] = *U; in ContractX1d()
25 if (data.t_id_x < Q_1D) { in ContractX1d()
27 *V += B[i + data.t_id_x * P_1D] * data.slice[i]; // Contract x direction in ContractX1d()
36 inline __device__ void ContractTransposeX1d(SharedData_Hip &data, const CeedScalar *U, const CeedSc… in ContractTransposeX1d() argument
38 data.slice[data.t_id_x] = *U; in ContractTransposeX1d()
41 if (data.t_id_x < P_1D) { in ContractTransposeX1d()
43 *V += B[data.t_id_x + i * P_1D] * data.slice[i]; // Contract x direction in ContractTransposeX1d()
52 inline __device__ void Interp1d(SharedData_Hip &data, const CeedScalar *__restrict__ r_U, const Cee… in Interp1d() argument
54 ContractX1d<NUM_COMP, P_1D, Q_1D>(data, &r_U[comp], c_B, &r_V[comp]); in Interp1d()
[all …]
|
| H A D | hip-shared-basis-nontensor.h | 22 SharedData_Hip data; in __launch_bounds__() local
23 data.t_id_x = threadIdx.x; in __launch_bounds__()
24 data.t_id_y = threadIdx.y; in __launch_bounds__()
25 data.t_id_z = threadIdx.z; in __launch_bounds__()
26 data.t_id = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x; in __launch_bounds__()
27 data.slice = slice + data.t_id_z * BASIS_T_1D; in __launch_bounds__()
34 LoadMatrix<BASIS_P, BASIS_Q>(data, c_B, s_B); in __launch_bounds__()
39 …ReadElementStrided1d<BASIS_NUM_COMP, BASIS_P>(data, elem, 1, BASIS_P * num_elem, BASIS_P, d_U, r_U… in __launch_bounds__()
40 InterpNonTensor<BASIS_NUM_COMP, BASIS_P, BASIS_Q, BASIS_T_1D>(data, r_U, s_B, r_V); in __launch_bounds__()
41 …WriteElementStrided1d<BASIS_NUM_COMP, BASIS_Q>(data, elem, 1, BASIS_Q * num_elem, BASIS_Q, r_V, d_… in __launch_bounds__()
[all …]
|
| H A D | hip-shared-basis-tensor.h | 22 SharedData_Hip data; in __launch_bounds__() local
23 data.t_id_x = threadIdx.x; in __launch_bounds__()
24 data.t_id_y = threadIdx.y; in __launch_bounds__()
25 data.t_id_z = threadIdx.z; in __launch_bounds__()
26 data.t_id = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x; in __launch_bounds__()
27 data.slice = slice + data.t_id_z * BASIS_T_1D * (BASIS_DIM > 1 ? BASIS_T_1D : 1); in __launch_bounds__()
34 LoadMatrix<BASIS_P_1D, BASIS_Q_1D>(data, c_B, s_B); in __launch_bounds__()
40 …ReadElementStrided1d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * num_elem, BASIS_P_1D,… in __launch_bounds__()
41 Interp1d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, s_B, r_V); in __launch_bounds__()
42 …WriteElementStrided1d<BASIS_NUM_COMP, BASIS_Q_1D>(data, elem, 1, BASIS_Q_1D * num_elem, BASIS_Q_1D… in __launch_bounds__()
[all …]
|
| H A D | hip-shared-basis-tensor-at-points-templates.h | 44 inline __device__ void InterpAtPoints1d(SharedData_Hip &data, const CeedInt p, const CeedScalar *__… in InterpAtPoints1d() argument
53 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()
66 inline __device__ void InterpTransposeAtPoints1d(SharedData_Hip &data, const CeedInt p, const CeedS… in InterpTransposeAtPoints1d() argument
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()
91 inline __device__ void GradAtPoints1d(SharedData_Hip &data, const CeedInt p, const CeedScalar *__re… in GradAtPoints1d() argument
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()
[all …]
|
| H A D | hip-shared-basis-tensor-flattened-templates.h | 20 inline __device__ void ContractX2dFlattened(SharedData_Hip &data, const int t_id_x, const int t_id_… in ContractX2dFlattened() argument
23 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()
37 inline __device__ void ContractY2dFlattened(SharedData_Hip &data, const int t_id_x, const int t_id_… in ContractY2dFlattened() argument
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()
54 inline __device__ void ContractTransposeY2dFlattened(SharedData_Hip &data, const int t_id_x, const … in ContractTransposeY2dFlattened() argument
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()
71 inline __device__ void ContractTransposeX2dFlattened(SharedData_Hip &data, const int t_id_x, const … in ContractTransposeX2dFlattened() argument
[all …]
|
| H A D | hip-gen-templates.h | 16 inline __device__ void LoadMatrix(SharedData_Hip &data, const CeedScalar *__restrict__ d_B, CeedSca… in LoadMatrix() argument
17 for (CeedInt i = data.t_id; i < P * Q; i += blockDim.x * blockDim.y * blockDim.z) B[i] = d_B[i]; in LoadMatrix()
28 inline __device__ void ReadPoint(SharedData_Hip &data, const CeedInt elem, const CeedInt p, const C… in ReadPoint() argument
41 inline __device__ void WritePoint(SharedData_Hip &data, const CeedInt elem, const CeedInt p, const … in WritePoint() argument
60 inline __device__ void SetEVecStandard1d_Single(SharedData_Hip &data, const CeedInt n, const CeedSc… in SetEVecStandard1d_Single() argument
64 if (data.t_id_x == target_node) { in SetEVecStandard1d_Single()
73 inline __device__ void ReadLVecStandard1d(SharedData_Hip &data, const CeedInt num_nodes, const Ceed… in ReadLVecStandard1d() argument
75 if (data.t_id_x < P_1D) { in ReadLVecStandard1d()
76 const CeedInt node = data.t_id_x; in ReadLVecStandard1d()
87 inline __device__ void ReadLVecStrided1d(SharedData_Hip &data, const CeedInt elem, const CeedScalar… in ReadLVecStrided1d() argument
[all …]
|
| H A D | hip-shared-basis-tensor-at-points.h | 28 SharedData_Hip data; in __launch_bounds__() local
29 data.t_id_x = threadIdx.x; in __launch_bounds__()
30 data.t_id_y = threadIdx.y; in __launch_bounds__()
31 data.t_id_z = threadIdx.z; in __launch_bounds__()
32 data.t_id = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x; in __launch_bounds__()
33 data.slice = slice + data.t_id_z * BASIS_T_1D * (BASIS_DIM > 1 ? BASIS_T_1D : 1); in __launch_bounds__()
42 LoadMatrix<BASIS_P_1D, BASIS_Q_1D>(data, c_B, s_B); in __launch_bounds__()
49 …ReadElementStrided1d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * num_elem, BASIS_P_1D,… in __launch_bounds__()
50 Interp1d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, s_B, r_C); in __launch_bounds__()
52 …ReadElementStrided2d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * BASIS_P_1D * num_elem… in __launch_bounds__()
[all …]
|
| H A D | hip-shared-basis-read-write-templates.h | 16 inline __device__ void LoadMatrix(SharedData_Hip &data, const CeedScalar *__restrict__ d_B, CeedSca… in LoadMatrix() argument
17 for (CeedInt i = data.t_id; i < P * Q; i += blockDim.x * blockDim.y * blockDim.z) B[i] = d_B[i]; in LoadMatrix()
28 inline __device__ void ReadElementStrided1d(SharedData_Hip &data, const CeedInt elem, const CeedInt… in ReadElementStrided1d() argument
30 if (data.t_id_x < P_1D) { in ReadElementStrided1d()
31 const CeedInt node = data.t_id_x; in ReadElementStrided1d()
44 inline __device__ void WriteElementStrided1d(SharedData_Hip &data, const CeedInt elem, const CeedIn… in WriteElementStrided1d() argument
46 if (data.t_id_x < P_1D) { in WriteElementStrided1d()
47 const CeedInt node = data.t_id_x; in WriteElementStrided1d()
57 inline __device__ void SumElementStrided1d(SharedData_Hip &data, const CeedInt elem, const CeedInt … in SumElementStrided1d() argument
59 if (data.t_id_x < P_1D) { in SumElementStrided1d()
[all …]
|
| /libCEED/backends/hip-ref/ |
| H A D | ceed-hip-ref-basis.c | 29 CeedBasis_Hip *data; in CeedBasisApplyCore_Hip() local
32 CeedCallBackend(CeedBasisGetData(basis, &data)); in CeedBasisApplyCore_Hip()
51 …void *interp_args[] = {(void *)&num_elem, (void *)&is_transpose, &data->d_interp_1d, &d_u,… in CeedBasisApplyCore_Hip()
54 CeedCallBackend(CeedRunKernel_Hip(ceed, data->Interp, num_elem, block_size, interp_args)); in CeedBasisApplyCore_Hip()
57 … *grad_args[] = {(void *)&num_elem, (void *)&is_transpose, &data->d_interp_1d, &data->d_grad_1… in CeedBasisApplyCore_Hip()
60 CeedCallBackend(CeedRunKernel_Hip(ceed, data->Grad, num_elem, block_size, grad_args)); in CeedBasisApplyCore_Hip()
63 …CeedCheck(data->d_q_weight_1d, ceed, CEED_ERROR_BACKEND, "%s not supported; q_weights_1d not set",… in CeedBasisApplyCore_Hip()
64 void *weight_args[] = {(void *)&num_elem, (void *)&data->d_q_weight_1d, &d_v}; in CeedBasisApplyCore_Hip()
68 …CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Weight, num_elem, block_size_x, block_size_y, 1, … in CeedBasisApplyCore_Hip()
109 CeedBasis_Hip *data; in CeedBasisApplyAtPointsCore_Hip() local
[all …]
|
| H A D | ceed-hip-ref-qfunction.c | 25 CeedQFunction_Hip *data; in CeedQFunctionApply_Hip() local
32 CeedCallBackend(CeedQFunctionGetData(qf, &data)); in CeedQFunctionApply_Hip()
39 CeedCallBackend(CeedVectorGetArrayRead(U[i], CEED_MEM_DEVICE, &data->fields.inputs[i])); in CeedQFunctionApply_Hip()
42 CeedCallBackend(CeedVectorGetArrayWrite(V[i], CEED_MEM_DEVICE, &data->fields.outputs[i])); in CeedQFunctionApply_Hip()
46 CeedCallBackend(CeedQFunctionGetInnerContextData(qf, CEED_MEM_DEVICE, &data->d_c)); in CeedQFunctionApply_Hip()
49 void *args[] = {&data->d_c, (void *)&Q, &data->fields}; in CeedQFunctionApply_Hip()
51 …CeedCallBackend(CeedRunKernel_Hip(ceed, data->QFunction, CeedDivUpInt(Q, block_size), block_size, … in CeedQFunctionApply_Hip()
55 CeedCallBackend(CeedVectorRestoreArrayRead(U[i], &data->fields.inputs[i])); in CeedQFunctionApply_Hip()
58 CeedCallBackend(CeedVectorRestoreArray(V[i], &data->fields.outputs[i])); in CeedQFunctionApply_Hip()
62 CeedCallBackend(CeedQFunctionRestoreInnerContextData(qf, &data->d_c)); in CeedQFunctionApply_Hip()
[all …]
|
| /libCEED/backends/cuda-ref/ |
| H A D | ceed-cuda-ref-basis.c | 30 CeedBasis_Cuda *data; in CeedBasisApplyCore_Cuda() local
33 CeedCallBackend(CeedBasisGetData(basis, &data)); in CeedBasisApplyCore_Cuda()
51 …void *interp_args[] = {(void *)&num_elem, (void *)&is_transpose, &data->d_interp_1d, &d_u,… in CeedBasisApplyCore_Cuda()
54 CeedCallBackend(CeedRunKernel_Cuda(ceed, data->Interp, num_elem, block_size, interp_args)); in CeedBasisApplyCore_Cuda()
57 … *grad_args[] = {(void *)&num_elem, (void *)&is_transpose, &data->d_interp_1d, &data->d_grad_1… in CeedBasisApplyCore_Cuda()
60 CeedCallBackend(CeedRunKernel_Cuda(ceed, data->Grad, num_elem, block_size, grad_args)); in CeedBasisApplyCore_Cuda()
63 …CeedCheck(data->d_q_weight_1d, ceed, CEED_ERROR_BACKEND, "%s not supported; q_weights_1d not set",… in CeedBasisApplyCore_Cuda()
64 void *weight_args[] = {(void *)&num_elem, (void *)&data->d_q_weight_1d, &d_v}; in CeedBasisApplyCore_Cuda()
68 …CeedCallBackend(CeedRunKernelDim_Cuda(ceed, data->Weight, num_elem, block_size_x, block_size_y, 1,… in CeedBasisApplyCore_Cuda()
110 CeedBasis_Cuda *data; in CeedBasisApplyAtPointsCore_Cuda() local
[all …]
|
| H A D | ceed-cuda-ref-qfunction.c | 25 CeedQFunction_Cuda *data; in CeedQFunctionApply_Cuda() local
32 CeedCallBackend(CeedQFunctionGetData(qf, &data)); in CeedQFunctionApply_Cuda()
38 CeedCallBackend(CeedVectorGetArrayRead(U[i], CEED_MEM_DEVICE, &data->fields.inputs[i])); in CeedQFunctionApply_Cuda()
41 CeedCallBackend(CeedVectorGetArrayWrite(V[i], CEED_MEM_DEVICE, &data->fields.outputs[i])); in CeedQFunctionApply_Cuda()
45 CeedCallBackend(CeedQFunctionGetInnerContextData(qf, CEED_MEM_DEVICE, &data->d_c)); in CeedQFunctionApply_Cuda()
48 void *args[] = {&data->d_c, (void *)&Q, &data->fields}; in CeedQFunctionApply_Cuda()
49 CeedCallBackend(CeedRunKernelAutoblockCuda(ceed, data->QFunction, Q, args)); in CeedQFunctionApply_Cuda()
53 CeedCallBackend(CeedVectorRestoreArrayRead(U[i], &data->fields.inputs[i])); in CeedQFunctionApply_Cuda()
56 CeedCallBackend(CeedVectorRestoreArray(V[i], &data->fields.outputs[i])); in CeedQFunctionApply_Cuda()
60 CeedCallBackend(CeedQFunctionRestoreInnerContextData(qf, &data->d_c)); in CeedQFunctionApply_Cuda()
[all …]
|
| /libCEED/backends/hip-shared/ |
| H A D | ceed-hip-shared-basis.c | 98 CeedBasis_Hip_shared *data; in CeedBasisApplyTensorCore_Hip_shared() local
102 CeedCallBackend(CeedBasisGetData(basis, &data)); in CeedBasisApplyTensorCore_Hip_shared()
119 CeedInt block_size = data->block_sizes[0]; in CeedBasisApplyTensorCore_Hip_shared()
121 …CeedCheck(data->d_interp_1d, ceed, CEED_ERROR_BACKEND, "%s not supported; interp_1d not set", Ceed… in CeedBasisApplyTensorCore_Hip_shared()
125 void *interp_args[] = {(void *)&num_elem, &data->d_interp_1d, &d_u, &d_v}; in CeedBasisApplyTensorCore_Hip_shared()
134 …CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, apply_add ? data->InterpTransposeAdd : data->Inte… in CeedBasisApplyTensorCore_Hip_shared()
137 …CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, data->Interp, NULL, grid, thread_1d, 1, elems_per… in CeedBasisApplyTensorCore_Hip_shared()
146 …CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, apply_add ? data->InterpTransposeAdd : data->Inte… in CeedBasisApplyTensorCore_Hip_shared()
149 …CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, data->Interp, NULL, grid, thread_1d, thread_1d, e… in CeedBasisApplyTensorCore_Hip_shared()
157 …CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, apply_add ? data->InterpTransposeAdd : data->Inte… in CeedBasisApplyTensorCore_Hip_shared()
[all …]
|
| /libCEED/examples/solids/src/ |
| H A D | setup-libceed.c | 37 PetscErrorCode CeedDataDestroy(CeedInt level, CeedData data) { in CeedDataDestroy() argument
41 CeedVectorDestroy(&data->x_ceed); in CeedDataDestroy()
42 CeedVectorDestroy(&data->y_ceed); in CeedDataDestroy()
43 CeedVectorDestroy(&data->geo_data); in CeedDataDestroy()
44 for (CeedInt i = 0; i < SOLIDS_MAX_NUMBER_FIELDS; i++) CeedVectorDestroy(&data->stored_fields[i]); in CeedDataDestroy()
45 CeedVectorDestroy(&data->geo_data_diagnostic); in CeedDataDestroy()
46 CeedVectorDestroy(&data->true_soln); in CeedDataDestroy()
48 CeedElemRestrictionDestroy(&data->elem_restr_x); in CeedDataDestroy()
49 CeedElemRestrictionDestroy(&data->elem_restr_u); in CeedDataDestroy()
50 CeedElemRestrictionDestroy(&data->elem_restr_geo_data_i); in CeedDataDestroy()
[all …]
|
| /libCEED/backends/cuda-shared/ |
| H A D | ceed-cuda-shared-basis.c | 31 CeedBasis_Cuda_shared *data; in CeedBasisApplyTensorCore_Cuda_shared() local
35 CeedCallBackend(CeedBasisGetData(basis, &data)); in CeedBasisApplyTensorCore_Cuda_shared()
53 …CeedCheck(data->d_interp_1d, ceed, CEED_ERROR_BACKEND, "%s not supported; interp_1d not set", Ceed… in CeedBasisApplyTensorCore_Cuda_shared()
58 void *interp_args[] = {(void *)&num_elem, &data->d_interp_1d, &d_u, &d_v}; in CeedBasisApplyTensorCore_Cuda_shared()
67 …eedCallBackend(CeedRunKernelDimShared_Cuda(ceed, apply_add ? data->InterpTransposeAdd : data->Inte… in CeedBasisApplyTensorCore_Cuda_shared()
70 …CeedCallBackend(CeedRunKernelDimShared_Cuda(ceed, data->Interp, NULL, grid, thread_1d, 1, elems_pe… in CeedBasisApplyTensorCore_Cuda_shared()
80 …eedCallBackend(CeedRunKernelDimShared_Cuda(ceed, apply_add ? data->InterpTransposeAdd : data->Inte… in CeedBasisApplyTensorCore_Cuda_shared()
83 …CeedCallBackend(CeedRunKernelDimShared_Cuda(ceed, data->Interp, NULL, grid, thread_1d, thread_1d, … in CeedBasisApplyTensorCore_Cuda_shared()
92 …eedCallBackend(CeedRunKernelDimShared_Cuda(ceed, apply_add ? data->InterpTransposeAdd : data->Inte… in CeedBasisApplyTensorCore_Cuda_shared()
95 …CeedCallBackend(CeedRunKernelDimShared_Cuda(ceed, data->Interp, NULL, grid, thread_1d, thread_1d, … in CeedBasisApplyTensorCore_Cuda_shared()
[all …]
|
| /libCEED/backends/hip-gen/ |
| H A D | ceed-hip-gen-operator.c | 62 CeedOperator_Hip_gen *data; in CeedOperatorApplyAddCore_Hip_gen() local
69 CeedCallBackend(CeedOperatorGetData(op, &data)); in CeedOperatorApplyAddCore_Hip_gen()
80 data->fields.inputs[i] = NULL; in CeedOperatorApplyAddCore_Hip_gen()
88 if (is_active) data->fields.inputs[i] = input_arr; in CeedOperatorApplyAddCore_Hip_gen()
89 else CeedCallBackend(CeedVectorGetArrayRead(vec, CEED_MEM_DEVICE, &data->fields.inputs[i])); in CeedOperatorApplyAddCore_Hip_gen()
98 data->fields.outputs[i] = NULL; in CeedOperatorApplyAddCore_Hip_gen()
106 if (is_active) data->fields.outputs[i] = output_arr; in CeedOperatorApplyAddCore_Hip_gen()
107 else CeedCallBackend(CeedVectorGetArray(vec, CEED_MEM_DEVICE, &data->fields.outputs[i])); in CeedOperatorApplyAddCore_Hip_gen()
119 CeedCallBackend(CeedVectorGetArrayRead(vec, CEED_MEM_DEVICE, &data->points.coords)); in CeedOperatorApplyAddCore_Hip_gen()
123 if (num_elem != data->points.num_elem) { in CeedOperatorApplyAddCore_Hip_gen()
[all …]
|