smyalygames/FiniteVolumeGPU
1
0
Fork 0
mirror of https://github.com/smyalygames/FiniteVolumeGPU.git synced 2025-11-29 17:28:03 +01:00
Code Issues Actions Packages Projects Releases Wiki Activity

Refactoring - broke 2nd order

Browse Source
This commit is contained in:
André R. Brodtkorb
2018-11-09 11:46:34 +01:00
parent 812ebcc9ba
commit e98ae0a979
21 changed files with 499023 additions and 611 deletions
Download Patch File Download Diff File Expand all files Collapse all files

33
GPUSimulators/cuda/EE2D_KP07_dimsplit.cu
View File

@@ -157,8 +157,8 @@ __global__ void KP07DimsplitKernel(
//Read into shared memory
readBlock<w, h, gc, 1, 1>( rho0_ptr_, rho0_pitch_, Q[0], nx_, ny_, boundary_conditions_);
readBlock<w, h, gc, 1, -1>(rho_u0_ptr_, rho_u0_pitch_, Q[1], nx_, ny_, boundary_conditions_);
readBlock<w, h, gc, -1, 1>(rho_v0_ptr_, rho_v0_pitch_, Q[2], nx_, ny_, boundary_conditions_);
readBlock<w, h, gc, -1, 1>(rho_u0_ptr_, rho_u0_pitch_, Q[1], nx_, ny_, boundary_conditions_);
readBlock<w, h, gc, 1, -1>(rho_v0_ptr_, rho_v0_pitch_, Q[2], nx_, ny_, boundary_conditions_);
readBlock<w, h, gc, 1, 1>( E0_ptr_, E0_pitch_, Q[3], nx_, ny_, boundary_conditions_);
__syncthreads();
@@ -226,33 +226,16 @@ __global__ void KP07DimsplitKernel(
Q[3][j][i] -= g_*rho_v*dt_;
__syncthreads();
}
//This is the RK2-part
if (getOrder(step_order_) == 2) {
const int i = threadIdx.x + gc;
const int j = threadIdx.y + gc;
const int tx = blockDim.x*blockIdx.x + i;
const int ty = blockDim.y*blockIdx.y + j;
const float q1 = ((float*) ((char*) rho1_ptr_ + rho1_pitch_*ty))[tx];
const float q2 = ((float*) ((char*) rho_u1_ptr_ + rho_u1_pitch_*ty))[tx];
const float q3 = ((float*) ((char*) rho_v1_ptr_ + rho_v1_pitch_*ty))[tx];
const float q4 = ((float*) ((char*) E1_ptr_ + E1_pitch_*ty))[tx];
Q[0][j][i] = 0.5f*( Q[0][j][i] + q1 );
Q[1][j][i] = 0.5f*( Q[1][j][i] + q2 );
Q[2][j][i] = 0.5f*( Q[2][j][i] + q3 );
Q[3][j][i] = 0.5f*( Q[3][j][i] + q4 );
__syncthreads();
}
}
// Write to main memory for all internal cells
writeBlock<w, h, gc>( rho1_ptr_, rho1_pitch_, Q[0], nx_, ny_);
writeBlock<w, h, gc>(rho_u1_ptr_, rho_u1_pitch_, Q[1], nx_, ny_);
writeBlock<w, h, gc>(rho_v1_ptr_, rho_v1_pitch_, Q[2], nx_, ny_);
writeBlock<w, h, gc>( E1_ptr_, E1_pitch_, Q[3], nx_, ny_);
const int step = getStep(step_order_);
const int order = getOrder(step_order_);
writeBlock<w, h, gc>( rho1_ptr_, rho1_pitch_, Q[0], nx_, ny_, step, order);
writeBlock<w, h, gc>(rho_u1_ptr_, rho_u1_pitch_, Q[1], nx_, ny_, step, order);
writeBlock<w, h, gc>(rho_v1_ptr_, rho_v1_pitch_, Q[2], nx_, ny_, step, order);
writeBlock<w, h, gc>( E1_ptr_, E1_pitch_, Q[3], nx_, ny_, step, order);
}
} // extern "C"

26
GPUSimulators/cuda/SWE2D_FORCE.cu
View File

@@ -87,6 +87,8 @@ __global__ void FORCEKernel(
float dx_, float dy_, float dt_,
float g_,
int boundary_conditions_,
//Input h^n
float* h0_ptr_, int h0_pitch_,
float* hu0_ptr_, int hu0_pitch_,
@@ -106,15 +108,9 @@ __global__ void FORCEKernel(
__shared__ float F[3][h+2][w+2];
//Read into shared memory
readBlock<w, h, gc>( h0_ptr_, h0_pitch_, Q[0], nx_+2, ny_+2);
readBlock<w, h, gc>(hu0_ptr_, hu0_pitch_, Q[1], nx_+2, ny_+2);
readBlock<w, h, gc>(hv0_ptr_, hv0_pitch_, Q[2], nx_+2, ny_+2);
__syncthreads();
//Set boundary conditions
noFlowBoundary<w, h, gc, 1, 1>(Q[0], nx_, ny_);
noFlowBoundary<w, h, gc, -1, 1>(Q[1], nx_, ny_);
noFlowBoundary<w, h, gc, 1, -1>(Q[2], nx_, ny_);
readBlock<w, h, gc, 1, 1>( h0_ptr_, h0_pitch_, Q[0], nx_, ny_, boundary_conditions_);
readBlock<w, h, gc, -1, 1>(hu0_ptr_, hu0_pitch_, Q[1], nx_, ny_, boundary_conditions_);
readBlock<w, h, gc, 1, -1>(hv0_ptr_, hv0_pitch_, Q[2], nx_, ny_, boundary_conditions_);
__syncthreads();
//Compute flux along x, and evolve
@@ -123,12 +119,6 @@ __global__ void FORCEKernel(
evolveF<w, h, gc, vars>(Q, F, dx_, dt_);
__syncthreads();
//Set boundary conditions
noFlowBoundary<w, h, gc, 1, 1>(Q[0], nx_, ny_);
noFlowBoundary<w, h, gc, -1, 1>(Q[1], nx_, ny_);
noFlowBoundary<w, h, gc, 1, -1>(Q[2], nx_, ny_);
__syncthreads();
//Compute flux along y, and evolve
computeFluxG(Q, F, g_, dy_, dt_);
__syncthreads();
@@ -136,9 +126,9 @@ __global__ void FORCEKernel(
__syncthreads();
//Write to main memory
writeBlock<w, h, gc>( h1_ptr_, h1_pitch_, Q[0], nx_, ny_);
writeBlock<w, h, gc>(hu1_ptr_, hu1_pitch_, Q[1], nx_, ny_);
writeBlock<w, h, gc>(hv1_ptr_, hv1_pitch_, Q[2], nx_, ny_);
writeBlock<w, h, gc>( h1_ptr_, h1_pitch_, Q[0], nx_, ny_, 0, 1);
writeBlock<w, h, gc>(hu1_ptr_, hu1_pitch_, Q[1], nx_, ny_, 0, 1);
writeBlock<w, h, gc>(hv1_ptr_, hv1_pitch_, Q[2], nx_, ny_, 0, 1);
}
} // extern "C"

27
GPUSimulators/cuda/SWE2D_HLL.cu
View File

@@ -103,6 +103,8 @@ __global__ void HLLKernel(
float dx_, float dy_, float dt_,
float g_,
int boundary_conditions_,
//Input h^n
float* h0_ptr_, int h0_pitch_,
float* hu0_ptr_, int hu0_pitch_,
@@ -123,16 +125,9 @@ __global__ void HLLKernel(
__shared__ float F[3][h+2][w+2];
//Read into shared memory
readBlock<w, h, gc>( h0_ptr_, h0_pitch_, Q[0], nx_+2, ny_+2);
readBlock<w, h, gc>(hu0_ptr_, hu0_pitch_, Q[1], nx_+2, ny_+2);
readBlock<w, h, gc>(hv0_ptr_, hv0_pitch_, Q[2], nx_+2, ny_+2);
__syncthreads();
//Set boundary conditions
noFlowBoundary<w, h, gc, 1, 1>(Q[0], nx_, ny_);
noFlowBoundary<w, h, gc, -1, 1>(Q[1], nx_, ny_);
noFlowBoundary<w, h, gc, 1, -1>(Q[2], nx_, ny_);
__syncthreads();
readBlock<w, h, gc, 1, 1>( h0_ptr_, h0_pitch_, Q[0], nx_, ny_, boundary_conditions_);
readBlock<w, h, gc, -1, 1>(hu0_ptr_, hu0_pitch_, Q[1], nx_, ny_, boundary_conditions_);
readBlock<w, h, gc, 1, -1>(hv0_ptr_, hv0_pitch_, Q[2], nx_, ny_, boundary_conditions_);
//Compute F flux
computeFluxF(Q, F, g_);
@@ -141,12 +136,6 @@ __global__ void HLLKernel(
evolveF<w, h, gc, vars>(Q, F, dx_, dt_);
__syncthreads();
//Set boundary conditions
noFlowBoundary<w, h, gc, 1, 1>(Q[0], nx_, ny_);
noFlowBoundary<w, h, gc, -1, 1>(Q[1], nx_, ny_);
noFlowBoundary<w, h, gc, 1, -1>(Q[2], nx_, ny_);
__syncthreads();
//Compute G flux
computeFluxG(Q, F, g_);
__syncthreads();
@@ -155,9 +144,9 @@ __global__ void HLLKernel(
__syncthreads();
// Write to main memory for all internal cells
writeBlock<w, h, gc>( h1_ptr_, h1_pitch_, Q[0], nx_, ny_);
writeBlock<w, h, gc>(hu1_ptr_, hu1_pitch_, Q[1], nx_, ny_);
writeBlock<w, h, gc>(hv1_ptr_, hv1_pitch_, Q[2], nx_, ny_);
writeBlock<w, h, gc>( h1_ptr_, h1_pitch_, Q[0], nx_, ny_, 0, 1);
writeBlock<w, h, gc>(hu1_ptr_, hu1_pitch_, Q[1], nx_, ny_, 0, 1);
writeBlock<w, h, gc>(hv1_ptr_, hv1_pitch_, Q[2], nx_, ny_, 0, 1);
}
} // extern "C"

38
GPUSimulators/cuda/SWE2D_HLL2.cu
View File

@@ -130,7 +130,8 @@ __global__ void HLL2Kernel(
float theta_,
int step_,
int step_order_,
int boundary_conditions_,
//Input h^n
float* h0_ptr_, int h0_pitch_,
@@ -153,19 +154,12 @@ __global__ void HLL2Kernel(
__shared__ float F[3][h+4][w+4];
//Read into shared memory
readBlock<w, h, gc>( h0_ptr_, h0_pitch_, Q[0], nx_+2, ny_+2);
readBlock<w, h, gc>(hu0_ptr_, hu0_pitch_, Q[1], nx_+2, ny_+2);
readBlock<w, h, gc>(hv0_ptr_, hv0_pitch_, Q[2], nx_+2, ny_+2);
__syncthreads();
//Set boundary conditions
noFlowBoundary<w, h, gc, 1, 1>(Q[0], nx_, ny_);
noFlowBoundary<w, h, gc, -1, 1>(Q[1], nx_, ny_);
noFlowBoundary<w, h, gc, 1, -1>(Q[2], nx_, ny_);
__syncthreads();
readBlock<w, h, gc, 1, 1>( h0_ptr_, h0_pitch_, Q[0], nx_, ny_, boundary_conditions_);
readBlock<w, h, gc, -1, 1>(hu0_ptr_, hu0_pitch_, Q[1], nx_, ny_, boundary_conditions_);
readBlock<w, h, gc, 1, -1>(hv0_ptr_, hv0_pitch_, Q[2], nx_, ny_, boundary_conditions_);
//Step 0 => evolve x first, then y
if (step_ == 0) {
if (getStep(step_order_) == 0) {
//Compute fluxes along the x axis and evolve
minmodSlopeX<w, h, gc, vars>(Q, Qx, theta_);
__syncthreads();
@@ -174,12 +168,6 @@ __global__ void HLL2Kernel(
evolveF<w, h, gc, vars>(Q, F, dx_, dt_);
__syncthreads();
//Set boundary conditions
noFlowBoundary<w, h, gc, 1, 1>(Q[0], nx_, ny_);
noFlowBoundary<w, h, gc, -1, 1>(Q[1], nx_, ny_);
noFlowBoundary<w, h, gc, 1, -1>(Q[2], nx_, ny_);
__syncthreads();
//Compute fluxes along the y axis and evolve
minmodSlopeY<w, h, gc, vars>(Q, Qx, theta_);
__syncthreads();
@@ -198,12 +186,6 @@ __global__ void HLL2Kernel(
evolveG<w, h, gc, vars>(Q, F, dy_, dt_);
__syncthreads();
//Set boundary conditions
noFlowBoundary<w, h, gc, 1, 1>(Q[0], nx_, ny_);
noFlowBoundary<w, h, gc, -1, 1>(Q[1], nx_, ny_);
noFlowBoundary<w, h, gc, 1, -1>(Q[2], nx_, ny_);
__syncthreads();
//Compute fluxes along the x axis and evolve
minmodSlopeX<w, h, gc, vars>(Q, Qx, theta_);
__syncthreads();
@@ -217,9 +199,11 @@ __global__ void HLL2Kernel(
// Write to main memory for all internal cells
writeBlock<w, h, 2>( h1_ptr_, h1_pitch_, Q[0], nx_, ny_);
writeBlock<w, h, 2>(hu1_ptr_, hu1_pitch_, Q[1], nx_, ny_);
writeBlock<w, h, 2>(hv1_ptr_, hv1_pitch_, Q[2], nx_, ny_);
const int step = getStep(step_order_);
const int order = getOrder(step_order_);
writeBlock<w, h, 2>( h1_ptr_, h1_pitch_, Q[0], nx_, ny_, step, order);
writeBlock<w, h, 2>(hu1_ptr_, hu1_pitch_, Q[1], nx_, ny_, step, order);
writeBlock<w, h, 2>(hv1_ptr_, hv1_pitch_, Q[2], nx_, ny_, step, order);
}
} // extern "C"

50
GPUSimulators/cuda/SWE2D_KP07.cu
View File

@@ -140,7 +140,8 @@ __global__ void KP07Kernel(
float theta_,
int step_,
int step_order_,
int boundary_conditions_,
//Input h^n
float* h0_ptr_, int h0_pitch_,
@@ -167,9 +168,6 @@ __global__ void KP07Kernel(
//Shared memory variables
__shared__ float Q[3][h+4][w+4];
//The following slightly wastes memory, but enables us to reuse the
//funcitons in common.opencl
__shared__ float Qx[3][h+2][w+2];
__shared__ float Qy[3][h+2][w+2];
__shared__ float F[3][h+1][w+1];
@@ -178,17 +176,9 @@ __global__ void KP07Kernel(
//Read into shared memory
readBlock<w, h, gc>( h0_ptr_, h0_pitch_, Q[0], nx_+2, ny_+2);
readBlock<w, h, gc>(hu0_ptr_, hu0_pitch_, Q[1], nx_+2, ny_+2);
readBlock<w, h, gc>(hv0_ptr_, hv0_pitch_, Q[2], nx_+2, ny_+2);
__syncthreads();
//Fix boundary conditions
noFlowBoundary<w, h, gc, 1, 1>(Q[0], nx_, ny_);
noFlowBoundary<w, h, gc, -1, 1>(Q[1], nx_, ny_);
noFlowBoundary<w, h, gc, 1, -1>(Q[2], nx_, ny_);
__syncthreads();
readBlock<w, h, gc, 1, 1>( h0_ptr_, h0_pitch_, Q[0], nx_, ny_, boundary_conditions_);
readBlock<w, h, gc, -1, 1>(hu0_ptr_, hu0_pitch_, Q[1], nx_, ny_, boundary_conditions_);
readBlock<w, h, gc, 1, -1>(hv0_ptr_, hv0_pitch_, Q[2], nx_, ny_, boundary_conditions_);
//Reconstruct slopes along x and axis
@@ -218,32 +208,18 @@ __global__ void KP07Kernel(
float* const h_row = (float*) ((char*) h1_ptr_ + h1_pitch_*tj);
float* const hu_row = (float*) ((char*) hu1_ptr_ + hu1_pitch_*tj);
float* const hv_row = (float*) ((char*) hv1_ptr_ + hv1_pitch_*tj);
if (step_ == 0) {
//First step of RK2 ODE integrator
if (getOrder(step_order_) == 2 && getStep(step_order_) == 1) {
//Write to main memory
h_row[ti] = 0.5f*(h_row[ti] + h1);
hu_row[ti] = 0.5f*(hu_row[ti] + hu1);
hv_row[ti] = 0.5f*(hv_row[ti] + hv1);
}
else {
h_row[ti] = h1;
hu_row[ti] = hu1;
hv_row[ti] = hv1;
}
else if (step_ == 1) {
//Second step of RK2 ODE integrator
//First read Q^n
const float h_a = h_row[ti];
const float hu_a = hu_row[ti];
const float hv_a = hv_row[ti];
//Compute Q^n+1
const float h_b = 0.5f*(h_a + h1);
const float hu_b = 0.5f*(hu_a + hu1);
const float hv_b = 0.5f*(hv_a + hv1);
//Write to main memory
h_row[ti] = h_b;
hu_row[ti] = hu_b;
hv_row[ti] = hv_b;
}
}
}
} //extern "C"

42
GPUSimulators/cuda/SWE2D_KP07_dimsplit.cu
View File

@@ -121,7 +121,8 @@ __global__ void KP07DimsplitKernel(
float theta_,
int step_,
int step_order_,
int boundary_conditions_,
//Input h^n
float* h0_ptr_, int h0_pitch_,
@@ -147,22 +148,14 @@ __global__ void KP07DimsplitKernel(
//Read into shared memory
readBlock<w, h, gc>( h0_ptr_, h0_pitch_, Q[0], nx_+2, ny_+2);
readBlock<w, h, gc>(hu0_ptr_, hu0_pitch_, Q[1], nx_+2, ny_+2);
readBlock<w, h, gc>(hv0_ptr_, hv0_pitch_, Q[2], nx_+2, ny_+2);
__syncthreads();
//Fix boundary conditions
noFlowBoundary<w, h, gc, 1, 1>(Q[0], nx_, ny_);
noFlowBoundary<w, h, gc, -1, 1>(Q[1], nx_, ny_);
noFlowBoundary<w, h, gc, 1, -1>(Q[2], nx_, ny_);
__syncthreads();
readBlock<w, h, gc, 1, 1>( h0_ptr_, h0_pitch_, Q[0], nx_, ny_, boundary_conditions_);
readBlock<w, h, gc, -1, 1>(hu0_ptr_, hu0_pitch_, Q[1], nx_, ny_, boundary_conditions_);
readBlock<w, h, gc, 1, -1>(hv0_ptr_, hv0_pitch_, Q[2], nx_, ny_, boundary_conditions_);
//Step 0 => evolve x first, then y
if (step_ == 0) {
if (getStep(step_order_) == 0) {
//Compute fluxes along the x axis and evolve
minmodSlopeX<w, h, gc, vars>(Q, Qx, theta_);
__syncthreads();
@@ -171,18 +164,9 @@ __global__ void KP07DimsplitKernel(
evolveF<w, h, gc, vars>(Q, F, dx_, dt_);
__syncthreads();
//Set boundary conditions
noFlowBoundary<w, h, gc, 1, 1>(Q[0], nx_, ny_);
noFlowBoundary<w, h, gc, -1, 1>(Q[1], nx_, ny_);
noFlowBoundary<w, h, gc, 1, -1>(Q[2], nx_, ny_);
__syncthreads();
//Compute fluxes along the y axis and evolve
minmodSlopeY<w, h, gc, vars>(Q, Qx, theta_);
__syncthreads();
computeFluxG(Q, Qx, F, g_, dy_, dt_);
__syncthreads();
evolveG<w, h, gc, vars>(Q, F, dy_, dt_);
@@ -198,12 +182,6 @@ __global__ void KP07DimsplitKernel(
evolveG<w, h, gc, vars>(Q, F, dy_, dt_);
__syncthreads();
//Set boundary conditions
noFlowBoundary<w, h, gc, 1, 1>(Q[0], nx_, ny_);
noFlowBoundary<w, h, gc, -1, 1>(Q[1], nx_, ny_);
noFlowBoundary<w, h, gc, 1, -1>(Q[2], nx_, ny_);
__syncthreads();
//Compute fluxes along the x axis and evolve
minmodSlopeX<w, h, gc, vars>(Q, Qx, theta_);
__syncthreads();
@@ -215,9 +193,11 @@ __global__ void KP07DimsplitKernel(
// Write to main memory for all internal cells
writeBlock<w, h, gc>( h1_ptr_, h1_pitch_, Q[0], nx_, ny_);
writeBlock<w, h, gc>(hu1_ptr_, hu1_pitch_, Q[1], nx_, ny_);
writeBlock<w, h, gc>(hv1_ptr_, hv1_pitch_, Q[2], nx_, ny_);
const int step = getStep(step_order_);
const int order = getOrder(step_order_);
writeBlock<w, h, gc>( h1_ptr_, h1_pitch_, Q[0], nx_, ny_, step, order);
writeBlock<w, h, gc>(hu1_ptr_, hu1_pitch_, Q[1], nx_, ny_, step, order);
writeBlock<w, h, gc>(hv1_ptr_, hv1_pitch_, Q[2], nx_, ny_, step, order);
}
} // extern "C"

21
GPUSimulators/cuda/SWE2D_LxF.cu
View File

@@ -104,6 +104,8 @@ void LxFKernel(
float dx_, float dy_, float dt_,
float g_,
int boundary_conditions_,
//Input h^n
float* h0_ptr_, int h0_pitch_,
float* hu0_ptr_, int hu0_pitch_,
@@ -123,16 +125,9 @@ void LxFKernel(
__shared__ float G[3][h+1][w ];
//Read from global memory
readBlock<w, h, gc>( h0_ptr_, h0_pitch_, Q[0], nx_+2, ny_+2);
readBlock<w, h, gc>(hu0_ptr_, hu0_pitch_, Q[1], nx_+2, ny_+2);
readBlock<w, h, gc>(hv0_ptr_, hv0_pitch_, Q[2], nx_+2, ny_+2);
__syncthreads();
//Set boundary conditions
noFlowBoundary<w, h, gc, 1, 1>(Q[0], nx_, ny_);
noFlowBoundary<w, h, gc, -1, 1>(Q[1], nx_, ny_);
noFlowBoundary<w, h, gc, 1, -1>(Q[2], nx_, ny_);
__syncthreads();
readBlock<w, h, gc, 1, 1>( h0_ptr_, h0_pitch_, Q[0], nx_, ny_, boundary_conditions_);
readBlock<w, h, gc, 1, -1>(hu0_ptr_, hu0_pitch_, Q[1], nx_, ny_, boundary_conditions_);
readBlock<w, h, gc, -1, 1>(hv0_ptr_, hv0_pitch_, Q[2], nx_, ny_, boundary_conditions_);
//Compute fluxes along the x and y axis
computeFluxF<w, h>(Q, F, g_, dx_, dt_);
@@ -154,9 +149,9 @@ void LxFKernel(
__syncthreads();
//Write to main memory
writeBlock<w, h, gc>( h1_ptr_, h1_pitch_, Q[0], nx_, ny_);
writeBlock<w, h, gc>(hu1_ptr_, hu1_pitch_, Q[1], nx_, ny_);
writeBlock<w, h, gc>(hv1_ptr_, hv1_pitch_, Q[2], nx_, ny_);
writeBlock<w, h, gc>( h1_ptr_, h1_pitch_, Q[0], nx_, ny_, 0, 1);
writeBlock<w, h, gc>(hu1_ptr_, hu1_pitch_, Q[1], nx_, ny_, 0, 1);
writeBlock<w, h, gc>(hv1_ptr_, hv1_pitch_, Q[2], nx_, ny_, 0, 1);
}
} // extern "C"

40
GPUSimulators/cuda/SWE2D_WAF.cu
View File

@@ -103,7 +103,10 @@ extern "C" {
__global__ void WAFKernel(
int nx_, int ny_,
float dx_, float dy_, float dt_,
float g_, int step_,
float g_,
int step_order_,
int boundary_conditions_,
//Input h^n
float* h0_ptr_, int h0_pitch_,
@@ -127,34 +130,21 @@ __global__ void WAFKernel(
//Read into shared memory Q from global memory
readBlock<w, h, gc>( h0_ptr_, h0_pitch_, Q[0], nx_+2, ny_+2);
readBlock<w, h, gc>(hu0_ptr_, hu0_pitch_, Q[1], nx_+2, ny_+2);
readBlock<w, h, gc>(hv0_ptr_, hv0_pitch_, Q[2], nx_+2, ny_+2);
__syncthreads();
//Set boundary conditions
noFlowBoundary<w, h, gc, 1, 1>(Q[0], nx_, ny_);
noFlowBoundary<w, h, gc, -1, 1>(Q[1], nx_, ny_);
noFlowBoundary<w, h, gc, 1, -1>(Q[2], nx_, ny_);
readBlock<w, h, gc, 1, 1>( h0_ptr_, h0_pitch_, Q[0], nx_, ny_, boundary_conditions_);
readBlock<w, h, gc, -1, 1>(hu0_ptr_, hu0_pitch_, Q[1], nx_, ny_, boundary_conditions_);
readBlock<w, h, gc, 1, -1>(hv0_ptr_, hv0_pitch_, Q[2], nx_, ny_, boundary_conditions_);
__syncthreads();
//Step 0 => evolve x first, then y
if (step_ == 0) {
if (getStep(step_order_) == 0) {
//Compute fluxes along the x axis and evolve
computeFluxF(Q, F, g_, dx_, dt_);
__syncthreads();
evolveF<w, h, gc, vars>(Q, F, dx_, dt_);
__syncthreads();
//Fix boundary conditions
noFlowBoundary<w, h, gc, 1, 1>(Q[0], nx_, ny_);
noFlowBoundary<w, h, gc, -1, 1>(Q[1], nx_, ny_);
noFlowBoundary<w, h, gc, 1, -1>(Q[2], nx_, ny_);
__syncthreads();
//Compute fluxes along the y axis and evolve
computeFluxG(Q, F, g_, dy_, dt_);
__syncthreads();
@@ -169,12 +159,6 @@ __global__ void WAFKernel(
evolveG<w, h, gc, vars>(Q, F, dy_, dt_);
__syncthreads();
//Fix boundary conditions
noFlowBoundary<w, h, gc, 1, 1>(Q[0], nx_, ny_);
noFlowBoundary<w, h, gc, -1, 1>(Q[1], nx_, ny_);
noFlowBoundary<w, h, gc, 1, -1>(Q[2], nx_, ny_);
__syncthreads();
//Compute fluxes along the x axis and evolve
computeFluxF(Q, F, g_, dx_, dt_);
__syncthreads();
@@ -185,9 +169,11 @@ __global__ void WAFKernel(
// Write to main memory for all internal cells
writeBlock<w, h, 2>( h1_ptr_, h1_pitch_, Q[0], nx_, ny_);
writeBlock<w, h, 2>(hu1_ptr_, hu1_pitch_, Q[1], nx_, ny_);
writeBlock<w, h, 2>(hv1_ptr_, hv1_pitch_, Q[2], nx_, ny_);
const int step = getStep(step_order_);
const int order = getOrder(step_order_);
writeBlock<w, h, gc>( h1_ptr_, h1_pitch_, Q[0], nx_, ny_, step, order);
writeBlock<w, h, gc>(hu1_ptr_, hu1_pitch_, Q[1], nx_, ny_, step, order);
writeBlock<w, h, gc>(hv1_ptr_, hv1_pitch_, Q[2], nx_, ny_, step, order);
}
} // extern "C"

77
GPUSimulators/cuda/common.h
View File

@@ -134,14 +134,19 @@ inline __device__ BoundaryCondition getBCWest(int bc_) {
template<int block_width, int block_height, int ghost_cells>
/**
* Alter the index l so that it gives periodic boundary conditions when reading
*/
template<int ghost_cells>
inline __device__ int handlePeriodicBoundaryX(int k, int nx_, int boundary_conditions_) {
const int gc_pad = 2*ghost_cells;
//West boundary: add an offset to read from east of domain
if ((k < gc_pad)
&& getBCWest(boundary_conditions_) == Periodic) {
k += (nx_+2*ghost_cells - 2*gc_pad);
}
//East boundary: subtract an offset to read from west of domain
else if ((k >= nx_+2*ghost_cells-gc_pad)
&& getBCEast(boundary_conditions_) == Periodic) {
k -= (nx_+2*ghost_cells - 2*gc_pad);
@@ -150,14 +155,19 @@ inline __device__ int handlePeriodicBoundaryX(int k, int nx_, int boundary_condi
return k;
}
template<int block_width, int block_height, int ghost_cells>
/**
* Alter the index l so that it gives periodic boundary conditions when reading
*/
template<int ghost_cells>
inline __device__ int handlePeriodicBoundaryY(int l, int ny_, int boundary_conditions_) {
const int gc_pad = 2*ghost_cells;
//South boundary: add an offset to read from north of domain
if ((l < gc_pad)
&& getBCSouth(boundary_conditions_) == Periodic) {
l += (ny_+2*ghost_cells - 2*gc_pad);
}
//North boundary: subtract an offset to read from south of domain
else if ((l >= ny_+2*ghost_cells-gc_pad)
&& getBCNorth(boundary_conditions_) == Periodic) {
l -= (ny_+2*ghost_cells - 2*gc_pad);
@@ -165,12 +175,36 @@ inline __device__ int handlePeriodicBoundaryY(int l, int ny_, int boundary_condi
return l;
}
template<int block_width, int block_height, int ghost_cells, int sign_x, int sign_y>
inline __device__ int handleReflectiveBoundary(
float Q[block_height+2*ghost_cells][block_width+2*ghost_cells],
const int nx_, const int ny_,
const int boundary_conditions_) {
//Handle reflective boundary conditions
if (getBCNorth(boundary_conditions_) == Reflective) {
bcNorthReflective<block_width, block_height, ghost_cells, sign_y>(Q, nx_, ny_);
__syncthreads();
}
if (getBCSouth(boundary_conditions_) == Reflective) {
bcSouthReflective<block_width, block_height, ghost_cells, sign_y>(Q, nx_, ny_);
__syncthreads();
}
if (getBCEast(boundary_conditions_) == Reflective) {
bcEastReflective<block_width, block_height, ghost_cells, sign_x>(Q, nx_, ny_);
__syncthreads();
}
if (getBCWest(boundary_conditions_) == Reflective) {
bcWestReflective<block_width, block_height, ghost_cells, sign_x>(Q, nx_, ny_);
__syncthreads();
}
}
/**
* Reads a block of data with ghost cells
*/
template<int block_width, int block_height, int ghost_cells, int sign_north_south, int sign_east_west>
template<int block_width, int block_height, int ghost_cells, int sign_x, int sign_y>
inline __device__ void readBlock(float* ptr_, int pitch_,
float Q[block_height+2*ghost_cells][block_width+2*ghost_cells],
const int nx_, const int ny_,
@@ -183,13 +217,13 @@ inline __device__ void readBlock(float* ptr_, int pitch_,
//Loop over all variables
for (int j=threadIdx.y; j<block_height+2*ghost_cells; j+=block_height) {
//Handle periodic boundary conditions here
int l = handlePeriodicBoundaryY<block_width, block_height, ghost_cells>(by + j, ny_, boundary_conditions_);
int l = handlePeriodicBoundaryY<ghost_cells>(by + j, ny_, boundary_conditions_);
l = min(l, ny_+2*ghost_cells-1);
float* row = (float*) ((char*) ptr_ + pitch_*l);
for (int i=threadIdx.x; i<block_width+2*ghost_cells; i+=block_width) {
//Handle periodic boundary conditions here
int k = handlePeriodicBoundaryX<block_width, block_height, ghost_cells>(bx + i, nx_, boundary_conditions_);
int k = handlePeriodicBoundaryX<ghost_cells>(bx + i, nx_, boundary_conditions_);
k = min(k, nx_+2*ghost_cells-1);
//Read from global memory
@@ -198,23 +232,7 @@ inline __device__ void readBlock(float* ptr_, int pitch_,
}
__syncthreads();
//Handle reflective boundary conditions
if (getBCNorth(boundary_conditions_) == Reflective) {
bcNorthReflective<block_width, block_height, ghost_cells, sign_north_south>(Q, nx_, ny_);
__syncthreads();
}
if (getBCSouth(boundary_conditions_) == Reflective) {
bcSouthReflective<block_width, block_height, ghost_cells, sign_north_south>(Q, nx_, ny_);
__syncthreads();
}
if (getBCEast(boundary_conditions_) == Reflective) {
bcEastReflective<block_width, block_height, ghost_cells, sign_east_west>(Q, nx_, ny_);
__syncthreads();
}
if (getBCWest(boundary_conditions_) == Reflective) {
bcWestReflective<block_width, block_height, ghost_cells, sign_east_west>(Q, nx_, ny_);
__syncthreads();
}
handleReflectiveBoundary<block_width, block_height, ghost_cells, sign_x, sign_y>(Q, nx_, ny_, boundary_conditions_);
}
@@ -226,7 +244,8 @@ inline __device__ void readBlock(float* ptr_, int pitch_,
template<int block_width, int block_height, int ghost_cells>
inline __device__ void writeBlock(float* ptr_, int pitch_,
float shmem[block_height+2*ghost_cells][block_width+2*ghost_cells],
const int width, const int height) {
const int width, const int height,
int rk_step_, int rk_order_) {
//Index of cell within domain
const int ti = blockDim.x*blockIdx.x + threadIdx.x + ghost_cells;
@@ -239,7 +258,14 @@ inline __device__ void writeBlock(float* ptr_, int pitch_,
const int ty = threadIdx.y + ghost_cells;
float* const row = (float*) ((char*) ptr_ + pitch_*tj);
row[ti] = shmem[ty][tx];
//Handle runge-kutta timestepping here
if (rk_order_ == 2 && rk_step_ == 1) {
row[ti] = 0.5f*(row[ti] + shmem[ty][tx]);
}
else {
row[ti] = shmem[ty][tx];
}
}
}
@@ -255,6 +281,7 @@ inline __device__ void writeBlock(float* ptr_, int pitch_,
// West boundary
template<int block_width, int block_height, int ghost_cells, int sign>
__device__ void bcWestReflective(float Q[block_height+2*ghost_cells][block_width+2*ghost_cells], const int nx_, const int ny_) {