smyalygames/FiniteVolumeGPU
1
0
Fork 0
mirror of https://github.com/smyalygames/FiniteVolumeGPU.git synced 2026-01-14 15:48:43 +01:00
Code Issues Actions Packages Projects Releases Wiki Activity

Fixed order again

Browse Source
This commit is contained in:
André R. Brodtkorb
2018-11-15 16:47:13 +01:00
parent dcb849b705
commit 7592ad5b9f
22 changed files with 758 additions and 619 deletions
Download Patch File Download Diff File Expand all files Collapse all files

277
GPUSimulators/cuda/common.h
View File

@@ -137,19 +137,21 @@ inline __device__ BoundaryCondition getBCWest(int bc_) {
/**
* Alter the index l so that it gives periodic boundary conditions when reading
*/
template<int ghost_cells>
template<int gc_x>
inline __device__ int handlePeriodicBoundaryX(int k, int nx_, int boundary_conditions_) {
const int gc_pad = 2*ghost_cells;
const int gc_pad = gc_x;
//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);
if (gc_x > 0) {
if ((k < gc_pad)
&& getBCWest(boundary_conditions_) == Periodic) {
k += (nx_+2*gc_x - 2*gc_pad);
}
//East boundary: subtract an offset to read from west of domain
else if ((k >= nx_+2*gc_x-gc_pad)
&& getBCEast(boundary_conditions_) == Periodic) {
k -= (nx_+2*gc_x - 2*gc_pad);
}
}
return k;
@@ -158,45 +160,49 @@ inline __device__ int handlePeriodicBoundaryX(int k, int nx_, int boundary_condi
/**
* Alter the index l so that it gives periodic boundary conditions when reading
*/
template<int ghost_cells>
template<int gc_y>
inline __device__ int handlePeriodicBoundaryY(int l, int ny_, int boundary_conditions_) {
const int gc_pad = 2*ghost_cells;
const int gc_pad = gc_y;
//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);
if (gc_y > 0) {
if ((l < gc_pad)
&& getBCSouth(boundary_conditions_) == Periodic) {
l += (ny_+2*gc_y - 2*gc_pad);
}
//North boundary: subtract an offset to read from south of domain
else if ((l >= ny_+2*gc_y-gc_pad)
&& getBCNorth(boundary_conditions_) == Periodic) {
l -= (ny_+2*gc_y - 2*gc_pad);
}
}
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],
template<int w, int h, int gc_x, int gc_y, int sign_x, int sign_y>
inline __device__
void handleReflectiveBoundary(
float Q[h+2*gc_y][w+2*gc_x],
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_);
bcNorthReflective<w, h, gc_x, gc_y, sign_y>(Q, nx_, ny_);
__syncthreads();
}
if (getBCSouth(boundary_conditions_) == Reflective) {
bcSouthReflective<block_width, block_height, ghost_cells, sign_y>(Q, nx_, ny_);
bcSouthReflective<w, h, gc_x, gc_y, sign_y>(Q, nx_, ny_);
__syncthreads();
}
if (getBCEast(boundary_conditions_) == Reflective) {
bcEastReflective<block_width, block_height, ghost_cells, sign_x>(Q, nx_, ny_);
bcEastReflective<w, h, gc_x, gc_y, sign_x>(Q, nx_, ny_);
__syncthreads();
}
if (getBCWest(boundary_conditions_) == Reflective) {
bcWestReflective<block_width, block_height, ghost_cells, sign_x>(Q, nx_, ny_);
bcWestReflective<w, h, gc_x, gc_y, sign_x>(Q, nx_, ny_);
__syncthreads();
}
}
@@ -204,9 +210,9 @@ inline __device__ int handleReflectiveBoundary(
/**
* Reads a block of data with ghost cells
*/
template<int block_width, int block_height, int ghost_cells, int sign_x, int sign_y>
template<int w, int h, int gc_x, int gc_y, 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],
float Q[h+2*gc_y][w+2*gc_x],
const int nx_, const int ny_,
const int boundary_conditions_) {
//Index of block within domain
@@ -215,16 +221,16 @@ inline __device__ void readBlock(float* ptr_, int pitch_,
//Read into shared memory
//Loop over all variables
for (int j=threadIdx.y; j<block_height+2*ghost_cells; j+=block_height) {
for (int j=threadIdx.y; j<h+2*gc_y; j+=h) {
//Handle periodic boundary conditions here
int l = handlePeriodicBoundaryY<ghost_cells>(by + j, ny_, boundary_conditions_);
l = min(l, ny_+2*ghost_cells-1);
int l = handlePeriodicBoundaryY<gc_y>(by + j, ny_, boundary_conditions_);
l = min(l, ny_+2*gc_y-1);
float* row = (float*) ((char*) ptr_ + pitch_*l);
for (int i=threadIdx.x; i<block_width+2*ghost_cells; i+=block_width) {
for (int i=threadIdx.x; i<w+2*gc_x; i+=w) {
//Handle periodic boundary conditions here
int k = handlePeriodicBoundaryX<ghost_cells>(bx + i, nx_, boundary_conditions_);
k = min(k, nx_+2*ghost_cells-1);
int k = handlePeriodicBoundaryX<gc_x>(bx + i, nx_, boundary_conditions_);
k = min(k, nx_+2*gc_x-1);
//Read from global memory
Q[j][i] = row[k];
@@ -232,7 +238,7 @@ inline __device__ void readBlock(float* ptr_, int pitch_,
}
__syncthreads();
handleReflectiveBoundary<block_width, block_height, ghost_cells, sign_x, sign_y>(Q, nx_, ny_, boundary_conditions_);
handleReflectiveBoundary<w, h, gc_x, gc_y, sign_x, sign_y>(Q, nx_, ny_, boundary_conditions_);
}
@@ -241,45 +247,68 @@ inline __device__ void readBlock(float* ptr_, int pitch_,
/**
* Writes a block of data to global memory for the shallow water equations.
*/
template<int block_width, int block_height, int ghost_cells>
template<int w, int h, int gc_x, int gc_y>
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,
float shmem[h+2*gc_y][w+2*gc_x],
const int nx_, const int ny_,
int rk_step_, int rk_order_) {
//Index of cell within domain
const int ti = blockDim.x*blockIdx.x + threadIdx.x + ghost_cells;
const int tj = blockDim.y*blockIdx.y + threadIdx.y + ghost_cells;
const int ti = blockDim.x*blockIdx.x + threadIdx.x + gc_x;
const int tj = blockDim.y*blockIdx.y + threadIdx.y + gc_y;
//Only write internal cells
if (ti < width+ghost_cells && tj < height+ghost_cells) {
if (ti < nx_+gc_x && tj < ny_+gc_y) {
//Index of thread within block
const int tx = threadIdx.x + ghost_cells;
const int ty = threadIdx.y + ghost_cells;
const int tx = threadIdx.x + gc_x;
const int ty = threadIdx.y + gc_y;
float* const row = (float*) ((char*) ptr_ + pitch_*tj);
//Handle runge-kutta timestepping here
row[ti] = shmem[ty][tx];
/**
* SSPRK1 (forward Euler)
* u^1 = u^n + dt*f(u^n)
*/
if (rk_order_ == 1) {
row[ti] = shmem[ty][tx];
}
/**
* SSPRK2
* u^1 = u^n + dt*f(u^n)
* u^n+1 = 1/2*u^n + 1/2*(u^1 + dt*f(u^1))
*
*/
else if (rk_order_ == 2) {
if (rk_step_ == 0) {
row[ti] = shmem[ty][tx];
}
else if (rk_step_ == 1) {
row[ti] = 0.5f*row[ti] + 0.5f*shmem[ty][tx];
}
}
/**
* SSPRK3
* u^1 = u^n + dt*f(u^n)
* u^2 = 3/4 * u^n + 1/4 * (u^1 + dt*f(u^1))
* u^n+1 = 1/3 * u^n + 2/3 * (u^2 + dt*f(u^2))
* FIXME: This is not correct now, need a temporary to hold intermediate step u^2
*/
/*
if (rk_order_ == 2 && rk_step_ == 1) {
row[ti] = 0.5f*(row[ti] + shmem[ty][tx]);
else if (rk_order_ == 3) {
if (rk_step_ == 0) {
row[ti] = shmem[ty][tx];
}
else if (rk_step_ == 1) {
row[ti] = 0.75f*row[ti] + 0.25f*shmem[ty][tx];
}
else if (rk_step_ == 2) {
const float t = 1.0f / 3.0f; //Not representable in base 2
row[ti] = t*row[ti] + (1.0f-t)*shmem[ty][tx];
}
}
else {
row[ti] = shmem[ty][tx];
}*/
}
}
@@ -297,25 +326,26 @@ 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_) {
for (int j=threadIdx.y; j<block_height+2*ghost_cells; j+= block_height) {
const int i = threadIdx.x + ghost_cells;
template<int w, int h, int gc_x, int gc_y, int sign>
__device__ void bcWestReflective(float Q[h+2*gc_y][w+2*gc_x],
const int nx_, const int ny_) {
for (int j=threadIdx.y; j<h+2*gc_y; j+=h) {
const int i = threadIdx.x + gc_x;
const int ti = blockDim.x*blockIdx.x + i;
if (ti == ghost_cells) {
if (gc_x >= 1 && ti == gc_x) {
Q[j][i-1] = sign*Q[j][i];
}
if (ghost_cells >= 2 && ti == ghost_cells + 1) {
if (gc_x >= 2 && ti == gc_x + 1) {
Q[j][i-3] = sign*Q[j][i];
}
if (ghost_cells >= 3 && ti == ghost_cells + 2) {
if (gc_x >= 3 && ti == gc_x + 2) {
Q[j][i-5] = sign*Q[j][i];
}
if (ghost_cells >= 4 && ti == ghost_cells + 3) {
if (gc_x >= 4 && ti == gc_x + 3) {
Q[j][i-7] = sign*Q[j][i];
}
if (ghost_cells >= 5 && ti == ghost_cells + 4) {
if (gc_x >= 5 && ti == gc_x + 4) {
Q[j][i-9] = sign*Q[j][i];
}
}
@@ -323,25 +353,26 @@ __device__ void bcWestReflective(float Q[block_height+2*ghost_cells][block_width
// East boundary
template<int block_width, int block_height, int ghost_cells, int sign>
__device__ void bcEastReflective(float Q[block_height+2*ghost_cells][block_width+2*ghost_cells], const int nx_, const int ny_) {
for (int j=threadIdx.y; j<block_height+2*ghost_cells; j+= block_height) {
const int i = threadIdx.x + ghost_cells;
template<int w, int h, int gc_x, int gc_y, int sign>
__device__ void bcEastReflective(float Q[h+2*gc_y][w+2*gc_x],
const int nx_, const int ny_) {
for (int j=threadIdx.y; j<h+2*gc_y; j+=h) {
const int i = threadIdx.x + gc_x;
const int ti = blockDim.x*blockIdx.x + i;
if (ti == nx_ + ghost_cells - 1) {
if (gc_x >= 1 && ti == nx_ + gc_x - 1) {
Q[j][i+1] = sign*Q[j][i];
}
if (ghost_cells >= 2 && ti == nx_ + ghost_cells - 2) {
if (gc_x >= 2 && ti == nx_ + gc_x - 2) {
Q[j][i+3] = sign*Q[j][i];
}
if (ghost_cells >= 3 && ti == nx_ + ghost_cells - 3) {
if (gc_x >= 3 && ti == nx_ + gc_x - 3) {
Q[j][i+5] = sign*Q[j][i];
}
if (ghost_cells >= 4 && ti == nx_ + ghost_cells - 4) {
if (gc_x >= 4 && ti == nx_ + gc_x - 4) {
Q[j][i+7] = sign*Q[j][i];
}
if (ghost_cells >= 5 && ti == nx_ + ghost_cells - 5) {
if (gc_x >= 5 && ti == nx_ + gc_x - 5) {
Q[j][i+9] = sign*Q[j][i];
}
}
@@ -349,25 +380,26 @@ __device__ void bcEastReflective(float Q[block_height+2*ghost_cells][block_width
// South boundary
template<int block_width, int block_height, int ghost_cells, int sign>
__device__ void bcSouthReflective(float Q[block_height+2*ghost_cells][block_width+2*ghost_cells], const int nx_, const int ny_) {
for (int i=threadIdx.x; i<block_width+2*ghost_cells; i+= block_width) {
const int j = threadIdx.y + ghost_cells;
template<int w, int h, int gc_x, int gc_y, int sign>
__device__ void bcSouthReflective(float Q[h+2*gc_y][w+2*gc_x],
const int nx_, const int ny_) {
for (int i=threadIdx.x; i<w+2*gc_x; i+=w) {
const int j = threadIdx.y + gc_y;
const int tj = blockDim.y*blockIdx.y + j;
if (tj == ghost_cells) {
if (gc_y >= 1 && tj == gc_y) {
Q[j-1][i] = sign*Q[j][i];
}
if (ghost_cells >= 2 && tj == ghost_cells + 1) {
if (gc_y >= 2 && tj == gc_y + 1) {
Q[j-3][i] = sign*Q[j][i];
}
if (ghost_cells >= 3 && tj == ghost_cells + 2) {
if (gc_y >= 3 && tj == gc_y + 2) {
Q[j-5][i] = sign*Q[j][i];
}
if (ghost_cells >= 4 && tj == ghost_cells + 3) {
if (gc_y >= 4 && tj == gc_y + 3) {
Q[j-7][i] = sign*Q[j][i];
}
if (ghost_cells >= 5 && tj == ghost_cells + 4) {
if (gc_y >= 5 && tj == gc_y + 4) {
Q[j-9][i] = sign*Q[j][i];
}
}
@@ -377,25 +409,25 @@ __device__ void bcSouthReflective(float Q[block_height+2*ghost_cells][block_widt
// North boundary
template<int block_width, int block_height, int ghost_cells, int sign>
__device__ void bcNorthReflective(float Q[block_height+2*ghost_cells][block_width+2*ghost_cells], const int nx_, const int ny_) {
for (int i=threadIdx.x; i<block_width+2*ghost_cells; i+= block_width) {
const int j = threadIdx.y + ghost_cells;
template<int w, int h, int gc_x, int gc_y, int sign>
__device__ void bcNorthReflective(float Q[h+2*gc_y][w+2*gc_x], const int nx_, const int ny_) {
for (int i=threadIdx.x; i<w+2*gc_x; i+=w) {
const int j = threadIdx.y + gc_y;
const int tj = blockDim.y*blockIdx.y + j;
if (tj == ny_ + ghost_cells - 1) {
if (gc_y >= 1 && tj == ny_ + gc_y - 1) {
Q[j+1][i] = sign*Q[j][i];
}
if (ghost_cells >= 2 && tj == ny_ + ghost_cells - 2) {
if (gc_y >= 2 && tj == ny_ + gc_y - 2) {
Q[j+3][i] = sign*Q[j][i];
}
if (ghost_cells >= 3 && tj == ny_ + ghost_cells - 3) {
if (gc_y >= 3 && tj == ny_ + gc_y - 3) {
Q[j+5][i] = sign*Q[j][i];
}
if (ghost_cells >= 4 && tj == ny_ + ghost_cells - 4) {
if (gc_y >= 4 && tj == ny_ + gc_y - 4) {
Q[j+7][i] = sign*Q[j][i];
}
if (ghost_cells >= 5 && tj == ny_ + ghost_cells - 5) {
if (gc_y >= 5 && tj == ny_ + gc_y - 5) {
Q[j+9][i] = sign*Q[j][i];
}
}
@@ -422,13 +454,13 @@ __device__ void bcNorthReflective(float Q[block_height+2*ghost_cells][block_widt
template<int block_width, int block_height, int ghost_cells, int vars>
__device__ void evolveF(float Q[vars][block_height+2*ghost_cells][block_width+2*ghost_cells],
float F[vars][block_height+2*ghost_cells][block_width+2*ghost_cells],
template<int w, int h, int gc_x, int gc_y, int vars>
__device__ void evolveF(float Q[vars][h+2*gc_y][w+2*gc_x],
float F[vars][h+2*gc_y][w+2*gc_x],
const float dx_, const float dt_) {
for (int var=0; var < vars; ++var) {
for (int j=threadIdx.y; j<block_height+2*ghost_cells; j+=block_height) {
for (int i=threadIdx.x+ghost_cells; i<block_width+ghost_cells; i+=block_width) {
for (int j=threadIdx.y; j<h+2*gc_y; j+=h) {
for (int i=threadIdx.x+gc_x; i<w+gc_x; i+=w) {
Q[var][j][i] = Q[var][j][i] + (F[var][j][i-1] - F[var][j][i]) * dt_ / dx_;
}
}
@@ -443,13 +475,13 @@ __device__ void evolveF(float Q[vars][block_height+2*ghost_cells][block_width+2*
/**
* Evolves the solution in time along the y axis (dimensional splitting)
*/
template<int block_width, int block_height, int ghost_cells, int vars>
__device__ void evolveG(float Q[vars][block_height+2*ghost_cells][block_width+2*ghost_cells],
float G[vars][block_height+2*ghost_cells][block_width+2*ghost_cells],
template<int w, int h, int gc_x, int gc_y, int vars>
__device__ void evolveG(float Q[vars][h+2*gc_y][w+2*gc_x],
float G[vars][h+2*gc_y][w+2*gc_x],
const float dy_, const float dt_) {
for (int var=0; var < vars; ++var) {
for (int j=threadIdx.y+ghost_cells; j<block_height+ghost_cells; j+=block_height) {
for (int i=threadIdx.x; i<block_width+2*ghost_cells; i+=block_width) {
for (int j=threadIdx.y+gc_y; j<h+gc_y; j+=h) {
for (int i=threadIdx.x; i<w+2*gc_x; i+=w) {
Q[var][j][i] = Q[var][j][i] + (G[var][j-1][i] - G[var][j][i]) * dt_ / dy_;
}
}
@@ -478,6 +510,55 @@ __device__ void memset(float Q[vars][shmem_height][shmem_width], float value) {
template <unsigned int threads>
__device__ void reduce_max(float* data, unsigned int n) {
__shared__ float sdata[threads];
unsigned int tid = threadIdx.x;
//Reduce to "threads" elements
sdata[tid] = FLT_MIN;
for (unsigned int i=tid; i<n; i += threads) {
sdata[tid] = max(sdata[tid], dt_ctx.L[i]);
}
__syncthreads();
//Now, reduce all elements into a single element
if (threads >= 512) {
if (tid < 256) {
sdata[tid] = max(sdata[tid], sdata[tid + 256]);
}
__syncthreads();
}
if (threads >= 256) {
if (tid < 128) {
sdata[tid] = max(sdata[tid], sdata[tid + 128]);
}
__syncthreads();
}
if (threads >= 128) {
if (tid < 64) {
sdata[tid] = max(sdata[tid], sdata[tid + 64]);
}
__syncthreads();
}
if (tid < 32) {
volatile float* sdata_volatile = sdata;
if (threads >= 64) {
sdata_volatile[tid] = max(sdata_volatile[tid], sdata_volatile[tid + 32]);
}
if (tid < 16) {
if (threads >= 32) sdata_volatile[tid] = max(sdata_volatile[tid], sdata_volatile[tid + 16]);
if (threads >= 16) sdata_volatile[tid] = max(sdata_volatile[tid], sdata_volatile[tid + 8]);
if (threads >= 8) sdata_volatile[tid] = max(sdata_volatile[tid], sdata_volatile[tid + 4]);
if (threads >= 4) sdata_volatile[tid] = max(sdata_volatile[tid], sdata_volatile[tid + 2]);
if (threads >= 2) sdata_volatile[tid] = max(sdata_volatile[tid], sdata_volatile[tid + 1]);
}
if (tid == 0) {
return sdata_volatile[0];
}
}
}