@@ -10,7 +10,7 @@ void CUDA_CHECK(cudaError CUDerr)
10
10
write_text_to_log_file (" Cuda error in file " + std::string (__FILE__) + " in line " + std::to_string (__LINE__) + " " + std::string (cudaGetErrorString (CUDerr)));
11
11
12
12
fprintf (stderr, " Cuda error in file '%s' in line %i : %s.\n " , __FILE__, __LINE__, cudaGetErrorString (CUDerr));
13
-
13
+
14
14
write_text_to_log_file (" Cuda error in file " + std::string (__FILE__) + " in line " + std::to_string (__LINE__) + " " + std::string (cudaGetErrorString (CUDerr)));
15
15
exit (EXIT_FAILURE);
16
16
@@ -29,7 +29,7 @@ XBGPUParam waveinitGPU(XBGPUParam Param, std::vector<Wavebndparam> wavebnd)
29
29
ny = Param.ny ;
30
30
31
31
32
-
32
+
33
33
if (Param.dtheta > 0.0 )
34
34
{
35
35
Param.ntheta = round ((Param.thetamax - Param.thetamin ) / Param.dtheta );
@@ -43,7 +43,7 @@ XBGPUParam waveinitGPU(XBGPUParam Param, std::vector<Wavebndparam> wavebnd)
43
43
Param.dtheta = (Param.thetamax - Param.thetamin ) / Param.ntheta ;
44
44
}
45
45
46
-
46
+
47
47
ntheta = Param.ntheta ;
48
48
dtheta = Param.dtheta ;
49
49
@@ -60,7 +60,7 @@ XBGPUParam waveinitGPU(XBGPUParam Param, std::vector<Wavebndparam> wavebnd)
60
60
{
61
61
nwavbnd = ceil (Param.rtlength / Param.dtbc )+1 ; // +1 needed here
62
62
}
63
-
63
+
64
64
theta = (DECNUM *)malloc (ntheta*sizeof (DECNUM));
65
65
66
66
Stfile = (double *)malloc (ntheta*ny*nwavbnd*sizeof (double ));
@@ -122,7 +122,7 @@ XBGPUParam waveinitGPU(XBGPUParam Param, std::vector<Wavebndparam> wavebnd)
122
122
{
123
123
// readXbbndstep(nx, ny, ntheta, Param.wavebndfile.c_str(), 1, Trepold, qfile, Stfile);
124
124
readXbbndstep (Param, wavebnd, 0 , Trep, qfile, Stfile);
125
-
125
+
126
126
}
127
127
if (Param.wavebndtype == 3 )
128
128
{
@@ -327,7 +327,7 @@ XBGPUParam waveinitGPU(XBGPUParam Param, std::vector<Wavebndparam> wavebnd)
327
327
328
328
// Clac Stat
329
329
330
-
330
+
331
331
for (int i = 0 ; i < ntheta; i++) // ! Fill St
332
332
{
333
333
// St[i]=Stold[i];
@@ -375,7 +375,7 @@ XBGPUParam waveinitGPU(XBGPUParam Param, std::vector<Wavebndparam> wavebnd)
375
375
}
376
376
377
377
378
- // run dispersion relation
378
+ // run dispersion relation
379
379
380
380
return Param;
381
381
@@ -521,17 +521,17 @@ void wavebnd(XBGPUParam Param, std::vector<Wavebndparam> wavebndvec)
521
521
522
522
double difft = wavebndvec[WAVstepinbnd].time - totaltime;
523
523
524
-
524
+
525
525
if (difft < 0.0 )
526
526
{
527
527
WAVstepinbnd++;
528
528
529
-
529
+
530
530
if (Param.wavebndtype == 2 )
531
531
{
532
532
// Read new STfile and qfile XBeach style
533
533
readXbbndstep (Param, wavebndvec, WAVstepinbnd - 1 , Trep, qfile, Stfile);
534
-
534
+
535
535
}
536
536
537
537
if (Param.wavebndtype == 3 )
@@ -551,7 +551,7 @@ void wavebnd(XBGPUParam Param, std::vector<Wavebndparam> wavebndvec)
551
551
int nfHR, ndHR;
552
552
553
553
makjonswap (Param, wavebndvec, WAVstepinbnd - 1 , nfHR, ndHR, HRfreq, HRdir, HRSpec);
554
-
554
+
555
555
// Then generate wave group timeseries based on that spectra
556
556
// void GenWGnLBW(XBGPUParam Param, int nf, int ndir, double * HRfreq, double * HRdir, double * HRSpec, float Trep, double * qfile, double * Stfile)
557
557
GenWGnLBW (Param, nfHR, ndHR, HRfreq, HRdir, HRSpec, Trep, qfile, Stfile);
@@ -741,7 +741,7 @@ void wavebnd(XBGPUParam Param, std::vector<Wavebndparam> wavebndvec)
741
741
{
742
742
St[ni + i*ny] = interptime (Stnew[ni + i*ny], Stold[ni + i*ny], timenext, timesincelast);
743
743
}
744
-
744
+
745
745
}
746
746
747
747
if (Param.flow == 1 )
@@ -879,13 +879,15 @@ void wavestep(XBGPUParam Param)
879
879
CUDA_CHECK (cudaMalloc ((void **)&yadvec_g, nx*ny*ntheta*sizeof (DECNUM)));
880
880
CUDA_CHECK (cudaMalloc ((void **)&thetaadvec_g, nx*ny*ntheta*sizeof (DECNUM)));
881
881
882
- xadvecupwind2 << <gridDim , blockDim , 0 >> >(nx, ny, ntheta, Param.dtheta , Param.dx , dt, wci_g, ee_g, cg_g, cxsth_g, uu_g, xadvec_g);
882
+ // xadvecupwind2 << <gridDim, blockDim, 0 >> >(nx, ny, ntheta, Param.dtheta, Param.dx, dt, wci_g, ee_g, cg_g, cxsth_g, uu_g, xadvec_g);
883
+ xadvecupwind2SD << <gridDim , blockDim , 0 >> >(nx, ny, ntheta, Param.dtheta , Param.dx , dt, thetamean_g, wci_g, ee_g, cg_g, cxsth_g, uu_g, xadvec_g);
883
884
// CUT_CHECK_ERROR("eulerupwind xadvec execution failed\n");
884
885
CUDA_CHECK (cudaThreadSynchronize ());
885
886
886
887
887
888
888
- yadvecupwind2 << <gridDim , blockDim , 0 >> >(nx, ny, ntheta, Param.dtheta , Param.dx , dt, wci_g, ee_g, cg_g, sxnth_g, vv_g, yadvec_g);
889
+ // yadvecupwind2 << <gridDim, blockDim, 0 >> >(nx, ny, ntheta, Param.dtheta, Param.dx, dt, wci_g, ee_g, cg_g, sxnth_g, vv_g, yadvec_g);
890
+ yadvecupwind2SD << <gridDim , blockDim , 0 >> >(nx, ny, ntheta, Param.dtheta , Param.dx , dt, thetamean_g, wci_g, ee_g, cg_g, sxnth_g, vv_g, yadvec_g);
889
891
// CUT_CHECK_ERROR("eulerupwind yadvec execution failed\n");
890
892
CUDA_CHECK (cudaThreadSynchronize ());
891
893
@@ -978,7 +980,7 @@ void wavestep(XBGPUParam Param)
978
980
979
981
980
982
981
- //
983
+ //
982
984
// Total dissipation from breaking and bottom friction
983
985
//
984
986
@@ -1017,7 +1019,7 @@ void wavestep(XBGPUParam Param)
1017
1019
1018
1020
// thetaadvecuw<<<gridDim, blockDim, 0>>>(nx,ny,ntheta,dtheta,eect_g,thetaadvec_g);
1019
1021
// //CUT_CHECK_ERROR("eulerupwind thetaadvecuw execution failed\n");
1020
- // CUDA_CHECK( cudaThreadSynchronize() );
1022
+ // CUDA_CHECK( cudaThreadSynchronize() );
1021
1023
1022
1024
thetaadvecuw2ho << <gridDim , blockDim , 0 >> >(nx, ny, ntheta, Param.dtheta , Param.dx , dt, Param.wci , rr_g, ctheta_g, thetaadvec_g);
1023
1025
// CUT_CHECK_ERROR("eulerupwind thetaadvec execution failed\n");
@@ -1044,9 +1046,9 @@ void wavestep(XBGPUParam Param)
1044
1046
// CUDA_CHECK( cudaMemcpy(D_g, uu, nx*ny*sizeof(DECNUM ), cudaMemcpyHostToDevice) );
1045
1047
1046
1048
1047
- //
1049
+ //
1048
1050
// Distribution of dissipation over directions and frequencies
1049
- //
1051
+ //
1050
1052
dissipation << <gridDim , blockDim , 0 >> >(nx, ny, ntheta, Param.dtheta , Param.eps , dt, Param.g , Param.beta , wci_g, hh_g, ee_g, D_g, E_g, rr_g, c_g, cxsth_g, sxnth_g, uu_g, vv_g, DR_g, R_g);
1051
1053
// CUT_CHECK_ERROR("dissipation execution failed\n");
1052
1054
CUDA_CHECK (cudaThreadSynchronize ());
@@ -1062,9 +1064,9 @@ void wavestep(XBGPUParam Param)
1062
1064
1063
1065
1064
1066
1065
- //
1067
+ //
1066
1068
// Compute mean wave direction
1067
- //
1069
+ //
1068
1070
1069
1071
meandir << <gridDim , blockDim , 0 >> >(nx, ny, ntheta, Param.rho , Param.g , Param.dtheta , ee_g, theta_g, thetamean_g, E_g, H_g);
1070
1072
// CUT_CHECK_ERROR("meandir execution failed\n");
@@ -1095,7 +1097,7 @@ void wavestep(XBGPUParam Param)
1095
1097
// CUT_CHECK_ERROR("radstress execution failed\n");
1096
1098
CUDA_CHECK (cudaThreadSynchronize ());
1097
1099
1098
- //
1100
+ //
1099
1101
// Wave forces
1100
1102
//
1101
1103
wavforce << <gridDim , blockDim , 0 >> >(nx, ny, ntheta, Param.dx , Param.dtheta , Sxx_g, Sxy_g, Syy_g, Fx_g, Fy_g, hh_g);
0 commit comments