OilfieldToolsNet/LBPM
4
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

gpu color clenaup

Browse Source
This commit is contained in:
James E McClure 2018-09-22 17:15:00 -04:00
parent 86de2527c4
commit 379b731069
3 changed files with 5 additions and 91 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

60
cpu/D3Q19.cpp
View File

@ -12,36 +12,6 @@ extern "C" void ScaLBL_D3Q19_Pack(int q, int *list, int start, int count, double
}
}
/*extern "C" void ScaLBL_D3Q19_Unpack(int q, int Cqx, int Cqy, int Cqz, int *list, int start, int count,
double *recvbuf, double *dist, int Nx, int Ny, int Nz){
//....................................................................................
// Unack distribution from the recv buffer
// Distribution q matche Cqx, Cqy, Cqz
// swap rule means that the distributions in recvbuf are OPPOSITE of q
// dist may be even or odd distributions stored by stream layout
//....................................................................................
int i,j,k,n,nn,idx;
int N = Nx*Ny*Nz;
for (idx=0; idx<count; idx++){
// Get the value from the list -- note that n is the index is from the send (non-local) process
n = list[idx];
// Get the 3-D indices
k = n/(Nx*Ny);
j = (n-Nx*Ny*k)/Nx;
i = n-Nx*Ny*k-Nx*j;
// Streaming for the non-local distribution
i += Cqx;
j += Cqy;
k += Cqz;
nn = k*Nx*Ny+j*Nx+i;
// unpack the distribution to the proper location
dist[q*N+nn] = recvbuf[start+idx];
}
}*/
extern "C" void ScaLBL_D3Q19_Unpack(int q, int *list, int start, int count,
double *recvbuf, double *dist, int N){
//....................................................................................
@ -61,36 +31,6 @@ extern "C" void ScaLBL_D3Q19_Unpack(int q, int *list, int start, int count,
}
}
/*
extern "C" void ScaLBL_D3Q19_MapRecv(int q, int Cqx, int Cqy, int Cqz, int *list, int start, int count,
int *d3q19_recvlist, int Nx, int Ny, int Nz){
//....................................................................................
// Map the recieve distributions to
// Distribution q matche Cqx, Cqy, Cqz
// swap rule means that the distributions in recvbuf are OPPOSITE of q
// dist may be even or odd distributions stored by stream layout
//....................................................................................
int i,j,k,n,nn,idx;
int N = Nx*Ny*Nz;
for (idx=0; idx<count; idx++){
// Get the value from the list -- note that n is the index is from the send (non-local) process
n = list[idx];
// Get the 3-D indices
k = n/(Nx*Ny);
j = (n-Nx*Ny*k)/Nx;
i = n-Nx*Ny*k-Nx*j;
// Streaming for the non-local distribution
i += Cqx;
j += Cqy;
k += Cqz;
// compute 1D index for the neighbor and save
nn = k*Nx*Ny+j*Nx+i;
d3q19_recvlist[start+idx] = nn;
}
}
*/
extern "C" void ScaLBL_D3Q19_AA_Init(double *f_even, double *f_odd, int Np)
{
int n;

32
gpu/D3Q19.cu
View File

@ -127,7 +127,7 @@ __global__ void deviceReduceKernel(double *in, double* out, int N) {
out[blockIdx.x]=sum;
}
__global__ void dvc_ScaLBL_D3Q19_Pack(int q, int *list, int start, int count, double *sendbuf, double *dist, int N){
__global__ void dvc_ScaLBL_D3Q19_Pack(int q, int *list, int start, int count, double *sendbuf, double *dist, int N){
//....................................................................................
// Pack distribution q into the send buffer for the listed lattice sites
// dist may be even or odd distributions stored by stream layout
@ -161,36 +161,6 @@ __global__ void dvc_ScaLBL_D3Q19_Unpack(int q, int *list, int start, int count
}
}
}
/*
__global__ void dvc_ScaLBL_D3Q19_MapRecv(int q, int Cqx, int Cqy, int Cqz, int *list, int start, int count,
int *d3q19_recvlist, int Nx, int Ny, int Nz){
//....................................................................................
// Unack distribution from the recv buffer
// Distribution q matche Cqx, Cqy, Cqz
// swap rule means that the distributions in recvbuf are OPPOSITE of q
// dist may be even or odd distributions stored by stream layout
//....................................................................................
int i,j,k,n,nn,idx;
int N = Nx*Ny*Nz;
idx = blockIdx.x*blockDim.x + threadIdx.x;
if (idx<count){
// Get the value from the list -- note that n is the index is from the send (non-local) process
n = list[idx];
// Get the 3-D indices
k = n/(Nx*Ny);
j = (n-Nx*Ny*k)/Nx;
i = n-Nx*Ny*k-Nx*j;
// Streaming for the non-local distribution
i += Cqx;
j += Cqy;
k += Cqz;
nn = k*Nx*Ny+j*Nx+i;
// unpack the distribution to the proper location
d3q19_recvlist[start+idx]=nn;
}
}
*/
__global__ void dvc_ScaLBL_D3Q19_Init(char *ID, double *f_even, double *f_odd, int Nx, int Ny, int Nz)
{

4
models/ColorModel.cpp
View File

@ -408,6 +408,7 @@ void ScaLBL_ColorModel::Run(){
ScaLBL_Comm->BiSendD3Q7AA(Aq,Bq); //READ FROM NORMAL
ScaLBL_D3Q7_AAodd_PhaseField(NeighborList, dvcMap, Aq, Bq, Den, Phi, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
ScaLBL_Comm->BiRecvD3Q7AA(Aq,Bq); //WRITE INTO OPPOSITE
ScaLBL_DeviceBarrier();
ScaLBL_D3Q7_AAodd_PhaseField(NeighborList, dvcMap, Aq, Bq, Den, Phi, 0, ScaLBL_Comm->LastExterior(), Np);
// Perform the collision operation
@ -423,6 +424,7 @@ void ScaLBL_ColorModel::Run(){
alpha, beta, Fx, Fy, Fz, Nx, Nx*Ny, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
ScaLBL_Comm_Regular->RecvHalo(Phi);
ScaLBL_Comm->RecvD3Q19AA(fq); //WRITE INTO OPPOSITE
ScaLBL_DeviceBarrier();
// Set BCs
if (BoundaryCondition == 3){
ScaLBL_Comm->D3Q19_Pressure_BC_z(NeighborList, fq, din, timestep);
@ -442,6 +444,7 @@ void ScaLBL_ColorModel::Run(){
ScaLBL_Comm->BiSendD3Q7AA(Aq,Bq); //READ FROM NORMAL
ScaLBL_D3Q7_AAeven_PhaseField(dvcMap, Aq, Bq, Den, Phi, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
ScaLBL_Comm->BiRecvD3Q7AA(Aq,Bq); //WRITE INTO OPPOSITE
ScaLBL_DeviceBarrier();
ScaLBL_D3Q7_AAeven_PhaseField(dvcMap, Aq, Bq, Den, Phi, 0, ScaLBL_Comm->LastExterior(), Np);
// Perform the collision operation
@ -456,6 +459,7 @@ void ScaLBL_ColorModel::Run(){
alpha, beta, Fx, Fy, Fz, Nx, Nx*Ny, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
ScaLBL_Comm_Regular->RecvHalo(Phi);
ScaLBL_Comm->RecvD3Q19AA(fq); //WRITE INTO OPPOSITE
ScaLBL_DeviceBarrier();
// Set boundary conditions
if (BoundaryCondition == 3){
ScaLBL_Comm->D3Q19_Pressure_BC_z(NeighborList, fq, din, timestep);