// apply the Discontinuous Galerkin approximation for computing
// the time derivative of the field
void dtField(Field* f){
MacroCell mcell[f->macromesh.nbelems];
for(int ie=0;ie<f->macromesh.nbelems;ie++){
mcell[ie].field=f;
mcell[ie].first_cell=ie;
mcell[ie].last_cell_p1=ie+1;
}
#ifdef _WITH_PTHREAD
// we will have only one flying thread per
// macrocell
pthread_t tmcell[f->macromesh.nbelems];
int status;
// launch a thread for each macro cell
// computation of the inter subcell fluxes
for(int ie=0;ie<f->macromesh.nbelems;ie++){
status=pthread_create (&(tmcell[ie]), // thread
NULL, // default attributes
DGMacroCellInterface, // called function
(void*) (mcell+ie)); // function params
assert(status==0);
//DGMacroCellInterface((void*) (mcell+ie));
}
// wait the end of the threads before next step
for (int ie=0;ie<f->macromesh.nbelems;ie++){
pthread_join(tmcell[ie], NULL);
}
for(int ie=0;ie<f->macromesh.nbelems;ie++){
status=pthread_create (&(tmcell[ie]), // thread
NULL, // default attributes
DGSubCellInterface, // called function
(void*) (mcell+ie)); // function params
assert(status==0);
//DGSubCellInterface((void*) (mcell+ie));
}
// wait the end of the threads before next step
for (int ie=0;ie<f->macromesh.nbelems;ie++){
pthread_join(tmcell[ie], NULL);
}
for(int ie=0;ie<f->macromesh.nbelems;ie++){
status=pthread_create (&(tmcell[ie]), // thread
NULL, // default attributes
DGVolume, // called function
(void*) (mcell+ie)); // function params
assert(status==0);
//DGVolume((void*) (mcell+ie));
}
// wait the end of the threads before next step
for (int ie=0;ie<f->macromesh.nbelems;ie++){
pthread_join(tmcell[ie], NULL);
}
//DisplayField(f);
//assert(1==2);
for(int ie=0;ie<f->macromesh.nbelems;ie++){
status=pthread_create (&(tmcell[ie]), // thread
NULL, // default attributes
DGMass, // called function
(void*) (mcell+ie)); // function params
assert(status==0);
//DGMass((void*) (mcell+ie));
}
// wait the end of the threads before next step
for (int ie=0;ie<f->macromesh.nbelems;ie++){
pthread_join(tmcell[ie], NULL);
}
#else
#ifdef _OPENMP
#pragma omp parallel for schedule(dynamic, 1)
#endif
for(int ie=0; ie < f->macromesh.nbelems; ++ie) {
DGMacroCellInterface((void*) (mcell+ie));
DGSubCellInterface((void*) (mcell+ie));
DGVolume((void*) (mcell+ie));
DGMass((void*) (mcell+ie));
}
#endif
}
int TestfieldSubCellDGVol()
{
int test = true;
field f;
init_empty_field(&f);
f.model.cfl = 0.05;
f.model.m = 1; // only one conservative variable
f.model.NumFlux = TransNumFlux;
f.model.BoundaryFlux = TestTransBoundaryFlux;
f.model.InitData = TestTransInitData;
f.model.ImposedData = TestTransImposedData;
f.varindex = GenericVarindex;
f.interp.interp_param[0] = 1; // _M
f.interp.interp_param[1] = 2; // x direction degree
f.interp.interp_param[2] = 2; // y direction degree
f.interp.interp_param[3] = 2; // z direction degree
f.interp.interp_param[4] = 2; // x direction refinement
f.interp.interp_param[5] = 2; // y direction refinement
f.interp.interp_param[6] = 1; // z direction refinement
ReadMacroMesh(&f.macromesh, "../test/testcube.msh");
//ReadMacroMesh(&f.macromesh,"test/testdisque.msh");
BuildConnectivity(&f.macromesh);
PrintMacroMesh(&f.macromesh);
//AffineMapMacroMesh(&f.macromesh);
PrintMacroMesh(&f.macromesh);
Initfield(&f);
CheckMacroMesh(&f.macromesh, f.interp.interp_param + 1);
real tnow = 0.0;
for(int ie = 0;ie < f.macromesh.nbelems; ie++)
DGMacroCellInterfaceSlow((void*) (f.mcell+ie), &f, f.wn, f.dtwn);
for(int ie = 0; ie < f.macromesh.nbelems; ie++) {
DGSubCellInterface((void*) (f.mcell+ie), &f, f.wn, f.dtwn);
DGVolume((void*) (f.mcell+ie), &f, f.wn, f.dtwn);
DGMass((void*) (f.mcell+ie), &f, f.dtwn);
DGSource((void*) (f.mcell+ie), &f, tnow, f.wn, f.dtwn);
}
/* DGMacroCellInterfaceSlow(&f); */
/* DGSubCellInterface(&f); */
/* DGVolume(&f); */
/* DGMass(&f); */
Displayfield(&f);
/* Plotfield(0, false, &f, NULL, "visu.msh"); */
/* Plotfield(0, true, &f, "error", "error.msh"); */
// test the time derivative with the exact solution
int *raf = f.interp.interp_param + 4;
int *deg = f.interp.interp_param + 1;
for(int i=0;
i < f.model.m * f.macromesh.nbelems * NPG(raf, deg);
i++) {
test = test && fabs(4 * f.wn[i] - pow(f.dtwn[i] , 2)) < 1e-2;
assert(test);
}
return test;
}
int TestKernelVolume(void){
bool test=true;
if(!cldevice_is_acceptable(nplatform_cl, ndevice_cl)) {
printf("OpenCL device not acceptable.\n");
return true;
}
field f;
init_empty_field(&f);
f.model.cfl = 0.05;
f.model.m = 1; // only one conservative variable
f.model.NumFlux = TransNumFlux2d;
f.model.BoundaryFlux = TransBoundaryFlux2d;
f.model.InitData = TransInitData2d;
f.model.ImposedData = TransImposedData2d;
f.varindex = GenericVarindex;
f.interp.interp_param[0] = 1; // _M
f.interp.interp_param[1] = 1; // x direction degree
f.interp.interp_param[2] = 1; // y direction degree
f.interp.interp_param[3] = 0; // z direction degree
f.interp.interp_param[4] = 3; // x direction refinement
f.interp.interp_param[5] = 3; // y direction refinement
f.interp.interp_param[6] = 1; // z direction refinement
ReadMacroMesh(&f.macromesh,"../test/testmacromesh.msh");
//ReadMacroMesh(&f.macromesh,"test/testcube.msh");
Detect2DMacroMesh(&f.macromesh);
assert(f.macromesh.is2d);
BuildConnectivity(&f.macromesh);
//PrintMacroMesh(&f.macromesh);
//AffineMapMacroMesh(&f.macromesh);
Initfield(&f);
CopyfieldtoGPU(&f);
/* // set dtwn to 1 for testing */
/* void* chkptr; */
/* cl_int status; */
/* chkptr=clEnqueueMapBuffer(f.cli.commandqueue, */
/* f.dtwn_cl, // buffer to copy from */
/* CL_TRUE, // block until the buffer is available */
/* CL_MAP_WRITE, */
/* 0, // offset */
/* sizeof(real)*(f.wsize), // buffersize */
/* 0,NULL,NULL, // events management */
/* &status); */
/* assert(status == CL_SUCCESS); */
/* assert(chkptr == f.dtwn); */
/* for(int i=0;i<f.wsize;i++){ */
/* f.dtwn[i]=1; */
/* } */
/* status=clEnqueueUnmapMemObject (f.cli.commandqueue, */
/* f.dtwn_cl, */
/* f.dtwn, */
/* 0,NULL,NULL); */
/* assert(status == CL_SUCCESS); */
/* status=clFinish(f.cli.commandqueue); */
/* assert(status == CL_SUCCESS); */
clFinish(f.cli.commandqueue);
for(int ie = 0; ie < f.macromesh.nbelems; ++ie) {
DGVolume_CL(f.mcell + ie, &f, f.wn_cl + ie, 0, NULL, NULL);
clFinish(f.cli.commandqueue);
}
clFinish(f.cli.commandqueue);
CopyfieldtoCPU(&f);
//Displayfield(&f);
// save the dtwn pointer
real *dtwn_cl = f.dtwn;
// malloc a new dtwn.
f.dtwn = calloc(f.wsize, sizeof(real));
for(int ie = 0; ie < f.macromesh.nbelems; ++ie) {
MacroCell *mcell = f.mcell + ie;
real *dtwmc = f.dtwn + mcell->woffset;
real *wmc = f.wn + mcell->woffset;
DGVolume(f.mcell + ie, &f, wmc, dtwmc);
}
//Displayfield(&f);
//check that the results are the same
real maxerr = 0.0;
//printf("\nDifference\tC\t\tOpenCL\n");
for(int i = 0; i < f.wsize; ++i) {
real err = fabs(f.dtwn[i] - dtwn_cl[i]);
//printf("%f\t%f\t%f\n", err, f.dtwn[i], dtwn_cl[i]);
maxerr = fmax(err, maxerr);
}
printf("max error: %f\n",maxerr);
real tolerance;
if(sizeof(real) == sizeof(double))
tolerance = 1e-8;
else
tolerance = 1e-4;
test = (maxerr < tolerance);
return test;
}