smyalygames/FiniteVolumeGPU
1
0
Fork 0
mirror of https://github.com/smyalygames/FiniteVolumeGPU.git synced 2025-07-05 08:00:59 +02:00
Code Issues Actions Packages Projects Releases Wiki Activity

Refactoring

Browse Source
This commit is contained in:
André R. Brodtkorb 2018-11-01 15:17:13 +01:00
parent 064027fc0b
commit 2b899d1c80
10 changed files with 234 additions and 194 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
GPUSimulators/Common.py
View File

@ -53,6 +53,9 @@ class Timer(object):
self.secs = self.end - self.start
self.msecs = self.secs * 1000 # millisecs
self.logger.log(self.log_level, "%s: %f ms", self.tag, self.msecs)
def elapsed(self):
return time.time() - self.start

2
GPUSimulators/EE2D_KP07_dimsplit.py
View File

@ -57,7 +57,7 @@ class EE2D_KP07_dimsplit (Simulator.BaseSimulator):
dx, dy, dt, \
gamma, \
theta=1.3, \
block_width=16, block_height=16):
block_width=8, block_height=4):
# Call super constructor
super().__init__(context, \

79
GPUSimulators/Simulator.py
View File

@ -82,6 +82,9 @@ class BaseSimulator:
#Keep track of simulation time
self.t = 0.0;
#Log progress every n seconds during simulation
self.log_every = 5
def __str__(self):
@ -101,6 +104,8 @@ class BaseSimulator:
with Common.Timer(self.__class__.__name__ + ".simulateEuler") as t:
# Compute number of timesteps to perform
n = int(t_end / self.dt + 1)
next_print = self.log_every
for i in range(0, n):
# Compute timestep for "this" iteration
@ -112,6 +117,12 @@ class BaseSimulator:
# Step with forward Euler
self.stepEuler(local_dt)
#Print info
if (t.elapsed() >= next_print):
self.logger.info("%s simulated %d of %d steps (Euler)", self, i, n)
next_print += self.log_every
self.logger.info("%s simulated %f seconds to %f with %d steps (Euler)", self, t_end, self.t, n)
return self.t, n
@ -121,19 +132,27 @@ class BaseSimulator:
Requires that the stepRK functionality is implemented in the subclasses
"""
def simulateRK(self, t_end, order):
# Compute number of timesteps to perform
n = int(t_end / self.dt + 1)
for i in range(0, n):
# Compute timestep for "this" iteration
local_dt = np.float32(min(self.dt, t_end-i*self.dt))
with Common.Timer(self.__class__.__name__ + ".simulateRK") as t:
# Compute number of timesteps to perform
n = int(t_end / self.dt + 1)
next_print = self.log_every
# Stop if end reached (should not happen)
if (local_dt <= 0.0):
break
# Perform all the Runge-Kutta substeps
self.stepRK(local_dt, order)
for i in range(0, n):
# Compute timestep for "this" iteration
local_dt = np.float32(min(self.dt, t_end-i*self.dt))
# Stop if end reached (should not happen)
if (local_dt <= 0.0):
break
# Perform all the Runge-Kutta substeps
self.stepRK(local_dt, order)
#Print info
if (t.elapsed() >= next_print):
self.logger.info("%s simulated %d of %d steps (RK2)", self, i, n)
next_print += self.log_every
self.logger.info("%s simulated %f seconds to %f with %d steps (RK2)", self, t_end, self.t, n)
return self.t, n
@ -143,22 +162,30 @@ class BaseSimulator:
Requires that the stepDimsplitX and stepDimsplitY functionality is implemented in the subclasses
"""
def simulateDimsplit(self, t_end):
# Compute number of timesteps to perform
n = int(t_end / (2.0*self.dt) + 1)
for i in range(0, n):
# Compute timestep for "this" iteration
local_dt = np.float32(0.5*min(2*self.dt, t_end-2*i*self.dt))
with Common.Timer(self.__class__.__name__ + ".simulateDimsplit") as t:
# Compute number of timesteps to perform
n = int(t_end / (2.0*self.dt) + 1)
# Stop if end reached (should not happen)
if (local_dt <= 0.0):
break
next_print = self.log_every
for i in range(0, n):
# Compute timestep for "this" iteration
local_dt = np.float32(0.5*min(2*self.dt, t_end-2*i*self.dt))
# Stop if end reached (should not happen)
if (local_dt <= 0.0):
break
# Perform the dimensional split substeps
self.stepDimsplitXY(local_dt)
self.stepDimsplitYX(local_dt)
#Print info
if (t.elapsed() >= next_print):
self.logger.info("%s simulated %d of %d steps (Dimsplit)", self, i, n)
next_print += self.log_every
# Perform the dimensional split substeps
self.stepDimsplitXY(local_dt)
self.stepDimsplitYX(local_dt)
self.logger.info("%s simulated %f seconds to %f with %d steps (dimsplit)", self, t_end, self.t, 2*n)
self.logger.info("%s simulated %f seconds to %f with %d steps (Dimsplit)", self, t_end, self.t, 2*n)
return self.t, 2*n

202
GPUSimulators/cuda/EE2D_KP07_dimsplit.cu
View File

@ -29,101 +29,89 @@ __device__
void computeFluxF(float Q[4][BLOCK_HEIGHT+4][BLOCK_WIDTH+4],
float Qx[4][BLOCK_HEIGHT+2][BLOCK_WIDTH+2],
float F[4][BLOCK_HEIGHT+1][BLOCK_WIDTH+1],
const float gamma_, const float dx_, const float dt_) {
//Index of thread within block
const int tx = threadIdx.x;
const int ty = threadIdx.y;
{
int j=ty;
const int l = j + 2; //Skip ghost cells
for (int i=tx; i<BLOCK_WIDTH+1; i+=BLOCK_WIDTH) {
const int k = i + 1;
// Reconstruct point values of Q at the left and right hand side
// of the cell for both the left (i) and right (i+1) cell
const float4 Q_rl = make_float4(Q[0][l][k+1] - 0.5f*Qx[0][j][i+1],
Q[1][l][k+1] - 0.5f*Qx[1][j][i+1],
Q[2][l][k+1] - 0.5f*Qx[2][j][i+1],
Q[4][l][k+1] - 0.5f*Qx[4][j][i+1]);
const float4 Q_rr = make_float4(Q[0][l][k+1] + 0.5f*Qx[0][j][i+1],
Q[1][l][k+1] + 0.5f*Qx[1][j][i+1],
Q[2][l][k+1] + 0.5f*Qx[2][j][i+1],
Q[4][l][k+1] + 0.5f*Qx[4][j][i+1]);
const float4 Q_ll = make_float4(Q[0][l][k] - 0.5f*Qx[0][j][i],
Q[1][l][k] - 0.5f*Qx[1][j][i],
Q[2][l][k] - 0.5f*Qx[2][j][i],
Q[4][l][k] - 0.5f*Qx[4][j][i]);
const float4 Q_lr = make_float4(Q[0][l][k] + 0.5f*Qx[0][j][i],
Q[1][l][k] + 0.5f*Qx[1][j][i],
Q[2][l][k] + 0.5f*Qx[2][j][i],
Q[4][l][k] + 0.5f*Qx[4][j][i]);
//Evolve half a timestep (predictor step)
const float4 Q_r_bar = Q_rl + dt_/(2.0f*dx_) * (F_func(Q_rl, gamma_) - F_func(Q_rr, gamma_));
const float4 Q_l_bar = Q_lr + dt_/(2.0f*dx_) * (F_func(Q_ll, gamma_) - F_func(Q_lr, gamma_));
const float gamma_, const float dx_, const float dt_) {
int j=threadIdx.y;
const int l = j + 2; //Skip ghost cells
for (int i=threadIdx.x; i<BLOCK_WIDTH+1; i+=BLOCK_WIDTH) {
const int k = i + 1;
// Reconstruct point values of Q at the left and right hand side
// of the cell for both the left (i) and right (i+1) cell
const float4 Q_rl = make_float4(Q[0][l][k+1] - 0.5f*Qx[0][j][i+1],
Q[1][l][k+1] - 0.5f*Qx[1][j][i+1],
Q[2][l][k+1] - 0.5f*Qx[2][j][i+1],
Q[3][l][k+1] - 0.5f*Qx[3][j][i+1]);
const float4 Q_rr = make_float4(Q[0][l][k+1] + 0.5f*Qx[0][j][i+1],
Q[1][l][k+1] + 0.5f*Qx[1][j][i+1],
Q[2][l][k+1] + 0.5f*Qx[2][j][i+1],
Q[3][l][k+1] + 0.5f*Qx[3][j][i+1]);
const float4 Q_ll = make_float4(Q[0][l][k] - 0.5f*Qx[0][j][i],
Q[1][l][k] - 0.5f*Qx[1][j][i],
Q[2][l][k] - 0.5f*Qx[2][j][i],
Q[3][l][k] - 0.5f*Qx[3][j][i]);
const float4 Q_lr = make_float4(Q[0][l][k] + 0.5f*Qx[0][j][i],
Q[1][l][k] + 0.5f*Qx[1][j][i],
Q[2][l][k] + 0.5f*Qx[2][j][i],
Q[3][l][k] + 0.5f*Qx[3][j][i]);
//Evolve half a timestep (predictor step)
const float4 Q_r_bar = Q_rl + dt_/(2.0f*dx_) * (F_func(Q_rl, gamma_) - F_func(Q_rr, gamma_));
const float4 Q_l_bar = Q_lr + dt_/(2.0f*dx_) * (F_func(Q_ll, gamma_) - F_func(Q_lr, gamma_));
// Compute flux based on prediction
const float4 flux = CentralUpwindFlux(Q_l_bar, Q_r_bar, gamma_);
//Write to shared memory
F[0][j][i] = flux.x;
F[1][j][i] = flux.y;
F[2][j][i] = flux.z;
F[3][j][i] = flux.w;
}
}
// Compute flux based on prediction
const float4 flux = CentralUpwindFlux(Q_l_bar, Q_r_bar, gamma_);
//Write to shared memory
F[0][j][i] = flux.x;
F[1][j][i] = flux.y;
F[2][j][i] = flux.z;
F[3][j][i] = flux.w;
}
}
__device__
void computeFluxG(float Q[4][BLOCK_HEIGHT+4][BLOCK_WIDTH+4],
float Qy[4][BLOCK_HEIGHT+2][BLOCK_WIDTH+2],
float G[4][BLOCK_HEIGHT+1][BLOCK_WIDTH+1],
const float gamma_, const float dy_, const float dt_) {
//Index of thread within block
const int tx = threadIdx.x;
const int ty = threadIdx.y;
for (int j=ty; j<BLOCK_HEIGHT+1; j+=BLOCK_HEIGHT) {
const float gamma_, const float dy_, const float dt_) {
int i=threadIdx.x;
const int k = i + 2; //Skip ghost cells
for (int j=threadIdx.y; j<BLOCK_HEIGHT+1; j+=BLOCK_HEIGHT) {
const int l = j + 1;
{
int i=tx;
const int k = i + 2; //Skip ghost cells
// Reconstruct point values of Q at the left and right hand side
// of the cell for both the left (i) and right (i+1) cell
//NOte that hu and hv are swapped ("transposing" the domain)!
const float4 Q_rl = make_float4(Q[0][l+1][k] - 0.5f*Qy[0][j+1][i],
Q[2][l+1][k] - 0.5f*Qy[2][j+1][i],
Q[1][l+1][k] - 0.5f*Qy[1][j+1][i],
Q[3][l+1][k] - 0.5f*Qy[3][j+1][i]);
const float4 Q_rr = make_float4(Q[0][l+1][k] + 0.5f*Qy[0][j+1][i],
Q[2][l+1][k] + 0.5f*Qy[2][j+1][i],
Q[1][l+1][k] + 0.5f*Qy[1][j+1][i],
Q[3][l+1][k] + 0.5f*Qy[3][j+1][i]);
const float4 Q_ll = make_float4(Q[0][l][k] - 0.5f*Qy[0][j][i],
Q[2][l][k] - 0.5f*Qy[2][j][i],
Q[1][l][k] - 0.5f*Qy[1][j][i],
Q[3][l][k] - 0.5f*Qy[3][j][i]);
const float4 Q_lr = make_float4(Q[0][l][k] + 0.5f*Qy[0][j][i],
Q[2][l][k] + 0.5f*Qy[2][j][i],
Q[1][l][k] + 0.5f*Qy[1][j][i],
Q[3][l][k] + 0.5f*Qy[3][j][i]);
//Evolve half a timestep (predictor step)
const float4 Q_r_bar = Q_rl + dt_/(2.0f*dy_) * (F_func(Q_rl, gamma_) - F_func(Q_rr, gamma_));
const float4 Q_l_bar = Q_lr + dt_/(2.0f*dy_) * (F_func(Q_ll, gamma_) - F_func(Q_lr, gamma_));
// Compute flux based on prediction
const float4 flux = make_float4(0.01, 0.01, 0.01, 0.01);//CentralUpwindFlux(Q_l_bar, Q_r_bar, gamma_);
//Write to shared memory
//Note that we here swap hu and hv back to the original
G[0][j][i] = flux.x;
G[1][j][i] = flux.z;
G[2][j][i] = flux.y;
G[3][j][i] = flux.w;
}
// Reconstruct point values of Q at the left and right hand side
// of the cell for both the left (i) and right (i+1) cell
//NOte that hu and hv are swapped ("transposing" the domain)!
const float4 Q_rl = make_float4(Q[0][l+1][k] - 0.5f*Qy[0][j+1][i],
Q[2][l+1][k] - 0.5f*Qy[2][j+1][i],
Q[1][l+1][k] - 0.5f*Qy[1][j+1][i],
Q[3][l+1][k] - 0.5f*Qy[3][j+1][i]);
const float4 Q_rr = make_float4(Q[0][l+1][k] + 0.5f*Qy[0][j+1][i],
Q[2][l+1][k] + 0.5f*Qy[2][j+1][i],
Q[1][l+1][k] + 0.5f*Qy[1][j+1][i],
Q[3][l+1][k] + 0.5f*Qy[3][j+1][i]);
const float4 Q_ll = make_float4(Q[0][l][k] - 0.5f*Qy[0][j][i],
Q[2][l][k] - 0.5f*Qy[2][j][i],
Q[1][l][k] - 0.5f*Qy[1][j][i],
Q[3][l][k] - 0.5f*Qy[3][j][i]);
const float4 Q_lr = make_float4(Q[0][l][k] + 0.5f*Qy[0][j][i],
Q[2][l][k] + 0.5f*Qy[2][j][i],
Q[1][l][k] + 0.5f*Qy[1][j][i],
Q[3][l][k] + 0.5f*Qy[3][j][i]);
//Evolve half a timestep (predictor step)
const float4 Q_r_bar = Q_rl + dt_/(2.0f*dy_) * (F_func(Q_rl, gamma_) - F_func(Q_rr, gamma_));
const float4 Q_l_bar = Q_lr + dt_/(2.0f*dy_) * (F_func(Q_ll, gamma_) - F_func(Q_lr, gamma_));
// Compute flux based on prediction
const float4 flux = CentralUpwindFlux(Q_l_bar, Q_r_bar, gamma_);
//Write to shared memory
//Note that we here swap hu and hv back to the original
G[0][j][i] = flux.x;
G[1][j][i] = flux.z;
G[2][j][i] = flux.y;
G[3][j][i] = flux.w;
}
}
@ -158,6 +146,7 @@ __global__ void KP07DimsplitKernel(
const unsigned int w = BLOCK_WIDTH;
const unsigned int h = BLOCK_HEIGHT;
const unsigned int gc = 2;
const unsigned int vars = 4;
//Shared memory variables
__shared__ float Q[4][h+4][w+4];
@ -167,10 +156,10 @@ __global__ void KP07DimsplitKernel(
//Read into shared memory
readBlock<w, h, gc>( rho0_ptr_, rho0_pitch_, Q[0], nx_+2, ny_+2);
readBlock<w, h, gc>(rho_u0_ptr_, rho_u0_pitch_, Q[1], nx_+2, ny_+2);
readBlock<w, h, gc>(rho_v0_ptr_, rho_v0_pitch_, Q[2], nx_+2, ny_+2);
readBlock<w, h, gc>( E0_ptr_, E0_pitch_, Q[3], nx_+2, ny_+2);
readBlock<w, h, gc>( rho0_ptr_, rho0_pitch_, Q[0], nx_+4, ny_+4);
readBlock<w, h, gc>(rho_u0_ptr_, rho_u0_pitch_, Q[1], nx_+4, ny_+4);
readBlock<w, h, gc>(rho_v0_ptr_, rho_v0_pitch_, Q[2], nx_+4, ny_+4);
readBlock<w, h, gc>( E0_ptr_, E0_pitch_, Q[3], nx_+4, ny_+4);
__syncthreads();
@ -181,40 +170,47 @@ __global__ void KP07DimsplitKernel(
noFlowBoundary<w, h, gc, 1, 1>(Q[3], nx_, ny_);
__syncthreads();
//Step 0 => evolve x first, then y
if (step_ == 0) {
//Compute fluxes along the x axis and evolve
minmodSlopeX(Q, Qx, theta_);
minmodSlopeX<w, h, gc, vars>(Q, Qx, theta_);
__syncthreads();
computeFluxF(Q, Qx, F, gamma_, dx_, dt_);
__syncthreads();
evolveF2(Q, F, nx_, ny_, dx_, dt_);
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_);
noFlowBoundary<w, h, gc, 1, 1>(Q[3], nx_, ny_);
__syncthreads();
//Compute fluxes along the y axis and evolve
minmodSlopeY(Q, Qx, theta_);
minmodSlopeY<w, h, gc, vars>(Q, Qx, theta_);
__syncthreads();
computeFluxG(Q, Qx, F, gamma_, dy_, dt_);
__syncthreads();
evolveG2(Q, F, nx_, ny_, dy_, dt_);
evolveG<w, h, gc, vars>(Q, F, dy_, dt_);
__syncthreads();
}
//Step 1 => evolve y first, then x
else {
//Compute fluxes along the y axis and evolve
minmodSlopeY(Q, Qx, theta_);
minmodSlopeY<w, h, gc, vars>(Q, Qx, theta_);
__syncthreads();
computeFluxG(Q, Qx, F, gamma_, dy_, dt_);
__syncthreads();
evolveG2(Q, F, nx_, ny_, dy_, dt_);
evolveG<w, h, gc, vars>(Q, F, dy_, dt_);
__syncthreads();
//Set boundary conditions
@ -225,14 +221,14 @@ __global__ void KP07DimsplitKernel(
__syncthreads();
//Compute fluxes along the x axis and evolve
minmodSlopeX(Q, Qx, theta_);
minmodSlopeX<w, h, gc, vars>(Q, Qx, theta_);
__syncthreads();
computeFluxF(Q, Qx, F, gamma_, dx_, dt_);
__syncthreads();
evolveF2(Q, F, nx_, ny_, dx_, dt_);
evolveF<w, h, gc, vars>(Q, F, dx_, dt_);
__syncthreads();
}
// Write to main memory for all internal cells
writeBlock<w, h, gc>( rho1_ptr_, rho1_pitch_, Q[0], nx_, ny_);

6
GPUSimulators/cuda/EulerCommon.h
View File

@ -61,17 +61,17 @@ __device__ float4 F_func(const float4 Q, float P) {
/**
* Central upwind flux function
*/
__device__ float4 CentralUpwindFlux(const float4 Qm, float4 Qp, const float gamma) {
__device__ float4 CentralUpwindFlux(const float4 Qm, const float4 Qp, const float gamma) {
const float Pp = pressure(Qp, gamma);
const float4 Fp = F_func(Qp, Pp);
const float up = Qp.y / Qp.x; // rho*u / rho
const float cp = sqrt(gamma*Pp*Qp.x); // sqrt(gamma*P/rho)
const float cp = sqrt(gamma*Pp/Qp.x); // sqrt(gamma*P/rho)
const float Pm = pressure(Qm, gamma);
const float4 Fm = F_func(Qm, Pm);
const float um = Qm.y / Qm.x; // rho*u / rho
const float cm = sqrt(gamma*Pm/Qm.x); // sqrt(g*h)
const float cm = sqrt(gamma*Pm/Qm.x); // sqrt(gamma*P/rho)
const float am = min(min(um-cm, up-cp), 0.0f); // largest negative wave speed
const float ap = max(max(um+cm, up+cp), 0.0f); // largest positive wave speed

8
GPUSimulators/cuda/SWE2D_HLL2.cu
View File

@ -181,7 +181,7 @@ __global__ void HLL2Kernel(
//Step 0 => evolve x first, then y
if (step_ == 0) {
//Compute fluxes along the x axis and evolve
minmodSlopeX(Q, Qx, theta_);
minmodSlopeX<w, h, gc, vars>(Q, Qx, theta_);
__syncthreads();
computeFluxF(Q, Qx, F, g_, dx_, dt_);
__syncthreads();
@ -195,7 +195,7 @@ __global__ void HLL2Kernel(
__syncthreads();
//Compute fluxes along the y axis and evolve
minmodSlopeY(Q, Qx, theta_);
minmodSlopeY<w, h, gc, vars>(Q, Qx, theta_);
__syncthreads();
computeFluxG(Q, Qx, F, g_, dy_, dt_);
__syncthreads();
@ -205,7 +205,7 @@ __global__ void HLL2Kernel(
//Step 1 => evolve y first, then x
else {
//Compute fluxes along the y axis and evolve
minmodSlopeY(Q, Qx, theta_);
minmodSlopeY<w, h, gc, vars>(Q, Qx, theta_);
__syncthreads();
computeFluxG(Q, Qx, F, g_, dy_, dt_);
__syncthreads();
@ -219,7 +219,7 @@ __global__ void HLL2Kernel(
__syncthreads();
//Compute fluxes along the x axis and evolve
minmodSlopeX(Q, Qx, theta_);
minmodSlopeX<w, h, gc, vars>(Q, Qx, theta_);
__syncthreads();
computeFluxF(Q, Qx, F, g_, dx_, dt_);
__syncthreads();

4
GPUSimulators/cuda/SWE2D_KP07.cu
View File

@ -159,8 +159,8 @@ __global__ void KP07Kernel(
//Reconstruct slopes along x and axis
minmodSlopeX(Q, Qx, theta_);
minmodSlopeY(Q, Qy, theta_);
minmodSlopeX<w, h, gc, vars>(Q, Qx, theta_);
minmodSlopeY<w, h, gc, vars>(Q, Qy, theta_);
__syncthreads();

70
GPUSimulators/cuda/SWE2D_KP07_dimsplit.cu
View File

@ -38,39 +38,37 @@ void computeFluxF(float Q[3][BLOCK_HEIGHT+4][BLOCK_WIDTH+4],
const int tx = threadIdx.x;
const int ty = threadIdx.y;
{
int j=ty;
const int l = j + 2; //Skip ghost cells
for (int i=tx; i<BLOCK_WIDTH+1; i+=BLOCK_WIDTH) {
const int k = i + 1;
// Reconstruct point values of Q at the left and right hand side
// of the cell for both the left (i) and right (i+1) cell
const float3 Q_rl = make_float3(Q[0][l][k+1] - 0.5f*Qx[0][j][i+1],
Q[1][l][k+1] - 0.5f*Qx[1][j][i+1],
Q[2][l][k+1] - 0.5f*Qx[2][j][i+1]);
const float3 Q_rr = make_float3(Q[0][l][k+1] + 0.5f*Qx[0][j][i+1],
Q[1][l][k+1] + 0.5f*Qx[1][j][i+1],
Q[2][l][k+1] + 0.5f*Qx[2][j][i+1]);
const float3 Q_ll = make_float3(Q[0][l][k] - 0.5f*Qx[0][j][i],
Q[1][l][k] - 0.5f*Qx[1][j][i],
Q[2][l][k] - 0.5f*Qx[2][j][i]);
const float3 Q_lr = make_float3(Q[0][l][k] + 0.5f*Qx[0][j][i],
Q[1][l][k] + 0.5f*Qx[1][j][i],
Q[2][l][k] + 0.5f*Qx[2][j][i]);
//Evolve half a timestep (predictor step)
const float3 Q_r_bar = Q_rl + dt_/(2.0f*dx_) * (F_func(Q_rl, g_) - F_func(Q_rr, g_));
const float3 Q_l_bar = Q_lr + dt_/(2.0f*dx_) * (F_func(Q_ll, g_) - F_func(Q_lr, g_));
int j=ty;
const int l = j + 2; //Skip ghost cells
for (int i=tx; i<BLOCK_WIDTH+1; i+=BLOCK_WIDTH) {
const int k = i + 1;
// Reconstruct point values of Q at the left and right hand side
// of the cell for both the left (i) and right (i+1) cell
const float3 Q_rl = make_float3(Q[0][l][k+1] - 0.5f*Qx[0][j][i+1],
Q[1][l][k+1] - 0.5f*Qx[1][j][i+1],
Q[2][l][k+1] - 0.5f*Qx[2][j][i+1]);
const float3 Q_rr = make_float3(Q[0][l][k+1] + 0.5f*Qx[0][j][i+1],
Q[1][l][k+1] + 0.5f*Qx[1][j][i+1],
Q[2][l][k+1] + 0.5f*Qx[2][j][i+1]);
const float3 Q_ll = make_float3(Q[0][l][k] - 0.5f*Qx[0][j][i],
Q[1][l][k] - 0.5f*Qx[1][j][i],
Q[2][l][k] - 0.5f*Qx[2][j][i]);
const float3 Q_lr = make_float3(Q[0][l][k] + 0.5f*Qx[0][j][i],
Q[1][l][k] + 0.5f*Qx[1][j][i],
Q[2][l][k] + 0.5f*Qx[2][j][i]);
//Evolve half a timestep (predictor step)
const float3 Q_r_bar = Q_rl + dt_/(2.0f*dx_) * (F_func(Q_rl, g_) - F_func(Q_rr, g_));
const float3 Q_l_bar = Q_lr + dt_/(2.0f*dx_) * (F_func(Q_ll, g_) - F_func(Q_lr, g_));
// Compute flux based on prediction
const float3 flux = CentralUpwindFlux(Q_l_bar, Q_r_bar, g_);
//Write to shared memory
F[0][j][i] = flux.x;
F[1][j][i] = flux.y;
F[2][j][i] = flux.z;
}
// Compute flux based on prediction
const float3 flux = CentralUpwindFlux(Q_l_bar, Q_r_bar, g_);
//Write to shared memory
F[0][j][i] = flux.x;
F[1][j][i] = flux.y;
F[2][j][i] = flux.z;
}
}
@ -178,7 +176,7 @@ __global__ void KP07DimsplitKernel(
//Step 0 => evolve x first, then y
if (step_ == 0) {
//Compute fluxes along the x axis and evolve
minmodSlopeX(Q, Qx, theta_);
minmodSlopeX<w, h, gc, vars>(Q, Qx, theta_);
__syncthreads();
computeFluxF(Q, Qx, F, g_, dx_, dt_);
__syncthreads();
@ -194,7 +192,7 @@ __global__ void KP07DimsplitKernel(
//Compute fluxes along the y axis and evolve
minmodSlopeY(Q, Qx, theta_);
minmodSlopeY<w, h, gc, vars>(Q, Qx, theta_);
__syncthreads();
computeFluxG(Q, Qx, F, g_, dy_, dt_);
@ -205,7 +203,7 @@ __global__ void KP07DimsplitKernel(
//Step 1 => evolve y first, then x
else {
//Compute fluxes along the y axis and evolve
minmodSlopeY(Q, Qx, theta_);
minmodSlopeY<w, h, gc, vars>(Q, Qx, theta_);
__syncthreads();
computeFluxG(Q, Qx, F, g_, dy_, dt_);
__syncthreads();
@ -219,7 +217,7 @@ __global__ void KP07DimsplitKernel(
__syncthreads();
//Compute fluxes along the x axis and evolve
minmodSlopeX(Q, Qx, theta_);
minmodSlopeX<w, h, gc, vars>(Q, Qx, theta_);
__syncthreads();
computeFluxF(Q, Qx, F, g_, dx_, dt_);
__syncthreads();

15
GPUSimulators/cuda/common.h
View File

@ -293,6 +293,21 @@ __device__ void evolveG(float Q[vars][block_height+2*ghost_cells][block_width+2*
/**
* Helper function for debugging etc.
*/
template<int shmem_width, int shmem_height, int vars>
__device__ void memset(float Q[vars][shmem_height][shmem_width], float value) {
for (int k=0; k<vars; ++k) {
for (int j=threadIdx.y; j<shmem_height; ++j) {
for (int i=threadIdx.x; i<shmem_width; ++i) {
Q[k][j][i] = value;
}
}
}
}

39
GPUSimulators/cuda/limiters.h
View File

@ -46,22 +46,22 @@ __device__ __inline__ float minmodSlope(float left, float center, float right, f
/**
* Reconstructs a minmod slope for a whole block along the abscissa
*/
__device__ void minmodSlopeX(float Q[3][BLOCK_HEIGHT+4][BLOCK_WIDTH+4],
float Qx[3][BLOCK_HEIGHT+2][BLOCK_WIDTH+2],
template<int block_width, int block_height, int ghost_cells, int vars>
__device__ void minmodSlopeX(float Q[vars][block_height+2*ghost_cells][block_width+2*ghost_cells],
float Qx[vars][block_height+2*(ghost_cells-1)][block_width+2*(ghost_cells-1)],
const float theta_) {
//Index of thread within block
const int tx = threadIdx.x;
const int ty = threadIdx.y;
const int j = ty;
const int l = j + ghost_cells; //Skip ghost cells
//Reconstruct slopes along x axis
{
const int j = ty;
const int l = j + 2; //Skip ghost cells
for (int i=tx; i<BLOCK_WIDTH+2; i+=BLOCK_WIDTH) {
const int k = i + 1;
for (int p=0; p<3; ++p) {
Qx[p][j][i] = minmodSlope(Q[p][l][k-1], Q[p][l][k], Q[p][l][k+1], theta_);
}
for (int i=tx; i<block_width+2*(ghost_cells-1); i+=block_width) {
const int k = i + 1;
for (int p=0; p<vars; ++p) {
Qx[p][j][i] = minmodSlope(Q[p][l][k-1], Q[p][l][k], Q[p][l][k+1], theta_);
}
}
}
@ -70,21 +70,22 @@ __device__ void minmodSlopeX(float Q[3][BLOCK_HEIGHT+4][BLOCK_WIDTH+4],
/**
* Reconstructs a minmod slope for a whole block along the ordinate
*/
__device__ void minmodSlopeY(float Q[3][BLOCK_HEIGHT+4][BLOCK_WIDTH+4],
float Qy[3][BLOCK_HEIGHT+2][BLOCK_WIDTH+2],
template<int block_width, int block_height, int ghost_cells, int vars>
__device__ void minmodSlopeY(float Q[vars][block_height+2*ghost_cells][block_width+2*ghost_cells],
float Qy[vars][block_height+2*(ghost_cells-1)][block_width+2*(ghost_cells-1)],
const float theta_) {
//Index of thread within block
const int tx = threadIdx.x;
const int ty = threadIdx.y;
for (int j=ty; j<BLOCK_HEIGHT+2; j+=BLOCK_HEIGHT) {
const int i = tx;
const int k = i + ghost_cells; //Skip ghost cells
//Reconstruct slopes along y axis
for (int j=ty; j<block_height+2*(ghost_cells-1); j+=block_height) {
const int l = j + 1;
{
const int i = tx;
const int k = i + 2; //Skip ghost cells
for (int p=0; p<3; ++p) {
Qy[p][j][i] = minmodSlope(Q[p][l-1][k], Q[p][l][k], Q[p][l+1][k], theta_);
}
for (int p=0; p<vars; ++p) {
Qy[p][j][i] = minmodSlope(Q[p][l-1][k], Q[p][l][k], Q[p][l+1][k], theta_);
}
}
}