Skip to content

GitLab

Commit 9e0b6851 authored by ccfd's avatar ccfd
Browse files

Modify a bug of boundary scheme

parent acb9b19b
Hide whitespace changes
Inline Side-by-side
Showing with 240 additions and 427 deletions
obj/OCFD_mpi.o
View file @ 9e0b6851
No preview for this file type
obj/OCFD_mpi_dev.o
View file @ 9e0b6851
No preview for this file type
obj/OCFD_split.o
View file @ 9e0b6851
No preview for this file type
obj/OCFD_time.o
View file @ 9e0b6851
No preview for this file type
obj/commen_kernel.o
View file @ 9e0b6851
No preview for this file type
obj/cuda_commen.o
View file @ 9e0b6851
No preview for this file type
obj/cuda_utility.o
View file @ 9e0b6851
No preview for this file type
obj/libocfd.a
View file @ 9e0b6851
No preview for this file type
obj/parameters.o
View file @ 9e0b6851
No preview for this file type
obj/parameters_d.o
View file @ 9e0b6851
No preview for this file type
obj/test.o
View file @ 9e0b6851
No preview for this file type
obj/utility.o
View file @ 9e0b6851
No preview for this file type
opencfd-scu.out
View file @ 9e0b6851
No preview for this file type
opencfd.o
View file @ 9e0b6851
No preview for this file type
src/OCFD_IO_mpi.c
View file @ 9e0b6851
...@@ -7,7 +7,7 @@ ...@@ -7,7 +7,7 @@
#include "utility.h" #include "utility.h"
#include "OCFD_mpi.h" #include "OCFD_mpi.h"
#include "io_warp.h" #include "io_warp.h"
__device__ __constant__ unsigned int ny_d,ny_2lap_d;
//--------------------------------------------------- //---------------------------------------------------
#ifdef __cplusplus #ifdef __cplusplus
extern "C"{ extern "C"{
......
src/OCFD_bound_Scheme.cu
View file @ 9e0b6851
...@@ -76,187 +76,111 @@ extern "C"{ ...@@ -76,187 +76,111 @@ extern "C"{
#endif #endif
#define D0bound1(coords)\
if(coords == 0){\
*tmp = (stencil[-ka1+1] - stencil[-ka1]);\
return 0;\
}else if(coords == 1){\
*tmp = (stencil[-ka1+1] - stencil[-ka1-1])*0.5;\
return 0;\
}else if(coords == 2){\
*tmp = (stencil[-ka1-2] - 8.0*stencil[-ka1-1] + 8.0*stencil[-ka1+1] - stencil[-ka1+2])/12.0;\
return 0;\
}else if(coords == 3){\
*tmp = (stencil[-ka1+3] - stencil[-ka1-3]\
-9.0*(stencil[-ka1+2] - stencil[-ka1-2])\
+45.0*(stencil[-ka1+1] - stencil[-ka1-1]))/60.0;\
return 0;\
}\
#define D0bound2(coords)\
if(coords == -1){\
*tmp = (stencil[-ka1] - stencil[-ka1-1]);\
return 0;\
}else if(coords == -2){\
*tmp = (stencil[-ka1+1] - stencil[-ka1-1])*0.5;\
return 0;\
}else if(coords == -3){\
*tmp = (stencil[-ka1-2] - 8.0*stencil[-ka1-1] + 8.0*stencil[-ka1+1] - stencil[-ka1+2])/12.0;\
return 0;\
}else if(coords == -4){\
*tmp = (stencil[-ka1+3] - stencil[-ka1-3]\
-9.0*(stencil[-ka1+2] - stencil[-ka1-2])\
+45.0*(stencil[-ka1+1] - stencil[-ka1-1]))/60.0;\
return 0;\
}\
// =========================================================================================================== // // =========================================================================================================== //
__device__ int OCFD_D0bound_scheme_kernel(REAL* tmp, dim3 flagxyzb, dim3 coords, REAL *stencil, int ka1, cudaJobPackage job){ __device__ int OCFD_D0bound_scheme_kernel(REAL* tmp, dim3 flagxyzb, dim3 coords, REAL *stencil, int ka1, cudaJobPackage job){
int tmp1;
switch(flagxyzb.y){ switch(flagxyzb.y){
case 1: case 1:
{ {
if(coords.x == 0){ D0bound1(coords.x)
*tmp = (stencil[-ka1+1] - stencil[-ka1]);
return 0;
}else if(coords.x == 1){
*tmp = (stencil[-ka1+1] - stencil[-ka1-1])*0.5;
return 0;
}else if(coords.x == 2){
*tmp = (stencil[-ka1-2] - 8.0*stencil[-ka1-1] + 8.0*stencil[-ka1+1] - stencil[-ka1+2])/12.0;
return 0;
}else if(coords.x == 3){
*tmp = (stencil[-ka1+3] - stencil[-ka1-3]
-9.0*(stencil[-ka1+2] - stencil[-ka1-2])
+45.0*(stencil[-ka1+1] - stencil[-ka1-1]))/60.0;
return 0;
}
} }
break; break;
case 2: case 2:
{ {
if(coords.y == 0){ D0bound1(coords.y)
*tmp = (stencil[-ka1+1] - stencil[-ka1]);
return 0;
}else if(coords.y == 1){
*tmp = (stencil[-ka1+1] - stencil[-ka1-1])*0.5;
return 0;
}else if(coords.y == 2){
*tmp = (stencil[-ka1-2] - 8.0*stencil[-ka1-1] + 8.0*stencil[-ka1+1] - stencil[-ka1+2])/12.0;
return 0;
}else if(coords.y == 3){
*tmp = (stencil[-ka1+3] - stencil[-ka1-3]
-9.0*(stencil[-ka1+2] - stencil[-ka1-2])
+45.0*(stencil[-ka1+1] - stencil[-ka1-1]))/60.0;
return 0;
}
} }
break; break;
case 3: case 3:
{ {
if(coords.z == 0){ D0bound1(coords.z)
*tmp = (stencil[-ka1+1] - stencil[-ka1]);
return 0;
}else if(coords.z == 1){
*tmp = (stencil[-ka1+1] - stencil[-ka1-1])*0.5;
return 0;
}else if(coords.z == 2){
*tmp = (stencil[-ka1-2] - 8.0*stencil[-ka1-1] + 8.0*stencil[-ka1+1] - stencil[-ka1+2])/12.0;
return 0;
}else if(coords.z == 3){
*tmp = (stencil[-ka1+3] - stencil[-ka1-3]
-9.0*(stencil[-ka1+2] - stencil[-ka1-2])
+45.0*(stencil[-ka1+1] - stencil[-ka1-1]))/60.0;
return 0;
}
} }
break; break;
case 4: case 4:
{ {
if(coords.x == job.end.x-job.start.x-1){ tmp1 = coords.x + job.start.x - job.end.x;
D0bound2(tmp1)
*tmp = (stencil[-ka1] - stencil[-ka1-1]);
return 0;
}else if(coords.x == job.end.x-job.start.x-2){
*tmp = (stencil[-ka1+1] - stencil[-ka1-1])*0.5;
return 0;
}else if(coords.x == job.end.x-job.start.x-3){
*tmp = (stencil[-ka1-2] - 8.0*stencil[-ka1-1] + 8.0*stencil[-ka1+1] - stencil[-ka1+2])/12.0;
return 0;
}else if(coords.x == job.end.x-job.start.x-4){
*tmp = (stencil[-ka1+3] - stencil[-ka1-3]
-9.0*(stencil[-ka1+2] - stencil[-ka1-2])
+45.0*(stencil[-ka1+1] - stencil[-ka1-1]))/60.0;
return 0;
}
} }
break; break;
case 5: case 5:
{ {
if(coords.y == job.end.y-job.start.y-1){ tmp1 = coords.y + job.start.y - job.end.y;
D0bound2(tmp1)
*tmp = (stencil[-ka1] - stencil[-ka1-1]);
return 0;
}else if(coords.y == job.end.y-job.start.y-2){
*tmp = (stencil[-ka1+1] - stencil[-ka1-1])*0.5;
return 0;
}else if(coords.y == job.end.y-job.start.y-3){
*tmp = (stencil[-ka1-2] - 8.0*stencil[-ka1-1] + 8.0*stencil[-ka1+1] - stencil[-ka1+2])/12.0;
return 0;
}else if(coords.y == job.end.y-job.start.y-4){
*tmp = (stencil[-ka1+3] - stencil[-ka1-3]
-9.0*(stencil[-ka1+2] - stencil[-ka1-2])
+45.0*(stencil[-ka1+1] - stencil[-ka1-1]))/60.0;
return 0;
}
} }
break; break;
case 6: case 6:
{ {
if(coords.z == job.end.z-job.start.z-1){ tmp1 = coords.z + job.start.z - job.end.z;
D0bound2(tmp1)
*tmp = (stencil[-ka1] - stencil[-ka1-1]); }
break;
return 0;
}else if(coords.z == job.end.z-job.start.z-2){
*tmp = (stencil[-ka1+1] - stencil[-ka1-1])*0.5;
return 0;
}else if(coords.z == job.end.z-job.start.z-3){
*tmp = (stencil[-ka1-2] - 8.0*stencil[-ka1-1] + 8.0*stencil[-ka1+1] - stencil[-ka1+2])/12.0; case 7:
return 0; {
D0bound1(coords.x)
tmp1 = coords.x + job.start.x - job.end.x;
D0bound2(tmp1)
}
break;
}else if(coords.z == job.end.z-job.start.z-4){ case 8:
*tmp = (stencil[-ka1+3] - stencil[-ka1-3] {
-9.0*(stencil[-ka1+2] - stencil[-ka1-2]) D0bound1(coords.y)
+45.0*(stencil[-ka1+1] - stencil[-ka1-1]))/60.0; tmp1 = coords.y + job.start.y - job.end.y;
D0bound2(tmp1)
}
break;
return 0; case 9:
} {
D0bound1(coords.z)
tmp1 = coords.z + job.start.z - job.end.z;
D0bound2(tmp1)
} }
break; break;
} }
return 1; return 1;
} }
...@@ -429,29 +353,48 @@ void OCFD_Dz0_bound(cudaField f , cudaField fx , cudaJobPackage job_in , dim3 bl ...@@ -429,29 +353,48 @@ void OCFD_Dz0_bound(cudaField f , cudaField fx , cudaJobPackage job_in , dim3 bl
void OCFD_bound(dim3 *flagxyzb, int boundp, int boundm, cudaJobPackage job){ void OCFD_bound(dim3 *flagxyzb, int boundp, int boundm, cudaJobPackage job){
// eyes on field WITH LAPs // eyes on field WITH LAPs
dim3 size; dim3 size;
int flag = 0;
jobsize(&job, &size); jobsize(&job, &size);
switch(flagxyzb->x){ switch(flagxyzb->x){
case 1: case 1:
case 4: case 4:
{ {
if(npx == 0 && job.start.x == LAP && boundp == 1) flagxyzb->y = 1; if(npx == 0 && job.start.x == LAP && boundp == 1){
if(npx == NPX0-1 && job.end.x == nx_lap && boundm == 1) flagxyzb->y = 4; flagxyzb->y = 1;
flag = 1;
}
if(npx == NPX0-1 && job.end.x == nx_lap && boundm == 1){
flagxyzb->y = 4;
if(flag == 1) flagxyzb->y = 7;
}
} }
break; break;
case 2: case 2:
case 5: case 5:
{ {
if(npy == 0 && job.start.y == LAP && boundp == 1) flagxyzb->y = 2; if(npy == 0 && job.start.y == LAP && boundp == 1){
if(npy == NPY0-1 && job.end.y == ny_lap && boundm == 1) flagxyzb->y = 5; flagxyzb->y = 2;
flag = 1;
}
if(npy == NPY0-1 && job.end.y == ny_lap && boundm == 1){
flagxyzb->y = 5;
if(flag == 1) flagxyzb->y = 8;
}
} }
break; break;
case 3: case 3:
case 6: case 6:
{ {
if(npz == 0 && job.start.z == LAP && boundp == 1) flagxyzb->y = 3; if(npz == 0 && job.start.z == LAP && boundp == 1){
if(npz == NPZ0-1 && job.end.z == nz_lap && boundm == 1) flagxyzb->y = 6; flagxyzb->y = 3;
flag = 1;
}
if(npz == NPZ0-1 && job.end.z == nz_lap && boundm == 1){
flagxyzb->y = 6;
if(flag == 1) flagxyzb->y = 9;
}
} }
break; break;
} }
...@@ -725,337 +668,207 @@ __device__ REAL OCFD_weno5_kernel_M_lift_plus(REAL *stencil){ ...@@ -725,337 +668,207 @@ __device__ REAL OCFD_weno5_kernel_M_lift_plus(REAL *stencil){
//tmp = (2.0*stencil[-ka1+1] + 5.0*stencil[-ka1+2] + stencil[-ka1+3])/6.0; //tmp = (2.0*stencil[-ka1+1] + 5.0*stencil[-ka1+2] + stencil[-ka1+3])/6.0;
//tmp = (11.0*stencil[-ka1] - 7.0*stencil[-ka1-1] + 2.0*stencil[-ka1-2])/6.0; //tmp = (11.0*stencil[-ka1] - 7.0*stencil[-ka1-1] + 2.0*stencil[-ka1-2])/6.0;
#define boundp1(coords)\
if(coords <= -ka1){\
if(coords == 0){\
*tmp = stencil[-ka1+1];\
}\
if(coords == 1){\
*tmp = OCFD_weno5_kernel_P_lift(&stencil[-ka1-2]);\
}\
if(coords == 2){\
*tmp = OCFD_weno5_kernel_P_lift_plus(&stencil[-ka1-2]);\
}\
if(coords == 3){\
*tmp = OCFD_weno5_kernel_P(&stencil[-ka1-2]);\
}\
return 0;\
}\
#define boundp2(coords)\
if(coords > -kb1){\
if(coords == -1){\
*tmp = stencil[-ka1] + 0.5*minmod2(stencil[-ka1] - stencil[-ka1-1], stencil[-ka1] - stencil[-ka1-1]);\
}\
if(coords == -2){\
*tmp = OCFD_weno5_kernel_P_right_plus(&stencil[-ka1-2]);\
}\
if(coords <= -3){\
*tmp = OCFD_weno5_kernel_P(&stencil[-ka1-2]);\
}\
return 0;\
}\
__device__ int OCFD_bound_scheme_kernel_p(REAL* tmp, dim3 flagxyzb, dim3 coords, REAL *stencil, int ka1, int kb1, cudaJobPackage job){ __device__ int OCFD_bound_scheme_kernel_p(REAL* tmp, dim3 flagxyzb, dim3 coords, REAL *stencil, int ka1, int kb1, cudaJobPackage job){
int tmp1;
switch(flagxyzb.y){ switch(flagxyzb.y){
case 1: case 1:
{ {
if(coords.x <= -ka1){ boundp1(coords.x)
if(coords.x == 0){
*tmp = stencil[-ka1+1];
}
if(coords.x == 1){
//*tmp = stencil[-ka1] + 0.5*minmod2(stencil[-ka1+1] - stencil[-ka1], stencil[-ka1+1] - stencil[-ka1]);
*tmp = OCFD_weno5_kernel_P_lift(&stencil[-ka1-2]);
}
if(coords.x == 2){
//*tmp = stencil[-ka1] + 0.5*minmod2(stencil[-ka1+1] - stencil[-ka1], stencil[-ka1] - stencil[-ka1-1]);
*tmp = OCFD_weno5_kernel_P_lift_plus(&stencil[-ka1-2]);
}
if(coords.x == 3){
//*tmp = stencil[-ka1] + 0.5*minmod2(stencil[-ka1+1] - stencil[-ka1], stencil[-ka1] - stencil[-ka1-1]);
*tmp = OCFD_weno5_kernel_P(&stencil[-ka1-2]);
}
return 0;
}
} }
break; break;
case 2: case 2:
{ {
if(coords.y <= -ka1){ boundp1(coords.y)
if(coords.y == 0){
*tmp = stencil[-ka1+1];
}
if(coords.y == 1){
//*tmp = stencil[-ka1] + 0.5*minmod2(stencil[-ka1+1] - stencil[-ka1], stencil[-ka1+1] - stencil[-ka1]);
*tmp = OCFD_weno5_kernel_P_lift(&stencil[-ka1-2]);
}
if(coords.y == 2){
//*tmp = stencil[-ka1] + 0.5*minmod2(stencil[-ka1+1] - stencil[-ka1], stencil[-ka1] - stencil[-ka1-1]);
*tmp = OCFD_weno5_kernel_P_lift_plus(&stencil[-ka1-2]);
}
if(coords.y == 3){
//*tmp = stencil[-ka1] + 0.5*minmod2(stencil[-ka1+1] - stencil[-ka1], stencil[-ka1] - stencil[-ka1-1]);
*tmp = OCFD_weno5_kernel_P(&stencil[-ka1-2]);
}
return 0;
}
} }
break; break;
case 3: case 3:
{ {
if(coords.z <= -ka1){ boundp1(coords.z)
if(coords.z == 0){
*tmp = stencil[-ka1+1];
}
if(coords.z == 1){
//*tmp = stencil[-ka1] + 0.5*minmod2(stencil[-ka1+1] - stencil[-ka1], stencil[-ka1+1] - stencil[-ka1]);
*tmp = OCFD_weno5_kernel_P_lift(&stencil[-ka1-2]);
}
if(coords.z == 2){
//*tmp = stencil[-ka1] + 0.5*minmod2(stencil[-ka1+1] - stencil[-ka1], stencil[-ka1] - stencil[-ka1-1]);
*tmp = OCFD_weno5_kernel_P_lift_plus(&stencil[-ka1-2]);
}
if(coords.z == 3){
//*tmp = stencil[-ka1] + 0.5*minmod2(stencil[-ka1+1] - stencil[-ka1], stencil[-ka1] - stencil[-ka1-1]);
*tmp = OCFD_weno5_kernel_P(&stencil[-ka1-2]);
}
return 0;
}
} }
break; break;
case 4: case 4:
{ {
if(coords.x > job.end.x-job.start.x-kb1){ tmp1 = coords.x + job.start.x - job.end.x;
boundp2(tmp1)
if(coords.x == job.end.x-job.start.x-1){
//*tmp = OCFD_weno5_kernel_P_right(&stencil[-ka1-2]);
*tmp = stencil[-ka1] + 0.5*minmod2(stencil[-ka1] - stencil[-ka1-1], stencil[-ka1] - stencil[-ka1-1]);
}
if(coords.x == job.end.x-job.start.x-2){
*tmp = OCFD_weno5_kernel_P_right_plus(&stencil[-ka1-2]);
//*tmp = stencil[-ka1] + 0.5*minmod2(stencil[-ka1+1] - stencil[-ka1], stencil[-ka1] - stencil[-ka1-1]);
}
if(coords.x <= job.end.x-job.start.x-3){
*tmp = OCFD_weno5_kernel_P(&stencil[-ka1-2]);
//*tmp = stencil[-ka1] + 0.5*minmod2(stencil[-ka1+1] - stencil[-ka1], stencil[-ka1] - stencil[-ka1-1]);
}
return 0;
}
} }
break; break;
case 5: case 5:
{ {
if(coords.y > job.end.y-job.start.y-kb1){ tmp1 = coords.y + job.start.y - job.end.y;
boundp2(tmp1)
if(coords.y == job.end.y-job.start.y-1){
//*tmp = OCFD_weno5_kernel_P_right(&stencil[-ka1-2]);
*tmp = stencil[-ka1] + 0.5*minmod2(stencil[-ka1] - stencil[-ka1-1], stencil[-ka1] - stencil[-ka1-1]);
}
if(coords.y == job.end.y-job.start.y-2){
*tmp = OCFD_weno5_kernel_P_right_plus(&stencil[-ka1-2]);
//*tmp = stencil[-ka1] + 0.5*minmod2(stencil[-ka1+1] - stencil[-ka1], stencil[-ka1] - stencil[-ka1-1]);
}
if(coords.y <= job.end.y-job.start.y-3){
*tmp = OCFD_weno5_kernel_P(&stencil[-ka1-2]);
//*tmp = stencil[-ka1] + 0.5*minmod2(stencil[-ka1+1] - stencil[-ka1], stencil[-ka1] - stencil[-ka1-1]);
}
return 0;
}
} }
break; break;
case 6: case 6:
{ {
if(coords.z > job.end.z-job.start.z-kb1){ tmp1 = coords.z + job.start.z - job.end.z;
boundp2(tmp1)
if(coords.z == job.end.z-job.start.z-1){ }
//*tmp = OCFD_weno5_kernel_P_right(&stencil[-ka1-2]); break;
*tmp = stencil[-ka1] + 0.5*minmod2(stencil[-ka1] - stencil[-ka1-1], stencil[-ka1] - stencil[-ka1-1]);
}
if(coords.z == job.end.z-job.start.z-2){ case 7:
*tmp = OCFD_weno5_kernel_P_right_plus(&stencil[-ka1-2]); {
//*tmp = stencil[-ka1] + 0.5*minmod2(stencil[-ka1+1] - stencil[-ka1], stencil[-ka1] - stencil[-ka1-1]); boundp1(coords.x)
} tmp1 = coords.x + job.start.x - job.end.x;
boundp2(tmp1)
}
break;
if(coords.z <= job.end.z-job.start.z-3){ case 8:
*tmp = OCFD_weno5_kernel_P(&stencil[-ka1-2]); {
//*tmp = stencil[-ka1] + 0.5*minmod2(stencil[-ka1+1] - stencil[-ka1], stencil[-ka1] - stencil[-ka1-1]); boundp1(coords.y)
} tmp1 = coords.y + job.start.y - job.end.y;
boundp2(tmp1)
}
break;
return 0; case 9:
} {
boundp1(coords.z)
tmp1 = coords.z + job.start.z - job.end.z;
boundp2(tmp1)
} }
break; break;
} }
return 1; return 1;
} }
#define boundm1(coords)\
if(coords < -ka1){\
if(coords == 0){\
*tmp = OCFD_weno5_kernel_M_lift(&stencil[-ka1-2]);\
}\
if(coords == 1){\
*tmp = OCFD_weno5_kernel_M_lift_plus(&stencil[-ka1-2]);\
}\
if(coords >= 2){\
*tmp = OCFD_weno5_kernel_M(&stencil[-ka1-2]);\
}\
return 0;\
}\
#define boundm2(coords)\
if(coords >= -kb1-1){\
if(coords == -1){\
*tmp = stencil[-ka1-1];\
}\
if(coords == -2){\
*tmp = stencil[-ka1] - 0.5*minmod2(stencil[-ka1] - stencil[-ka1-1], stencil[-ka1] - stencil[-ka1-1]);\
}\
if(coords == -3){\
*tmp = OCFD_weno5_kernel_M_right_plus(&stencil[-ka1-2]);\
}\
if(coords == -4){\
*tmp = OCFD_weno5_kernel_M(&stencil[-ka1-2]);\
}\
return 0;\
}\
__device__ int OCFD_bound_scheme_kernel_m(REAL* tmp, dim3 flagxyzb, dim3 coords, REAL *stencil, int ka1, int kb1, cudaJobPackage job){ __device__ int OCFD_bound_scheme_kernel_m(REAL* tmp, dim3 flagxyzb, dim3 coords, REAL *stencil, int ka1, int kb1, cudaJobPackage job){
int tmp1;
switch(flagxyzb.y){ switch(flagxyzb.y){
case 1: case 1:
{ {
if(coords.x < -ka1){ boundm1(coords.x)
if(coords.x == 0){
*tmp = OCFD_weno5_kernel_M_lift(&stencil[-ka1-2]);
//*tmp = stencil[-ka1] - 0.5*minmod2(stencil[-ka1+1] - stencil[-ka1], stencil[-ka1+1] - stencil[-ka1]);
}
if(coords.x == 1){
*tmp = OCFD_weno5_kernel_M_lift_plus(&stencil[-ka1-2]);
//*tmp = stencil[-ka1] - 0.5*minmod2(stencil[-ka1+1] - stencil[-ka1], stencil[-ka1] - stencil[-ka1-1]);
}
if(coords.x >= 2){
*tmp = OCFD_weno5_kernel_M(&stencil[-ka1-2]);
//*tmp = stencil[-ka1] - 0.5*minmod2(stencil[-ka1+1] - stencil[-ka1], stencil[-ka1] - stencil[-ka1-1]);
}
return 0;
}
} }
break; break;
case 2: case 2:
{ {
if(coords.y < -ka1){ boundm1(coords.y)
if(coords.y == 0){
*tmp = OCFD_weno5_kernel_M_lift(&stencil[-ka1-2]);
//*tmp = stencil[-ka1] - 0.5*minmod2(stencil[-ka1+1] - stencil[-ka1], stencil[-ka1+1] - stencil[-ka1]);
}
if(coords.y == 1){
*tmp = OCFD_weno5_kernel_M_lift_plus(&stencil[-ka1-2]);
//*tmp = stencil[-ka1] - 0.5*minmod2(stencil[-ka1+1] - stencil[-ka1], stencil[-ka1] - stencil[-ka1-1]);
}
if(coords.y >= 2){
*tmp = OCFD_weno5_kernel_M(&stencil[-ka1-2]);
//*tmp = stencil[-ka1] - 0.5*minmod2(stencil[-ka1+1] - stencil[-ka1], stencil[-ka1] - stencil[-ka1-1]);
}
return 0;
}
} }
break; break;
case 3: case 3:
{ {
if(coords.z < -ka1){ boundm1(coords.z)
if(coords.z == 0){
*tmp = OCFD_weno5_kernel_M_lift(&stencil[-ka1-2]);
//*tmp = stencil[-ka1] - 0.5*minmod2(stencil[-ka1+1] - stencil[-ka1], stencil[-ka1+1] - stencil[-ka1]);
}
if(coords.z == 1){
*tmp = OCFD_weno5_kernel_M_lift_plus(&stencil[-ka1-2]);
//*tmp = stencil[-ka1] - 0.5*minmod2(stencil[-ka1+1] - stencil[-ka1], stencil[-ka1] - stencil[-ka1-1]);
}
if(coords.z >= 2){
*tmp = OCFD_weno5_kernel_M(&stencil[-ka1-2]);
//*tmp = stencil[-ka1] - 0.5*minmod2(stencil[-ka1+1] - stencil[-ka1], stencil[-ka1] - stencil[-ka1-1]);
}
return 0;
}
} }
break; break;
case 4: case 4:
{ {
{ tmp1 = coords.x + job.start.x - job.end.x;
if(coords.x >= job.end.x-job.start.x-kb1-1){ boundm2(tmp1)
if(coords.x == job.end.x-job.start.x-1){
*tmp = stencil[-ka1-1];
}
if(coords.x == job.end.x-job.start.x-2){
//*tmp = OCFD_weno5_kernel_M_right(&stencil[-ka1-2]);
*tmp = stencil[-ka1] - 0.5*minmod2(stencil[-ka1] - stencil[-ka1-1], stencil[-ka1] - stencil[-ka1-1]);
}
if(coords.x == job.end.x-job.start.x-3){
*tmp = OCFD_weno5_kernel_M_right_plus(&stencil[-ka1-2]);
//*tmp = stencil[-ka1] - 0.5*minmod2(stencil[-ka1+1] - stencil[-ka1], stencil[-ka1] - stencil[-ka1-1]);
}
if(coords.x == job.end.x-job.start.x-4){
//*tmp = stencil[-ka1] - 0.5*minmod2(stencil[-ka1+1] - stencil[-ka1], stencil[-ka1] - stencil[-ka1-1]);
*tmp = OCFD_weno5_kernel_M(&stencil[-ka1-2]);
}
return 0;
}
}
} }
break; break;
case 5: case 5:
{ {
{ tmp1 = coords.y + job.start.y - job.end.y;
if(coords.y >= job.end.y-job.start.y-kb1-1){ boundm2(tmp1)
if(coords.y == job.end.y-job.start.y-1){
*tmp = stencil[-ka1-1];
}
if(coords.y == job.end.y-job.start.y-2){
//*tmp = OCFD_weno5_kernel_M_right(&stencil[-ka1-2]);
*tmp = stencil[-ka1] - 0.5*minmod2(stencil[-ka1] - stencil[-ka1-1], stencil[-ka1] - stencil[-ka1-1]);
}
if(coords.y == job.end.y-job.start.y-3){
*tmp = OCFD_weno5_kernel_M_right_plus(&stencil[-ka1-2]);
//*tmp = stencil[-ka1] - 0.5*minmod2(stencil[-ka1+1] - stencil[-ka1], stencil[-ka1] - stencil[-ka1-1]);
}
if(coords.y == job.end.y-job.start.y-4){
//*tmp = stencil[-ka1] - 0.5*minmod2(stencil[-ka1+1] - stencil[-ka1], stencil[-ka1] - stencil[-ka1-1]);
*tmp = OCFD_weno5_kernel_M(&stencil[-ka1-2]);
}
return 0;
}
}
} }
break; break;
case 6: case 6:
{ {
{ tmp1 = coords.z + job.start.z - job.end.z;
if(coords.z >= job.end.z-job.start.z-kb1-1){ boundm2(tmp1)
if(coords.z == job.end.z-job.start.z-1){
*tmp = stencil[-ka1-1];
}
if(coords.z == job.end.z-job.start.z-2){
//*tmp = OCFD_weno5_kernel_M_right(&stencil[-ka1-2]);
*tmp = stencil[-ka1] - 0.5*minmod2(stencil[-ka1] - stencil[-ka1-1], stencil[-ka1] - stencil[-ka1-1]);
}
if(coords.z == job.end.z-job.start.z-3){
*tmp = OCFD_weno5_kernel_M_right_plus(&stencil[-ka1-2]);
//*tmp = stencil[-ka1] - 0.5*minmod2(stencil[-ka1+1] - stencil[-ka1], stencil[-ka1] - stencil[-ka1-1]);
}
if(coords.z == job.end.z-job.start.z-4){
//*tmp = stencil[-ka1] - 0.5*minmod2(stencil[-ka1+1] - stencil[-ka1], stencil[-ka1] - stencil[-ka1-1]);
*tmp = OCFD_weno5_kernel_M(&stencil[-ka1-2]);
}
return 0;
}
}
} }
break; break;
case 7:
{
boundm1(coords.x)
tmp1 = coords.x + job.start.x - job.end.x;
boundm2(tmp1)
}
break;
case 8:
{
boundm1(coords.y)
tmp1 = coords.y + job.start.y - job.end.y;
boundm2(tmp1)
}
break;
case 9:
{
boundm1(coords.z)
tmp1 = coords.z + job.start.z - job.end.z;
boundm2(tmp1)
}
break;
} }
return 1; return 1;
......
src/OCFD_init.cu
View file @ 9e0b6851
...@@ -315,8 +315,8 @@ void opencfd_mem_init_boundary(){ ...@@ -315,8 +315,8 @@ void opencfd_mem_init_boundary(){
pu_dist_upper = (REAL*)malloc_me(sizeof(REAL)*ny*nx); pu_dist_upper = (REAL*)malloc_me(sizeof(REAL)*ny*nx);
pfx = (REAL*)malloc_me(sizeof(REAL)*nx); pfx = (REAL*)malloc_me(sizeof(REAL)*nx);
pgz = (REAL*)malloc_me(sizeof(REAL)*ny); pgz = (REAL*)malloc_me(sizeof(REAL)*ny);
TM = (REAL*)malloc_me(sizeof(REAL)*MTMAX); TM = (REAL*)malloc_me(sizeof(REAL)*abs(MTMAX));
fait = (REAL*)malloc_me(sizeof(REAL)*MTMAX); fait = (REAL*)malloc_me(sizeof(REAL)*abs(MTMAX));
new_cudaField(&pu2d_inlet_d, ny, nz, 5); new_cudaField(&pu2d_inlet_d, ny, nz, 5);
new_cudaField(&pu2d_upper_d, nx, ny, 5); new_cudaField(&pu2d_upper_d, nx, ny, 5);
...@@ -330,8 +330,8 @@ void opencfd_mem_init_boundary(){ ...@@ -330,8 +330,8 @@ void opencfd_mem_init_boundary(){
pfx = (REAL*)malloc_me(sizeof(REAL)*nx); pfx = (REAL*)malloc_me(sizeof(REAL)*nx);
pgz = (REAL*)malloc_me(sizeof(REAL)*nz); pgz = (REAL*)malloc_me(sizeof(REAL)*nz);
TM = (REAL*)malloc_me(sizeof(REAL)*MTMAX); TM = (REAL*)malloc_me(sizeof(REAL)*abs(MTMAX));
fait = (REAL*)malloc_me(sizeof(REAL)*MTMAX); fait = (REAL*)malloc_me(sizeof(REAL)*abs(MTMAX));
//ptmp = (REAL*)malloc_me_Host(sizeof(REAL)*10*nx*nz); //ptmp = (REAL*)malloc_me_Host(sizeof(REAL)*10*nx*nz);
new_cudaField(&pub1_d, ny, 4, 1); new_cudaField(&pub1_d, ny, 4, 1);
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment