From f44167d2ef8ef7d449389c25d7741509feb1f50f Mon Sep 17 00:00:00 2001 From: James McClure Date: Wed, 24 Mar 2021 07:57:27 -0400 Subject: [PATCH] add hip version of Lee model --- hip/FreeLee.cu | 1483 ++++++++++++++++++++++++++---------------------- 1 file changed, 794 insertions(+), 689 deletions(-) diff --git a/hip/FreeLee.cu b/hip/FreeLee.cu index 558bd2f1..09bc8689 100644 --- a/hip/FreeLee.cu +++ b/hip/FreeLee.cu @@ -1,11 +1,12 @@ #include +#include #include "hip/hip_runtime.h" +#define STOKES + #define NBLOCKS 1024 #define NTHREADS 256 -#define STOKES - __global__ void dvc_ScaLBL_D3Q19_FreeLeeModel_TwoFluid_Init(double *gqbar, double *mu_phi, double *ColorGrad, double Fx, double Fy, double Fz, int Np) { int n; @@ -186,10 +187,16 @@ __global__ void dvc_ScaLBL_D3Q7_AAodd_FreeLeeModel_PhaseField(int *neighborList, } } -__global__ void dvc_ScaLBL_D3Q7_AAeven_FreeLeeModel_PhaseField(int *Map, double *hq, double *Den, double *Phi, - double rhoA, double rhoB, int start, int finish, int Np){ - int idx,n; - double fq,phi; +__global__ void dvc_ScaLBL_D3Q7_AAodd_FreeLee_PhaseField(int *neighborList, int *Map, double *hq, double *Den, double *Phi, double *ColorGrad, double *Vel, + double rhoA, double rhoB, double tauM, double W, int start, int finish, int Np){ + + int n,idx,nr1,nr2,nr3,nr4,nr5,nr6; + double h0,h1,h2,h3,h4,h5,h6; + double nx,ny,nz,C; + double ux,uy,uz; + double phi; + double M = 2.0/9.0*(tauM-0.5);//diffusivity (or mobility) for the phase field D3Q7 + double factor = 1.0; // for (int n=start; n>>( gqbar, mu_phi, ColorGrad, Fx, Fy, Fz, Np); + hipError_t err = hipGetLastError(); + if (hipSuccess != err){ + printf("CUDA error in ScaLBL_D3Q19_FreeLeeModel_TwoFluid_Init: %s \n",hipGetErrorString(err)); + } } extern "C" void ScaLBL_D3Q19_FreeLeeModel_SingleFluid_Init(double *gqbar, double Fx, double Fy, double Fz, int Np){ - + + dvc_ScaLBL_D3Q19_FreeLeeModel_SingleFluid_Init<<>>( gqbar, Fx, Fy, Fz, Np); + hipError_t err = hipGetLastError(); + if (hipSuccess != err){ + printf("CUDA error in ScaLBL_D3Q19_FreeLeeModel_SingleFluid_Init: %s \n",hipGetErrorString(err)); + } } extern "C" void ScaLBL_FreeLeeModel_PhaseField_Init(int *Map, double *Phi, double *Den, double *hq, double *ColorGrad, double rhoA, double rhoB, double tauM, double W, int start, int finish, int Np){ -} -extern "C" void ScaLBL_D3Q7_AAodd_FreeLeeModel_PhaseField(int *neighborList, int *Map, double *hq, double *Den, double *Phi, - double rhoA, double rhoB, int start, int finish, int Np){ + dvc_ScaLBL_FreeLeeModel_PhaseField_Init<<>>(Map, Phi, Den, hq, ColorGrad, rhoA, rhoB, tauM, W, start, finish, Np); + hipError_t err = hipGetLastError(); + if (hipSuccess != err){ + printf("CUDA error in ScaLBL_FreeLeeModel_PhaseField_Init: %s \n",hipGetErrorString(err)); + } + } - -extern "C" void ScaLBL_D3Q7_AAeven_FreeLeeModel_PhaseField(int *Map, double *hq, double *Den, double *Phi, - double rhoA, double rhoB, int start, int finish, int Np){ - +extern "C" void ScaLBL_D3Q7_AAodd_FreeLee_PhaseField(int *neighborList, int *Map, double *hq, double *Den, double *Phi, double *ColorGrad, double *Vel, + double rhoA, double rhoB, double tauM, double W, int start, int finish, int Np) +{ + hipFuncSetCacheConfig(dvc_ScaLBL_D3Q7_AAodd_FreeLee_PhaseField, hipFuncCachePreferL1); + dvc_ScaLBL_D3Q7_AAodd_FreeLee_PhaseField<<>>(neighborList, Map, hq, Den, Phi, ColorGrad, Vel, + rhoA, rhoB, tauM, W, start, finish, Np); + hipError_t err = hipGetLastError(); + if (hipSuccess != err){ + printf("CUDA error in ScaLBL_D3Q7_AAodd_FreeLee_PhaseField: %s \n",hipGetErrorString(err)); + } } -extern "C" void ScaLBL_D3Q19_AAodd_FreeLeeModel(int *neighborList, int *Map, double *dist, double *hq, double *Den, double *Phi, double *mu_phi, double *Vel, double *Pressure, double *ColorGrad, - double rhoA, double rhoB, double tauA, double tauB, double tauM, double kappa, double beta, double W, double Fx, double Fy, double Fz, +extern "C" void ScaLBL_D3Q7_AAeven_FreeLee_PhaseField( int *Map, double *hq, double *Den, double *Phi, double *ColorGrad, double *Vel, + double rhoA, double rhoB, double tauM, double W, int start, int finish, int Np){ + + hipFuncSetCacheConfig(dvc_ScaLBL_D3Q7_AAeven_FreeLee_PhaseField, hipFuncCachePreferL1); + dvc_ScaLBL_D3Q7_AAeven_FreeLee_PhaseField<<>>( Map, hq, Den, Phi, ColorGrad, Vel, rhoA, rhoB, tauM, W, start, finish, Np); + hipError_t err = hipGetLastError(); + if (hipSuccess != err){ + printf("CUDA error in ScaLBL_D3Q7_AAeven_FreeLee_PhaseField: %s \n",hipGetErrorString(err)); + } +} + + +extern "C" void ScaLBL_D3Q7_ComputePhaseField(int *Map, double *hq, double *Den, double *Phi, double rhoA, double rhoB, int start, int finish, int Np){ + + hipFuncSetCacheConfig(dvc_ScaLBL_D3Q7_ComputePhaseField, hipFuncCachePreferL1); + dvc_ScaLBL_D3Q7_ComputePhaseField<<>>( Map, hq, Den, Phi, rhoA, rhoB, start, finish, Np); + hipError_t err = hipGetLastError(); + if (hipSuccess != err){ + printf("CUDA error in ScaLBL_D3Q7_ComputePhaseField: %s \n",hipGetErrorString(err)); + } +} + + +extern "C" void ScaLBL_D3Q19_AAodd_FreeLeeModel(int *neighborList, int *Map, double *dist, double *Den, double *Phi, double *mu_phi, double *Vel, double *Pressure, double *ColorGrad, + double rhoA, double rhoB, double tauA, double tauB, double kappa, double beta, double W, double Fx, double Fy, double Fz, int strideY, int strideZ, int start, int finish, int Np){ -} - + hipFuncSetCacheConfig(dvc_ScaLBL_D3Q19_AAodd_FreeLeeModel, hipFuncCachePreferL1); + dvc_ScaLBL_D3Q19_AAodd_FreeLeeModel<<>>(neighborList, Map, dist, Den, Phi, mu_phi, Vel, Pressure, ColorGrad, + rhoA, rhoB, tauA, tauB, kappa, beta, W, Fx, Fy, Fz, strideY, strideZ, start, finish, Np); + hipError_t err = hipGetLastError(); + if (hipSuccess != err){ + printf("CUDA error in ScaLBL_D3Q19_AAodd_FreeLeeModel: %s \n",hipGetErrorString(err)); + } +} -extern "C" void ScaLBL_D3Q19_AAeven_FreeLeeModel(int *Map, double *dist, double *hq, double *Den, double *Phi, double *mu_phi, double *Vel, double *Pressure, double *ColorGrad, - double rhoA, double rhoB, double tauA, double tauB, double tauM, double kappa, double beta, double W, double Fx, double Fy, double Fz, +extern "C" void ScaLBL_D3Q19_AAeven_FreeLeeModel(int *Map, double *dist, double *Den, double *Phi, double *mu_phi, double *Vel, double *Pressure, double *ColorGrad, + double rhoA, double rhoB, double tauA, double tauB, double kappa, double beta, double W, double Fx, double Fy, double Fz, int strideY, int strideZ, int start, int finish, int Np){ + hipFuncSetCacheConfig(dvc_ScaLBL_D3Q19_AAeven_FreeLeeModel, hipFuncCachePreferL1); + dvc_ScaLBL_D3Q19_AAeven_FreeLeeModel<<>>(Map, dist, Den, Phi, mu_phi, Vel, Pressure, ColorGrad, + rhoA, rhoB, tauA, tauB, kappa, beta, W, Fx, Fy, Fz, strideY, strideZ, start, finish, Np); + hipError_t err = hipGetLastError(); + if (hipSuccess != err){ + printf("CUDA error in ScaLBL_D3Q19_AAeven_FreeLeeModel: %s \n",hipGetErrorString(err)); + } + } extern "C" void ScaLBL_D3Q19_AAodd_FreeLeeModel_SingleFluid_BGK(int *neighborList, double *dist, double *Vel, double *Pressure, double tau, double rho0, double Fx, double Fy, double Fz, int start, int finish, int Np){ - + hipFuncSetCacheConfig(dvc_ScaLBL_D3Q19_AAodd_FreeLeeModel_SingleFluid_BGK, hipFuncCachePreferL1); + dvc_ScaLBL_D3Q19_AAodd_FreeLeeModel_SingleFluid_BGK<<>>(neighborList, dist, Vel, Pressure, + tau, rho0, Fx, Fy, Fz, start, finish, Np); + hipError_t err = hipGetLastError(); + if (hipSuccess != err){ + printf("CUDA error in ScaLBL_D3Q19_AAodd_FreeLeeModel_SingleFluid_BGK: %s \n",hipGetErrorString(err)); + } } extern "C" void ScaLBL_D3Q19_AAeven_FreeLeeModel_SingleFluid_BGK(double *dist, double *Vel, double *Pressure, double tau, double rho0, double Fx, double Fy, double Fz, int start, int finish, int Np){ - + + hipFuncSetCacheConfig(dvc_ScaLBL_D3Q19_AAeven_FreeLeeModel_SingleFluid_BGK, hipFuncCachePreferL1); + dvc_ScaLBL_D3Q19_AAeven_FreeLeeModel_SingleFluid_BGK<<>>(dist, Vel, Pressure, + tau, rho0, Fx, Fy, Fz, start, finish, Np); + hipError_t err = hipGetLastError(); + if (hipSuccess != err){ + printf("CUDA error in ScaLBL_D3Q19_AAeven_FreeLeeModel_SingleFluid_BGK: %s \n",hipGetErrorString(err)); + } } + +extern "C" void ScaLBL_D3Q9_MGTest(int *Map, double *Phi,double *ColorGrad,int strideY, int strideZ, int start, int finish, int Np){ +} \ No newline at end of file