diff --git a/src/gpu/core/GPU/CudaMemoryManager.cpp b/src/gpu/core/GPU/CudaMemoryManager.cpp index 312e379c476b54e868e97e2d1b5fb88311bb4158..a6adb34c165019372512873d153f45681c01ea11 100644 --- a/src/gpu/core/GPU/CudaMemoryManager.cpp +++ b/src/gpu/core/GPU/CudaMemoryManager.cpp @@ -1863,82 +1863,6 @@ void CudaMemoryManager::cudaFreePressX1(int lev) checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->QpressX1.RhoBC )); checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->QpressX1.deltaVz)); } -//Propeller Velocity -void CudaMemoryManager::cudaAllocVeloPropeller(int lev) -{ - unsigned int mem_size_Propeller_k = sizeof(int)*parameter->getParH(lev)->propellerBC.numberOfBCnodes; - unsigned int mem_size_Propeller_q = sizeof(real)*parameter->getParH(lev)->propellerBC.numberOfBCnodes; - - //Host - //checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->propellerBC.q27[0]), parameter->getD3Qxx()*mem_size_Propeller_q )); - checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->propellerBC.k), mem_size_Propeller_k )); - checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->propellerBC.Vx), mem_size_Propeller_q )); - checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->propellerBC.Vy), mem_size_Propeller_q )); - checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->propellerBC.Vz), mem_size_Propeller_q )); - checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->propellerBC.RhoBC), mem_size_Propeller_q )); - - //Device - //checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->propellerBC.q27[0]), parameter->getD3Qxx()*mem_size_Propeller_q )); - checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->propellerBC.k), mem_size_Propeller_k )); - checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->propellerBC.Vx), mem_size_Propeller_q )); - checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->propellerBC.Vy), mem_size_Propeller_q )); - checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->propellerBC.Vz), mem_size_Propeller_q )); - checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->propellerBC.RhoBC), mem_size_Propeller_q )); - - ////////////////////////////////////////////////////////////////////////// - double tmp = (double)mem_size_Propeller_k + 4. * (double)mem_size_Propeller_q; - setMemsizeGPU(tmp, false); -} -void CudaMemoryManager::cudaCopyVeloPropeller(int lev) -{ - unsigned int mem_size_Propeller_k = sizeof(int)*parameter->getParH(lev)->propellerBC.numberOfBCnodes; - unsigned int mem_size_Propeller_q = sizeof(real)*parameter->getParH(lev)->propellerBC.numberOfBCnodes; - - //checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->propellerBC.q27[0], parameter->getParH(lev)->propellerBC.q27[0], parameter->getD3Qxx()* mem_size_Propeller_q, cudaMemcpyHostToDevice)); - checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->propellerBC.k, parameter->getParH(lev)->propellerBC.k, mem_size_Propeller_k, cudaMemcpyHostToDevice)); - checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->propellerBC.Vx, parameter->getParH(lev)->propellerBC.Vx, mem_size_Propeller_q, cudaMemcpyHostToDevice)); - checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->propellerBC.Vy, parameter->getParH(lev)->propellerBC.Vy, mem_size_Propeller_q, cudaMemcpyHostToDevice)); - checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->propellerBC.Vz, parameter->getParH(lev)->propellerBC.Vz, mem_size_Propeller_q, cudaMemcpyHostToDevice)); - checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->propellerBC.RhoBC, parameter->getParH(lev)->propellerBC.RhoBC, mem_size_Propeller_q, cudaMemcpyHostToDevice)); -} -void CudaMemoryManager::cudaFreeVeloPropeller(int lev) -{ - //checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->propellerBC.q27[0] )); - checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->propellerBC.k )); - checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->propellerBC.Vx )); - checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->propellerBC.Vy )); - checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->propellerBC.Vz )); - checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->propellerBC.RhoBC )); -} -//Measure Points -//void CudaMemoryManager::cudaAllocMeasurePoints(int lev, int i) -//{ -// //Host -// checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->MP[i].Vx), parameter->getParH(lev)->memSizerealMP )); -// checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->MP[i].Vy), parameter->getParH(lev)->memSizerealMP )); -// checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->MP[i].Vz), parameter->getParH(lev)->memSizerealMP )); -// checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->MP[i].Rho), parameter->getParH(lev)->memSizerealMP )); -// -// //Device -// checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->MP[i].Vx), parameter->getParD(lev)->memSizerealMP )); -// checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->MP[i].Vy), parameter->getParD(lev)->memSizerealMP )); -// checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->MP[i].Vz), parameter->getParD(lev)->memSizerealMP )); -// checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->MP[i].Rho), parameter->getParD(lev)->memSizerealMP )); -//} -//void CudaMemoryManager::cudaCopyMeasurePoints(int lev, int i) -//{ -// checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->MP[i].Vx, parameter->getParH(lev)->MP[i].Vx, parameter->getParH(lev)->memSizerealMP, cudaMemcpyHostToDevice)); -// checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->MP[i].Vy, parameter->getParH(lev)->MP[i].Vy, parameter->getParH(lev)->memSizerealMP, cudaMemcpyHostToDevice)); -// checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->MP[i].Vz, parameter->getParH(lev)->MP[i].Vz, parameter->getParH(lev)->memSizerealMP, cudaMemcpyHostToDevice)); -// checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->MP[i].Rho, parameter->getParH(lev)->MP[i].Rho, parameter->getParH(lev)->memSizerealMP, cudaMemcpyHostToDevice)); -//} -//void CudaMemoryManager::cudaFreeMeasurePoints(int lev, int i) -//{ -// checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->MP[i].Vx )); -// checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->MP[i].Vy )); -// checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->MP[i].Vz )); -// checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->MP[i].Rho )); -//} void CudaMemoryManager::cudaAllocMeasurePointsIndex(int lev) { //Host diff --git a/src/gpu/core/GPU/CudaMemoryManager.h b/src/gpu/core/GPU/CudaMemoryManager.h index dcdfecd6aad6ace700bc5e83e1d6e3f14e1f1bcc..89fed68436e4d40be2bf1abc9959825301a41488 100644 --- a/src/gpu/core/GPU/CudaMemoryManager.h +++ b/src/gpu/core/GPU/CudaMemoryManager.h @@ -38,9 +38,6 @@ public: void setMemsizeGPU(double admem, bool reset); double getMemsizeGPU(); - //void cudaAllocFull(int lev); //DEPRECATED: related to full matrix - //void cudaFreeFull(int lev); //DEPRECATED: related to full matrix - void cudaCopyPrint(int lev); void cudaCopyMedianPrint(int lev); @@ -251,14 +248,6 @@ public: void cudaCopyPressX1(int lev); void cudaFreePressX1(int lev); - void cudaAllocVeloPropeller(int lev); - void cudaCopyVeloPropeller(int lev); - void cudaFreeVeloPropeller(int lev); - - void cudaAllocMeasurePoints(int lev, int i); - void cudaCopyMeasurePoints(int lev, int i); - void cudaFreeMeasurePoints(int lev, int i); - void cudaAllocMeasurePointsIndex(int lev); void cudaCopyMeasurePointsIndex(int lev); void cudaCopyMeasurePointsToHost(int lev); diff --git a/src/gpu/core/GPU/GPU_Interface.h b/src/gpu/core/GPU/GPU_Interface.h index 72e839bcf741b9966a4e38a7b78c1085b20132e2..2f25f5f15cb5f9334c1cef070ba9c32a7c11c68b 100644 --- a/src/gpu/core/GPU/GPU_Interface.h +++ b/src/gpu/core/GPU/GPU_Interface.h @@ -933,21 +933,6 @@ void QADPressIncompDev27( unsigned int numberOfThreads, unsigned long long numberOfLBnodes, bool isEvenTimestep); -void PropVelo( unsigned int numberOfThreads, - unsigned int* neighborX, - unsigned int* neighborY, - unsigned int* neighborZ, - real* rho, - real* ux, - real* uy, - real* uz, - int* k_Q, - unsigned int size_Prop, - unsigned long long numberOfLBnodes, - unsigned int* bcMatD, - real* DD, - bool EvenOrOdd); - void ScaleCF27( real* DC, real* DF, unsigned int* neighborCX, diff --git a/src/gpu/core/GPU/GPU_Kernels.cuh b/src/gpu/core/GPU/GPU_Kernels.cuh index ab8ab0fa3564de523e8fc8e6985f6d98c09fb404..6e007fb6ced0b4301b0f212d4b0199ae4988b728 100644 --- a/src/gpu/core/GPU/GPU_Kernels.cuh +++ b/src/gpu/core/GPU/GPU_Kernels.cuh @@ -1264,21 +1264,6 @@ __global__ void QADPressIncomp27( real* DD, unsigned long long numberOfLBnodes, bool isEvenTimestep); -//Propeller BC -__global__ void PropellerBC(unsigned int* neighborX, - unsigned int* neighborY, - unsigned int* neighborZ, - real* rho, - real* ux, - real* uy, - real* uz, - int* k_Q, - unsigned int size_Prop, - unsigned long long numberOfLBnodes, - unsigned int* bcMatD, - real* DD, - bool EvenOrOdd); - //coarse to fine diff --git a/src/gpu/core/GPU/LBMKernel.cu b/src/gpu/core/GPU/LBMKernel.cu index 403ed748f3d525bb5eaefe35b44ef26391230fcd..a870d43bb36811441afa654a01f4eeb7671ba28e 100644 --- a/src/gpu/core/GPU/LBMKernel.cu +++ b/src/gpu/core/GPU/LBMKernel.cu @@ -2688,41 +2688,6 @@ void QPrecursorDevDistributions( LBMSimulationParameter* parameterDevice, } ////////////////////////////////////////////////////////////////////////// -extern "C" void PropVelo( - unsigned int numberOfThreads, - unsigned int* neighborX, - unsigned int* neighborY, - unsigned int* neighborZ, - real* rho, - real* ux, - real* uy, - real* uz, - int* k_Q, - unsigned int size_Prop, - unsigned long long numberOfLBnodes, - unsigned int* bcMatD, - real* DD, - bool EvenOrOdd) -{ - vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(numberOfThreads, size_Prop); - - PropellerBC<<< grid.grid, grid.threads >>>( - neighborX, - neighborY, - neighborZ, - rho, - ux, - uy, - uz, - k_Q, - size_Prop, - numberOfLBnodes, - bcMatD, - DD, - EvenOrOdd); - getLastCudaError("PropellerBC execution failed"); -} -////////////////////////////////////////////////////////////////////////// void ScaleCF27( real* DC, real* DF, diff --git a/src/gpu/core/GPU/VelocityBCs27.cu b/src/gpu/core/GPU/VelocityBCs27.cu index d1798e84490ffed3e803717e22cc8dbbaccf2222..437c19d81e513c70fbdddbb8e20819933d1c42d0 100644 --- a/src/gpu/core/GPU/VelocityBCs27.cu +++ b/src/gpu/core/GPU/VelocityBCs27.cu @@ -5444,401 +5444,3 @@ __global__ void QVelDevice27( } } //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - -//////////////////////////////////////////////////////////////////////////////// -__global__ void PropellerBC( - unsigned int* neighborX, - unsigned int* neighborY, - unsigned int* neighborZ, - real* rho, - real* ux, - real* uy, - real* uz, - int* k_Q, - unsigned int size_Prop, - unsigned long long numberOfLBnodes, - unsigned int* bcMatD, - real* DD, - bool EvenOrOdd) -{ - //////////////////////////////////////////////////////////////////////////////// - const unsigned x = threadIdx.x; // Globaler x-Index - const unsigned y = blockIdx.x; // Globaler y-Index - const unsigned z = blockIdx.y; // Globaler z-Index - - const unsigned nx = blockDim.x; - const unsigned ny = gridDim.x; - - const unsigned k = nx*(ny*z + y) + x; - ////////////////////////////////////////////////////////////////////////// - - if(k<size_Prop) - { - //////////////////////////////////////////////////////////////////////////////// - Distributions27 D; - if (EvenOrOdd==true) - { - D.f[dP00] = &DD[dP00 * numberOfLBnodes]; - D.f[dM00] = &DD[dM00 * numberOfLBnodes]; - D.f[d0P0] = &DD[d0P0 * numberOfLBnodes]; - D.f[d0M0] = &DD[d0M0 * numberOfLBnodes]; - D.f[d00P] = &DD[d00P * numberOfLBnodes]; - D.f[d00M] = &DD[d00M * numberOfLBnodes]; - D.f[dPP0] = &DD[dPP0 * numberOfLBnodes]; - D.f[dMM0] = &DD[dMM0 * numberOfLBnodes]; - D.f[dPM0] = &DD[dPM0 * numberOfLBnodes]; - D.f[dMP0] = &DD[dMP0 * numberOfLBnodes]; - D.f[dP0P] = &DD[dP0P * numberOfLBnodes]; - D.f[dM0M] = &DD[dM0M * numberOfLBnodes]; - D.f[dP0M] = &DD[dP0M * numberOfLBnodes]; - D.f[dM0P] = &DD[dM0P * numberOfLBnodes]; - D.f[d0PP] = &DD[d0PP * numberOfLBnodes]; - D.f[d0MM] = &DD[d0MM * numberOfLBnodes]; - D.f[d0PM] = &DD[d0PM * numberOfLBnodes]; - D.f[d0MP] = &DD[d0MP * numberOfLBnodes]; - D.f[d000] = &DD[d000 * numberOfLBnodes]; - D.f[dPPP] = &DD[dPPP * numberOfLBnodes]; - D.f[dMMP] = &DD[dMMP * numberOfLBnodes]; - D.f[dPMP] = &DD[dPMP * numberOfLBnodes]; - D.f[dMPP] = &DD[dMPP * numberOfLBnodes]; - D.f[dPPM] = &DD[dPPM * numberOfLBnodes]; - D.f[dMMM] = &DD[dMMM * numberOfLBnodes]; - D.f[dPMM] = &DD[dPMM * numberOfLBnodes]; - D.f[dMPM] = &DD[dMPM * numberOfLBnodes]; - } - else - { - D.f[dM00] = &DD[dP00 * numberOfLBnodes]; - D.f[dP00] = &DD[dM00 * numberOfLBnodes]; - D.f[d0M0] = &DD[d0P0 * numberOfLBnodes]; - D.f[d0P0] = &DD[d0M0 * numberOfLBnodes]; - D.f[d00M] = &DD[d00P * numberOfLBnodes]; - D.f[d00P] = &DD[d00M * numberOfLBnodes]; - D.f[dMM0] = &DD[dPP0 * numberOfLBnodes]; - D.f[dPP0] = &DD[dMM0 * numberOfLBnodes]; - D.f[dMP0] = &DD[dPM0 * numberOfLBnodes]; - D.f[dPM0] = &DD[dMP0 * numberOfLBnodes]; - D.f[dM0M] = &DD[dP0P * numberOfLBnodes]; - D.f[dP0P] = &DD[dM0M * numberOfLBnodes]; - D.f[dM0P] = &DD[dP0M * numberOfLBnodes]; - D.f[dP0M] = &DD[dM0P * numberOfLBnodes]; - D.f[d0MM] = &DD[d0PP * numberOfLBnodes]; - D.f[d0PP] = &DD[d0MM * numberOfLBnodes]; - D.f[d0MP] = &DD[d0PM * numberOfLBnodes]; - D.f[d0PM] = &DD[d0MP * numberOfLBnodes]; - D.f[d000] = &DD[d000 * numberOfLBnodes]; - D.f[dMMM] = &DD[dPPP * numberOfLBnodes]; - D.f[dPPM] = &DD[dMMP * numberOfLBnodes]; - D.f[dMPM] = &DD[dPMP * numberOfLBnodes]; - D.f[dPMM] = &DD[dMPP * numberOfLBnodes]; - D.f[dMMP] = &DD[dPPM * numberOfLBnodes]; - D.f[dPPP] = &DD[dMMM * numberOfLBnodes]; - D.f[dMPP] = &DD[dPMM * numberOfLBnodes]; - D.f[dPMP] = &DD[dMPM * numberOfLBnodes]; - } - ////////////////////////////////////////////////////////////////////////// - unsigned int KQK = k_Q[k]; - unsigned int BC = bcMatD[KQK]; - if( (BC != GEO_SOLID) && (BC != GEO_VOID)) - { - ////////////////////////////////////////////////////////////////////////// - real vx1 = ux[k]; - real vx2 = uy[k]; - real vx3 = uz[k]; - //real vx1 = -c1o100; - //real vx2 = zero; - //real vx3 = zero; - ////////////////////////////////////////////////////////////////////////// - //index - ////////////////////////////////////////////////////////////////////////// - unsigned int kzero= KQK; - unsigned int ke = KQK; - unsigned int kw = neighborX[KQK]; - unsigned int kn = KQK; - unsigned int ks = neighborY[KQK]; - unsigned int kt = KQK; - unsigned int kb = neighborZ[KQK]; - unsigned int ksw = neighborY[kw]; - unsigned int kne = KQK; - unsigned int kse = ks; - unsigned int knw = kw; - unsigned int kbw = neighborZ[kw]; - unsigned int kte = KQK; - unsigned int kbe = kb; - unsigned int ktw = kw; - unsigned int kbs = neighborZ[ks]; - unsigned int ktn = KQK; - unsigned int kbn = kb; - unsigned int kts = ks; - unsigned int ktse = ks; - unsigned int kbnw = kbw; - unsigned int ktnw = kw; - unsigned int kbse = kbs; - unsigned int ktsw = ksw; - unsigned int kbne = kb; - unsigned int ktne = KQK; - unsigned int kbsw = neighborZ[ksw]; - ////////////////////////////////////////////////////////////////////////// - real f_E, f_W, f_N, f_S, f_T, f_B, f_NE, f_SW, f_SE, f_NW, f_TE, f_BW, f_BE, - f_TW, f_TN, f_BS, f_BN, f_TS, f_TNE, f_TSW, f_TSE, f_TNW, f_BNE, f_BSW, f_BSE, f_BNW, f_ZERO; - - f_ZERO= (D.f[d000])[kzero]; - f_E = (D.f[dP00])[ke ]; - f_W = (D.f[dM00])[kw ]; - f_N = (D.f[d0P0])[kn ]; - f_S = (D.f[d0M0])[ks ]; - f_T = (D.f[d00P])[kt ]; - f_B = (D.f[d00M])[kb ]; - f_NE = (D.f[dPP0])[kne ]; - f_SW = (D.f[dMM0])[ksw ]; - f_SE = (D.f[dPM0])[kse ]; - f_NW = (D.f[dMP0])[knw ]; - f_TE = (D.f[dP0P])[kte ]; - f_BW = (D.f[dM0M])[kbw ]; - f_BE = (D.f[dP0M])[kbe ]; - f_TW = (D.f[dM0P])[ktw ]; - f_TN = (D.f[d0PP])[ktn ]; - f_BS = (D.f[d0MM])[kbs ]; - f_BN = (D.f[d0PM])[kbn ]; - f_TS = (D.f[d0MP])[kts ]; - f_TNE = (D.f[dPPP])[ktne ]; - f_BSW = (D.f[dMMM])[kbsw ]; - f_BNE = (D.f[dPPM])[kbne ]; - f_TSW = (D.f[dMMP])[ktsw ]; - f_TSE = (D.f[dPMP])[ktse ]; - f_BNW = (D.f[dMPM])[kbnw ]; - f_BSE = (D.f[dPMM])[kbse ]; - f_TNW = (D.f[dMPP])[ktnw ]; - //f_W = (D.f[dP00])[ke ]; - //f_E = (D.f[dM00])[kw ]; - //f_S = (D.f[d0P0])[kn ]; - //f_N = (D.f[d0M0])[ks ]; - //f_B = (D.f[d00P])[kt ]; - //f_T = (D.f[d00M])[kb ]; - //f_SW = (D.f[dPP0])[kne ]; - //f_NE = (D.f[dMM0])[ksw ]; - //f_NW = (D.f[dPM0])[kse ]; - //f_SE = (D.f[dMP0])[knw ]; - //f_BW = (D.f[dP0P])[kte ]; - //f_TE = (D.f[dM0M])[kbw ]; - //f_TW = (D.f[dP0M])[kbe ]; - //f_BE = (D.f[dM0P])[ktw ]; - //f_BS = (D.f[d0PP])[ktn ]; - //f_TN = (D.f[d0MM])[kbs ]; - //f_TS = (D.f[d0PM])[kbn ]; - //f_BN = (D.f[d0MP])[kts ]; - //f_BSW = (D.f[dPPP])[ktne ]; - //f_TNE = (D.f[dMMM])[kbsw ]; - //f_TSW = (D.f[dPPM])[kbne ]; - //f_BNE = (D.f[dMMP])[ktsw ]; - //f_BNW = (D.f[dPMP])[ktse ]; - //f_TSE = (D.f[dMPM])[kbnw ]; - //f_TNW = (D.f[dPMM])[kbse ]; - //f_BSE = (D.f[dMPP])[ktnw ]; - ////////////////////////////////////////////////////////////////////////////////// - real vxo1, vxo2, vxo3, drho; - drho = /*zero;*/f_TSE + f_TNW + f_TNE + f_TSW + f_BSE + f_BNW + f_BNE + f_BSW + - f_BN + f_TS + f_TN + f_BS + f_BE + f_TW + f_TE + f_BW + f_SE + f_NW + f_NE + f_SW + - f_T + f_B + f_N + f_S + f_E + f_W + f_ZERO; - - vxo1 = (((f_TSE - f_BNW) - (f_TNW - f_BSE)) + ((f_TNE - f_BSW) - (f_TSW - f_BNE)) + - ((f_BE - f_TW) + (f_TE - f_BW)) + ((f_SE - f_NW) + (f_NE - f_SW)) + - (f_E - f_W) )/ (c1o1 + drho); - - - vxo2 = ((-(f_TSE - f_BNW) + (f_TNW - f_BSE)) + ((f_TNE - f_BSW) - (f_TSW - f_BNE)) + - ((f_BN - f_TS) + (f_TN - f_BS)) + (-(f_SE - f_NW) + (f_NE - f_SW)) + - (f_N - f_S) )/ (c1o1 + drho); - - vxo3 = (((f_TSE - f_BNW) + (f_TNW - f_BSE)) + ((f_TNE - f_BSW) + (f_TSW - f_BNE)) + - (-(f_BN - f_TS) + (f_TN - f_BS)) + ((f_TE - f_BW) - (f_BE - f_TW)) + - (f_T - f_B) )/ (c1o1 + drho); - - real cusq=c3o2*(vxo1*vxo1+vxo2*vxo2+vxo3*vxo3); - //vx1 = vx1 * two - vxo1; - //vx2 = vx2 * two - vxo2; - //vx3 = vx3 * two - vxo3; - real cusq2=c3o2*(vx1*vx1+vx2*vx2+vx3*vx3); - - //f_ZERO = ((one+drho) * ( c8over27 *(one+(-cusq2)))) - c8over27; - //f_E = ((one+drho) * ( c2over27 *(one+three*( vx1 )+c9over2*( vx1 )*( vx1 )-cusq2))) - c2over27 ; - //f_W = ((one+drho) * ( c2over27 *(one+three*(-vx1 )+c9over2*(-vx1 )*(-vx1 )-cusq2))) - c2over27 ; - //f_N = ((one+drho) * ( c2over27 *(one+three*( vx2 )+c9over2*( vx2 )*( vx2 )-cusq2))) - c2over27 ; - //f_S = ((one+drho) * ( c2over27 *(one+three*( -vx2 )+c9over2*( -vx2 )*( -vx2 )-cusq2))) - c2over27 ; - //f_T = ((one+drho) * ( c2over27 *(one+three*( vx3)+c9over2*( vx3)*( vx3)-cusq2))) - c2over27 ; - //f_B = ((one+drho) * ( c2over27 *(one+three*( -vx3)+c9over2*( -vx3)*( -vx3)-cusq2))) - c2over27 ; - //f_NE = ((one+drho) * ( c1over54 *(one+three*( vx1+vx2 )+c9over2*( vx1+vx2 )*( vx1+vx2 )-cusq2))) - c1over54 ; - //f_SW = ((one+drho) * ( c1over54 *(one+three*(-vx1-vx2 )+c9over2*(-vx1-vx2 )*(-vx1-vx2 )-cusq2))) - c1over54 ; - //f_SE = ((one+drho) * ( c1over54 *(one+three*( vx1-vx2 )+c9over2*( vx1-vx2 )*( vx1-vx2 )-cusq2))) - c1over54 ; - //f_NW = ((one+drho) * ( c1over54 *(one+three*(-vx1+vx2 )+c9over2*(-vx1+vx2 )*(-vx1+vx2 )-cusq2))) - c1over54 ; - //f_TE = ((one+drho) * ( c1over54 *(one+three*( vx1 +vx3)+c9over2*( vx1 +vx3)*( vx1 +vx3)-cusq2))) - c1over54 ; - //f_BW = ((one+drho) * ( c1over54 *(one+three*(-vx1 -vx3)+c9over2*(-vx1 -vx3)*(-vx1 -vx3)-cusq2))) - c1over54 ; - //f_BE = ((one+drho) * ( c1over54 *(one+three*( vx1 -vx3)+c9over2*( vx1 -vx3)*( vx1 -vx3)-cusq2))) - c1over54 ; - //f_TW = ((one+drho) * ( c1over54 *(one+three*(-vx1 +vx3)+c9over2*(-vx1 +vx3)*(-vx1 +vx3)-cusq2))) - c1over54 ; - //f_TN = ((one+drho) * ( c1over54 *(one+three*( vx2+vx3)+c9over2*( vx2+vx3)*( vx2+vx3)-cusq2))) - c1over54 ; - //f_BS = ((one+drho) * ( c1over54 *(one+three*( -vx2-vx3)+c9over2*( -vx2-vx3)*( -vx2-vx3)-cusq2))) - c1over54 ; - //f_BN = ((one+drho) * ( c1over54 *(one+three*( vx2-vx3)+c9over2*( vx2-vx3)*( vx2-vx3)-cusq2))) - c1over54 ; - //f_TS = ((one+drho) * ( c1over54 *(one+three*( -vx2+vx3)+c9over2*( -vx2+vx3)*( -vx2+vx3)-cusq2))) - c1over54 ; - //f_TNE = ((one+drho) * ( c1over216*(one+three*( vx1+vx2+vx3)+c9over2*( vx1+vx2+vx3)*( vx1+vx2+vx3)-cusq2))) - c1over216; - //f_BSW = ((one+drho) * ( c1over216*(one+three*(-vx1-vx2-vx3)+c9over2*(-vx1-vx2-vx3)*(-vx1-vx2-vx3)-cusq2))) - c1over216; - //f_BNE = ((one+drho) * ( c1over216*(one+three*( vx1+vx2-vx3)+c9over2*( vx1+vx2-vx3)*( vx1+vx2-vx3)-cusq2))) - c1over216; - //f_TSW = ((one+drho) * ( c1over216*(one+three*(-vx1-vx2+vx3)+c9over2*(-vx1-vx2+vx3)*(-vx1-vx2+vx3)-cusq2))) - c1over216; - //f_TSE = ((one+drho) * ( c1over216*(one+three*( vx1-vx2+vx3)+c9over2*( vx1-vx2+vx3)*( vx1-vx2+vx3)-cusq2))) - c1over216; - //f_BNW = ((one+drho) * ( c1over216*(one+three*(-vx1+vx2-vx3)+c9over2*(-vx1+vx2-vx3)*(-vx1+vx2-vx3)-cusq2))) - c1over216; - //f_BSE = ((one+drho) * ( c1over216*(one+three*( vx1-vx2-vx3)+c9over2*( vx1-vx2-vx3)*( vx1-vx2-vx3)-cusq2))) - c1over216; - //f_TNW = ((one+drho) * ( c1over216*(one+three*(-vx1+vx2+vx3)+c9over2*(-vx1+vx2+vx3)*(-vx1+vx2+vx3)-cusq2))) - c1over216; - f_ZERO = f_ZERO + ((c1o1+drho) * (- c8o27* (-cusq) + c8o27* (-cusq2))); - f_E = f_E + ((c1o1+drho) * (- c2o27* (c3o1*( vxo1 )+c9o2*( vxo1 )*( vxo1 )-cusq) + c2o27* (c3o1*( vx1 )+c9o2*( vx1 )*( vx1 )-cusq2))); - f_W = f_W + ((c1o1+drho) * (- c2o27* (c3o1*(-vxo1 )+c9o2*(-vxo1 )*(-vxo1 )-cusq) + c2o27* (c3o1*(-vx1 )+c9o2*(-vx1 )*(-vx1 )-cusq2))); - f_N = f_N + ((c1o1+drho) * (- c2o27* (c3o1*( vxo2 )+c9o2*( vxo2 )*( vxo2 )-cusq) + c2o27* (c3o1*( vx2 )+c9o2*( vx2 )*( vx2 )-cusq2))); - f_S = f_S + ((c1o1+drho) * (- c2o27* (c3o1*( -vxo2 )+c9o2*( -vxo2 )*( -vxo2 )-cusq) + c2o27* (c3o1*( -vx2 )+c9o2*( -vx2 )*( -vx2 )-cusq2))); - f_T = f_T + ((c1o1+drho) * (- c2o27* (c3o1*( vxo3)+c9o2*( vxo3)*( vxo3)-cusq) + c2o27* (c3o1*( vx3)+c9o2*( vx3)*( vx3)-cusq2))); - f_B = f_B + ((c1o1+drho) * (- c2o27* (c3o1*( -vxo3)+c9o2*( -vxo3)*( -vxo3)-cusq) + c2o27* (c3o1*( -vx3)+c9o2*( -vx3)*( -vx3)-cusq2))); - f_NE = f_NE + ((c1o1+drho) * (- c1o54* (c3o1*( vxo1+vxo2 )+c9o2*( vxo1+vxo2 )*( vxo1+vxo2 )-cusq) + c1o54* (c3o1*( vx1+vx2 )+c9o2*( vx1+vx2 )*( vx1+vx2 )-cusq2))); - f_SW = f_SW + ((c1o1+drho) * (- c1o54* (c3o1*(-vxo1-vxo2 )+c9o2*(-vxo1-vxo2 )*(-vxo1-vxo2 )-cusq) + c1o54* (c3o1*(-vx1-vx2 )+c9o2*(-vx1-vx2 )*(-vx1-vx2 )-cusq2))); - f_SE = f_SE + ((c1o1+drho) * (- c1o54* (c3o1*( vxo1-vxo2 )+c9o2*( vxo1-vxo2 )*( vxo1-vxo2 )-cusq) + c1o54* (c3o1*( vx1-vx2 )+c9o2*( vx1-vx2 )*( vx1-vx2 )-cusq2))); - f_NW = f_NW + ((c1o1+drho) * (- c1o54* (c3o1*(-vxo1+vxo2 )+c9o2*(-vxo1+vxo2 )*(-vxo1+vxo2 )-cusq) + c1o54* (c3o1*(-vx1+vx2 )+c9o2*(-vx1+vx2 )*(-vx1+vx2 )-cusq2))); - f_TE = f_TE + ((c1o1+drho) * (- c1o54* (c3o1*( vxo1 +vxo3)+c9o2*( vxo1 +vxo3)*( vxo1 +vxo3)-cusq) + c1o54* (c3o1*( vx1 +vx3)+c9o2*( vx1 +vx3)*( vx1 +vx3)-cusq2))); - f_BW = f_BW + ((c1o1+drho) * (- c1o54* (c3o1*(-vxo1 -vxo3)+c9o2*(-vxo1 -vxo3)*(-vxo1 -vxo3)-cusq) + c1o54* (c3o1*(-vx1 -vx3)+c9o2*(-vx1 -vx3)*(-vx1 -vx3)-cusq2))); - f_BE = f_BE + ((c1o1+drho) * (- c1o54* (c3o1*( vxo1 -vxo3)+c9o2*( vxo1 -vxo3)*( vxo1 -vxo3)-cusq) + c1o54* (c3o1*( vx1 -vx3)+c9o2*( vx1 -vx3)*( vx1 -vx3)-cusq2))); - f_TW = f_TW + ((c1o1+drho) * (- c1o54* (c3o1*(-vxo1 +vxo3)+c9o2*(-vxo1 +vxo3)*(-vxo1 +vxo3)-cusq) + c1o54* (c3o1*(-vx1 +vx3)+c9o2*(-vx1 +vx3)*(-vx1 +vx3)-cusq2))); - f_TN = f_TN + ((c1o1+drho) * (- c1o54* (c3o1*( vxo2+vxo3)+c9o2*( vxo2+vxo3)*( vxo2+vxo3)-cusq) + c1o54* (c3o1*( vx2+vx3)+c9o2*( vx2+vx3)*( vx2+vx3)-cusq2))); - f_BS = f_BS + ((c1o1+drho) * (- c1o54* (c3o1*( -vxo2-vxo3)+c9o2*( -vxo2-vxo3)*( -vxo2-vxo3)-cusq) + c1o54* (c3o1*( -vx2-vx3)+c9o2*( -vx2-vx3)*( -vx2-vx3)-cusq2))); - f_BN = f_BN + ((c1o1+drho) * (- c1o54* (c3o1*( vxo2-vxo3)+c9o2*( vxo2-vxo3)*( vxo2-vxo3)-cusq) + c1o54* (c3o1*( vx2-vx3)+c9o2*( vx2-vx3)*( vx2-vx3)-cusq2))); - f_TS = f_TS + ((c1o1+drho) * (- c1o54* (c3o1*( -vxo2+vxo3)+c9o2*( -vxo2+vxo3)*( -vxo2+vxo3)-cusq) + c1o54* (c3o1*( -vx2+vx3)+c9o2*( -vx2+vx3)*( -vx2+vx3)-cusq2))); - f_TNE = f_TNE + ((c1o1+drho) * (- c1o216*(c3o1*( vxo1+vxo2+vxo3)+c9o2*( vxo1+vxo2+vxo3)*( vxo1+vxo2+vxo3)-cusq) + c1o216*(c3o1*( vx1+vx2+vx3)+c9o2*( vx1+vx2+vx3)*( vx1+vx2+vx3)-cusq2))); - f_BSW = f_BSW + ((c1o1+drho) * (- c1o216*(c3o1*(-vxo1-vxo2-vxo3)+c9o2*(-vxo1-vxo2-vxo3)*(-vxo1-vxo2-vxo3)-cusq) + c1o216*(c3o1*(-vx1-vx2-vx3)+c9o2*(-vx1-vx2-vx3)*(-vx1-vx2-vx3)-cusq2))); - f_BNE = f_BNE + ((c1o1+drho) * (- c1o216*(c3o1*( vxo1+vxo2-vxo3)+c9o2*( vxo1+vxo2-vxo3)*( vxo1+vxo2-vxo3)-cusq) + c1o216*(c3o1*( vx1+vx2-vx3)+c9o2*( vx1+vx2-vx3)*( vx1+vx2-vx3)-cusq2))); - f_TSW = f_TSW + ((c1o1+drho) * (- c1o216*(c3o1*(-vxo1-vxo2+vxo3)+c9o2*(-vxo1-vxo2+vxo3)*(-vxo1-vxo2+vxo3)-cusq) + c1o216*(c3o1*(-vx1-vx2+vx3)+c9o2*(-vx1-vx2+vx3)*(-vx1-vx2+vx3)-cusq2))); - f_TSE = f_TSE + ((c1o1+drho) * (- c1o216*(c3o1*( vxo1-vxo2+vxo3)+c9o2*( vxo1-vxo2+vxo3)*( vxo1-vxo2+vxo3)-cusq) + c1o216*(c3o1*( vx1-vx2+vx3)+c9o2*( vx1-vx2+vx3)*( vx1-vx2+vx3)-cusq2))); - f_BNW = f_BNW + ((c1o1+drho) * (- c1o216*(c3o1*(-vxo1+vxo2-vxo3)+c9o2*(-vxo1+vxo2-vxo3)*(-vxo1+vxo2-vxo3)-cusq) + c1o216*(c3o1*(-vx1+vx2-vx3)+c9o2*(-vx1+vx2-vx3)*(-vx1+vx2-vx3)-cusq2))); - f_BSE = f_BSE + ((c1o1+drho) * (- c1o216*(c3o1*( vxo1-vxo2-vxo3)+c9o2*( vxo1-vxo2-vxo3)*( vxo1-vxo2-vxo3)-cusq) + c1o216*(c3o1*( vx1-vx2-vx3)+c9o2*( vx1-vx2-vx3)*( vx1-vx2-vx3)-cusq2))); - f_TNW = f_TNW + ((c1o1+drho) * (- c1o216*(c3o1*(-vxo1+vxo2+vxo3)+c9o2*(-vxo1+vxo2+vxo3)*(-vxo1+vxo2+vxo3)-cusq) + c1o216*(c3o1*(-vx1+vx2+vx3)+c9o2*(-vx1+vx2+vx3)*(-vx1+vx2+vx3)-cusq2))); - - (D.f[d000])[kzero] = f_ZERO; - (D.f[dP00])[ke ] = f_E ; // f_W ;// - (D.f[dM00])[kw ] = f_W ; // f_E ;// - (D.f[d0P0])[kn ] = f_N ; // f_S ;// - (D.f[d0M0])[ks ] = f_S ; // f_N ;// - (D.f[d00P])[kt ] = f_T ; // f_B ;// - (D.f[d00M])[kb ] = f_B ; // f_T ;// - (D.f[dPP0])[kne ] = f_NE ; // f_SW ;// - (D.f[dMM0])[ksw ] = f_SW ; // f_NE ;// - (D.f[dPM0])[kse ] = f_SE ; // f_NW ;// - (D.f[dMP0])[knw ] = f_NW ; // f_SE ;// - (D.f[dP0P])[kte ] = f_TE ; // f_BW ;// - (D.f[dM0M])[kbw ] = f_BW ; // f_TE ;// - (D.f[dP0M])[kbe ] = f_BE ; // f_TW ;// - (D.f[dM0P])[ktw ] = f_TW ; // f_BE ;// - (D.f[d0PP])[ktn ] = f_TN ; // f_BS ;// - (D.f[d0MM])[kbs ] = f_BS ; // f_TN ;// - (D.f[d0PM])[kbn ] = f_BN ; // f_TS ;// - (D.f[d0MP])[kts ] = f_TS ; // f_BN ;// - (D.f[dPPP])[ktne ] = f_TNE ; // f_BSW ;// - (D.f[dMMM])[kbsw ] = f_BSW ; // f_BNE ;// - (D.f[dPPM])[kbne ] = f_BNE ; // f_BNW ;// - (D.f[dMMP])[ktsw ] = f_TSW ; // f_BSE ;// - (D.f[dPMP])[ktse ] = f_TSE ; // f_TSW ;// - (D.f[dMPM])[kbnw ] = f_BNW ; // f_TNE ;// - (D.f[dPMM])[kbse ] = f_BSE ; // f_TNW ;// - (D.f[dMPP])[ktnw ] = f_TNW ; // f_TSE ;// - - ////////////////////////////////////////////////////////////////////////// - ////(D.f[d000])[kzero] = c8over27* (drho-cu_sq); - //(D.f[dP00])[ke ] = three*c2over27* ( vx1 ); //six - //(D.f[dM00])[kw ] = three*c2over27* (-vx1 ); //six - //(D.f[d0P0])[kn ] = three*c2over27* ( vx2 ); //six - //(D.f[d0M0])[ks ] = three*c2over27* ( -vx2 ); //six - //(D.f[d00P])[kt ] = three*c2over27* ( vx3); //six - //(D.f[d00M])[kb ] = three*c2over27* ( -vx3); //six - //(D.f[dPP0])[kne ] = three*c1over54* ( vx1+vx2 ); //six - //(D.f[dMM0])[ksw ] = three*c1over54* (-vx1-vx2 ); //six - //(D.f[dPM0])[kse ] = three*c1over54* ( vx1-vx2 ); //six - //(D.f[dMP0])[knw ] = three*c1over54* (-vx1+vx2 ); //six - //(D.f[dP0P])[kte ] = three*c1over54* ( vx1 +vx3); //six - //(D.f[dM0M])[kbw ] = three*c1over54* (-vx1 -vx3); //six - //(D.f[dP0M])[kbe ] = three*c1over54* ( vx1 -vx3); //six - //(D.f[dM0P])[ktw ] = three*c1over54* (-vx1 +vx3); //six - //(D.f[d0PP])[ktn ] = three*c1over54* ( vx2+vx3); //six - //(D.f[d0MM])[kbs ] = three*c1over54* ( -vx2-vx3); //six - //(D.f[d0PM])[kbn ] = three*c1over54* ( vx2-vx3); //six - //(D.f[d0MP])[kts ] = three*c1over54* ( -vx2+vx3); //six - //(D.f[dPPP])[ktne ] = three*c1over216*( vx1+vx2+vx3); //six - //(D.f[dMMM])[kbsw ] = three*c1over216*(-vx1-vx2-vx3); //six - //(D.f[dPPM])[kbne ] = three*c1over216*( vx1+vx2-vx3); //six - //(D.f[dMMP])[ktsw ] = three*c1over216*(-vx1-vx2+vx3); //six - //(D.f[dPMP])[ktse ] = three*c1over216*( vx1-vx2+vx3); //six - //(D.f[dMPM])[kbnw ] = three*c1over216*(-vx1+vx2-vx3); //six - //(D.f[dPMM])[kbse ] = three*c1over216*( vx1-vx2-vx3); //six - //(D.f[dMPP])[ktnw ] = three*c1over216*(-vx1+vx2+vx3); //six - //(D.f[d000])[kzero] = c8over27* (drho-cu_sq); - //(D.f[dP00])[ke ] = c2over27* (drho+three*( vx1 )+c9over2*( vx1 )*( vx1 )-cu_sq); - //(D.f[dM00])[kw ] = c2over27* (drho+three*(-vx1 )+c9over2*(-vx1 )*(-vx1 )-cu_sq); - //(D.f[d0P0])[kn ] = c2over27* (drho+three*( vx2 )+c9over2*( vx2 )*( vx2 )-cu_sq); - //(D.f[d0M0])[ks ] = c2over27* (drho+three*( -vx2 )+c9over2*( -vx2 )*( -vx2 )-cu_sq); - //(D.f[d00P])[kt ] = c2over27* (drho+three*( vx3)+c9over2*( vx3)*( vx3)-cu_sq); - //(D.f[d00M])[kb ] = c2over27* (drho+three*( -vx3)+c9over2*( -vx3)*( -vx3)-cu_sq); - //(D.f[dPP0])[kne ] = c1over54* (drho+three*( vx1+vx2 )+c9over2*( vx1+vx2 )*( vx1+vx2 )-cu_sq); - //(D.f[dMM0])[ksw ] = c1over54* (drho+three*(-vx1-vx2 )+c9over2*(-vx1-vx2 )*(-vx1-vx2 )-cu_sq); - //(D.f[dPM0])[kse ] = c1over54* (drho+three*( vx1-vx2 )+c9over2*( vx1-vx2 )*( vx1-vx2 )-cu_sq); - //(D.f[dMP0])[knw ] = c1over54* (drho+three*(-vx1+vx2 )+c9over2*(-vx1+vx2 )*(-vx1+vx2 )-cu_sq); - //(D.f[dP0P])[kte ] = c1over54* (drho+three*( vx1 +vx3)+c9over2*( vx1 +vx3)*( vx1 +vx3)-cu_sq); - //(D.f[dM0M])[kbw ] = c1over54* (drho+three*(-vx1 -vx3)+c9over2*(-vx1 -vx3)*(-vx1 -vx3)-cu_sq); - //(D.f[dP0M])[kbe ] = c1over54* (drho+three*( vx1 -vx3)+c9over2*( vx1 -vx3)*( vx1 -vx3)-cu_sq); - //(D.f[dM0P])[ktw ] = c1over54* (drho+three*(-vx1 +vx3)+c9over2*(-vx1 +vx3)*(-vx1 +vx3)-cu_sq); - //(D.f[d0PP])[ktn ] = c1over54* (drho+three*( vx2+vx3)+c9over2*( vx2+vx3)*( vx2+vx3)-cu_sq); - //(D.f[d0MM])[kbs ] = c1over54* (drho+three*( -vx2-vx3)+c9over2*( -vx2-vx3)*( -vx2-vx3)-cu_sq); - //(D.f[d0PM])[kbn ] = c1over54* (drho+three*( vx2-vx3)+c9over2*( vx2-vx3)*( vx2-vx3)-cu_sq); - //(D.f[d0MP])[kts ] = c1over54* (drho+three*( -vx2+vx3)+c9over2*( -vx2+vx3)*( -vx2+vx3)-cu_sq); - //(D.f[dPPP])[ktne ] = c1over216*(drho+three*( vx1+vx2+vx3)+c9over2*( vx1+vx2+vx3)*( vx1+vx2+vx3)-cu_sq); - //(D.f[dMMM])[kbsw ] = c1over216*(drho+three*(-vx1-vx2-vx3)+c9over2*(-vx1-vx2-vx3)*(-vx1-vx2-vx3)-cu_sq); - //(D.f[dPPM])[kbne ] = c1over216*(drho+three*( vx1+vx2-vx3)+c9over2*( vx1+vx2-vx3)*( vx1+vx2-vx3)-cu_sq); - //(D.f[dMMP])[ktsw ] = c1over216*(drho+three*(-vx1-vx2+vx3)+c9over2*(-vx1-vx2+vx3)*(-vx1-vx2+vx3)-cu_sq); - //(D.f[dPMP])[ktse ] = c1over216*(drho+three*( vx1-vx2+vx3)+c9over2*( vx1-vx2+vx3)*( vx1-vx2+vx3)-cu_sq); - //(D.f[dMPM])[kbnw ] = c1over216*(drho+three*(-vx1+vx2-vx3)+c9over2*(-vx1+vx2-vx3)*(-vx1+vx2-vx3)-cu_sq); - //(D.f[dPMM])[kbse ] = c1over216*(drho+three*( vx1-vx2-vx3)+c9over2*( vx1-vx2-vx3)*( vx1-vx2-vx3)-cu_sq); - //(D.f[dMPP])[ktnw ] = c1over216*(drho+three*(-vx1+vx2+vx3)+c9over2*(-vx1+vx2+vx3)*(-vx1+vx2+vx3)-cu_sq); - } - } -} -////////////////////////////////////////////////////////////////////////// - - diff --git a/src/gpu/core/Init/PositionReader.cpp b/src/gpu/core/Init/PositionReader.cpp deleted file mode 100644 index c94ea46acb9aa1d9b86feb77ebbc58c2aa47e431..0000000000000000000000000000000000000000 --- a/src/gpu/core/Init/PositionReader.cpp +++ /dev/null @@ -1,214 +0,0 @@ -#include "PositionReader.h" - -#include "Parameter/Parameter.h" - -#include <basics/utilities/UbFileInputASCII.h> - -using namespace vf::lbm::dir; - -////////////////////////////////////////////////////////////////////////// -void PositionReader::readFilePropellerCylinderForAlloc(Parameter* para) -{ - UbFileInputASCII in(para->getpropellerCylinder()); - int test = 0, count = 0; - int maxlevel = in.readInteger(); - in.readLine(); - - for (int level = 0; level < maxlevel; level++) - { - para->getParH(level)->propellerBC.numberOfBCnodes = in.readInteger(); - para->getParD(level)->propellerBC.numberOfBCnodes = para->getParH(level)->propellerBC.numberOfBCnodes; - in.readLine(); - if (level == para->getFine()) - { - for(uint u=0; u<para->getParH(level)->propellerBC.numberOfBCnodes; u++) - { - test = in.readInteger(); - if (para->getParH(level)->typeOfGridNode[test] == GEO_FLUID) - { - count++; - } - //////////////////////////////////////////////////////////////////////// - //for(unsigned int ix3=0; ix3<para->getParH(level)->nz; ix3++) - //{ - // for(unsigned int ix2=0; ix2<para->getParH(level)->ny; ix2++) - // { - // for(unsigned int ix1=0; ix1<para->getParH(level)->nx; ix1++) - // { - // unsigned int m = para->getParH(level)->nx*(para->getParH(level)->ny*ix3 + ix2) + ix1; - // if (para->getParH(level)->k[m] == test) - // { - // if(para->getParH(level)->geo[m] == 1) - // { - // count++; - // } - // } - // } - // } - //} - //count++; - //////////////////////////////////////////////////////////////////////// - in.readDouble(); - in.readDouble(); - in.readDouble(); - in.readLine(); - } - } - else - { - for(uint u=0; u<para->getParH(level)->propellerBC.numberOfBCnodes; u++) - { - in.readInteger(); - in.readDouble(); - in.readDouble(); - in.readDouble(); - in.readLine(); - } - } - para->getParH(level)->propellerBC.numberOfBCnodes = count; - para->getParD(level)->propellerBC.numberOfBCnodes = para->getParH(level)->propellerBC.numberOfBCnodes; - } -} -////////////////////////////////////////////////////////////////////////// - - -////////////////////////////////////////////////////////////////////////// -void PositionReader::readFilePropellerCylinder(Parameter* para) -{ - UbFileInputASCII in(para->getpropellerCylinder()); - int test = 0, count = 0; - int maxlevel = in.readInteger(); - in.readLine(); - - for (int level = 0; level < maxlevel; level++) - { - int allnodes = in.readInteger(); - in.readLine(); - if (level == para->getFine()) - { - for(int u=0; u<allnodes; u++) - { - test = in.readInteger(); - //////////////////////////////////////////////////////////////////////// - if (para->getParH(level)->typeOfGridNode[test] == GEO_FLUID) - { - para->getParH(level)->propellerBC.k[count] = test; - para->getParH(level)->propellerBC.Vx[count] = (real)in.readDouble(); - para->getParH(level)->propellerBC.Vy[count] = (real)in.readDouble(); - para->getParH(level)->propellerBC.Vz[count] = (real)in.readDouble(); - para->getParH(level)->propellerBC.RhoBC[count] = 0.0f; - count++; - } - else - { - in.readDouble(); - in.readDouble(); - in.readDouble(); - } - //para->getParH(level)->propellerBC.k[count] = test; - //para->getParH(level)->propellerBC.Vx[count] = (real)in.readDouble(); - //para->getParH(level)->propellerBC.Vy[count] = (real)in.readDouble(); - //para->getParH(level)->propellerBC.Vz[count] = (real)in.readDouble(); - //para->getParH(level)->propellerBC.Vx[count] = 0.07f; - //para->getParH(level)->propellerBC.Vy[count] = 0.0f; - //para->getParH(level)->propellerBC.Vz[count] = 0.0f; - in.readLine(); - } - } - else - { - for(int u=0; u<allnodes; u++) - { - in.readInteger(); - in.readDouble(); - in.readDouble(); - in.readDouble(); - in.readLine(); - } - } - printf("allnodes = %d, count = %d\n", allnodes, count); - } -} -////////////////////////////////////////////////////////////////////////// - - -////////////////////////////////////////////////////////////////////////// -void PositionReader::definePropellerQs(Parameter* para) -{ - ////////////////////////////////////////////////////////////////// - //preprocessing - real* QQ = para->getParH(para->getFine())->propellerBC.q27[0]; - unsigned int sizeQ = para->getParH(para->getFine())->propellerBC.numberOfBCnodes; - QforBoundaryConditions Q; - Q.q27[dP00 ] = &QQ[dP00 *sizeQ]; - Q.q27[dM00 ] = &QQ[dM00 *sizeQ]; - Q.q27[d0P0 ] = &QQ[d0P0 *sizeQ]; - Q.q27[d0M0 ] = &QQ[d0M0 *sizeQ]; - Q.q27[d00P ] = &QQ[d00P *sizeQ]; - Q.q27[d00M ] = &QQ[d00M *sizeQ]; - Q.q27[dPP0 ] = &QQ[dPP0 *sizeQ]; - Q.q27[dMM0 ] = &QQ[dMM0 *sizeQ]; - Q.q27[dPM0 ] = &QQ[dPM0 *sizeQ]; - Q.q27[dMP0 ] = &QQ[dMP0 *sizeQ]; - Q.q27[dP0P ] = &QQ[dP0P *sizeQ]; - Q.q27[dM0M ] = &QQ[dM0M *sizeQ]; - Q.q27[dP0M ] = &QQ[dP0M *sizeQ]; - Q.q27[dM0P ] = &QQ[dM0P *sizeQ]; - Q.q27[d0PP ] = &QQ[d0PP *sizeQ]; - Q.q27[d0MM ] = &QQ[d0MM *sizeQ]; - Q.q27[d0PM ] = &QQ[d0PM *sizeQ]; - Q.q27[d0MP ] = &QQ[d0MP *sizeQ]; - Q.q27[d000] = &QQ[d000*sizeQ]; - Q.q27[dPPP ] = &QQ[dPPP *sizeQ]; - Q.q27[dMMP ] = &QQ[dMMP *sizeQ]; - Q.q27[dPMP ] = &QQ[dPMP *sizeQ]; - Q.q27[dMPP ] = &QQ[dMPP *sizeQ]; - Q.q27[dPPM ] = &QQ[dPPM *sizeQ]; - Q.q27[dMMM ] = &QQ[dMMM *sizeQ]; - Q.q27[dPMM ] = &QQ[dPMM *sizeQ]; - Q.q27[dMPM ] = &QQ[dMPM *sizeQ]; - ////////////////////////////////////////////////////////////////// - for(uint u=0; u<para->getParH(para->getFine())->propellerBC.numberOfBCnodes; u++) - { - for (size_t dir = dP00; dir<=dMMM; dir++) - { - if ((dir==dP00) || - (dir==dPP0) || (dir==dPM0) || (dir==dP0P) || (dir==dP0M) || - (dir==dPPP)|| (dir==dPPM)|| (dir==dPMP)|| (dir==dPMM)) - { - Q.q27[dir][u] = 1.0f; - } - else - { - Q.q27[dir][u] = -1.0f; - } - } - } - ////////////////////////////////////////////////////////////////// -} -////////////////////////////////////////////////////////////////////////// - - -////////////////////////////////////////////////////////////////////////// -void PositionReader::readMeasurePoints( Parameter* para ) -{ - UbFileInputASCII in(para->getmeasurePoints()); - int numberOfAllNodes = in.readInteger(); - in.readLine(); - int tempLevel; - MeasurePoints tempMP; - //printf("done, init the values...\n"); - for (int u = 0; u < numberOfAllNodes; u++) - { - tempMP.name = in.readString(); - //printf("done, read the name...\n"); - tempMP.k = in.readInteger(); - //printf("done, read k...\n"); - tempLevel = in.readInteger(); - //printf("done, read level...\n"); - in.readLine(); - //printf("done, read the values...\n"); - para->getParH(tempLevel)->MP.push_back(tempMP); - //printf("done, put it into a vector...\n"); - } -} diff --git a/src/gpu/core/Init/PositionReader.h b/src/gpu/core/Init/PositionReader.h deleted file mode 100644 index b12a17c8148f4b5b1d5b656e0c9ea04d88eee89f..0000000000000000000000000000000000000000 --- a/src/gpu/core/Init/PositionReader.h +++ /dev/null @@ -1,15 +0,0 @@ -#ifndef POSITION_READER_H -#define POSITION_READER_H - -class Parameter; - -class PositionReader -{ -public: - static void readFilePropellerCylinderForAlloc(Parameter* para); - static void readFilePropellerCylinder(Parameter* para); - static void definePropellerQs(Parameter* para); - static void readMeasurePoints(Parameter* para); -}; - -#endif \ No newline at end of file diff --git a/src/gpu/core/Init/VfReader.cpp b/src/gpu/core/Init/VfReader.cpp deleted file mode 100644 index 8a567200863cfa5ff78c24b04a619cfccf6fa036..0000000000000000000000000000000000000000 --- a/src/gpu/core/Init/VfReader.cpp +++ /dev/null @@ -1,74 +0,0 @@ -#include "Init/VfReader.h" - -#include "Parameter/Parameter.h" -#include "Init/PositionReader.h" -#include "GPU/CudaMemoryManager.h" - -//////////////////////////////////////////////////////////////////////////////// -void readPropellerCylinder(Parameter* para, CudaMemoryManager* cudaMemoryManager) -{ - PositionReader::readFilePropellerCylinderForAlloc(para); - - cudaMemoryManager->cudaAllocVeloPropeller(para->getFine()); - - PositionReader::readFilePropellerCylinder(para); - //PositionReader::definePropellerQs(para); - - cudaMemoryManager->cudaCopyVeloPropeller(para->getFine()); -} - -//////////////////////////////////////////////////////////////////////////////// -void readMeasurePoints(Parameter* para, CudaMemoryManager* cudaMemoryManager) -{ - //read measure points from file - PositionReader::readMeasurePoints(para); - //printf("done, reading the file...\n"); - //level loop - for (int lev = 0; lev <= para->getMaxLevel(); lev++) - { - //set Memory Size and malloc of the indices and macroscopic values per level - para->getParH(lev)->numberOfValuesMP = (unsigned int)para->getParH(lev)->MP.size()*(unsigned int)para->getclockCycleForMP()/((unsigned int)para->getTimestepForMP()); - para->getParD(lev)->numberOfValuesMP = para->getParH(lev)->numberOfValuesMP; - - para->getParH(lev)->numberOfPointskMP = (int)para->getParH(lev)->MP.size(); - para->getParD(lev)->numberOfPointskMP = para->getParH(lev)->numberOfPointskMP; - - para->getParH(lev)->memSizeIntkMP = sizeof(unsigned int)*(int)para->getParH(lev)->MP.size(); - para->getParD(lev)->memSizeIntkMP = para->getParH(lev)->memSizeIntkMP; - - para->getParH(lev)->memSizerealkMP = sizeof(real)*para->getParH(lev)->numberOfValuesMP; - para->getParD(lev)->memSizerealkMP = para->getParH(lev)->memSizerealkMP; - - printf("Level: %d, numberOfValuesMP: %d, memSizeIntkMP: %d, memSizerealkMP: %d\n",lev,para->getParH(lev)->numberOfValuesMP,para->getParH(lev)->memSizeIntkMP, para->getParD(lev)->memSizerealkMP); - - cudaMemoryManager->cudaAllocMeasurePointsIndex(lev); - - //loop over all measure points per level - for(int index = 0; index < (int)para->getParH(lev)->MP.size(); index++) - { - //set indices - para->getParH(lev)->kMP[index] = para->getParH(lev)->MP[index].k; - } - //loop over all measure points per level times MPClockCycle - for(int index = 0; index < (int)para->getParH(lev)->numberOfValuesMP; index++) - { - //init values - para->getParH(lev)->VxMP[index] = (real)0.0; - para->getParH(lev)->VyMP[index] = (real)0.0; - para->getParH(lev)->VzMP[index] = (real)0.0; - para->getParH(lev)->RhoMP[index] = (real)0.0; - } - - //copy indices-arrays - cudaMemoryManager->cudaCopyMeasurePointsIndex(lev); - } -} -//////////////////////////////////////////////////////////////////////////////// - - - - - - - - diff --git a/src/gpu/core/Init/VfReader.h b/src/gpu/core/Init/VfReader.h deleted file mode 100644 index f8f792375f58edd61b9ec4e876245ad3a334dd81..0000000000000000000000000000000000000000 --- a/src/gpu/core/Init/VfReader.h +++ /dev/null @@ -1,13 +0,0 @@ -#ifndef VF_READER_H -#define VF_READER_H - - -class CudaMemoryManager; -class Parameter; - - -void readPropellerCylinder(Parameter* para, CudaMemoryManager* cudaMemoryManager); - -void readMeasurePoints(Parameter* para, CudaMemoryManager* cudaMemoryManager); - -#endif diff --git a/src/gpu/core/LBM/Simulation.cpp b/src/gpu/core/LBM/Simulation.cpp index 3566c26f8849e586bcdadbc6c916f0815b8d54cf..93896f1f8065cd63b81499c1e810619d8128024c 100644 --- a/src/gpu/core/LBM/Simulation.cpp +++ b/src/gpu/core/LBM/Simulation.cpp @@ -25,8 +25,8 @@ #include "Utilities/Buffer2D.hpp" #include "StringUtilities/StringUtil.h" ////////////////////////////////////////////////////////////////////////// -#include "Init/InitLattice.h" -#include "Init/VfReader.h" +#include "PreProcessor/InitLattice.h" +#include "PreProcessor/ReaderMeasurePoints.h" ////////////////////////////////////////////////////////////////////////// #include "FindQ/FindQ.h" #include "FindQ/DefineBCs.h" @@ -245,7 +245,7 @@ void Simulation::init(GridProvider &gridProvider, BoundaryConditionFactory *bcFa ////////////////////////////////////////////////////////////////////////// if (para->getUseMeasurePoints()) { VF_LOG_INFO("read measure points"); - readMeasurePoints(para.get(), cudaMemoryManager.get()); + ReaderMeasurePoints::readMeasurePoints(para.get(), cudaMemoryManager.get()); } ////////////////////////////////////////////////////////////////////////// diff --git a/src/gpu/core/Parameter/Parameter.cpp b/src/gpu/core/Parameter/Parameter.cpp index 4177aa2fc9bb5475f025700c8217a36ab4387eca..14aacb94a6da732d581b9820b2d34390b1472aff 100644 --- a/src/gpu/core/Parameter/Parameter.cpp +++ b/src/gpu/core/Parameter/Parameter.cpp @@ -397,7 +397,6 @@ void Parameter::initGridPaths(){ this->setnumberNodes(gridPath + "numberNodes.dat"); this->setLBMvsSI(gridPath + "LBMvsSI.dat"); this->setmeasurePoints(gridPath + "measurePoints.dat"); - this->setpropellerValues(gridPath + "propellerValues.dat"); this->setcpTop(gridPath + "cpTop.dat"); this->setcpBottom(gridPath + "cpBottom.dat"); this->setcpBottom2(gridPath + "cpBottom2.dat"); @@ -957,10 +956,6 @@ void Parameter::setIsOutflowNormal(bool isOutflowNormal) { this->isOutflowNormal = isOutflowNormal; } -void Parameter::setIsProp(bool isProp) -{ - this->isProp = isProp; -} void Parameter::setIsCp(bool isCp) { this->isCp = isCp; @@ -1288,18 +1283,6 @@ void Parameter::setperiodicBcValues(std::string periodicBcValues) { this->periodicBcValues = periodicBcValues; } -void Parameter::setpropellerQs(std::string propellerQs) -{ - this->propellerQs = propellerQs; -} -void Parameter::setpropellerValues(std::string propellerValues) -{ - this->propellerValues = propellerValues; -} -void Parameter::setpropellerCylinder(std::string propellerCylinder) -{ - this->propellerCylinder = propellerCylinder; -} void Parameter::setmeasurePoints(std::string measurePoints) { this->measurePoints = measurePoints; @@ -1356,8 +1339,6 @@ void Parameter::setObj(std::string str, bool isObj) { if (str == "geo") { this->setIsGeo(isObj); - } else if (str == "prop") { - this->setIsProp(isObj); } else if (str == "cp") { this->setIsCp(isObj); } else if (str == "geoNormal") { @@ -2246,18 +2227,6 @@ std::string Parameter::getperiodicBcValues() { return this->periodicBcValues; } -std::string Parameter::getpropellerQs() -{ - return this->propellerQs; -} -std::string Parameter::getpropellerValues() -{ - return this->propellerValues; -} -std::string Parameter::getpropellerCylinder() -{ - return this->propellerCylinder; -} std::string Parameter::getmeasurePoints() { return this->measurePoints; @@ -2409,10 +2378,6 @@ bool Parameter::getCalcHighOrderMoments() { return this->isHighOrderMoments; } -bool Parameter::getIsProp() -{ - return this->isProp; -} bool Parameter::overWritingRestart(uint t) { return t == getTimeDoRestart(); diff --git a/src/gpu/core/Parameter/Parameter.h b/src/gpu/core/Parameter/Parameter.h index cc062e2d96b5e049c7a2e692860da026e49ab5b5..35c638878fbd235480132e363df0df0e8e858c08 100644 --- a/src/gpu/core/Parameter/Parameter.h +++ b/src/gpu/core/Parameter/Parameter.h @@ -373,7 +373,6 @@ struct LBMSimulationParameter { QforBoundaryConditions QInlet, QOutlet, QPeriodic; // DEPRECATED BCs that are not used any more unsigned int kInletQread, kOutletQread; // DEPRECATED - QforBoundaryConditions propellerBC; // DEPRECATED QforBoundaryConditions geometryBCnormalX, geometryBCnormalY, geometryBCnormalZ; // DEPRECATED QforBoundaryConditions inflowBCnormalX, inflowBCnormalY, inflowBCnormalZ; // DEPRECATED QforBoundaryConditions outflowBCnormalX, outflowBCnormalY, outflowBCnormalZ; // DEPRECATED @@ -571,9 +570,6 @@ public: void setwallBcValues(std::string wallBcValues); void setperiodicBcQs(std::string periodicBcQs); void setperiodicBcValues(std::string periodicBcValues); - void setpropellerCylinder(std::string propellerCylinder); - void setpropellerValues(std::string propellerValues); - void setpropellerQs(std::string propellerQs); void setmeasurePoints(std::string measurePoints); void setnumberNodes(std::string numberNodes); void setLBMvsSI(std::string LBMvsSI); @@ -599,7 +595,6 @@ public: void setIsGeoNormal(bool isGeoNormal); void setIsInflowNormal(bool isInflowNormal); void setIsOutflowNormal(bool isOutflowNormal); - void setIsProp(bool isProp); void setIsCp(bool isCp); void setConcFile(bool concFile); void setUseMeasurePoints(bool useMeasurePoints); @@ -785,9 +780,6 @@ public: std::string getwallBcValues(); std::string getperiodicBcQs(); std::string getperiodicBcValues(); - std::string getpropellerQs(); - std::string getpropellerCylinder(); - std::string getpropellerValues(); std::string getmeasurePoints(); std::string getnumberNodes(); std::string getLBMvsSI(); @@ -885,7 +877,6 @@ public: bool getIsGeoNormal(); bool getIsInflowNormal(); bool getIsOutflowNormal(); - bool getIsProp(); bool getIsCp(); bool getIsGeometryValues(); bool getCalc2ndOrderMoments(); @@ -1003,7 +994,6 @@ private: bool doCheckPoint{ false }; bool readGeo{ false }; bool isGeo; - bool isProp; bool isCp; bool GeometryValues{ false }; bool is2ndOrderMoments{ false }; @@ -1060,7 +1050,7 @@ private: std::string pressBcPos, pressBcQs, pressBcValue; std::string geomBoundaryBcQs, velBcQs; std::string geomBoundaryBcValues, velBcValues, pressBcValues, noSlipBcValues; - std::string propellerCylinder, propellerValues, propellerQs, measurePoints; + std::string measurePoints; std::string inletBcQs, inletBcValues; std::string outletBcQs, outletBcValues; std::string topBcQs, topBcValues; diff --git a/src/gpu/core/Init/InitLattice.cpp b/src/gpu/core/PreProcessor/InitLattice.cpp similarity index 99% rename from src/gpu/core/Init/InitLattice.cpp rename to src/gpu/core/PreProcessor/InitLattice.cpp index f8e00006e3b6438080c23b3a24bb9d30aa17fe48..1702e4e44fa6b04f59cb517d92df61f35104c9ba 100644 --- a/src/gpu/core/Init/InitLattice.cpp +++ b/src/gpu/core/PreProcessor/InitLattice.cpp @@ -30,7 +30,7 @@ //! \ingroup Init //! \author Martin Schoenherr //======================================================================================= -#include "Init/InitLattice.h" +#include "PreProcessor/InitLattice.h" #include "GPU/CudaMemoryManager.h" #include "GPU/GPU_Interface.h" diff --git a/src/gpu/core/Init/InitLattice.h b/src/gpu/core/PreProcessor/InitLattice.h similarity index 100% rename from src/gpu/core/Init/InitLattice.h rename to src/gpu/core/PreProcessor/InitLattice.h diff --git a/src/gpu/core/PreProcessor/PreProcessorFactory/PreProcessorFactoryImp.cpp b/src/gpu/core/PreProcessor/PreProcessorFactory/PreProcessorFactoryImp.cpp index 18a1d4e7023e92cc0ac3be5dc2a57d1de64ac5cc..10fdf377bfb738f48d8083d40aedb6db6e2f627e 100644 --- a/src/gpu/core/PreProcessor/PreProcessorFactory/PreProcessorFactoryImp.cpp +++ b/src/gpu/core/PreProcessor/PreProcessorFactory/PreProcessorFactoryImp.cpp @@ -32,12 +32,9 @@ #include "PreProcessor/PreProcessorImp.h" -#include "PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressibleD3Q7/InitAdvectionDiffusionCompressibleD3Q7.h" #include "PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressible/InitAdvectionDiffusionCompressible.h" #include "PreProcessor/PreProcessorStrategy/InitNavierStokesCompressible/InitNavierStokesCompressible.h" -#include "PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible.h" #include "PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressible/InitAdvectionDiffusionIncompressible.h" -#include "PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressibleD3Q7/InitAdvectionDiffusionIncompressibleD3Q7.h" #include "PreProcessor/PreProcessorStrategy/InitNavierStokesIncompressible/InitNavierStokesIncompressible.h" @@ -61,18 +58,9 @@ std::shared_ptr<PreProcessorStrategy> PreProcessorFactoryImp::makePreProcessorSt case InitNavierStokesCompressible: return InitNavierStokesCompressible::getNewInstance(para); break; - case InitK18K20NavierStokesCompressible: - return InitK18K20NavierStokesCompressible::getNewInstance(para); - break; - case InitAdvectionDiffusionIncompressibleD3Q7: - return InitAdvectionDiffusionIncompressibleD3Q7::getNewInstance(para); - break; case InitAdvectionDiffusionIncompressible: return InitAdvectionDiffusionIncompressible::getNewInstance(para); break; - case InitAdvectionDiffusionCompressibleD3Q7: - return InitAdvectionDiffusionCompressibleD3Q7::getNewInstance(para); - break; case InitAdvectionDiffusionCompressible: return InitAdvectionDiffusionCompressible::getNewInstance(para); break; diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressibleD3Q7/InitAdvectionDiffusionCompressibleD3Q7.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressibleD3Q7/InitAdvectionDiffusionCompressibleD3Q7.cu deleted file mode 100644 index 2db77de45d67a0f74d80c50b6d875aa39d7b2d6d..0000000000000000000000000000000000000000 --- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressibleD3Q7/InitAdvectionDiffusionCompressibleD3Q7.cu +++ /dev/null @@ -1,73 +0,0 @@ -//======================================================================================= -// ____ ____ __ ______ __________ __ __ __ __ -// \ \ | | | | | _ \ |___ ___| | | | | / \ | | -// \ \ | | | | | |_) | | | | | | | / \ | | -// \ \ | | | | | _ / | | | | | | / /\ \ | | -// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____ -// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______| -// \ \ | | ________________________________________________________________ -// \ \ | | | ______________________________________________________________| -// \ \| | | | __ __ __ __ ______ _______ -// \ | | |_____ | | | | | | | | | _ \ / _____) -// \ | | _____| | | | | | | | | | | \ \ \_______ -// \ | | | | |_____ | \_/ | | | | |_/ / _____ | -// \ _____| |__| |________| \_______/ |__| |______/ (_______/ -// -// This file is part of VirtualFluids. VirtualFluids is free software: you can -// redistribute it and/or modify it under the terms of the GNU General Public -// License as published by the Free Software Foundation, either version 3 of -// the License, or (at your option) any later version. -// -// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT -// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or -// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License -// for more details. -// -// You should have received a copy of the GNU General Public License along -// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>. -// -//! \author Martin Schoenherr -//======================================================================================= -#include "InitAdvectionDiffusionCompressibleD3Q7.h" - -#include "InitAdvectionDiffusionCompressibleD3Q7_Device.cuh" -#include "Parameter/Parameter.h" -#include <cuda_helper/CudaGrid.h> - -std::shared_ptr<InitAdvectionDiffusionCompressibleD3Q7> InitAdvectionDiffusionCompressibleD3Q7::getNewInstance(std::shared_ptr<Parameter> para) -{ - return std::shared_ptr<InitAdvectionDiffusionCompressibleD3Q7>(new InitAdvectionDiffusionCompressibleD3Q7(para)); -} - -void InitAdvectionDiffusionCompressibleD3Q7::init(int level) -{ - vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes); - - InitAdvectionDiffusionCompressibleD3Q7_Device <<< grid.grid, grid.threads >>>( - para->getParD(level)->neighborX, - para->getParD(level)->neighborY, - para->getParD(level)->neighborZ, - para->getParD(level)->typeOfGridNode, - para->getParD(level)->concentration, - para->getParD(level)->velocityX, - para->getParD(level)->velocityY, - para->getParD(level)->velocityZ, - para->getParD(level)->numberOfNodes, - para->getParD(level)->distributionsAD7.f[0], - para->getParD(level)->isEvenTimestep); - getLastCudaError("InitAdvectionDiffusionCompressibleD3Q7_Device execution failed"); -} - -bool InitAdvectionDiffusionCompressibleD3Q7::checkParameter() -{ - return false; -} - -InitAdvectionDiffusionCompressibleD3Q7::InitAdvectionDiffusionCompressibleD3Q7(std::shared_ptr<Parameter> para) -{ - this->para = para; -} - -InitAdvectionDiffusionCompressibleD3Q7::InitAdvectionDiffusionCompressibleD3Q7() -{ -} diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressibleD3Q7/InitAdvectionDiffusionCompressibleD3Q7.h b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressibleD3Q7/InitAdvectionDiffusionCompressibleD3Q7.h deleted file mode 100644 index 2c68fca2fde9d70075ceb1dc9f0e0f238e39fb40..0000000000000000000000000000000000000000 --- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressibleD3Q7/InitAdvectionDiffusionCompressibleD3Q7.h +++ /dev/null @@ -1,52 +0,0 @@ -//======================================================================================= -// ____ ____ __ ______ __________ __ __ __ __ -// \ \ | | | | | _ \ |___ ___| | | | | / \ | | -// \ \ | | | | | |_) | | | | | | | / \ | | -// \ \ | | | | | _ / | | | | | | / /\ \ | | -// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____ -// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______| -// \ \ | | ________________________________________________________________ -// \ \ | | | ______________________________________________________________| -// \ \| | | | __ __ __ __ ______ _______ -// \ | | |_____ | | | | | | | | | _ \ / _____) -// \ | | _____| | | | | | | | | | | \ \ \_______ -// \ | | | | |_____ | \_/ | | | | |_/ / _____ | -// \ _____| |__| |________| \_______/ |__| |______/ (_______/ -// -// This file is part of VirtualFluids. VirtualFluids is free software: you can -// redistribute it and/or modify it under the terms of the GNU General Public -// License as published by the Free Software Foundation, either version 3 of -// the License, or (at your option) any later version. -// -// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT -// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or -// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License -// for more details. -// -// You should have received a copy of the GNU General Public License along -// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>. -// -//! \author Martin Schoenherr -//======================================================================================= -#ifndef InitAdvectionDiffusionCompressibleD3Q7_H -#define InitAdvectionDiffusionCompressibleD3Q7_H - -#include "PreProcessor/PreProcessorStrategy/PreProcessorStrategy.h" - -#include <memory> - -class Parameter; - -class InitAdvectionDiffusionCompressibleD3Q7 : public PreProcessorStrategy -{ -public: - static std::shared_ptr<InitAdvectionDiffusionCompressibleD3Q7> getNewInstance(std::shared_ptr< Parameter> para); - void init(int level); - bool checkParameter(); - -private: - InitAdvectionDiffusionCompressibleD3Q7(); - InitAdvectionDiffusionCompressibleD3Q7(std::shared_ptr< Parameter> para); - std::shared_ptr< Parameter> para; -}; -#endif \ No newline at end of file diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressibleD3Q7/InitAdvectionDiffusionCompressibleD3Q7_Device.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressibleD3Q7/InitAdvectionDiffusionCompressibleD3Q7_Device.cu deleted file mode 100644 index b9265e83bb1202a1126a1f2bb342bae3f10e842a..0000000000000000000000000000000000000000 --- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressibleD3Q7/InitAdvectionDiffusionCompressibleD3Q7_Device.cu +++ /dev/null @@ -1,125 +0,0 @@ -//======================================================================================= -// ____ ____ __ ______ __________ __ __ __ __ -// \ \ | | | | | _ \ |___ ___| | | | | / \ | | -// \ \ | | | | | |_) | | | | | | | / \ | | -// \ \ | | | | | _ / | | | | | | / /\ \ | | -// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____ -// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______| -// \ \ | | ________________________________________________________________ -// \ \ | | | ______________________________________________________________| -// \ \| | | | __ __ __ __ ______ _______ -// \ | | |_____ | | | | | | | | | _ \ / _____) -// \ | | _____| | | | | | | | | | | \ \ \_______ -// \ | | | | |_____ | \_/ | | | | |_/ / _____ | -// \ _____| |__| |________| \_______/ |__| |______/ (_______/ -// -// This file is part of VirtualFluids. VirtualFluids is free software: you can -// redistribute it and/or modify it under the terms of the GNU General Public -// License as published by the Free Software Foundation, either version 3 of -// the License, or (at your option) any later version. -// -// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT -// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or -// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License -// for more details. -// -// You should have received a copy of the GNU General Public License along -// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>. -// -//! \author Martin Schoenherr -//======================================================================================= -#include "LBM/LB.h" -#include "lbm/constants/D3Q27.h" -#include <basics/constants/NumericConstants.h> - -using namespace vf::basics::constant; -using namespace vf::lbm::dir; -#include "math.h" - -__global__ void InitAdvectionDiffusionCompressibleD3Q7_Device(unsigned int* neighborX, - unsigned int* neighborY, - unsigned int* neighborZ, - unsigned int* geoD, - real* Conc, - real* ux, - real* uy, - real* uz, - unsigned int size_Mat, - real* DD7, - bool EvenOrOdd) -{ - //////////////////////////////////////////////////////////////////////////////// - const unsigned x = threadIdx.x; // Globaler x-Index - const unsigned y = blockIdx.x; // Globaler y-Index - const unsigned z = blockIdx.y; // Globaler z-Index - - const unsigned nx = blockDim.x; - const unsigned ny = gridDim.x; - - const unsigned k = nx*(ny*z + y) + x; - ////////////////////////////////////////////////////////////////////////// - - if (k<size_Mat) - { - //////////////////////////////////////////////////////////////////////////////// - unsigned int BC; - BC = geoD[k]; - - if (BC != GEO_SOLID && BC != GEO_VOID) - { - Distributions7 D7; - if (EvenOrOdd == true) - { - D7.f[0] = &DD7[0 * size_Mat]; - D7.f[1] = &DD7[1 * size_Mat]; - D7.f[2] = &DD7[2 * size_Mat]; - D7.f[3] = &DD7[3 * size_Mat]; - D7.f[4] = &DD7[4 * size_Mat]; - D7.f[5] = &DD7[5 * size_Mat]; - D7.f[6] = &DD7[6 * size_Mat]; - } - else - { - D7.f[0] = &DD7[0 * size_Mat]; - D7.f[2] = &DD7[1 * size_Mat]; - D7.f[1] = &DD7[2 * size_Mat]; - D7.f[4] = &DD7[3 * size_Mat]; - D7.f[3] = &DD7[4 * size_Mat]; - D7.f[6] = &DD7[5 * size_Mat]; - D7.f[5] = &DD7[6 * size_Mat]; - } - ////////////////////////////////////////////////////////////////////////// - real ConcD = Conc[k]; - real vx1 = ux[k]; - real vx2 = uy[k]; - real vx3 = uz[k]; - real lambdaD = -c3o1 + sqrt(c3o1); - real Diffusivity = c1o20; - real Lam = -(c1o2 + c1o1 / lambdaD); - real nue_d = Lam / c3o1; - real ae = Diffusivity / nue_d - c1o1; - real ux_sq = vx1 * vx1; - real uy_sq = vx2 * vx2; - real uz_sq = vx3 * vx3; - ////////////////////////////////////////////////////////////////////////// - //index - ////////////////////////////////////////////////////////////////////////// - unsigned int kzero = k; - unsigned int ke = k; - unsigned int kw = neighborX[k]; - unsigned int kn = k; - unsigned int ks = neighborY[k]; - unsigned int kt = k; - unsigned int kb = neighborZ[k]; - ////////////////////////////////////////////////////////////////////////// - - (D7.f[0])[kzero] = ConcD*(c1o3*(ae*(-c3o1)) - (ux_sq + uy_sq + uz_sq)); - (D7.f[1])[ke] = ConcD*(c1o6*(ae + c1o1) + c1o2*(ux_sq)+vx1*c1o2); - (D7.f[2])[kw] = ConcD*(c1o6*(ae + c1o1) + c1o2*(ux_sq)-vx1*c1o2); - (D7.f[3])[kn] = ConcD*(c1o6*(ae + c1o1) + c1o2*(uy_sq)+vx2*c1o2); - (D7.f[4])[ks] = ConcD*(c1o6*(ae + c1o1) + c1o2*(uy_sq)-vx2*c1o2); - (D7.f[5])[kt] = ConcD*(c1o6*(ae + c1o1) + c1o2*(uz_sq)+vx3*c1o2); - (D7.f[6])[kb] = ConcD*(c1o6*(ae + c1o1) + c1o2*(uz_sq)-vx3*c1o2); - } - } -} \ No newline at end of file diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressibleD3Q7/InitAdvectionDiffusionCompressibleD3Q7_Device.cuh b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressibleD3Q7/InitAdvectionDiffusionCompressibleD3Q7_Device.cuh deleted file mode 100644 index fca4430fe935cc082692d2da837d7e7ea3fe6bec..0000000000000000000000000000000000000000 --- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionCompressibleD3Q7/InitAdvectionDiffusionCompressibleD3Q7_Device.cuh +++ /dev/null @@ -1,50 +0,0 @@ -//======================================================================================= -// ____ ____ __ ______ __________ __ __ __ __ -// \ \ | | | | | _ \ |___ ___| | | | | / \ | | -// \ \ | | | | | |_) | | | | | | | / \ | | -// \ \ | | | | | _ / | | | | | | / /\ \ | | -// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____ -// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______| -// \ \ | | ________________________________________________________________ -// \ \ | | | ______________________________________________________________| -// \ \| | | | __ __ __ __ ______ _______ -// \ | | |_____ | | | | | | | | | _ \ / _____) -// \ | | _____| | | | | | | | | | | \ \ \_______ -// \ | | | | |_____ | \_/ | | | | |_/ / _____ | -// \ _____| |__| |________| \_______/ |__| |______/ (_______/ -// -// This file is part of VirtualFluids. VirtualFluids is free software: you can -// redistribute it and/or modify it under the terms of the GNU General Public -// License as published by the Free Software Foundation, either version 3 of -// the License, or (at your option) any later version. -// -// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT -// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or -// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License -// for more details. -// -// You should have received a copy of the GNU General Public License along -// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>. -// -//! \author Martin Schoenherr -//======================================================================================= -#ifndef InitAdvectionDiffusionCompressibleD3Q7_Device_H -#define InitAdvectionDiffusionCompressibleD3Q7_Device_H - -#include <DataTypes.h> -#include <curand.h> - -__global__ void InitAdvectionDiffusionCompressibleD3Q7_Device( - unsigned int* neighborX, - unsigned int* neighborY, - unsigned int* neighborZ, - unsigned int* geoD, - real* Conc, - real* ux, - real* uy, - real* uz, - unsigned int size_Mat, - real* DD7, - bool EvenOrOdd); - -#endif \ No newline at end of file diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressibleD3Q7/InitAdvectionDiffusionIncompressibleD3Q7.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressibleD3Q7/InitAdvectionDiffusionIncompressibleD3Q7.cu deleted file mode 100644 index 94d7da3bcdc6296bfaaa03c5ce81508c5c3c53e2..0000000000000000000000000000000000000000 --- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressibleD3Q7/InitAdvectionDiffusionIncompressibleD3Q7.cu +++ /dev/null @@ -1,73 +0,0 @@ -//======================================================================================= -// ____ ____ __ ______ __________ __ __ __ __ -// \ \ | | | | | _ \ |___ ___| | | | | / \ | | -// \ \ | | | | | |_) | | | | | | | / \ | | -// \ \ | | | | | _ / | | | | | | / /\ \ | | -// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____ -// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______| -// \ \ | | ________________________________________________________________ -// \ \ | | | ______________________________________________________________| -// \ \| | | | __ __ __ __ ______ _______ -// \ | | |_____ | | | | | | | | | _ \ / _____) -// \ | | _____| | | | | | | | | | | \ \ \_______ -// \ | | | | |_____ | \_/ | | | | |_/ / _____ | -// \ _____| |__| |________| \_______/ |__| |______/ (_______/ -// -// This file is part of VirtualFluids. VirtualFluids is free software: you can -// redistribute it and/or modify it under the terms of the GNU General Public -// License as published by the Free Software Foundation, either version 3 of -// the License, or (at your option) any later version. -// -// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT -// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or -// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License -// for more details. -// -// You should have received a copy of the GNU General Public License along -// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>. -// -//! \author Martin Schoenherr -//======================================================================================= -#include "InitAdvectionDiffusionIncompressibleD3Q7.h" - -#include "InitAdvectionDiffusionIncompressibleD3Q7_Device.cuh" -#include "Parameter/Parameter.h" -#include <cuda_helper/CudaGrid.h> - -std::shared_ptr<PreProcessorStrategy> InitAdvectionDiffusionIncompressibleD3Q7::getNewInstance(std::shared_ptr<Parameter> para) -{ - return std::shared_ptr<PreProcessorStrategy>(new InitAdvectionDiffusionIncompressibleD3Q7(para)); -} - -void InitAdvectionDiffusionIncompressibleD3Q7::init(int level) -{ - vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes); - - InitAdvectionDiffusionIncompressibleD3Q7_Device <<< grid.grid, grid.threads >>>( - para->getParD(level)->neighborX, - para->getParD(level)->neighborY, - para->getParD(level)->neighborZ, - para->getParD(level)->typeOfGridNode, - para->getParD(level)->concentration, - para->getParD(level)->velocityX, - para->getParD(level)->velocityY, - para->getParD(level)->velocityZ, - para->getParD(level)->numberOfNodes, - para->getParD(level)->distributionsAD.f[0], - para->getParD(level)->isEvenTimestep); - getLastCudaError("InitAdvectionDiffusionIncompressibleD3Q7_Device execution failed"); -} - -bool InitAdvectionDiffusionIncompressibleD3Q7::checkParameter() -{ - return false; -} - -InitAdvectionDiffusionIncompressibleD3Q7::InitAdvectionDiffusionIncompressibleD3Q7(std::shared_ptr<Parameter> para) -{ - this->para = para; -} - -InitAdvectionDiffusionIncompressibleD3Q7::InitAdvectionDiffusionIncompressibleD3Q7() -{ -} diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressibleD3Q7/InitAdvectionDiffusionIncompressibleD3Q7.h b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressibleD3Q7/InitAdvectionDiffusionIncompressibleD3Q7.h deleted file mode 100644 index 4ece3ee7dd8c889ec373bcc08a5e47c255b7be5b..0000000000000000000000000000000000000000 --- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressibleD3Q7/InitAdvectionDiffusionIncompressibleD3Q7.h +++ /dev/null @@ -1,53 +0,0 @@ -//======================================================================================= -// ____ ____ __ ______ __________ __ __ __ __ -// \ \ | | | | | _ \ |___ ___| | | | | / \ | | -// \ \ | | | | | |_) | | | | | | | / \ | | -// \ \ | | | | | _ / | | | | | | / /\ \ | | -// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____ -// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______| -// \ \ | | ________________________________________________________________ -// \ \ | | | ______________________________________________________________| -// \ \| | | | __ __ __ __ ______ _______ -// \ | | |_____ | | | | | | | | | _ \ / _____) -// \ | | _____| | | | | | | | | | | \ \ \_______ -// \ | | | | |_____ | \_/ | | | | |_/ / _____ | -// \ _____| |__| |________| \_______/ |__| |______/ (_______/ -// -// This file is part of VirtualFluids. VirtualFluids is free software: you can -// redistribute it and/or modify it under the terms of the GNU General Public -// License as published by the Free Software Foundation, either version 3 of -// the License, or (at your option) any later version. -// -// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT -// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or -// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License -// for more details. -// -// You should have received a copy of the GNU General Public License along -// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>. -// -//! \author Martin Schoenherr -//======================================================================================= -#ifndef InitAdvectionDiffusionIncompressibleD3Q7_H -#define InitAdvectionDiffusionIncompressibleD3Q7_H - -#include "PreProcessor/PreProcessorStrategy/PreProcessorStrategy.h" - -#include <memory> - -class Parameter; - -class InitAdvectionDiffusionIncompressibleD3Q7 : public PreProcessorStrategy -{ -public: - static std::shared_ptr<PreProcessorStrategy> getNewInstance(std::shared_ptr< Parameter> para); - void init(int level); - bool checkParameter(); - -private: - InitAdvectionDiffusionIncompressibleD3Q7(); - InitAdvectionDiffusionIncompressibleD3Q7(std::shared_ptr< Parameter> para); - std::shared_ptr< Parameter> para; -}; - -#endif \ No newline at end of file diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressibleD3Q7/InitAdvectionDiffusionIncompressibleD3Q7_Device.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressibleD3Q7/InitAdvectionDiffusionIncompressibleD3Q7_Device.cu deleted file mode 100644 index 441e8c7adcc18d891615d0374fe5bfaa56d92652..0000000000000000000000000000000000000000 --- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressibleD3Q7/InitAdvectionDiffusionIncompressibleD3Q7_Device.cu +++ /dev/null @@ -1,126 +0,0 @@ -//======================================================================================= -// ____ ____ __ ______ __________ __ __ __ __ -// \ \ | | | | | _ \ |___ ___| | | | | / \ | | -// \ \ | | | | | |_) | | | | | | | / \ | | -// \ \ | | | | | _ / | | | | | | / /\ \ | | -// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____ -// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______| -// \ \ | | ________________________________________________________________ -// \ \ | | | ______________________________________________________________| -// \ \| | | | __ __ __ __ ______ _______ -// \ | | |_____ | | | | | | | | | _ \ / _____) -// \ | | _____| | | | | | | | | | | \ \ \_______ -// \ | | | | |_____ | \_/ | | | | |_/ / _____ | -// \ _____| |__| |________| \_______/ |__| |______/ (_______/ -// -// This file is part of VirtualFluids. VirtualFluids is free software: you can -// redistribute it and/or modify it under the terms of the GNU General Public -// License as published by the Free Software Foundation, either version 3 of -// the License, or (at your option) any later version. -// -// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT -// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or -// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License -// for more details. -// -// You should have received a copy of the GNU General Public License along -// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>. -// -//! \author Martin Schoenherr -//======================================================================================= -#include "LBM/LB.h" -#include "lbm/constants/D3Q27.h" -#include <basics/constants/NumericConstants.h> - -using namespace vf::basics::constant; -using namespace vf::lbm::dir; -#include "math.h" - -__global__ void InitAdvectionDiffusionIncompressibleD3Q7_Device( - unsigned int* neighborX, - unsigned int* neighborY, - unsigned int* neighborZ, - unsigned int* geoD, - real* Conc, - real* ux, - real* uy, - real* uz, - unsigned int size_Mat, - real* DD7, - bool EvenOrOdd) -{ - //////////////////////////////////////////////////////////////////////////////// - const unsigned x = threadIdx.x; // Globaler x-Index - const unsigned y = blockIdx.x; // Globaler y-Index - const unsigned z = blockIdx.y; // Globaler z-Index - - const unsigned nx = blockDim.x; - const unsigned ny = gridDim.x; - - const unsigned k = nx*(ny*z + y) + x; - ////////////////////////////////////////////////////////////////////////// - - if(k<size_Mat) - { - //////////////////////////////////////////////////////////////////////////////// - unsigned int BC; - BC = geoD[k]; - - if( BC != GEO_SOLID && BC != GEO_VOID) - { - Distributions7 D7; - if (EvenOrOdd==true) - { - D7.f[0] = &DD7[0*size_Mat]; - D7.f[1] = &DD7[1*size_Mat]; - D7.f[2] = &DD7[2*size_Mat]; - D7.f[3] = &DD7[3*size_Mat]; - D7.f[4] = &DD7[4*size_Mat]; - D7.f[5] = &DD7[5*size_Mat]; - D7.f[6] = &DD7[6*size_Mat]; - } - else - { - D7.f[0] = &DD7[0*size_Mat]; - D7.f[2] = &DD7[1*size_Mat]; - D7.f[1] = &DD7[2*size_Mat]; - D7.f[4] = &DD7[3*size_Mat]; - D7.f[3] = &DD7[4*size_Mat]; - D7.f[6] = &DD7[5*size_Mat]; - D7.f[5] = &DD7[6*size_Mat]; - } - ////////////////////////////////////////////////////////////////////////// - real ConcD = Conc[k]; - real vx1 = ux[k]; - real vx2 = uy[k]; - real vx3 = uz[k]; - real lambdaD = -c3o1 + sqrt(c3o1); - real Diffusivity = c1o20; - real Lam = -(c1o2+c1o1/lambdaD); - real nue_d = Lam/c3o1; - real ae = Diffusivity/nue_d - c1o1; - real ux_sq = vx1 * vx1; - real uy_sq = vx2 * vx2; - real uz_sq = vx3 * vx3; - ////////////////////////////////////////////////////////////////////////// - //index - ////////////////////////////////////////////////////////////////////////// - unsigned int kzero= k; - unsigned int ke = k; - unsigned int kw = neighborX[k]; - unsigned int kn = k; - unsigned int ks = neighborY[k]; - unsigned int kt = k; - unsigned int kb = neighborZ[k]; - ////////////////////////////////////////////////////////////////////////// - - (D7.f[0])[kzero] = ConcD*(c1o3*(ae*(-c3o1))-(ux_sq+uy_sq+uz_sq)); - (D7.f[1])[ke ] = ConcD*(c1o6*(ae+c1o1)+c1o2*(ux_sq)+vx1*c1o2); - (D7.f[2])[kw ] = ConcD*(c1o6*(ae+c1o1)+c1o2*(ux_sq)-vx1*c1o2); - (D7.f[3])[kn ] = ConcD*(c1o6*(ae+c1o1)+c1o2*(uy_sq)+vx2*c1o2); - (D7.f[4])[ks ] = ConcD*(c1o6*(ae+c1o1)+c1o2*(uy_sq)-vx2*c1o2); - (D7.f[5])[kt ] = ConcD*(c1o6*(ae+c1o1)+c1o2*(uz_sq)+vx3*c1o2); - (D7.f[6])[kb ] = ConcD*(c1o6*(ae+c1o1)+c1o2*(uz_sq)-vx3*c1o2); - } - } -} \ No newline at end of file diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressibleD3Q7/InitAdvectionDiffusionIncompressibleD3Q7_Device.cuh b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressibleD3Q7/InitAdvectionDiffusionIncompressibleD3Q7_Device.cuh deleted file mode 100644 index 3d0efb3a80edc645020226ca7fe88256a6c5097e..0000000000000000000000000000000000000000 --- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitAdvectionDiffusionIncompressibleD3Q7/InitAdvectionDiffusionIncompressibleD3Q7_Device.cuh +++ /dev/null @@ -1,50 +0,0 @@ -//======================================================================================= -// ____ ____ __ ______ __________ __ __ __ __ -// \ \ | | | | | _ \ |___ ___| | | | | / \ | | -// \ \ | | | | | |_) | | | | | | | / \ | | -// \ \ | | | | | _ / | | | | | | / /\ \ | | -// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____ -// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______| -// \ \ | | ________________________________________________________________ -// \ \ | | | ______________________________________________________________| -// \ \| | | | __ __ __ __ ______ _______ -// \ | | |_____ | | | | | | | | | _ \ / _____) -// \ | | _____| | | | | | | | | | | \ \ \_______ -// \ | | | | |_____ | \_/ | | | | |_/ / _____ | -// \ _____| |__| |________| \_______/ |__| |______/ (_______/ -// -// This file is part of VirtualFluids. VirtualFluids is free software: you can -// redistribute it and/or modify it under the terms of the GNU General Public -// License as published by the Free Software Foundation, either version 3 of -// the License, or (at your option) any later version. -// -// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT -// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or -// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License -// for more details. -// -// You should have received a copy of the GNU General Public License along -// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>. -// -//! \author Martin Schoenherr -//======================================================================================= -#ifndef InitAdvectionDiffusionIncompressibleD3Q7_Device_H -#define InitAdvectionDiffusionIncompressibleD3Q7_Device_H - -#include <DataTypes.h> -#include <curand.h> - -__global__ void InitAdvectionDiffusionIncompressibleD3Q7_Device( - unsigned int* neighborX, - unsigned int* neighborY, - unsigned int* neighborZ, - unsigned int* geoD, - real* Conc, - real* ux, - real* uy, - real* uz, - unsigned int size_Mat, - real* DD7, - bool EvenOrOdd); - -#endif \ No newline at end of file diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible.cu deleted file mode 100644 index b3f298988f64f4726743d8429f6df52da8eb8227..0000000000000000000000000000000000000000 --- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible.cu +++ /dev/null @@ -1,73 +0,0 @@ -//======================================================================================= -// ____ ____ __ ______ __________ __ __ __ __ -// \ \ | | | | | _ \ |___ ___| | | | | / \ | | -// \ \ | | | | | |_) | | | | | | | / \ | | -// \ \ | | | | | _ / | | | | | | / /\ \ | | -// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____ -// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______| -// \ \ | | ________________________________________________________________ -// \ \ | | | ______________________________________________________________| -// \ \| | | | __ __ __ __ ______ _______ -// \ | | |_____ | | | | | | | | | _ \ / _____) -// \ | | _____| | | | | | | | | | | \ \ \_______ -// \ | | | | |_____ | \_/ | | | | |_/ / _____ | -// \ _____| |__| |________| \_______/ |__| |______/ (_______/ -// -// This file is part of VirtualFluids. VirtualFluids is free software: you can -// redistribute it and/or modify it under the terms of the GNU General Public -// License as published by the Free Software Foundation, either version 3 of -// the License, or (at your option) any later version. -// -// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT -// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or -// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License -// for more details. -// -// You should have received a copy of the GNU General Public License along -// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>. -// -//! \author Martin Schoenherr -//======================================================================================= -#include "InitK18K20NavierStokesCompressible.h" - -#include "InitK18K20NavierStokesCompressible_Device.cuh" -#include "Parameter/Parameter.h" -#include <cuda_helper/CudaGrid.h> - -std::shared_ptr<PreProcessorStrategy> InitK18K20NavierStokesCompressible::getNewInstance(std::shared_ptr<Parameter> para) -{ - return std::shared_ptr<PreProcessorStrategy>(new InitK18K20NavierStokesCompressible(para)); -} - -void InitK18K20NavierStokesCompressible::init(int level) -{ - vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes); - - InitK18K20NavierStokesCompressible_Device <<< grid.grid, grid.threads >>>( - para->getParD(level)->neighborX, - para->getParD(level)->neighborY, - para->getParD(level)->neighborZ, - para->getParD(level)->typeOfGridNode, - para->getParD(level)->rho, - para->getParD(level)->velocityX, - para->getParD(level)->velocityY, - para->getParD(level)->velocityZ, - para->getParD(level)->numberOfNodes, - para->getParD(level)->g6.g[0], - para->getParD(level)->isEvenTimestep); - getLastCudaError("InitK18K20NavierStokesCompressible_Device execution failed"); -} - -bool InitK18K20NavierStokesCompressible::checkParameter() -{ - return false; -} - -InitK18K20NavierStokesCompressible::InitK18K20NavierStokesCompressible(std::shared_ptr<Parameter> para) -{ - this->para = para; -} - -InitK18K20NavierStokesCompressible::InitK18K20NavierStokesCompressible() -{ -} diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible.h b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible.h deleted file mode 100644 index 403dc4f0cac889af75f40da0cf92f8b821692407..0000000000000000000000000000000000000000 --- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible.h +++ /dev/null @@ -1,53 +0,0 @@ -//======================================================================================= -// ____ ____ __ ______ __________ __ __ __ __ -// \ \ | | | | | _ \ |___ ___| | | | | / \ | | -// \ \ | | | | | |_) | | | | | | | / \ | | -// \ \ | | | | | _ / | | | | | | / /\ \ | | -// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____ -// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______| -// \ \ | | ________________________________________________________________ -// \ \ | | | ______________________________________________________________| -// \ \| | | | __ __ __ __ ______ _______ -// \ | | |_____ | | | | | | | | | _ \ / _____) -// \ | | _____| | | | | | | | | | | \ \ \_______ -// \ | | | | |_____ | \_/ | | | | |_/ / _____ | -// \ _____| |__| |________| \_______/ |__| |______/ (_______/ -// -// This file is part of VirtualFluids. VirtualFluids is free software: you can -// redistribute it and/or modify it under the terms of the GNU General Public -// License as published by the Free Software Foundation, either version 3 of -// the License, or (at your option) any later version. -// -// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT -// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or -// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License -// for more details. -// -// You should have received a copy of the GNU General Public License along -// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>. -// -//! \author Martin Schoenherr -//======================================================================================= -#ifndef InitK18K20NavierStokesCompressible_H -#define InitK18K20NavierStokesCompressible_H - -#include "PreProcessor/PreProcessorStrategy/PreProcessorStrategy.h" - -#include <memory> - -class Parameter; - -class InitK18K20NavierStokesCompressible : public PreProcessorStrategy -{ -public: - static std::shared_ptr<PreProcessorStrategy> getNewInstance(std::shared_ptr< Parameter> para); - void init(int level); - bool checkParameter(); - -private: - InitK18K20NavierStokesCompressible(); - InitK18K20NavierStokesCompressible(std::shared_ptr< Parameter> para); - std::shared_ptr< Parameter> para; -}; - -#endif \ No newline at end of file diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible_Device.cu b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible_Device.cu deleted file mode 100644 index 6beb4ad8574236857b16cebf035c2796312bdf82..0000000000000000000000000000000000000000 --- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible_Device.cu +++ /dev/null @@ -1,110 +0,0 @@ -//======================================================================================= -// ____ ____ __ ______ __________ __ __ __ __ -// \ \ | | | | | _ \ |___ ___| | | | | / \ | | -// \ \ | | | | | |_) | | | | | | | / \ | | -// \ \ | | | | | _ / | | | | | | / /\ \ | | -// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____ -// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______| -// \ \ | | ________________________________________________________________ -// \ \ | | | ______________________________________________________________| -// \ \| | | | __ __ __ __ ______ _______ -// \ | | |_____ | | | | | | | | | _ \ / _____) -// \ | | _____| | | | | | | | | | | \ \ \_______ -// \ | | | | |_____ | \_/ | | | | |_/ / _____ | -// \ _____| |__| |________| \_______/ |__| |______/ (_______/ -// -// This file is part of VirtualFluids. VirtualFluids is free software: you can -// redistribute it and/or modify it under the terms of the GNU General Public -// License as published by the Free Software Foundation, either version 3 of -// the License, or (at your option) any later version. -// -// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT -// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or -// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License -// for more details. -// -// You should have received a copy of the GNU General Public License along -// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>. -// -//! \author Martin Schoenherr -//======================================================================================= -#include "LBM/LB.h" -#include "lbm/constants/D3Q27.h" -#include <basics/constants/NumericConstants.h> - -using namespace vf::basics::constant; -using namespace vf::lbm::dir; -#include "math.h" - - -__global__ void InitK18K20NavierStokesCompressible_Device(unsigned int* neighborX, - unsigned int* neighborY, - unsigned int* neighborZ, - unsigned int* geoD, - real* rho, - real* ux, - real* uy, - real* uz, - unsigned int size_Mat, - real* G6, - bool EvenOrOdd) -{ - //////////////////////////////////////////////////////////////////////////////// - const unsigned x = threadIdx.x; // Globaler x-Index - const unsigned y = blockIdx.x; // Globaler y-Index - const unsigned z = blockIdx.y; // Globaler z-Index - - const unsigned nx = blockDim.x; - const unsigned ny = gridDim.x; - - const unsigned k = nx*(ny*z + y) + x; - ////////////////////////////////////////////////////////////////////////// - - if (k < size_Mat) - { - //////////////////////////////////////////////////////////////////////////////// - unsigned int BC; - BC = geoD[k]; - - if (BC != GEO_SOLID && BC != GEO_VOID) - { - Distributions6 D; - if (EvenOrOdd == true) - { - D.g[dP00] = &G6[dP00 *size_Mat]; - D.g[dM00] = &G6[dM00 *size_Mat]; - D.g[d0P0] = &G6[d0P0 *size_Mat]; - D.g[d0M0] = &G6[d0M0 *size_Mat]; - D.g[d00P] = &G6[d00P *size_Mat]; - D.g[d00M] = &G6[d00M *size_Mat]; - } - else - { - D.g[dM00] = &G6[dP00 *size_Mat]; - D.g[dP00] = &G6[dM00 *size_Mat]; - D.g[d0M0] = &G6[d0P0 *size_Mat]; - D.g[d0P0] = &G6[d0M0 *size_Mat]; - D.g[d00M] = &G6[d00P *size_Mat]; - D.g[d00P] = &G6[d00M *size_Mat]; - } - ////////////////////////////////////////////////////////////////////////// - //index - ////////////////////////////////////////////////////////////////////////// - // unsigned int kzero = k; - unsigned int ke = k; - unsigned int kw = neighborX[k]; - unsigned int kn = k; - unsigned int ks = neighborY[k]; - unsigned int kt = k; - unsigned int kb = neighborZ[k]; - ////////////////////////////////////////////////////////////////////////// - - (D.g[dP00])[ke] = 0.0f; - (D.g[dM00])[kw] = 0.0f; - (D.g[d0P0])[kn] = 0.0f; - (D.g[d0M0])[ks] = 0.0f; - (D.g[d00P])[kt] = 0.0f; - (D.g[d00M])[kb] = 0.0f; - } - } -} \ No newline at end of file diff --git a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible_Device.cuh b/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible_Device.cuh deleted file mode 100644 index 90ecb070533a0bf40522bdf246f59921149caaa3..0000000000000000000000000000000000000000 --- a/src/gpu/core/PreProcessor/PreProcessorStrategy/InitK18K20NavierStokesCompressible/InitK18K20NavierStokesCompressible_Device.cuh +++ /dev/null @@ -1,50 +0,0 @@ -//======================================================================================= -// ____ ____ __ ______ __________ __ __ __ __ -// \ \ | | | | | _ \ |___ ___| | | | | / \ | | -// \ \ | | | | | |_) | | | | | | | / \ | | -// \ \ | | | | | _ / | | | | | | / /\ \ | | -// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____ -// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______| -// \ \ | | ________________________________________________________________ -// \ \ | | | ______________________________________________________________| -// \ \| | | | __ __ __ __ ______ _______ -// \ | | |_____ | | | | | | | | | _ \ / _____) -// \ | | _____| | | | | | | | | | | \ \ \_______ -// \ | | | | |_____ | \_/ | | | | |_/ / _____ | -// \ _____| |__| |________| \_______/ |__| |______/ (_______/ -// -// This file is part of VirtualFluids. VirtualFluids is free software: you can -// redistribute it and/or modify it under the terms of the GNU General Public -// License as published by the Free Software Foundation, either version 3 of -// the License, or (at your option) any later version. -// -// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT -// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or -// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License -// for more details. -// -// You should have received a copy of the GNU General Public License along -// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>. -// -//! \author Martin Schoenherr -//======================================================================================= -#ifndef InitK18K20NavierStokesCompressible_Device_H -#define InitK18K20NavierStokesCompressible_Device_H - -#include <DataTypes.h> -#include <curand.h> - -__global__ void InitK18K20NavierStokesCompressible_Device( - unsigned int* neighborX, - unsigned int* neighborY, - unsigned int* neighborZ, - unsigned int* geoD, - real* rho, - real* ux, - real* uy, - real* uz, - unsigned int size_Mat, - real* G6, - bool EvenOrOdd); - -#endif \ No newline at end of file diff --git a/src/gpu/core/PreProcessor/PreProcessorType.h b/src/gpu/core/PreProcessor/PreProcessorType.h index 20230c4b08817352f60cb5567dd99fdd19ce6438..88c4f2bfb35abb19d80380035df0b6cfe6020bfa 100644 --- a/src/gpu/core/PreProcessor/PreProcessorType.h +++ b/src/gpu/core/PreProcessor/PreProcessorType.h @@ -35,10 +35,7 @@ enum PreProcessorType { InitNavierStokesIncompressible, InitNavierStokesCompressible, - InitK18K20NavierStokesCompressible, - InitAdvectionDiffusionIncompressibleD3Q7, InitAdvectionDiffusionIncompressible, - InitAdvectionDiffusionCompressibleD3Q7, InitAdvectionDiffusionCompressible }; #endif \ No newline at end of file diff --git a/src/gpu/core/PreProcessor/ReaderMeasurePoints.cpp b/src/gpu/core/PreProcessor/ReaderMeasurePoints.cpp new file mode 100644 index 0000000000000000000000000000000000000000..59835b1f6ff95de78ac8885e3e30a34626a80350 --- /dev/null +++ b/src/gpu/core/PreProcessor/ReaderMeasurePoints.cpp @@ -0,0 +1,80 @@ +#include "ReaderMeasurePoints.h" + +#include "Parameter/Parameter.h" +#include "GPU/CudaMemoryManager.h" + +#include <basics/utilities/UbFileInputASCII.h> + +using namespace vf::lbm::dir; + +////////////////////////////////////////////////////////////////////////// +void ReaderMeasurePoints::readMeasurePoints( Parameter* para ) +{ + UbFileInputASCII in(para->getmeasurePoints()); + int numberOfAllNodes = in.readInteger(); + in.readLine(); + int tempLevel; + MeasurePoints tempMP; + //printf("done, init the values...\n"); + for (int u = 0; u < numberOfAllNodes; u++) + { + tempMP.name = in.readString(); + //printf("done, read the name...\n"); + tempMP.k = in.readInteger(); + //printf("done, read k...\n"); + tempLevel = in.readInteger(); + //printf("done, read level...\n"); + in.readLine(); + //printf("done, read the values...\n"); + para->getParH(tempLevel)->MP.push_back(tempMP); + //printf("done, put it into a vector...\n"); + } +} + +//////////////////////////////////////////////////////////////////////////////// +void ReaderMeasurePoints::readMeasurePoints(Parameter* para, CudaMemoryManager* cudaMemoryManager) +{ + // read measure points from file + ReaderMeasurePoints::readMeasurePoints(para); + // printf("done, reading the file...\n"); + // level loop + for (int lev = 0; lev <= para->getMaxLevel(); lev++) { + // set Memory Size and malloc of the indices and macroscopic values per level + para->getParH(lev)->numberOfValuesMP = (unsigned int)para->getParH(lev)->MP.size() * + (unsigned int)para->getclockCycleForMP() / + ((unsigned int)para->getTimestepForMP()); + para->getParD(lev)->numberOfValuesMP = para->getParH(lev)->numberOfValuesMP; + + para->getParH(lev)->numberOfPointskMP = (int)para->getParH(lev)->MP.size(); + para->getParD(lev)->numberOfPointskMP = para->getParH(lev)->numberOfPointskMP; + + para->getParH(lev)->memSizeIntkMP = sizeof(unsigned int) * (int)para->getParH(lev)->MP.size(); + para->getParD(lev)->memSizeIntkMP = para->getParH(lev)->memSizeIntkMP; + + para->getParH(lev)->memSizerealkMP = sizeof(real) * para->getParH(lev)->numberOfValuesMP; + para->getParD(lev)->memSizerealkMP = para->getParH(lev)->memSizerealkMP; + + printf("Level: %d, numberOfValuesMP: %d, memSizeIntkMP: %d, memSizerealkMP: %d\n", lev, + para->getParH(lev)->numberOfValuesMP, para->getParH(lev)->memSizeIntkMP, para->getParD(lev)->memSizerealkMP); + + cudaMemoryManager->cudaAllocMeasurePointsIndex(lev); + + // loop over all measure points per level + for (int index = 0; index < (int)para->getParH(lev)->MP.size(); index++) { + // set indices + para->getParH(lev)->kMP[index] = para->getParH(lev)->MP[index].k; + } + // loop over all measure points per level times MPClockCycle + for (int index = 0; index < (int)para->getParH(lev)->numberOfValuesMP; index++) { + // init values + para->getParH(lev)->VxMP[index] = (real)0.0; + para->getParH(lev)->VyMP[index] = (real)0.0; + para->getParH(lev)->VzMP[index] = (real)0.0; + para->getParH(lev)->RhoMP[index] = (real)0.0; + } + + // copy indices-arrays + cudaMemoryManager->cudaCopyMeasurePointsIndex(lev); + } +} +//////////////////////////////////////////////////////////////////////////////// diff --git a/src/gpu/core/PreProcessor/ReaderMeasurePoints.h b/src/gpu/core/PreProcessor/ReaderMeasurePoints.h new file mode 100644 index 0000000000000000000000000000000000000000..62dd9446c0ad9d0e3b1d1e5c34eb80d53ff8038a --- /dev/null +++ b/src/gpu/core/PreProcessor/ReaderMeasurePoints.h @@ -0,0 +1,16 @@ +#ifndef POSITION_READER_H +#define POSITION_READER_H + +class Parameter; +class CudaMemoryManager; + +class ReaderMeasurePoints +{ +private: + static void readMeasurePoints(Parameter* para); + +public: + static void readMeasurePoints(Parameter* para, CudaMemoryManager* cudaMemoryManager); +}; + +#endif \ No newline at end of file