Frank14f/DynamisLab
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

第一轮分析工作暂存

Browse Source
This commit is contained in:
Frank14f
2026-06-09 18:46:59 +08:00
parent 4612928398
commit d1b9922c6b
111 changed files with 24287 additions and 1581 deletions
Download Patch File Download Diff File Expand all files Collapse all files
LegacyCelerisLab/kernels/D2Q9.cu
+101
View File
@@ -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;
}
LegacyCelerisLab/kernels/IO.cu
View File
LegacyCelerisLab/kernels/const.h
+10
View File
@@ -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
LegacyCelerisLab/kernels/kernel.cu
+234
View File
@@ -0,0 +1,234 @@
// 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, uint32_t *error_flag)
{
__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++)
{
if (isnan((double)g[i]) || isinf((double)g[i]))
{
atomicOr(error_flag, (uint32_t)1);
}
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++)
{
if (isnan((double)g[i]) || isinf((double)g[i]))
{
atomicOr(error_flag, (uint32_t)1);
}
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;
if (isnan((double)u) || isinf((double)u) || isnan((double)v) || isinf((double)v))
{
atomicOr(error_flag, (uint32_t)1);
}
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];
}
}
// __global__ void AddVortex(LBtype *f, int32_t *config)
// {
// __shared__ LBtype f_share[NT * NQ];
// int x, y, k;
// LBtype u, v, u_vor, v_vor;
// Index_lattice(x, y, k);
// int totalCells = NX * NY;
// for (int i = 0; i < NQ; i++)
// {
// f_share[threadIdx.x + i * NT] = f[k + i * totalCells];
// }
// __syncthreads();
// u = f_share[threadIdx.x + 1 * NT] - f_share[threadIdx.x + 3 * NT] + f_share[threadIdx.x + 5 * NT] - f_share[threadIdx.x + 6 * NT] - f_share[threadIdx.x + 7 * NT] + f_share[threadIdx.x + 8 * NT];
// v = f_share[threadIdx.x + 2 * NT] - f_share[threadIdx.x + 4 * NT] + f_share[threadIdx.x + 5 * NT] + f_share[threadIdx.x + 6 * NT] - f_share[threadIdx.x + 7 * NT] - f_share[threadIdx.x + 8 * NT];
// if type & V_TAYLOR
// {
// u_vor = -2 * PI * U0 * sin(2 * PI * x / NX) * sin(2 * PI * y / NY);
// v_vor = 2 * PI * U0 * cos(2 * PI * x / NX) * cos(2 * PI * y / NY);
// }
// else
// {
// u_vor = 0;
// v_vor = 0;
// }
// }
}
LegacyCelerisLab/kernels/macros.h
+37
View File
@@ -0,0 +1,37 @@
// 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.008
#define RHO 1.0
#define U0 0.02
// constants
#define PI 3.141592653589793238
#define FLUID 0b00000001
#define SOLID 0b00000010
#define GAS 0b00000100
#define INTERFACE 0b00001000
#define SENSOR 0b00010000
// vortex type
#define V_TAYLOR 0b00000001
// variables
#define N_OBJS 7
// #define N_SENS 2
LegacyCelerisLab/kernels/preproc.cu
+2
View File
@@ -0,0 +1,2 @@
#include "macros.h"
#include "const.h"