#include #include //#include #define NBLOCKS 1024 #define NTHREADS 256 __global__ void dvc_ScaLBL_D3Q7_AAodd_Poisson_ElectricPotential(int *neighborList,int *Map, double *dist, double *Psi, int start, int finish, int Np){ int n; double psi;//electric potential double fq; int nread; int idx; int S = Np/NBLOCKS/NTHREADS + 1; for (int s=0; s 10Np => odd part of dist) f1 = dist[nr1]; // reading the f1 data into register fq nr2 = neighborList[n+Np]; // neighbor 1 ( < 10Np => even part of dist) f2 = dist[nr2]; // reading the f2 data into register fq // q=3 nr3 = neighborList[n+2*Np]; // neighbor 4 f3 = dist[nr3]; // q = 4 nr4 = neighborList[n+3*Np]; // neighbor 3 f4 = dist[nr4]; // q=5 nr5 = neighborList[n+4*Np]; f5 = dist[nr5]; // q = 6 nr6 = neighborList[n+5*Np]; f6 = dist[nr6]; Ex = (f1-f2)*rlx*4.0;//NOTE the unit of electric field here is V/lu Ey = (f3-f4)*rlx*4.0;//factor 4.0 is D3Q7 lattice speed of sound Ez = (f5-f6)*rlx*4.0; ElectricField[n+0*Np] = Ex; ElectricField[n+1*Np] = Ey; ElectricField[n+2*Np] = Ez; // q = 0 dist[n] = f0*(1.0-rlx) + 0.25*(rlx*psi+rho_e); // q = 1 dist[nr2] = f1*(1.0-rlx) + 0.125*(rlx*psi+rho_e); // q = 2 dist[nr1] = f2*(1.0-rlx) + 0.125*(rlx*psi+rho_e); // q = 3 dist[nr4] = f3*(1.0-rlx) + 0.125*(rlx*psi+rho_e); // q = 4 dist[nr3] = f4*(1.0-rlx) + 0.125*(rlx*psi+rho_e); // q = 5 dist[nr6] = f5*(1.0-rlx) + 0.125*(rlx*psi+rho_e); // q = 6 dist[nr5] = f6*(1.0-rlx) + 0.125*(rlx*psi+rho_e); //........................................................................ } } } __global__ void dvc_ScaLBL_D3Q7_AAeven_Poisson(int *Map, double *dist, double *Den_charge, double *Psi, double *ElectricField, double tau, double epsilon_LB,int start, int finish, int Np){ int n; double psi;//electric potential double Ex,Ey,Ez;//electric field double rho_e;//local charge density double f0,f1,f2,f3,f4,f5,f6; double rlx=1.0/tau; int idx; int S = Np/NBLOCKS/NTHREADS + 1; for (int s=0; s>>(neighborList,Map,dist,Psi,start,finish,Np); cudaError_t err = cudaGetLastError(); if (cudaSuccess != err){ printf("CUDA error in ScaLBL_D3Q7_AAodd_Poisson_ElectricPotential: %s \n",cudaGetErrorString(err)); } //cudaProfilerStop(); } extern "C" void ScaLBL_D3Q7_AAeven_Poisson_ElectricPotential(int *Map, double *dist, double *Psi, int start, int finish, int Np){ //cudaProfilerStart(); dvc_ScaLBL_D3Q7_AAeven_Poisson_ElectricPotential<<>>(Map,dist,Psi,start,finish,Np); cudaError_t err = cudaGetLastError(); if (cudaSuccess != err){ printf("CUDA error in ScaLBL_D3Q7_AAeven_Poisson_ElectricPotential: %s \n",cudaGetErrorString(err)); } //cudaProfilerStop(); } extern "C" void ScaLBL_D3Q7_AAodd_Poisson(int *neighborList, int *Map, double *dist, double *Den_charge, double *Psi, double *ElectricField, double tau, double epsilon_LB,int start, int finish, int Np){ //cudaProfilerStart(); dvc_ScaLBL_D3Q7_AAodd_Poisson<<>>(neighborList,Map,dist,Den_charge,Psi,ElectricField,tau,epsilon_LB,start,finish,Np); cudaError_t err = cudaGetLastError(); if (cudaSuccess != err){ printf("CUDA error in ScaLBL_D3Q7_AAodd_Poisson: %s \n",cudaGetErrorString(err)); } //cudaProfilerStop(); } extern "C" void ScaLBL_D3Q7_AAeven_Poisson(int *Map, double *dist, double *Den_charge, double *Psi, double *ElectricField, double tau, double epsilon_LB,int start, int finish, int Np){ //cudaProfilerStart(); dvc_ScaLBL_D3Q7_AAeven_Poisson<<>>(Map,dist,Den_charge,Psi,ElectricField,tau,epsilon_LB,start,finish,Np); cudaError_t err = cudaGetLastError(); if (cudaSuccess != err){ printf("CUDA error in ScaLBL_D3Q7_AAeven_Poisson: %s \n",cudaGetErrorString(err)); } //cudaProfilerStop(); } extern "C" void ScaLBL_D3Q7_Poisson_Init(int *Map, double *dist, double *Psi, int start, int finish, int Np){ //cudaProfilerStart(); dvc_ScaLBL_D3Q7_Poisson_Init<<>>(Map,dist,Psi,start,finish,Np); cudaError_t err = cudaGetLastError(); if (cudaSuccess != err){ printf("CUDA error in ScaLBL_D3Q7_Poisson_Init: %s \n",cudaGetErrorString(err)); } //cudaProfilerStop(); }