reflect BC for D3Q7

This commit is contained in:
James McClure 2020-04-07 10:38:21 -04:00
parent af8b2d799a
commit 3b006fbc3c
4 changed files with 66 additions and 1 deletions

View File

@ -1286,7 +1286,16 @@ void ScaLBL_Communicator::BiRecvD3Q7AA(double *Aq, double *Bq){
ScaLBL_D3Q7_Unpack(5,dvcRecvDist_Z,0,recvCount_Z,recvbuf_Z,Aq,N);
ScaLBL_D3Q7_Unpack(5,dvcRecvDist_Z,recvCount_Z,recvCount_Z,recvbuf_Z,Bq,N);
}
if (BoundaryCondition == 5){
if (kproc == 0){
ScaLBL_D3Q7_Reflection_BC_z(dvcSendList_z, Aq, sendCount_z, N);
ScaLBL_D3Q7_Reflection_BC_z(dvcSendList_z, Bq, sendCount_z, N);
}
if (kproc == nprocz-1){
ScaLBL_D3Q7_Reflection_BC_Z(dvcSendList_Z, Aq, sendCount_Z, N);
ScaLBL_D3Q7_Reflection_BC_Z(dvcSendList_Z, Bq, sendCount_Z, N);
}
}
//...................................................................................
Lock=false; // unlock the communicator after communications complete
//...................................................................................

View File

@ -141,6 +141,10 @@ extern "C" void ScaLBL_D3Q19_Reflection_BC_z(int *list, double *dist, int count,
extern "C" void ScaLBL_D3Q19_Reflection_BC_Z(int *list, double *dist, int count, int Np);
extern "C" void ScaLBL_D3Q7_Reflection_BC_z(int *list, double *dist, int count, int Np);
extern "C" void ScaLBL_D3Q7_Reflection_BC_Z(int *list, double *dist, int count, int Np);
extern "C" void ScaLBL_SetSlice_z(double *Phi, double value, int Nx, int Ny, int Nz, int Slice);
extern "C" void ScaLBL_CopySlice_z(double *Phi, int Nx, int Ny, int Nz, int Source, int Destination);

View File

@ -72,6 +72,21 @@ extern "C" void ScaLBL_UnpackDenD3Q7(int *list, int count, double *recvbuf, int
}
}
extern "C" void ScaLBL_D3Q7_Reflection_BC_z(int *list, double *dist, int count, int Np){
for (int idx=0; idx<count; idx++){
n = list[idx];
double f5 = 0.222222222222222222222222 - dist[6*Np+n];
dist[6*Np+n] = f5;
}
}
extern "C" void ScaLBL_D3Q7_Reflection_BC_Z(int *list, double *dist, int count, int Np){
for (int idx=0; idx<count; idx++){
n = list[idx];
double f6 = 0.222222222222222222222222 - dist[5*Np+n];
dist[5*Np+n] = f6;
}
}
extern "C" void ScaLBL_D3Q7_Init(char *ID, double *f_even, double *f_odd, double *Den, int Nx, int Ny, int Nz)
{
int n,N;

View File

@ -79,6 +79,25 @@ __global__ void dvc_ScaLBL_D3Q7_Unpack(int q, int *list, int start, int count,
}
}
__global__ void dvc_ScaLBL_D3Q7_Reflection_BC_z(int *list, double *dist, int count, int Np){
int idx, n;
idx = blockIdx.x*blockDim.x + threadIdx.x;
if (idx < count){
n = list[idx];
double f5 = 0.222222222222222222222222 - dist[6*Np+n];
dist[6*Np+n] = f5;
}
}
__global__ void dvc_ScaLBL_D3Q7_Reflection_BC_Z(int *list, double *dist, int count, int Np){
int idx, n;
idx = blockIdx.x*blockDim.x + threadIdx.x;
if (idx < count){
n = list[idx];
double f6 = 0.222222222222222222222222 - dist[5*Np+n];
dist[5*Np+n] = f6;
}
}
__global__ void dvc_ScaLBL_D3Q7_Init(char *ID, double *f_even, double *f_odd, double *Den, int Nx, int Ny, int Nz)
{
int n,N;
@ -207,6 +226,24 @@ __global__ void dvc_ScaLBL_D3Q7_Density(char *ID, double *disteven, double *dis
}
}
extern "C" void ScaLBL_D3Q7_Reflection_BC_z(int *list, double *dist, int count, int Np){
int GRID = count / 512 + 1;
dvc_ScaLBL_D3Q7_Reflection_BC_z<<<GRID,512>>>(list, dist, count, Np);
cudaError_t err = cudaGetLastError();
if (cudaSuccess != err){
printf("CUDA error in ScaLBL_D3Q7_Reflection_BC_z (kernel): %s \n",cudaGetErrorString(err));
}
}
extern "C" void ScaLBL_D3Q7_Reflection_BC_Z(int *list, double *dist, int count, int Np){
int GRID = count / 512 + 1;
dvc_ScaLBL_D3Q7_Reflection_BC_Z<<<GRID,512>>>(list, dist, count, Np);
cudaError_t err = cudaGetLastError();
if (cudaSuccess != err){
printf("CUDA error in ScaLBL_D3Q7_Reflection_BC_Z (kernel): %s \n",cudaGetErrorString(err));
}
}
extern "C" void ScaLBL_D3Q7_Unpack(int q, int *list, int start, int count, double *recvbuf, double *dist, int N){
int GRID = count / 512 + 1;
dvc_ScaLBL_D3Q7_Unpack <<<GRID,512 >>>(q, list, start, count, recvbuf, dist, N);