Skip to content

Commit f27810e

Browse files
committed
HYD_system: bug fixes in
-bilan gpu; - mixed gpu-cpu FP
1 parent a87030a commit f27810e

File tree

9 files changed

+142
-33
lines changed

9 files changed

+142
-33
lines changed

source_code/system_hydraulic/source_code/models/_Hyd_Model.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -265,8 +265,8 @@ void _Hyd_Model::init_solver_gpu(Hyd_Param_Global* global_params) {
265265
pManager->setSelectedDevice(scheme_info.selected_device); // Set GPU device to Use. Important: Has to be called after setExecutor. Default is the faster one.
266266
pManager->setSimulationLength(simulationLength); // Set Simulation Length
267267
pManager->setOutputFrequency(outputFrequency); // Set Output Frequency
268-
pManager->setFloatPrecision(model::floatPrecision::kDouble); // Set Precision
269-
268+
pManager->setFloatPrecision(model::floatPrecision::kDouble); // Set Precision
269+
//pManager->setFloatPrecision(model::floatPrecision::kSingle);
270270
//Create the domain
271271
CDomainCartesian* ourCartesianDomain = pManager->getDomain();
272272
ourCartesianDomain->setCellResolution(*myFloodplain->Param_FP.get_ptr_width_x(), *myFloodplain->Param_FP.get_ptr_width_y());

source_code/system_hydraulic/source_code/models/coupling/Hyd_Coupling_Management.cpp

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1421,7 +1421,9 @@ double Hyd_Coupling_Management::get_min_timestep_fp2fp(void) {
14211421

14221422

14231423
for (int i = 0; i < this->number_fp2fp; i++) {
1424-
min_t = min(this->coupling_fp2fp[i].list.get_min_timestep(), min_t);
1424+
if (this->coupling_fp2fp[i].list.get_min_timestep()>0) {
1425+
min_t = min(this->coupling_fp2fp[i].list.get_min_timestep(), min_t);
1426+
}
14251427
}
14261428

14271429

source_code/system_hydraulic/source_code/models/floodplain/elements/Hyd_Element_Floodplain_Type_Standard.cpp

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -254,15 +254,15 @@ void Hyd_Element_Floodplain_Type_Standard::calculate_v_out(const double sin_valu
254254
this->v_x_out = (this->v_x + this->element_x_minus->element_type->get_flowvelocity_vx())*0.5;
255255
}
256256
else {
257-
this->v_x_out = (this->v_x + 0.0)*0.5;
257+
this->v_x_out = (this->v_x)*0.5;
258258
}
259259

260260
//y-direction
261261
if (this->element_y_minus != NULL) {
262262
this->v_y_out = (this->v_y + this->element_y_minus->element_type->get_flowvelocity_vy())*0.5;
263263
}
264264
else {
265-
this->v_y_out = (this->v_y + 0.0)*0.5;
265+
this->v_y_out = (this->v_y)*0.5;
266266
}
267267

268268
//turning
@@ -1358,6 +1358,7 @@ void Hyd_Element_Floodplain_Type_Standard::calculate_wet_dry_duration(const doub
13581358
if(this->was_wet_flag==false && this->h_value>wet_boundary){
13591359
this->was_wet_flag=true;
13601360
}
1361+
13611362
if(this->wet_flag==false && this->h_value>wet_boundary){
13621363
this->wet_flag=true;
13631364
this->time_wet_start=time_point;
@@ -1369,6 +1370,7 @@ void Hyd_Element_Floodplain_Type_Standard::calculate_wet_dry_duration(const doub
13691370
this->wet_duration=this->wet_duration+(time_point-this->time_wet_start);
13701371
this->time_wet_start=time_point;
13711372
}
1373+
13721374
//set the duration
13731375
if(this->wet_flag==true && this->h_value<wet_boundary){
13741376
this->wet_flag=false;

source_code/system_hydraulic/source_code/solvergpu/source_code/CSchemeGodunov.cpp

Lines changed: 78 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -59,6 +59,7 @@ CSchemeGodunov::CSchemeGodunov(void)
5959
this->ulReductionGlobalSize = 0;
6060
this->ulReductionWorkgroupSize = 0;
6161
this->dAvgTimestep = 0.0;
62+
this->sum_deficit_bilan = 0.0;
6263

6364
this->dCurrentTime = 0.0;
6465
this->dCurrentTimestepMovAvg = 0.0;
@@ -572,10 +573,14 @@ void CSchemeGodunov::prepare1OMemory(){
572573
this->oclBufferCellManning = new COCLBuffer("Manning coefficients", oclModel, true, true);
573574
if (this->bUseOptimizedBoundary == false) {
574575
this->oclBufferCellBoundary = new COCLBuffer("Boundary Values", oclModel, false, true);
576+
this->oclBufferBilanDef = new COCLBuffer("BilanDef", oclModel, false, true);
577+
this->sum_deficit_bilan = 0.0;
575578
}
576579
else {
577580
this->oclBufferCouplingIDs = new COCLBuffer("Coupling IDs", oclModel, true, true);
578581
this->oclBufferCouplingValues = new COCLBuffer("Coupling Values", oclModel, false, true);
582+
this->oclBufferBilanDef = new COCLBuffer("BilanDef", oclModel, false, true);
583+
this->sum_deficit_bilan = 0.0;
579584
}
580585
this->oclBufferUsePoleni = new COCLBuffer("Poleni Booleans", oclModel, true, true);
581586
this->oclBuffer_opt_zxmax = new COCLBuffer("opt_zxmax Values", oclModel, true, true);
@@ -584,15 +589,19 @@ void CSchemeGodunov::prepare1OMemory(){
584589
this->oclBuffer_opt_cy = new COCLBuffer("opt_cy Values", oclModel, true, true);
585590
this->oclBufferCellBed = new COCLBuffer("Bed elevations", oclModel, true, true);
586591

592+
593+
587594
this->oclBufferCellStates->setPointer(pCellStates, ucFloatSize * 4 * pDomain->getCellCount());
588595
this->oclBufferCellStatesAlt->setPointer(pCellStates, ucFloatSize * 4 * pDomain->getCellCount());
589596
this->oclBufferCellManning->setPointer(pManningValues, ucFloatSize * pDomain->getCellCount());
590597
if (this->bUseOptimizedBoundary == false) {
591598
this->oclBufferCellBoundary->setPointer(pBoundaryValues, ucFloatSize * pDomain->getCellCount());
599+
this->oclBufferBilanDef->setPointer(pBoundaryValues, ucFloatSize * pDomain->getCellCount());
592600
}
593601
else {
594602
this->oclBufferCouplingIDs->setPointer(pCouplingIDs, sizeof(cl_ulong) * this->ulCouplingArraySize);
595603
this->oclBufferCouplingValues->setPointer(pCouplingValues, ucFloatSize * this->ulCouplingArraySize);
604+
this->oclBufferBilanDef->setPointer(pCouplingValues, ucFloatSize * this->ulCouplingArraySize);
596605
}
597606
this->oclBufferUsePoleni->setPointer(pPoleniValues, sizeof(sUsePoleni) * pDomain->getCellCount());
598607
this->oclBuffer_opt_zxmax->setPointer(pOpt_zxmax, ucFloatSize * pDomain->getCellCount());
@@ -606,10 +615,12 @@ void CSchemeGodunov::prepare1OMemory(){
606615
this->oclBufferCellManning->createBuffer();
607616
if (this->bUseOptimizedBoundary == false) {
608617
this->oclBufferCellBoundary->createBuffer();
618+
this->oclBufferBilanDef->createBuffer();
609619
}
610620
else {
611621
this->oclBufferCouplingIDs->createBuffer();
612622
this->oclBufferCouplingValues->createBuffer();
623+
this->oclBufferBilanDef->createBuffer();
613624
}
614625
this->oclBufferUsePoleni->createBuffer();
615626
this->oclBuffer_opt_zxmax->createBuffer();
@@ -628,7 +639,7 @@ void CSchemeGodunov::prepare1OMemory(){
628639
this->oclBufferTimeTarget = new COCLBuffer("Target time (sync)", oclModel, false, true, ucFloatSize, true);
629640
this->oclBufferTimestepMovAvg = new COCLBuffer("Timestep Moving Average", oclModel, false, true, ucFloatSize, true);
630641
this->oclBufferTimeHydrological = new COCLBuffer("Time (hydrological)", oclModel, false, true, ucFloatSize, true);
631-
this->oclBufferBilanDef= new COCLBuffer("BilanDef", oclModel, false, true, ucFloatSize, true);
642+
632643

633644
// We duplicate the time and timestep variables if we're using single-precision so we have copies in both formats
634645
if (cModel->getFloatPrecision() == model::floatPrecision::kSingle)
@@ -638,23 +649,23 @@ void CSchemeGodunov::prepare1OMemory(){
638649
*(this->oclBufferTimestepMovAvg->getHostBlock<float*>()) = 0.0f;
639650
*(this->oclBufferTimeHydrological->getHostBlock<float*>()) = 0.0f;
640651
*(this->oclBufferTimeTarget->getHostBlock<float*>()) = 0.0f;
641-
*(this->oclBufferBilanDef->getHostBlock<float*>()) = 0.0f;
652+
642653
}
643654
else {
644655
*(this->oclBufferTime->getHostBlock<double*>()) = this->dCurrentTime;
645656
*(this->oclBufferTimestep->getHostBlock<double*>()) = this->dCurrentTimestep;
646657
*(this->oclBufferTimestepMovAvg->getHostBlock<double*>()) = 0.0;
647658
*(this->oclBufferTimeHydrological->getHostBlock<double*>()) = 0.0;
648659
*(this->oclBufferTimeTarget->getHostBlock<double*>()) = 0.0;
649-
*(this->oclBufferBilanDef->getHostBlock<double*>()) = 0.0;
660+
650661
}
651662
//create the buffers
652663
this->oclBufferTimestep->createBuffer();
653664
this->oclBufferTime->createBuffer();
654665
this->oclBufferTimestepMovAvg->createBuffer();
655666
this->oclBufferTimeHydrological->createBuffer();
656667
this->oclBufferTimeTarget->createBuffer();
657-
this->oclBufferBilanDef->createBuffer();
668+
658669

659670
// --
660671
// Timestep reduction global array
@@ -791,6 +802,7 @@ void CSchemeGodunov::release1OResources()
791802
if (this->oclBufferCellBoundary != NULL) delete oclBufferCellBoundary;
792803
if (this->oclBufferCouplingIDs != NULL) delete oclBufferCouplingIDs;
793804
if (this->oclBufferCouplingValues != NULL) delete oclBufferCouplingValues;
805+
if (this->oclBufferBilanDef != NULL) delete this->oclBufferBilanDef;
794806
if (this->oclBufferUsePoleni != NULL) delete oclBufferUsePoleni;
795807
if (this->oclBuffer_opt_zxmax != NULL) delete oclBuffer_opt_zxmax;
796808
if (this->oclBuffer_opt_cx != NULL) delete oclBuffer_opt_cx;
@@ -820,6 +832,7 @@ void CSchemeGodunov::release1OResources()
820832
oclBufferCellBoundary = NULL;
821833
oclBufferCouplingIDs = NULL;
822834
oclBufferCouplingValues = NULL;
835+
this->oclBufferBilanDef = NULL;
823836
oclBufferUsePoleni = NULL;
824837
oclBuffer_opt_zxmax = NULL;
825838
oclBuffer_opt_cx = NULL;
@@ -853,10 +866,12 @@ void CSchemeGodunov::prepareSimulation()
853866
oclBufferCellManning->queueWriteAll();
854867
if (this->bUseOptimizedBoundary == false) {
855868
oclBufferCellBoundary->queueWriteAll();
869+
this->oclBufferBilanDef->queueWriteAll();
856870
}
857871
else {
858872
oclBufferCouplingIDs->queueWriteAll();
859873
oclBufferCouplingValues->queueWriteAll();
874+
this->oclBufferBilanDef->queueWriteAll();
860875
}
861876
oclBufferUsePoleni->queueWriteAll();
862877
oclBuffer_opt_zxmax->queueWriteAll();
@@ -867,7 +882,7 @@ void CSchemeGodunov::prepareSimulation()
867882
oclBufferTimestep->queueWriteAll();
868883
oclBufferTimestepMovAvg->queueWriteAll();
869884
oclBufferTimeHydrological->queueWriteAll();
870-
this->oclBufferBilanDef->queueWriteAll();
885+
871886

872887
this->pDomain->getDevice()->blockUntilFinished();
873888

@@ -1036,7 +1051,7 @@ void CSchemeGodunov::Threaded_runBatch()
10361051
if (this->dCurrentTime < dTargetTime - 1e-8) {
10371052
oclKernelResetCounters->scheduleExecution();
10381053
for (unsigned int i = 0; i < uiQueueAmount; i++) {
1039-
1054+
///here the clacluation is started
10401055
this->scheduleIteration();
10411056
uiIterationsSinceTargetChanged++;
10421057
ulCurrentCellsCalculated += this->pDomain->getCellCount();
@@ -1052,7 +1067,56 @@ void CSchemeGodunov::Threaded_runBatch()
10521067
oclBufferBatchSkipped->queueReadAll();
10531068
oclBufferBatchSuccessful->queueReadAll();
10541069
oclBufferBatchTimesteps->queueReadAll();
1055-
this->oclBufferBilanDef->queueReadAll();
1070+
this->sum_deficit_bilan = 0.0;
1071+
if (this->bUseOptimizedBoundary == false) {
1072+
this->oclBufferBilanDef->queueReadAll();
1073+
for (int i = 0; i < pDomain->getCellCount();i++) {
1074+
if (cModel->getFloatPrecision() == model::floatPrecision::kSingle)
1075+
{
1076+
float* data = (this->oclBufferBilanDef->getHostBlock<float*>());
1077+
1078+
this->sum_deficit_bilan = this->sum_deficit_bilan+data[i];
1079+
1080+
}
1081+
else {
1082+
1083+
double* data = (this->oclBufferBilanDef->getHostBlock<double*>());
1084+
1085+
this->sum_deficit_bilan = this->sum_deficit_bilan+data[i];
1086+
1087+
}
1088+
1089+
1090+
1091+
}
1092+
1093+
1094+
1095+
}
1096+
else {
1097+
this->oclBufferBilanDef->queueReadAll();
1098+
for (int i = 0; i < this->ulCouplingArraySize; i++) {
1099+
if (cModel->getFloatPrecision() == model::floatPrecision::kSingle)
1100+
{
1101+
float* data = (this->oclBufferBilanDef->getHostBlock<float*>());
1102+
1103+
this->sum_deficit_bilan = this->sum_deficit_bilan+ data[i];
1104+
1105+
}
1106+
else {
1107+
1108+
double* data = (this->oclBufferBilanDef->getHostBlock<double*>());
1109+
1110+
this->sum_deficit_bilan = this->sum_deficit_bilan+ data[i];
1111+
1112+
}
1113+
1114+
1115+
}
1116+
}
1117+
1118+
1119+
10561120

10571121
this->pDomain->getDevice()->blockUntilFinished();
10581122

@@ -1082,7 +1146,7 @@ void CSchemeGodunov::Threaded_runBatch()
10821146
this->bThreadTerminated = true;
10831147
}
10841148

1085-
//Runs the actual simulation until completion or error
1149+
//Runs the actual simulation until completion or error here the simulation is set together via the kernels this connects the clc-methods as kernels
10861150
void CSchemeGodunov::scheduleIteration() {
10871151
//std::cout << "start iteration" << std::endl;
10881152
COCLBuffer* bufferSrc = NULL;
@@ -1106,6 +1170,7 @@ void CSchemeGodunov::scheduleIteration() {
11061170
oclKernelTimestepReduction->assignArgument(0, bufferDst);
11071171

11081172
profiless
1173+
//here the timesteps is claculated depending to the scheme, e.g. function pro_cacheDisabled in CLSchemePromaides.clc/.h
11091174
oclKernelFullTimestep->scheduleExecution();
11101175
profilese
11111176

@@ -1115,11 +1180,14 @@ void CSchemeGodunov::scheduleIteration() {
11151180
}
11161181

11171182
profilebs
1183+
//here the´bounadries are set e.g. function pro_cacheDisabled in CLSchemePromaides.clc/.h
11181184
oclKernelBoundary->scheduleExecution();
11191185
profilebe
11201186

11211187
profilers
1122-
if (this->bDynamicTimestep) { oclKernelTimestepReduction->scheduleExecution(); }
1188+
if (this->bDynamicTimestep) {
1189+
oclKernelTimestepReduction->scheduleExecution();
1190+
}
11231191
profilere
11241192

11251193
profilets
@@ -1220,15 +1288,7 @@ void CSchemeGodunov::setDebugger(unsigned int debugX, unsigned int debugY) {
12201288
//Function to get the bilan deficit
12211289
double CSchemeGodunov::get_bilan_deficit(void) {
12221290
if (this->oclBufferBilanDef != NULL) {
1223-
1224-
1225-
if (cModel->getFloatPrecision() == model::floatPrecision::kSingle)
1226-
{
1227-
return *(this->oclBufferBilanDef->getHostBlock<float*>());
1228-
}
1229-
else {
1230-
return *(this->oclBufferBilanDef->getHostBlock<double*>());
1231-
}
1291+
return this->sum_deficit_bilan;
12321292
}
12331293
else {
12341294
return 0.0;

source_code/system_hydraulic/source_code/solvergpu/source_code/CSchemeGodunov.h

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -101,6 +101,7 @@ class CSchemeGodunov : public CScheme
101101
unsigned int uiDebugCellX; // Debug info cell X
102102
unsigned int uiDebugCellY; // Debug info cell Y
103103
unsigned int uiTimestepReductionWavefronts; // Number of wavefronts used in reduction
104+
double sum_deficit_bilan; //Sum of Bilan deficitin
104105

105106
// Private functions
106107
virtual void prepareCode(); // Prepare the code required

source_code/system_hydraulic/source_code/solvergpu/source_code/opencl/CLBoundaries.clc

Lines changed: 17 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -41,6 +41,9 @@ __kernel void bdy_Promaides (
4141
if (dLclTimestep <= 0.0 || pCellData.y <= -9999.0 || pCellData.x == -9999.0){
4242
return;
4343
}
44+
if(dRate==0.0){
45+
return;
46+
}
4447

4548
//if (ulIdx == 7116 ){
4649
// //printf("Postion: X: %f Y:%f \n",lIdxX*1.0 , lIdxY*1.0 );
@@ -56,11 +59,11 @@ __kernel void bdy_Promaides (
5659
//}
5760

5861
//calculate the deficit
59-
if(pCellData.x + dRate * dLclTimestep<dCellBedElev){
60-
61-
(*pBilanDef)=(*pBilanDef)+((dRate-((dCellBedElev-pCellData.x)/dLclTimestep))*dLclTimestep)*-1.0;
62-
//printf("deficit %f %f %f\n",(*pBilanDef), dLclTimestep, dRate);
63-
62+
if(dRate<0.0){
63+
if(pCellData.x + dRate * dLclTimestep<dCellBedElev){
64+
pBilanDef[ulIdx]=pBilanDef[ulIdx]+((dRate-((dCellBedElev-pCellData.x)/dLclTimestep))*dLclTimestep)*-1.0;
65+
//printf("deficit %f %f %f %f\n",pBilanDef[ulIdx], dLclTimestep, dRate, pCellData.x );
66+
}
6467
}
6568

6669
// Apply the value...
@@ -107,12 +110,18 @@ __kernel void bdy_Promaides_by_id (
107110
if (dLclTimestep <= 0.0){
108111
return;
109112
}
113+
if(dRate==0.0){
114+
return;
115+
}
110116

111117
//calculate the deficit
112-
if(pCellData.x + dRate * dLclTimestep<dCellBedElev){
118+
if(dRate<0.0){
119+
if(pCellData.x + dRate * dLclTimestep<dCellBedElev){
120+
121+
pBilanDef[lId]=pBilanDef[ulIdx]+((dRate-((dCellBedElev-pCellData.x)/dLclTimestep))*dLclTimestep)*-1.0;
122+
//printf("deficit %f %f %f %f\n",pBilanDef[lId], dLclTimestep, dRate, pCellData.x );
113123

114-
(*pBilanDef)=(*pBilanDef)+((dRate-((dCellBedElev-pCellData.x)/dLclTimestep))*dLclTimestep)*-1.0;
115-
//printf("deficit %f %f %f\n",(*pBilanDef), dLclTimestep, dRate);
124+
}
116125

117126
}
118127

source_code/system_hydraulic/source_code/solvergpu/source_code/opencl/CLSchemePromaides.clc

Lines changed: 6 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -15,6 +15,8 @@
1515
/*
1616
* Calculate everything without using LDS caching
1717
*/
18+
//______________________
19+
//calculation of the overfow for the four neihgboring elements; it is i decide her if Manning (normal flow) or poleni (weir-flow over 1d-dikeline) is used
1820
__kernel REQD_WG_SIZE_FULL_TS
1921
void pro_cacheDisabled (
2022
__constant cl_double * dTimestep, // Timestep
@@ -315,7 +317,8 @@ void pro_cacheDisabled (
315317
pCellStateDst[ ulIdx ] = pCellData;
316318
}
317319

318-
320+
//_________________________________________________
321+
//calculation the overflow via Manning-equation (normal flow)
319322
cl_double2 manning_Solver(
320323
cl_double timestep, // Current Time-step
321324
cl_double manning, // Manning of Main Cell
@@ -416,6 +419,8 @@ cl_double2 manning_Solver(
416419

417420
}
418421

422+
//_________________________________________________
423+
//calculation the overflow between elements via Poleni weir equation (1d-dike line element)
419424
cl_double2 poleni_Solver(
420425
cl_double timestep, // Current Time-step
421426
cl_double manning, // Manning of Main Cell

0 commit comments

Comments
 (0)