int TestKernelFlux() { 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] = f.model.m; f.interp.interp_param[1] = 2; // x direction degree f.interp.interp_param[2] = 2; // 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); Initfield(&f); CopyfieldtoGPU(&f); clFinish(f.cli.commandqueue); for(int ie = 0; ie < f.macromesh.nbelems; ++ie) { MacroCell *mcell = f.mcell + ie; DGFlux_CL(mcell, &f, 0, f.wn_cl + ie, 0, NULL, NULL); clFinish(f.cli.commandqueue); DGFlux_CL(mcell, &f, 1, f.wn_cl + ie, 0, NULL, NULL); clFinish(f.cli.commandqueue); if(!f.macromesh.is2d) { DGFlux_CL(mcell, &f, 2, f.wn_cl + ie, 0, NULL, NULL); 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 *wmc = f.wn + mcell->woffset; real *dtwmc = f.dtwn + mcell->woffset; DGSubCellInterface(f.mcell + ie, &f, wmc, dtwmc); //DGVolume((void*) &f.mcell[ie], &f, f.wn, f.dtwn); } //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) { printf("%f\t%f\t%f\n", f.dtwn[i] - dtwn_cl[i], f.dtwn[i], dtwn_cl[i]); maxerr = fmax(fabs(f.dtwn[i] - dtwn_cl[i]), 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; }
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; }
// 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 }