Skip to content

Commit 3843268

Browse files
sebastiangrimbergsebastiangrimberg
authored
Merge pull request #1265 from CEED/sjg/hcurl-restr-dev
CeedElemRestriction for H(curl)
2 parents c2bc9a8 + c16dd8e commit 3843268

51 files changed

Lines changed: 3104 additions & 538 deletions

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

‎AUTHORS‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -6,6 +6,7 @@ Jean-Sylvain Camier
66
Veselin Dobrev
77
Yohann Dudouit
88
Leila Ghaffari
9+
Sebastian Grimberg
910
Tzanio Kolev
1011
David Medina
1112
Will Pazner

‎backends/blocked/ceed-blocked-operator.c‎

Lines changed: 41 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -41,29 +41,54 @@ static int CeedOperatorSetupFields_Blocked(CeedQFunction qf, CeedOperator op, bo
4141
CeedCallBackend(CeedQFunctionFieldGetEvalMode(qf_fields[i], &eval_mode));
4242

4343
if (eval_mode != CEED_EVAL_WEIGHT) {
44+
Ceed ceed_rstr;
4445
CeedCallBackend(CeedOperatorFieldGetElemRestriction(op_fields[i], &r));
45-
CeedCallBackend(CeedElemRestrictionGetCeed(r, &ceed));
46+
CeedCallBackend(CeedElemRestrictionGetCeed(r, &ceed_rstr));
4647
CeedSize l_size;
4748
CeedInt num_elem, elem_size, comp_stride;
4849
CeedCallBackend(CeedElemRestrictionGetNumElements(r, &num_elem));
4950
CeedCallBackend(CeedElemRestrictionGetElementSize(r, &elem_size));
5051
CeedCallBackend(CeedElemRestrictionGetLVectorSize(r, &l_size));
5152
CeedCallBackend(CeedElemRestrictionGetNumComponents(r, &num_comp));
52-
53-
bool strided;
54-
CeedCallBackend(CeedElemRestrictionIsStrided(r, &strided));
55-
if (strided) {
56-
CeedInt strides[3];
57-
CeedCallBackend(CeedElemRestrictionGetStrides(r, &strides));
58-
CeedCallBackend(
59-
CeedElemRestrictionCreateBlockedStrided(ceed, num_elem, elem_size, blk_size, num_comp, l_size, strides, &blk_restr[i + start_e]));
60-
} else {
61-
const CeedInt *offsets = NULL;
62-
CeedCallBackend(CeedElemRestrictionGetOffsets(r, CEED_MEM_HOST, &offsets));
63-
CeedCallBackend(CeedElemRestrictionGetCompStride(r, &comp_stride));
64-
CeedCallBackend(CeedElemRestrictionCreateBlocked(ceed, num_elem, elem_size, blk_size, num_comp, comp_stride, l_size, CEED_MEM_HOST,
65-
CEED_COPY_VALUES, offsets, &blk_restr[i + start_e]));
66-
CeedCallBackend(CeedElemRestrictionRestoreOffsets(r, &offsets));
53+
CeedCallBackend(CeedElemRestrictionGetCompStride(r, &comp_stride));
54+
55+
CeedRestrictionType rstr_type;
56+
CeedCallBackend(CeedElemRestrictionGetType(r, &rstr_type));
57+
switch (rstr_type) {
58+
case CEED_RESTRICTION_STANDARD: {
59+
const CeedInt *offsets = NULL;
60+
CeedCallBackend(CeedElemRestrictionGetOffsets(r, CEED_MEM_HOST, &offsets));
61+
CeedCallBackend(CeedElemRestrictionCreateBlocked(ceed_rstr, num_elem, elem_size, blk_size, num_comp, comp_stride, l_size, CEED_MEM_HOST,
62+
CEED_COPY_VALUES, offsets, &blk_restr[i + start_e]));
63+
CeedCallBackend(CeedElemRestrictionRestoreOffsets(r, &offsets));
64+
} break;
65+
case CEED_RESTRICTION_ORIENTED: {
66+
const CeedInt *offsets = NULL;
67+
const bool *orients = NULL;
68+
CeedCallBackend(CeedElemRestrictionGetOffsets(r, CEED_MEM_HOST, &offsets));
69+
CeedCallBackend(CeedElemRestrictionGetOrientations(r, CEED_MEM_HOST, &orients));
70+
CeedCallBackend(CeedElemRestrictionCreateBlockedOriented(ceed_rstr, num_elem, elem_size, blk_size, num_comp, comp_stride, l_size,
71+
CEED_MEM_HOST, CEED_COPY_VALUES, offsets, orients, &blk_restr[i + start_e]));
72+
CeedCallBackend(CeedElemRestrictionRestoreOffsets(r, &offsets));
73+
CeedCallBackend(CeedElemRestrictionRestoreOrientations(r, &orients));
74+
} break;
75+
case CEED_RESTRICTION_CURL_ORIENTED: {
76+
const CeedInt *offsets = NULL;
77+
const CeedInt8 *curl_orients = NULL;
78+
CeedCallBackend(CeedElemRestrictionGetOffsets(r, CEED_MEM_HOST, &offsets));
79+
CeedCallBackend(CeedElemRestrictionGetCurlOrientations(r, CEED_MEM_HOST, &curl_orients));
80+
CeedCallBackend(CeedElemRestrictionCreateBlockedCurlOriented(ceed_rstr, num_elem, elem_size, blk_size, num_comp, comp_stride, l_size,
81+
CEED_MEM_HOST, CEED_COPY_VALUES, offsets, curl_orients,
82+
&blk_restr[i + start_e]));
83+
CeedCallBackend(CeedElemRestrictionRestoreOffsets(r, &offsets));
84+
CeedCallBackend(CeedElemRestrictionRestoreCurlOrientations(r, &curl_orients));
85+
} break;
86+
case CEED_RESTRICTION_STRIDED: {
87+
CeedInt strides[3];
88+
CeedCallBackend(CeedElemRestrictionGetStrides(r, &strides));
89+
CeedCallBackend(
90+
CeedElemRestrictionCreateBlockedStrided(ceed_rstr, num_elem, elem_size, blk_size, num_comp, l_size, strides, &blk_restr[i + start_e]));
91+
} break;
6792
}
6893
CeedCallBackend(CeedElemRestrictionCreateVector(blk_restr[i + start_e], NULL, &e_vecs_full[i + start_e]));
6994
}

‎backends/cuda-ref/ceed-cuda-ref-restriction.c‎

Lines changed: 8 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -209,7 +209,8 @@ static int CeedElemRestrictionOffset_Cuda(const CeedElemRestriction r, const Cee
209209
//------------------------------------------------------------------------------
210210
// Create restriction
211211
//------------------------------------------------------------------------------
212-
int CeedElemRestrictionCreate_Cuda(CeedMemType mem_type, CeedCopyMode copy_mode, const CeedInt *indices, CeedElemRestriction r) {
212+
int CeedElemRestrictionCreate_Cuda(CeedMemType mem_type, CeedCopyMode copy_mode, const CeedInt *indices, const bool *orients,
213+
const CeedInt8 *curl_orients, CeedElemRestriction r) {
213214
Ceed ceed;
214215
CeedCallBackend(CeedElemRestrictionGetCeed(r, &ceed));
215216
CeedElemRestriction_Cuda *impl;
@@ -222,6 +223,11 @@ int CeedElemRestrictionCreate_Cuda(CeedMemType mem_type, CeedCopyMode copy_mode,
222223
CeedInt strides[3] = {1, size, elem_size};
223224
CeedInt comp_stride = 1;
224225

226+
CeedRestrictionType rstr_type;
227+
CeedCallBackend(CeedElemRestrictionGetType(r, &rstr_type));
228+
CeedCheck(rstr_type != CEED_RESTRICTION_ORIENTED && rstr_type != CEED_RESTRICTION_CURL_ORIENTED, ceed, CEED_ERROR_BACKEND,
229+
"Backend does not implement CeedElemRestrictionCreateOriented or CeedElemRestrictionCreateCurlOriented");
230+
225231
// Stride data
226232
bool is_strided;
227233
CeedCallBackend(CeedElemRestrictionIsStrided(r, &is_strided));
@@ -323,6 +329,7 @@ int CeedElemRestrictionCreate_Cuda(CeedMemType mem_type, CeedCopyMode copy_mode,
323329
// Register backend functions
324330
CeedCallBackend(CeedSetBackendFunction(ceed, "ElemRestriction", r, "Apply", CeedElemRestrictionApply_Cuda));
325331
CeedCallBackend(CeedSetBackendFunction(ceed, "ElemRestriction", r, "ApplyUnsigned", CeedElemRestrictionApply_Cuda));
332+
CeedCallBackend(CeedSetBackendFunction(ceed, "ElemRestriction", r, "ApplyUnoriented", CeedElemRestrictionApply_Cuda));
326333
CeedCallBackend(CeedSetBackendFunction(ceed, "ElemRestriction", r, "GetOffsets", CeedElemRestrictionGetOffsets_Cuda));
327334
CeedCallBackend(CeedSetBackendFunction(ceed, "ElemRestriction", r, "Destroy", CeedElemRestrictionDestroy_Cuda));
328335
return CEED_ERROR_SUCCESS;

‎backends/cuda-ref/ceed-cuda-ref.h‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -113,7 +113,8 @@ CEED_INTERN int CeedGetCublasHandle_Cuda(Ceed ceed, cublasHandle_t *handle);
113113

114114
CEED_INTERN int CeedVectorCreate_Cuda(CeedSize n, CeedVector vec);
115115

116-
CEED_INTERN int CeedElemRestrictionCreate_Cuda(CeedMemType mem_type, CeedCopyMode copy_mode, const CeedInt *indices, CeedElemRestriction r);
116+
CEED_INTERN int CeedElemRestrictionCreate_Cuda(CeedMemType mem_type, CeedCopyMode copy_mode, const CeedInt *indices, const bool *orients,
117+
const CeedInt8 *curl_orients, CeedElemRestriction r);
117118

118119
CEED_INTERN int CeedBasisCreateTensorH1_Cuda(CeedInt dim, CeedInt P_1d, CeedInt Q_1d, const CeedScalar *interp_1d, const CeedScalar *grad_1d,
119120
const CeedScalar *q_ref_1d, const CeedScalar *q_weight_1d, CeedBasis basis);

‎backends/hip-ref/ceed-hip-ref-restriction.c‎

Lines changed: 8 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -207,7 +207,8 @@ static int CeedElemRestrictionOffset_Hip(const CeedElemRestriction r, const Ceed
207207
//------------------------------------------------------------------------------
208208
// Create restriction
209209
//------------------------------------------------------------------------------
210-
int CeedElemRestrictionCreate_Hip(CeedMemType mem_type, CeedCopyMode copy_mode, const CeedInt *indices, CeedElemRestriction r) {
210+
int CeedElemRestrictionCreate_Hip(CeedMemType mem_type, CeedCopyMode copy_mode, const CeedInt *indices, const bool *orients,
211+
const CeedInt8 *curl_orients, CeedElemRestriction r) {
211212
Ceed ceed;
212213
CeedCallBackend(CeedElemRestrictionGetCeed(r, &ceed));
213214
CeedElemRestriction_Hip *impl;
@@ -220,6 +221,11 @@ int CeedElemRestrictionCreate_Hip(CeedMemType mem_type, CeedCopyMode copy_mode,
220221
CeedInt strides[3] = {1, size, elem_size};
221222
CeedInt comp_stride = 1;
222223

224+
CeedRestrictionType rstr_type;
225+
CeedCallBackend(CeedElemRestrictionGetType(r, &rstr_type));
226+
CeedCheck(rstr_type != CEED_RESTRICTION_ORIENTED && rstr_type != CEED_RESTRICTION_CURL_ORIENTED, ceed, CEED_ERROR_BACKEND,
227+
"Backend does not implement CeedElemRestrictionCreateOriented or CeedElemRestrictionCreateCurlOriented");
228+
223229
// Stride data
224230
bool is_strided;
225231
CeedCallBackend(CeedElemRestrictionIsStrided(r, &is_strided));
@@ -321,6 +327,7 @@ int CeedElemRestrictionCreate_Hip(CeedMemType mem_type, CeedCopyMode copy_mode,
321327
// Register backend functions
322328
CeedCallBackend(CeedSetBackendFunction(ceed, "ElemRestriction", r, "Apply", CeedElemRestrictionApply_Hip));
323329
CeedCallBackend(CeedSetBackendFunction(ceed, "ElemRestriction", r, "ApplyUnsigned", CeedElemRestrictionApply_Hip));
330+
CeedCallBackend(CeedSetBackendFunction(ceed, "ElemRestriction", r, "ApplyUnoriented", CeedElemRestrictionApply_Hip));
324331
CeedCallBackend(CeedSetBackendFunction(ceed, "ElemRestriction", r, "GetOffsets", CeedElemRestrictionGetOffsets_Hip));
325332
CeedCallBackend(CeedSetBackendFunction(ceed, "ElemRestriction", r, "Destroy", CeedElemRestrictionDestroy_Hip));
326333
return CEED_ERROR_SUCCESS;

‎backends/hip-ref/ceed-hip-ref.h‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -118,7 +118,8 @@ CEED_INTERN int CeedGetHipblasHandle_Hip(Ceed ceed, hipblasHandle_t *handle);
118118

119119
CEED_INTERN int CeedVectorCreate_Hip(CeedSize n, CeedVector vec);
120120

121-
CEED_INTERN int CeedElemRestrictionCreate_Hip(CeedMemType mem_type, CeedCopyMode copy_mode, const CeedInt *indices, CeedElemRestriction r);
121+
CEED_INTERN int CeedElemRestrictionCreate_Hip(CeedMemType mem_type, CeedCopyMode copy_mode, const CeedInt *indices, const bool *orients,
122+
const CeedInt8 *curl_orients, CeedElemRestriction r);
122123

123124
CEED_INTERN int CeedBasisCreateTensorH1_Hip(CeedInt dim, CeedInt P_1d, CeedInt Q_1d, const CeedScalar *interp_1d, const CeedScalar *grad_1d,
124125
const CeedScalar *q_ref_1d, const CeedScalar *q_weight_1d, CeedBasis basis);

‎backends/magma/ceed-magma-restriction.c‎

Lines changed: 8 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -151,7 +151,8 @@ static int CeedElemRestrictionDestroy_Magma(CeedElemRestriction r) {
151151
return CEED_ERROR_SUCCESS;
152152
}
153153

154-
int CeedElemRestrictionCreate_Magma(CeedMemType mtype, CeedCopyMode cmode, const CeedInt *offsets, CeedElemRestriction r) {
154+
int CeedElemRestrictionCreate_Magma(CeedMemType mtype, CeedCopyMode cmode, const CeedInt *offsets, const bool *orients, const CeedInt8 *curl_orients,
155+
CeedElemRestriction r) {
155156
Ceed ceed;
156157
CeedCallBackend(CeedElemRestrictionGetCeed(r, &ceed));
157158

@@ -163,6 +164,11 @@ int CeedElemRestrictionCreate_Magma(CeedMemType mtype, CeedCopyMode cmode, const
163164
CeedCallBackend(CeedElemRestrictionGetElementSize(r, &elemsize));
164165
CeedInt size = elemsize * nelem;
165166

167+
CeedRestrictionType rstr_type;
168+
CeedCallBackend(CeedElemRestrictionGetType(r, &rstr_type));
169+
CeedCheck(rstr_type != CEED_RESTRICTION_ORIENTED && rstr_type != CEED_RESTRICTION_CURL_ORIENTED, ceed, CEED_ERROR_BACKEND,
170+
"Backend does not implement CeedElemRestrictionCreateOriented or CeedElemRestrictionCreateCurlOriented");
171+
166172
CeedElemRestriction_Magma *impl;
167173
CeedCallBackend(CeedCalloc(1, &impl));
168174

@@ -261,6 +267,7 @@ int CeedElemRestrictionCreate_Magma(CeedMemType mtype, CeedCopyMode cmode, const
261267
CeedCallBackend(CeedElemRestrictionSetELayout(r, layout));
262268
CeedCallBackend(CeedSetBackendFunction(ceed, "ElemRestriction", r, "Apply", CeedElemRestrictionApply_Magma));
263269
CeedCallBackend(CeedSetBackendFunction(ceed, "ElemRestriction", r, "ApplyUnsigned", CeedElemRestrictionApply_Magma));
270+
CeedCallBackend(CeedSetBackendFunction(ceed, "ElemRestriction", r, "ApplyUnoriented", CeedElemRestrictionApply_Magma));
264271
CeedCallBackend(CeedSetBackendFunction(ceed, "ElemRestriction", r, "GetOffsets", CeedElemRestrictionGetOffsets_Magma));
265272
CeedCallBackend(CeedSetBackendFunction(ceed, "ElemRestriction", r, "Destroy", CeedElemRestrictionDestroy_Magma));
266273
CeedCallBackend(CeedFree(&restriction_kernel_path));

‎backends/magma/ceed-magma.h‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -134,7 +134,8 @@ CEED_INTERN int CeedBasisCreateTensorH1_Magma(CeedInt dim, CeedInt P1d, CeedInt
134134
CEED_INTERN int CeedBasisCreateH1_Magma(CeedElemTopology topo, CeedInt dim, CeedInt ndof, CeedInt nqpts, const CeedScalar *interp,
135135
const CeedScalar *grad, const CeedScalar *qref, const CeedScalar *qweight, CeedBasis basis);
136136

137-
CEED_INTERN int CeedElemRestrictionCreate_Magma(CeedMemType mtype, CeedCopyMode cmode, const CeedInt *offsets, CeedElemRestriction r);
137+
CEED_INTERN int CeedElemRestrictionCreate_Magma(CeedMemType mtype, CeedCopyMode cmode, const CeedInt *indices, const bool *orients,
138+
const CeedInt8 *curl_orients, CeedElemRestriction r);
138139

139140
// comment the line below to use the default magma_is_devptr function
140141
#define magma_is_devptr magma_isdevptr

‎backends/occa/ceed-occa-elem-restriction.cpp‎

Lines changed: 9 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -297,14 +297,21 @@ int ElemRestriction::registerCeedFunction(Ceed ceed, CeedElemRestriction r, cons
297297
return CeedSetBackendFunction(ceed, "ElemRestriction", r, fname, f);
298298
}
299299

300-
int ElemRestriction::ceedCreate(CeedMemType memType, CeedCopyMode copyMode, const CeedInt *indicesInput, CeedElemRestriction r) {
300+
int ElemRestriction::ceedCreate(CeedMemType memType, CeedCopyMode copyMode, const CeedInt *indicesInput, const bool *orientsInput,
301+
const CeedInt8 *curlOrientsInput, CeedElemRestriction r) {
301302
Ceed ceed;
302303
CeedCallBackend(CeedElemRestrictionGetCeed(r, &ceed));
303304

304305
if ((memType != CEED_MEM_DEVICE) && (memType != CEED_MEM_HOST)) {
305306
return staticCeedError("Only HOST and DEVICE CeedMemType supported");
306307
}
307308

309+
CeedRestrictionType rstr_type;
310+
CeedCallBackend(CeedElemRestrictionGetType(r, &rstr_type));
311+
if ((rstr_type == CEED_RESTRICTION_ORIENTED) || (rstr_type == CEED_RESTRICTION_CURL_ORIENTED)) {
312+
return staticCeedError("(OCCA) Backend does not implement CeedElemRestrictionCreateOriented or CeedElemRestrictionCreateCurlOriented");
313+
}
314+
308315
ElemRestriction *elemRestriction = new ElemRestriction();
309316
CeedCallBackend(CeedElemRestrictionSetData(r, elemRestriction));
310317

@@ -317,17 +324,14 @@ int ElemRestriction::ceedCreate(CeedMemType memType, CeedCopyMode copyMode, cons
317324

318325
CeedOccaRegisterFunction(r, "Apply", ElemRestriction::ceedApply);
319326
CeedOccaRegisterFunction(r, "ApplyUnsigned", ElemRestriction::ceedApply);
327+
CeedOccaRegisterFunction(r, "ApplyUnoriented", ElemRestriction::ceedApply);
320328
CeedOccaRegisterFunction(r, "ApplyBlock", ElemRestriction::ceedApplyBlock);
321329
CeedOccaRegisterFunction(r, "GetOffsets", ElemRestriction::ceedGetOffsets);
322330
CeedOccaRegisterFunction(r, "Destroy", ElemRestriction::ceedDestroy);
323331

324332
return CEED_ERROR_SUCCESS;
325333
}
326334

327-
int ElemRestriction::ceedCreateBlocked(CeedMemType memType, CeedCopyMode copyMode, const CeedInt *indicesInput, CeedElemRestriction r) {
328-
return staticCeedError("(OCCA) Backend does not implement CeedElemRestrictionCreateBlocked");
329-
}
330-
331335
int ElemRestriction::ceedApply(CeedElemRestriction r, CeedTransposeMode tmode, CeedVector u, CeedVector v, CeedRequest *request) {
332336
ElemRestriction *elemRestriction = ElemRestriction::from(r);
333337
Vector *uVector = Vector::from(u);

‎backends/occa/ceed-occa-elem-restriction.hpp‎

Lines changed: 2 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -77,9 +77,8 @@ class ElemRestriction : public CeedObject {
7777
//---[ Ceed Callbacks ]-----------
7878
static int registerCeedFunction(Ceed ceed, CeedElemRestriction r, const char *fname, ceed::occa::ceedFunction f);
7979

80-
static int ceedCreate(CeedMemType memType, CeedCopyMode copyMode, const CeedInt *indicesInput, CeedElemRestriction r);
81-
82-
static int ceedCreateBlocked(CeedMemType memType, CeedCopyMode copyMode, const CeedInt *indicesInput, CeedElemRestriction r);
80+
static int ceedCreate(CeedMemType memType, CeedCopyMode copyMode, const CeedInt *indicesInput, const bool *orientsInput,
81+
const CeedInt8 *curlOrientsInput, CeedElemRestriction r);
8382

8483
static int ceedApply(CeedElemRestriction r, CeedTransposeMode tmode, CeedVector u, CeedVector v, CeedRequest *request);
8584

0 commit comments

Comments
 (0)