Major update: functional LBM and DRL
This commit is contained in:
Frank14f
@@ -0,0 +1,101 @@
|
||||
#include "macros.h"
|
||||
#include "const.h"
|
||||
|
||||
__device__ void Index_lattice(int &x, int &y, int &k) {
|
||||
// Only for D2
|
||||
x = threadIdx.x + NT * blockIdx.x;
|
||||
y = blockIdx.y;
|
||||
k = y * NX + x;
|
||||
}
|
||||
|
||||
__device__ void CollisionKernel(LBtype* g, LBtype* m) {
|
||||
// Only for D2Q9
|
||||
LBtype p, u, v;
|
||||
LBtype niu = 1.0 / (0.5 + 3 * VIS);
|
||||
|
||||
u = (g[1]+g[5]+g[8]-g[3]-g[6]-g[7])/RHO;
|
||||
v = (g[2]+g[5]+g[6]-g[4]-g[7]-g[8])/RHO;
|
||||
p = (g[0]+g[1]+g[2]+g[3]+g[4]+g[5]+g[6]+g[7]+g[8])/3.0;
|
||||
|
||||
m[0]= g[0] +g[1] +g[2] +g[3] +g[4] +g[5] +g[6] +g[7] +g[8];
|
||||
m[1]=-4*g[0] -g[1] -g[2] -g[3] -g[4]+2*g[5]+2*g[6]+2*g[7]+2*g[8];
|
||||
m[2]= 4*g[0]-2*g[1]-2*g[2]-2*g[3]-2*g[4] +g[5] +g[6] +g[7] +g[8];
|
||||
m[3]= g[1] -g[3] +g[5] -g[6] -g[7] +g[8];
|
||||
m[4]= -2*g[1] +2*g[3] +g[5] -g[6] -g[7] +g[8];
|
||||
m[5]= g[2] -g[4] +g[5] +g[6] -g[7] -g[8];
|
||||
m[6]= -2*g[2] +2*g[4] +g[5] +g[6] -g[7] -g[8];
|
||||
m[7]= g[1] -g[2] +g[3] -g[4];
|
||||
m[8]= g[5] -g[6] +g[7] -g[8];
|
||||
|
||||
m[0]=1.00*( 3*p -m[0]);
|
||||
m[1]=1.20*(-6*p +3*RHO*(u*u+v*v)-m[1]);
|
||||
m[2]=1.20*( 3*p -3*RHO*(u*u+v*v)-m[2]);
|
||||
m[3]=1.00*( RHO*u -m[3]);
|
||||
m[4]=1.20*(-RHO*u -m[4]);
|
||||
m[5]=1.00*( RHO*v -m[5]);
|
||||
m[6]=1.20*(-RHO*v -m[6]);
|
||||
m[7]= niu*( RHO*(u*u-v*v) -m[7]);
|
||||
m[8]= niu*( RHO*u*v -m[8]);
|
||||
|
||||
g[0]=g[0]+( m[0] -m[1] +m[2] )/ 9.0;
|
||||
g[1]=g[1]+(4*m[0] -m[1]-2*m[2]+6*m[3]-6*m[4] +9*m[7])/36.0;
|
||||
g[2]=g[2]+(4*m[0] -m[1]-2*m[2] +6*m[5]-6*m[6]-9*m[7])/36.0;
|
||||
g[3]=g[3]+(4*m[0] -m[1]-2*m[2]-6*m[3]+6*m[4] +9*m[7])/36.0;
|
||||
g[4]=g[4]+(4*m[0] -m[1]-2*m[2] -6*m[5]+6*m[6]-9*m[7])/36.0;
|
||||
g[5]=g[5]+(4*m[0]+2*m[1] +m[2]+6*m[3]+3*m[4]+6*m[5]+3*m[6]+9*m[8])/36.0;
|
||||
g[6]=g[6]+(4*m[0]+2*m[1] +m[2]-6*m[3]-3*m[4]+6*m[5]+3*m[6]-9*m[8])/36.0;
|
||||
g[7]=g[7]+(4*m[0]+2*m[1] +m[2]-6*m[3]-3*m[4]-6*m[5]-3*m[6]+9*m[8])/36.0;
|
||||
g[8]=g[8]+(4*m[0]+2*m[1] +m[2]+6*m[3]+3*m[4]-6*m[5]-3*m[6]-9*m[8])/36.0;
|
||||
}
|
||||
|
||||
__device__ void ParabolicInlet(LBtype* f, LBtype* f_neb, LBtype y) {
|
||||
LBtype p, u, v, yy;
|
||||
LBtype feq1, feq5, feq8, feqn1, feqn5, feqn8;
|
||||
|
||||
p=(f_neb[0]+f_neb[1]+f_neb[2]+f_neb[3]+f_neb[4]+f_neb[5]+f_neb[6]+f_neb[7]+f_neb[8])/3.0;
|
||||
yy=(y-0.5*(NY-1))/(NY-2.0);
|
||||
u=U0*1.5*(1-4*yy*yy);
|
||||
v=0.0;
|
||||
|
||||
feq1=(2*p+RHO*(2*u*u+2*u -v*v) )/ 6.0;
|
||||
feq5=( p+RHO*( u*u+3*u*v+u+v*v+v))/12.0;
|
||||
feq8=( p+RHO*( u*u-3*u*v+u+v*v-v))/12.0;
|
||||
|
||||
u=(f_neb[1]+f_neb[5]+f_neb[8]-f_neb[3]-f_neb[6]-f_neb[7])/RHO;
|
||||
v=(f_neb[2]+f_neb[5]+f_neb[6]-f_neb[4]-f_neb[7]-f_neb[8])/RHO;
|
||||
|
||||
feqn1=(2*p+RHO*(2*u*u+2*u -v*v) )/ 6.0;
|
||||
feqn5=( p+RHO*( u*u+3*u*v+u+v*v+v))/12.0;
|
||||
feqn8=( p+RHO*( u*u-3*u*v+u+v*v-v))/12.0;
|
||||
|
||||
f[1]=f_neb[1]-feqn1+feq1;
|
||||
f[5]=f_neb[5]-feqn5+feq5;
|
||||
f[8]=f_neb[8]-feqn8+feq8;
|
||||
}
|
||||
|
||||
__device__ void PressureOutlet(LBtype* f, LBtype* f_neb, LBtype y) {
|
||||
// Edit to Parabolic Outlet temporarily
|
||||
LBtype p, u, v, yy;
|
||||
LBtype feq3, feq6, feq7, feqn3, feqn6, feqn7;
|
||||
|
||||
p=0.0;
|
||||
|
||||
yy=(y-0.5*(NY-1))/(NY-2.0);
|
||||
u=U0*1.5*(1-4*yy*yy);
|
||||
v=0.0;
|
||||
|
||||
feq3=(2*p-RHO*(-2*u*u+2*u +v*v) )/ 6.0;
|
||||
feq6=( p+RHO*( u*u-3*u*v-u+v*v+v))/12.0;
|
||||
feq7=( p+RHO*( u*u+3*u*v-u+v*v-v))/12.0;
|
||||
|
||||
u=(f_neb[1]+f_neb[5]+f_neb[8]-f_neb[3]-f_neb[6]-f_neb[7])/RHO;
|
||||
v=(f_neb[2]+f_neb[5]+f_neb[6]-f_neb[4]-f_neb[7]-f_neb[8])/RHO;
|
||||
// p=(f_neb[0]+f_neb[1]+f_neb[2]+f_neb[3]+f_neb[4]+f_neb[5]+f_neb[6]+f_neb[7]+f_neb[8])/3.0;
|
||||
feqn3=(2*p-RHO*(-2*u*u+2*u +v*v) )/ 6.0;
|
||||
feqn6=( p+RHO*( u*u-3*u*v-u+v*v+v))/12.0;
|
||||
feqn7=( p+RHO*( u*u+3*u*v-u+v*v-v))/12.0;
|
||||
|
||||
f[3]=f_neb[3]-feqn3+feq3;
|
||||
f[6]=f_neb[6]-feqn6+feq6;
|
||||
f[7]=f_neb[7]-feqn7+feq7;
|
||||
}
|
||||
Binary file not shown.
@@ -0,0 +1,10 @@
|
||||
// CelerisLab/kernels/const.h
|
||||
|
||||
#ifndef CONST_H
|
||||
#define CONST_H
|
||||
|
||||
__constant__ int e[9][2] = {{0, 0}, {1, 0}, {0, 1}, {-1, 0}, {0, -1}, {1, 1}, {-1, 1}, {-1, -1}, {1, -1}};
|
||||
__constant__ int opp[9] = {0, 3, 4, 1, 2, 7, 8, 5, 6};
|
||||
__constant__ float w[9] = {4/9., 1/9., 1/9., 1/9., 1/9., 1/36., 1/36., 1/36., 1/36.};
|
||||
|
||||
#endif
|
||||
@@ -0,0 +1,190 @@
|
||||
// CelerisLab/kernels/kernel.cu
|
||||
|
||||
#include <stdio.h>
|
||||
#include <stdint.h>
|
||||
#include <cuda.h>
|
||||
|
||||
#include "macros.h"
|
||||
#include "const.h"
|
||||
#include "D2Q9.cu"
|
||||
|
||||
extern "C"
|
||||
{
|
||||
__global__ void OneStep(uint8_t *flag, LBtype *f, LBtype *f_temp, int32_t *indx, LBtype *delta, LBtype *action, LBtype *obs)
|
||||
{
|
||||
__shared__ LBtype f_share[NT * NQ];
|
||||
__shared__ LBtype obs_share[(N_OBJS * DIM > 0) ? N_OBJS * DIM : 1];
|
||||
|
||||
int x, y, k;
|
||||
LBtype g[NQ], m[NQ];
|
||||
Index_lattice(x, y, k); // Only for D2
|
||||
int totalCells = NX * NY;
|
||||
int id = indx[k];
|
||||
|
||||
for (int i = 0; i < NQ; i++)
|
||||
{
|
||||
f_share[threadIdx.x + i * NT] = f[k + i * totalCells];
|
||||
}
|
||||
for (int i = threadIdx.x; i < N_OBJS * DIM; i+=NT)
|
||||
{
|
||||
obs_share[i] = 0;
|
||||
}
|
||||
|
||||
__syncthreads();
|
||||
|
||||
for (int i = 0; i < NQ; i++)
|
||||
{
|
||||
g[i] = f_share[threadIdx.x + i * NT];
|
||||
}
|
||||
|
||||
if (flag[k] & FLUID)
|
||||
{
|
||||
CollisionKernel(g, m);
|
||||
|
||||
for (int i = 0; i < NQ; i++)
|
||||
{
|
||||
f_share[threadIdx.x + i * NT] = g[i];
|
||||
}
|
||||
}
|
||||
else if (flag[k] & SOLID)
|
||||
{
|
||||
if (x == 0)
|
||||
{
|
||||
for (int i = 0; i < NQ; i++)
|
||||
{
|
||||
m[i] = f_share[threadIdx.x + i * NT + 1];
|
||||
}
|
||||
ParabolicInlet(g, m, y);
|
||||
}
|
||||
else if (x == NX - 1)
|
||||
{
|
||||
for (int i = 0; i < NQ; i++)
|
||||
{
|
||||
m[i] = f_share[threadIdx.x + i * NT - 1];
|
||||
}
|
||||
PressureOutlet(g, m, y);
|
||||
}
|
||||
|
||||
for (int i = 0; i < NQ; i++)
|
||||
{
|
||||
f_share[threadIdx.x + i * NT] = g[i];
|
||||
}
|
||||
}
|
||||
|
||||
__syncthreads();
|
||||
|
||||
for (int i = 0; i < NQ; i++)
|
||||
{
|
||||
int x_neb = x + e[i][0];
|
||||
int y_neb = y + e[i][1];
|
||||
|
||||
if (y != 0 && y != NY - 1)
|
||||
{
|
||||
if ((y == 1 && y_neb == 0) || (y == NY - 2 && y_neb == NY - 1))
|
||||
{
|
||||
f_temp[k + opp[i] * totalCells] = f_share[threadIdx.x + i * NT];
|
||||
}
|
||||
else
|
||||
{
|
||||
int k_neb = ((y_neb * NX + x_neb) + totalCells) % totalCells;
|
||||
f_temp[k_neb + i * totalCells] = f_share[threadIdx.x + i * NT];
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
__syncthreads();
|
||||
|
||||
if (flag[k] & SOLID && flag[k] & INTERFACE)
|
||||
{
|
||||
LBtype Uw, Vw;
|
||||
int id_obj = *reinterpret_cast<int*>(&delta[id]);
|
||||
Uw = action[id_obj] * delta[id + 9];
|
||||
Vw = action[id_obj] * delta[id + 10];
|
||||
|
||||
int x_neb, y_neb, k_neb;
|
||||
for (int i = 1; i < 9; i++)
|
||||
{
|
||||
x_neb = x + e[i][0];
|
||||
y_neb = y + e[i][1];
|
||||
k_neb = x_neb + y_neb * NX;
|
||||
if (flag[k_neb] & FLUID)
|
||||
{
|
||||
LBtype q = delta[id + i];
|
||||
int k_neb2 = (y + 2 * e[i][1]) * NX + (x + 2 * e[i][0]);
|
||||
LBtype temp = 6 * w[i] * (e[i][0] * Uw + e[i][1] * Vw);
|
||||
f_temp[k_neb + i * totalCells] = (q * f_temp[k + opp[i] * totalCells] \
|
||||
+ (1 - q) * f_temp[k_neb + opp[i] * totalCells] \
|
||||
+ q * f_temp[k_neb2 + i * totalCells] + temp) / (1 + q);
|
||||
f_temp[k + i * totalCells] = temp * Uw;
|
||||
k_neb2 = (y - e[i][1]) * NX + (x - e[i][0]);
|
||||
f_temp[k_neb2 + i * totalCells] = temp * Vw;
|
||||
|
||||
temp = f_temp[k_neb + i * totalCells] + f_temp[k + opp[i] * totalCells];
|
||||
k_neb2 = (y - e[i][1]) * NX + (x - e[i][0]);
|
||||
atomicAdd(&obs_share[DIM * id_obj], -temp * e[i][0] + f_temp[k + i * totalCells]);
|
||||
atomicAdd(&obs_share[DIM * id_obj + 1], -temp * e[i][1] + f_temp[k_neb2 + i * totalCells]);
|
||||
}
|
||||
}
|
||||
}
|
||||
if (flag[k] & SENSOR)
|
||||
{
|
||||
LBtype u, v;
|
||||
u = (g[1]+g[5]+g[8]-g[3]-g[6]-g[7])/RHO;
|
||||
v = (g[2]+g[5]+g[6]-g[4]-g[7]-g[8])/RHO;
|
||||
atomicAdd(&obs_share[DIM * id], u);
|
||||
atomicAdd(&obs_share[DIM * id + 1], v);
|
||||
}
|
||||
|
||||
__syncthreads();
|
||||
|
||||
for (int i = threadIdx.x; i < N_OBJS * DIM; i+=NT)
|
||||
{
|
||||
atomicAdd(&obs[i], obs_share[i]);
|
||||
}
|
||||
}
|
||||
|
||||
__global__ void InitTubeFlow(uint8_t *flag, LBtype *f)
|
||||
{
|
||||
__shared__ LBtype f_share[NT * NQ];
|
||||
__shared__ uint8_t flag_share[NT];
|
||||
int x, y, k;
|
||||
LBtype u;
|
||||
Index_lattice(x, y, k);
|
||||
int totalCells = NX * NY;
|
||||
|
||||
flag_share[threadIdx.x] = flag[k];
|
||||
for (int i = 0; i < NQ; i++)
|
||||
{
|
||||
f_share[threadIdx.x + i * NT] = f[k + i * totalCells];
|
||||
}
|
||||
|
||||
__syncthreads();
|
||||
|
||||
u = U0 * 1.5 * (1 - 4 * (y - 0.5 * (NY - 1)) * (y - 0.5 * (NY - 1)) / ((NY - 2) * (NY - 2)));
|
||||
if (y == 0 || y == NY - 1 || x == 0 || x == NX - 1)
|
||||
{
|
||||
flag_share[threadIdx.x] = SOLID;
|
||||
for (int i = 0; i < NQ; i++)
|
||||
{
|
||||
f_share[threadIdx.x + i * NT] = 0;
|
||||
}
|
||||
}
|
||||
else
|
||||
{
|
||||
flag_share[threadIdx.x] = FLUID;
|
||||
for (int i = 0; i < NQ; i++)
|
||||
{
|
||||
f_share[threadIdx.x + i * NT] = w[i] * RHO * (3 * e[i][0] * u + \
|
||||
4.5 * e[i][0] * e[i][0] * u * u - 1.5 * u * u);
|
||||
}
|
||||
}
|
||||
|
||||
__syncthreads();
|
||||
|
||||
flag[k] = flag_share[threadIdx.x];
|
||||
for (int i = 0; i < NQ; i++)
|
||||
{
|
||||
f[k + i * totalCells] = f_share[threadIdx.x + i * NT];
|
||||
}
|
||||
}
|
||||
}
|
||||
File diff suppressed because it is too large
Load Diff
@@ -0,0 +1,34 @@
|
||||
// CelerisLab/kernels/macros.h
|
||||
|
||||
// cuda parameters
|
||||
#define MULT_GPU False
|
||||
#define NT 128
|
||||
#define X_1U 128
|
||||
#define Y_1U 32
|
||||
#define Z_1U 1
|
||||
|
||||
// flow parameters
|
||||
#define LBtype float
|
||||
#define UX 10
|
||||
#define UY 16
|
||||
#define UZ 1
|
||||
#define NX 1280
|
||||
#define NY 512
|
||||
#define NZ 1
|
||||
#define DIM 2
|
||||
#define NQ 9
|
||||
#define VIS 0.004
|
||||
#define RHO 1.0
|
||||
#define U0 0.01
|
||||
|
||||
// constants
|
||||
#define PI 3.141592653589793238
|
||||
#define FLUID 0b00000001
|
||||
#define SOLID 0b00000010
|
||||
#define GAS 0b00000100
|
||||
#define INTERFACE 0b00001000
|
||||
#define SENSOR 0b00010000
|
||||
|
||||
// variables
|
||||
#define N_OBJS 7
|
||||
// #define N_SENS 2
|
||||
@@ -0,0 +1,2 @@
|
||||
#include "macros.h"
|
||||
#include "const.h"
|
||||
Reference in New Issue
Block a user