first commit

head_hip/OCFD_boundary.h

+#ifndef __OCFD_BOUNDARY_H
+#define __OCFD_BOUNDARY_H
+
+#include "parameters.h"
+#ifdef __cplusplus
+extern "C"{
+#endif
+
+void OCFD_bc();
+
+#ifdef __cplusplus
+}
+#endif
+#endif
\ No newline at end of file

head_hip/OCFD_boundary_Liftbody3D.h

+#ifndef __OCFD_BOUNDARY_LIFTBODY3D_H
+#define __OCFD_BOUNDARY_LIFTBODY3D_H
+#include "parameters.h"
+
+#ifdef __cplusplus
+extern "C"{
+#endif
+
+void bc_user_Liftbody3d();
+void bc_user_Liftbody3d_simple();
+void bc_user_Liftbody3d_plus();
+
+#ifdef __cplusplus
+}
+#endif
+#endif
\ No newline at end of file

head_hip/OCFD_boundary_compression_conner.h

+#ifndef __OCFD_BOUNDARY_COMPRESSION_H
+#define __OCFD_BOUNDARY_COMPRESSION_H
+#include "parameters.h"
+
+#ifdef __cplusplus
+extern "C"{
+#endif
+
+void get_ht_multifrequancy(REAL HT, REAL TT, int MT_MAX, REAL beta);
+void bc_user_Compression_conner();
+
+#ifdef __cplusplus
+}
+#endif
+#endif
\ No newline at end of file

head_hip/OCFD_boundary_init.h

+#ifndef __OCFD_BOUNDARY_INIT_H
+#define __OCFD_BOUNDARY_INIT_H
+#include "parameters.h"
+#include "cuda_commen.h"
+
+#ifdef __cplusplus
+extern "C"{
+#endif
+
+// used in boudary_liftbody*********************************************************************
+extern REAL *pu2d_inlet; //[5][nz][ny]
+extern REAL *pu2d_upper; //[5][ny][nx]
+        
+extern REAL * pv_dist_wall; // [ny][nx]
+extern REAL *pv_dist_coeff; // [3][ny][nx]
+extern REAL * pu_dist_upper; // [ny][nx]
+
+extern cudaField *pu2d_inlet_d; //[5][nz][ny]
+extern cudaField *pu2d_upper_d; //[5][ny][nx]
+//extern cudaField *pv_dist_wall_d; // [ny][nx]
+extern cudaField *pv_dist_coeff_d; // [3][ny][nx]
+extern cudaField *pu_dist_upper_d; // [ny][nx]
+    
+void bc_user_Liftbody3d_init();
+//**********************************************************************************************
+
+// used in boundary_compressible_conner*********************************************************
+extern REAL *pub1; // [ny][4]
+extern REAL *pfx; // [nx]
+extern REAL *pgz; // [nz]
+extern REAL *TM; // [MTMAX]
+extern REAL *fait; // [MTMAX]
+
+extern cudaField *pub1_d; // [ny][4]
+extern cudaField *pfx_d; // [nx]
+extern cudaField *pgz_d; // [nz]
+
+void bc_user_Compression_conner_init();
+//**********************************************************************************************
+
+
+void get_fait_multifrequancy(int MT_MAX);
+
+void get_xy_blow_suction_multiwave(int NX, int MZ_MAX, REAL *xx,
+REAL *fx, REAL *gz, REAL DIST_BEGIN, REAL DIST_END);
+
+void get_xs_blow_suction_multiwave(int NX, int NZ, int MZ_MAX, REAL *xx,
+REAL *zz, REAL SL, REAL *fx, REAL *gz, REAL DIST_BEGIN, REAL DIST_END);
+
+
+#ifdef __cplusplus
+}
+#endif
+#endif

head_hip/OCFD_filtering.h

+#ifndef __OCFD_FILTERINT_H
+#define __OCFD_FILTERINT_H
+#include "parameters.h"
+#include "cuda_commen.h"
+
+
+#ifdef __cplusplus
+extern "C"{
+#endif
+void filtering(REAL *pf,REAL *pf0,REAL *pp);
+
+void filter_x3d(REAL *pf, REAL *pf0, REAL s0, int ib, int ie, int jb, int je, int kb, int ke);
+void filter_y3d(REAL *pf, REAL *pf0, REAL s0, int ib, int ie, int jb, int je, int kb, int ke);
+void filter_z3d(REAL *pf, REAL *pf0, REAL s0, int ib, int ie, int jb, int je, int kb, int ke);
+
+void filter_x3d_shock(cudaSoA *pf, cudaSoA *pf0, cudaField *pp, REAL s0, REAL rth, int ib, int ie, int jb, int je, int kb, int ke, int IF_Filter);
+void filter_y3d_shock(cudaSoA *pf, cudaSoA *pf0, cudaField *pp, REAL s0, REAL rth, int ib, int ie, int jb, int je, int kb, int ke, int IF_Filter);
+void filter_z3d_shock(cudaSoA *pf, cudaSoA *pf0, cudaField *pp, REAL s0, REAL rth, int ib, int ie, int jb, int je, int kb, int ke, int IF_Filter);
+
+void set_para_filtering();
+#ifdef __cplusplus
+}
+#endif
+#endif
\ No newline at end of file

head_hip/OCFD_flux_charteric.h

+#include "hip/hip_runtime.h"
+#ifndef __OCFD_FLUX_CHARTERIC_H
+#define __OCFD_FLUX_CHARTERIC_H
+#include "parameters.h"
+#include "cuda_commen.h"
+
+#ifdef __cplusplus
+extern "C"{
+#endif
+
+__global__ void OCFD_weno7_SYMBO_character_P_kernel(int WENO_LMT_FLAG, 
+dim3 flagxyzb, cudaSoA f, cudaSoA du, cudaField Ajac, cudaField u, cudaField v,
+cudaField w, cudaField cc, cudaField Ax, cudaField Ay, cudaField Az,
+cudaJobPackage job);
+
+
+__global__ void OCFD_weno7_SYMBO_character_M_kernel(int WENO_LMT_FLAG, 
+dim3 flagxyzb, cudaSoA f, cudaSoA du, cudaField Ajac, cudaField u, cudaField v,
+cudaField w, cudaField cc, cudaField Ax, cudaField Ay, cudaField Az,
+cudaJobPackage job);
+
+
+__global__ void OCFD_HybridAuto_character_P_kernel( 
+dim3 flagxyzb, cudaSoA f, cudaSoA du, cudaField Ajac, cudaField u, cudaField v,
+cudaField w, cudaField cc, cudaField Ax, cudaField Ay, cudaField Az,
+cudaField_int scheme, cudaJobPackage job);
+
+
+__global__ void OCFD_HybridAuto_character_M_kernel(
+dim3 flagxyzb, cudaSoA f, cudaSoA du, cudaField Ajac, cudaField u, cudaField v,
+cudaField w, cudaField cc, cudaField Ax, cudaField Ay, cudaField Az,
+cudaField_int scheme, cudaJobPackage job);
+
+
+__global__ void OCFD_HybridAuto_character_P_Jameson_kernel( 
+dim3 flagxyzb, cudaSoA f, cudaSoA du, cudaField Ajac, cudaField u, cudaField v,
+cudaField w, cudaField cc, cudaField Ax, cudaField Ay, cudaField Az,
+cudaField_int scheme, cudaJobPackage job);
+
+
+__global__ void OCFD_HybridAuto_character_M_Jameson_kernel(
+dim3 flagxyzb, cudaSoA f, cudaSoA du, cudaField Ajac, cudaField u, cudaField v,
+cudaField w, cudaField cc, cudaField Ax, cudaField Ay, cudaField Az,
+cudaField_int scheme, cudaJobPackage job);
+
+#ifdef __cplusplus
+}
+#endif
+#endif
\ No newline at end of file

head_hip/OCFD_init.h

+#ifndef __OCFD_INIT_H
+#define __OCFD_INIT_H
+#include "parameters.h"
+
+#ifdef __cplusplus
+extern "C"{
+#endif
+
+void init();
+
+void opencfd_mem_init_all();
+void opencfd_mem_finalize_all();
+
+void opencfd_mem_init();
+void opencfd_mem_finalize();
+void opencfd_para_init();
+
+void opencfd_mem_init_boundary();
+void opencfd_mem_finalize_boundary();
+
+void opencfd_mem_init_dev();
+void opencfd_mem_finalize_dev();
+void opencfd_para_init_dev();
+
+#ifdef __cplusplus
+}
+#endif
+#endif

head_hip/OCFD_mpi.h

+#ifndef __OCFD_MPI_H
+#define __OCFD_MPI_H
+#include "parameters.h"
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+void opencfd_mem_init_mpi();
+void opencfd_mem_finalize_mpi();
+void mpi_init(int * , char ***);
+void mpi_finalize();
+void part();
+static inline int idx2int(int i,int j,int k){
+    return(i + j*(nx+2*LAP) + k*(nx+2*LAP)*(ny+2*LAP));
+}
+int my_mod1(int i,int n);
+void New_MPI_datatype();
+void get_i_node(int i_global,int * node_i, int * i_local);
+void get_j_node(int j_global,int * node_j, int * j_local);
+void get_k_node(int k_global,int * node_k, int * k_local);
+int get_id(int npx1,int npy1,int npz1);
+
+
+void exchange_boundary_xyz(REAL *pf);
+void exchange_boundary_x(REAL *pf,int Iperiodic1);
+void exchange_boundary_y(REAL *pf,int Iperiodic1);
+void exchange_boundary_z(REAL *pf,int Iperiodic1);
+
+
+void exchange_boundary_x_standard(REAL *pf, int Iperiodic1);
+void exchange_boundary_y_standard(REAL *pf, int Iperiodic1);
+void exchange_boundary_z_standard(REAL *pf, int Iperiodic1);
+
+
+void exchange_boundary_x_deftype(REAL * pf);
+void exchange_boundary_y_deftype(REAL * pf);
+void exchange_boundary_z_deftype(REAL * pf);
+
+#ifdef __cplusplus
+}
+#endif 
+#endif
\ No newline at end of file

head_hip/OCFD_mpi_dev.h

+#ifndef __OCFD_MPI_DEV_H
+#define __OCFD_MPI_DEV_H
+#include "cuda_commen.h"
+
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+void exchange_boundary_xyz_dev(REAL *hostptr, cudaField * devptr);
+void exchange_boundary_x_dev(REAL *hostptr, cudaField * devptr, int Iperiodic1);
+void exchange_boundary_y_dev(REAL *hostptr, cudaField * devptr, int Iperiodic1);
+void exchange_boundary_z_dev(REAL *hostptr, cudaField * devptr, int Iperiodic1);
+
+
+void exchange_boundary_x_standard_dev(REAL *hostptr, cudaField * devptr, int Iperiodic1);
+void exchange_boundary_y_standard_dev(REAL *hostptr, cudaField * devptr, int Iperiodic1);
+void exchange_boundary_z_standard_dev(REAL *hostptr, cudaField * devptr, int Iperiodic1);
+
+void opencfd_mem_init_mpi_dev();
+void opencfd_mem_finalize_mpi_dev();
+
+void exchange_boundary_xyz_Async_packed_dev(REAL *hostptr, cudaField * devptr , hipStream_t *stream);
+void exchange_boundary_x_Async_packed_dev(REAL *hostptr , cudaField * devptr, int Iperiodic1 , hipStream_t *stream);
+void exchange_boundary_y_Async_packed_dev(REAL *hostptr , cudaField * devptr, int Iperiodic1 , hipStream_t *stream);
+void exchange_boundary_z_Async_packed_dev(REAL *hostptr , cudaField * devptr, int Iperiodic1 , hipStream_t *stream);
+
+void exchange_boundary_xyz_packed_dev(REAL *hostptr, cudaField * devptr);
+void exchange_boundary_x_packed_dev(REAL *hostptr , cudaField * devptr, int Iperiodic1);
+void exchange_boundary_y_packed_dev(REAL *hostptr , cudaField * devptr, int Iperiodic1);
+void exchange_boundary_z_packed_dev(REAL *hostptr , cudaField * devptr, int Iperiodic1);
+
+#ifdef __cplusplus
+}
+#endif
+#endif
\ No newline at end of file

head_hip/OCFD_split.h

+#include "hip/hip_runtime.h"
+#ifndef __OCFD_SPLIT_H
+#define __OCFD_SPLIT_H
+#include "cuda_commen.h"
+
+#ifdef __cplusplus
+extern "C"{
+#endif
+//__global__ void split_Jac3d_Stager_Warming_ker(cudaField d0, cudaField u0, cudaField v0, cudaField w0, cudaField cc0, cudaSoA fp, cudaSoA fm, cudaField Akx, cudaField Aky, cudaField Akz, REAL tmp0, REAL split_C1, REAL split_C3, cudaJobPackage job);
+
+typedef struct sw_split_
+{
+	cudaField d;
+	cudaField u;
+	cudaField v;
+	cudaField w;
+	cudaField cc;
+
+	cudaField Akx;
+	cudaField Aky;
+	cudaField Akz;
+	cudaField Aix;
+	cudaField Aiy;
+	cudaField Aiz;
+	cudaField Asx;
+	cudaField Asy;
+	cudaField Asz;
+} sw_split;
+
+__global__ void split_Jac3d_Stager_Warming_ker(sw_split sw, cudaSoA fp_x, cudaSoA fm_x, cudaSoA fp_y, cudaSoA fm_y, cudaSoA fp_z, cudaSoA fm_z, REAL tmp0, REAL split_C1, REAL split_C3, cudaJobPackage job);
+
+void Stager_Warming(cudaJobPackage job_in, cudaSoA *fp_x, cudaSoA *fm_x, cudaSoA *fp_y, cudaSoA *fm_y, cudaSoA *fp_z, cudaSoA *fm_z, hipStream_t *stream);
+
+typedef struct sw_split_out_
+{
+	cudaField d;
+	cudaField u;
+	cudaField v;
+	cudaField w;
+	cudaField cc;
+
+	cudaField Ax;
+	cudaField Ay;
+	cudaField Az;
+} sw_split_out;
+
+void Stager_Warming_out(cudaJobPackage job_in, cudaSoA *fp, cudaSoA *fm, hipStream_t *stream);
+
+#ifdef __cplusplus
+}
+#endif
+#endif
\ No newline at end of file

head_hip/OCFD_time.h

+#include "hip/hip_runtime.h"
+#ifndef __OCFD_TIME_H
+#define __OCFD_TIME_H
+
+#include "cuda_commen.h"
+
+#ifdef __cplusplus
+extern "C"{
+#endif
+
+void OCFD_time_advance(int KRK);
+void OCFD_time_advance_plus(int KRK);
+
+__global__ void OCFD_time_advance_ker1(cudaSoA f , cudaSoA fn , cudaSoA du , cudaJobPackage job);
+__global__ void OCFD_time_advance_ker2(cudaSoA f , cudaSoA fn , cudaSoA du , cudaJobPackage job);
+__global__ void OCFD_time_advance_ker3(cudaSoA f , cudaSoA fn , cudaSoA du , cudaSoA pf_lap , cudaJobPackage job);
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif
\ No newline at end of file

head_hip/OCFD_warp_shuffle.h

+#ifndef __OCFD_WARP_SHUFFLE_H
+#define __OCFD_WARP_SHUFFLE_H
+#include "cuda_commen.h"
+#include "hip/hip_runtime.h"
+#include "hip/hip_runtime.h"
+#include "hip/hip_runtime_api.h"
+#ifdef __cplusplus
+extern "C"{
+#endif
+
+
+#ifdef __NVCC__
+__device__ __forceinline__ double __shfl_up_double(double & val , unsigned char delta , unsigned char width ){
+    return __shfl_up_sync(0xffffffff , val , delta , width);
+}
+__device__ __forceinline__ double __shfl_down_double(double & val , unsigned char delta , unsigned char width ){
+    return __shfl_down_sync(0xffffffff , val , delta , width);
+}
+__device__ __forceinline__ double __shfl_double(double & val , unsigned char srcLane , unsigned char width){
+    return __shfl_sync(0xffffffff , val , srcLane , width);
+}
+
+__device__ __forceinline__ double __shfl_xor_double(double & val , unsigned char srcLane , unsigned char width){
+    return __shfl_xor_sync(0xffffffff , val , srcLane , width);
+}
+
+#else
+
+#define __shfl_up_double(val , delta , witdh) __shfl_up_double_( *( (int2*)(&val) ) , delta , witdh)
+__device__ __forceinline__ double __shfl_up_double_(int2 & val , unsigned char delta , unsigned char width ){
+    int2 out = *( (int2*)(&val) );
+    out.x = __shfl_up(out.x , delta , width);
+    out.y = __shfl_up(out.y , delta , width);
+    return ( *( (double*)(&out) ) );
+}
+
+#define __shfl_down_double(val , delta , witdh) __shfl_down_double_( *( (int2*)(&val) ) , delta , witdh)
+__device__ __forceinline__ double __shfl_down_double_(int2 & val , unsigned char delta , unsigned char width ){
+    int2 out = *( (int2*)(&val) );
+    out.x = __shfl_down(out.x , delta , width);
+    out.y = __shfl_down(out.y , delta , width);
+    return ( *( (double*)(&out) ) );
+}
+
+#define __shfl_double(val , delta , witdh) __shfl_double_( *( (int2*)(&val) ) , delta , witdh)
+__device__ __forceinline__ double __shfl_double_(int2 & val , unsigned char srcLane , unsigned char width){
+    int2 out = *( (int2*)(&val) );
+    out.x = __shfl(out.x , srcLane , width);
+    out.y = __shfl(out.y , srcLane , width);
+    return ( *( (double*)(&out) ) );
+}
+
+#define __shfl_xor_double(val , delta , witdh) __shfl_xor_double_( *( (int2*)(&val) ) , delta , witdh)
+__device__ __forceinline__ double __shfl_xor_double_(int2 & val , unsigned char srcLane , unsigned char width){
+    int2 out = *( (int2*)(&val) );
+    out.x = __shfl_xor(out.x , srcLane , width);
+    out.y = __shfl_xor(out.y , srcLane , width);
+    return ( *( (double*)(&out) ) );
+}
+
+#endif
+
+
+#ifdef __cplusplus
+}
+#endif
+#endif

head_hip/commen_kernel.h

+#include "hip/hip_runtime.h"
+#ifndef _COMMEN_KERNEL_H
+#define _COMMEN_KERNEL_H
+#include "cuda_commen.h"
+
+#ifdef __cplusplus
+extern "C"{
+#endif
+
+__global__ void cuda_mem_value_init(REAL value , REAL * ptr , unsigned int pitch , unsigned int size_x , unsigned int size_y , unsigned int size_z);
+void cuda_mem_value_init_warp(REAL value , REAL * ptr , unsigned int pitch , unsigned int size_x , unsigned int size_y , unsigned int size_z);
+
+// eyes on no-lap region
+__global__ void pri_to_cons_kernel(cudaSoA pcons , cudaField pd , cudaField pu , cudaField pv , cudaField pw , cudaField pT , cudaJobPackage job);
+void pri_to_cons_kernel_warp(cudaSoA *pcons , cudaField *pd , cudaField *pu , cudaField *pv , cudaField *pw , cudaField *pT , cudaJobPackage job_in , dim3 blockdim_in );
+
+__global__ void cons_to_pri_kernel(cudaSoA f, cudaField d , cudaField u , cudaField v , cudaField w , cudaField T , cudaField P , cudaJobPackage job);
+void get_duvwT();
+
+__global__ void get_Amu_kernal(cudaField Amu , cudaField T , cudaJobPackage job);
+void get_Amu();
+
+__global__ void sound_speed_kernel(cudaField T , cudaField cc , cudaJobPackage job);
+
+
+__global__ void YF_Pe_XF(cudaField yF , cudaField xF , cudaField AJac , cudaJobPackage job);
+__global__ void ZF_e_XF_P_YF(cudaField out , cudaField xF , cudaField yF , cudaJobPackage job);
+__global__ void ZF_e_XF_P_YF_LAP(cudaField out , cudaField xF , cudaField yF , cudaJobPackage job);
+__global__ void ZF_Pe_XF_P_YF(cudaField zF , cudaField xF , cudaField yF , cudaField AJac , cudaJobPackage job);
+
+
+
+
+/* ========================================= */
+// inline function
+// #include "config_parameters.h"
+// #include "cuda_commen.h"
+// #include "cuda_utility.h"
+// #include "parameters_d.h"
+// __device__ inline void cons_to_pri_dev_fun(cudaField & d , cudaField & u , cudaField & v , cudaField & w , cudaField & T , cudaField & P , REAL & f0 , REAL & f1 , REAL & f2 , REAL & f3 , REAL & f4 ){
+//         get_Field_LAP(d , x+LAP , y+LAP , z+LAP) = f0;
+
+//         REAL u = f1/f0;
+//         get_Field_LAP(u , x+LAP , y+LAP , z+LAP) = u;
+
+//         REAL v = f2/f0;
+//         get_Field_LAP(v , x+LAP , y+LAP , z+LAP) = v;
+
+//         REAL w = f3/f0;
+//         get_Field_LAP(w , x+LAP , y+LAP , z+LAP) = w;
+
+//         REAL tmp = f4 - 0.5*f0*(u*u + v*v + w*w);
+//         get_Field_LAP(T , x+LAP , y+LAP , z+LAP) = tmp/(f0*Cv_d);
+//         tmp = tmp/d1;
+//         get_Field_LAP(P , x+LAP , y+LAP , z+LAP) = tmp*(Gamma_d - 1.0);
+// }
+
+// __device__ inline void get_Amu_kernal(cudaField & Amu , REAL & t){
+//         get_Field(Amu , x,y,z) = amu_C0_d * sqrt(t * t * t) / (Tsb_d + t);
+// }
+
+
+// __device__ inline void sound_speed_kernel(cudaField & cc , REAL & t){
+//         get_Field_LAP(cc , x,y,z) = sqrt( get_Field_LAP(T , x,y,z) )/Ama_d;
+// }
+#ifdef __cplusplus
+}
+#endif
+#endif

head_hip/config_parameters.h

+#ifndef __READ_PARAMETERS_H
+#define __READ_PARAMETERS_H
+
+#include "stdio.h"
+
+#ifdef __cplusplus
+extern "C"{
+#endif
+
+// config macro
+#define DEBUG_MODE
+
+typedef double REAL;
+// boundary lap size
+#define LAP 4
+// unknown number
+#define NVARS 5                         
+
+#define hipWarpSize 64
+
+typedef struct scheme_choose{
+    char *invis;
+    char *vis;
+}SCHEME_CHOOSE;
+
+#define Negatie_T_limit 100
+#define NegT_Max 500
+
+typedef struct TYPE_NegT_{
+    int Num_NegT;
+    int NT_Point[NegT_Max][3];
+}TYPE_NegT;
+
+typedef struct configItem_
+{
+    char name[1000];
+    char value[1000];
+} configItem;
+
+
+int ExtarctItem(char *src, char *name, char *value);
+int ItemNUM(FILE *file, char *Item_name, int *NameNUM);
+int PartItem(char *src, char part[][1000]);
+void ModifyItem(char *name, char *buff);
+void RemovalNUM(char *buff);
+void SearchItem(FILE *file, configItem *List, int configNum);
+void Schemes_Choose_ID(SCHEME_CHOOSE *scheme);
+
+#ifdef __cplusplus
+}
+#endif
+#endif
+

head_hip/cuda_commen.h

+#ifndef _CUDA_HEAD_H
+#define _CUDA_HEAD_H
+
+#include "config_parameters.h"
+#include "hip/hip_runtime.h"
+#include "hip/hip_runtime.h"
+#include "mpi.h"
+
+#ifdef __cplusplus
+extern "C"{
+#endif
+
+
+#ifdef DEBUG_MODE
+#include "stdio.h"
+#define CUDA_CALL(CALL) \
+    {\
+        hipError_t error;\
+        error = CALL;\
+        if(error != hipSuccess){\
+            printf("In file \"%s\" , line %d ( \"%s\" ) , cuda call failed\n",__FILE__ , __LINE__ , __FUNCTION__);\
+            printf("Error code : %s  (%s)\n",hipGetErrorString(error) , hipGetErrorName(error));\
+            MPI_Abort(((MPI_Comm)0x44000000) , 1);\
+        }\
+    }
+#else
+#define CUDA_CALL(CALL) CALL;
+#endif
+
+#ifdef DEBUG_MODE
+#include "stdio.h"
+#define CUDA_LAUNCH(call)\
+    {\
+        hipError_t error;\
+        call;\
+        error = hipGetLastError();\
+        if( error != hipSuccess){\
+            printf("In file \"%s\" , line %d ( \"%s\" ) , cuda call failed\n",__FILE__ , __LINE__ , __FUNCTION__);\
+            printf("Error code : %s  (%s)\n",hipGetErrorString(error) , hipGetErrorName(error));\
+            MPI_Abort(((MPI_Comm)0x44000000) , 1);\
+        }\
+    }
+#else
+#define CUDA_LAUNCH(call)\
+    call;
+#endif
+
+typedef struct cudaField_ {
+    REAL * ptr;
+    unsigned int pitch;
+    #ifdef __HIPCC__
+    unsigned int dummy0;
+    #endif
+} cudaField;
+
+typedef struct cudaField_int_ {
+    int * ptr;
+    unsigned int pitch;
+    #ifdef __HIPCC__
+    unsigned int dummy0;
+    #endif
+} cudaField_int;
+
+typedef struct cudaSoA_{
+    REAL * ptr;
+    unsigned int pitch;
+    unsigned int length_Y;
+    unsigned int length_Z;
+    #ifdef __HIPCC__
+    unsigned int dummy0;
+    #endif
+} cudaSoA;
+
+
+typedef struct cudaFieldPack_{
+    REAL * ptr;
+} cudaFieldPack;
+
+
+extern uint32_t BlockDimX;
+extern uint32_t BlockDimY;
+extern uint32_t BlockDimZ;
+
+extern dim3 BlockDim_X;
+extern dim3 BlockDim_Y;
+extern dim3 BlockDim_Z;
+
+extern dim3 BlockDim;
+
+extern int MaxThreadsPerBlock;
+extern int MaxBlockDimX;
+extern int MaxBlockDimY;
+extern int MaxBlockDimZ;
+extern int MaxGridDimX;
+extern int MaxGridDimY;
+extern int MaxGridDimZ;
+
+
+void cuda_commen_init();
+
+
+typedef struct cudaJobPackage_
+{
+    dim3 start;
+    dim3 end;
+    #ifdef __cplusplus
+    __host__ __device__ cudaJobPackage_(){};
+    __host__ __device__ cudaJobPackage_(const dim3 & start_in , const dim3 & end_in):start(start_in),end(end_in){};
+    __host__ __device__ void setup(const dim3 & start_in , const dim3 & end_in){start = start_in ,end = end_in;};
+    #endif
+} cudaJobPackage;
+
+inline void jobsize(cudaJobPackage *job , dim3 *size){
+    size->x = job->end.x - job->start.x;
+    size->y = job->end.y - job->start.y;
+    size->z = job->end.z - job->start.z;
+}
+
+inline void job2size(cudaJobPackage *job , dim3 *size){
+    size->x = job->end.x;
+    size->y = job->end.y;
+    size->z = job->end.z;
+}
+inline void job2job(cudaJobPackage * job , cudaJobPackage * out){
+    out->start.x = job->start.x;
+    out->start.y = job->start.y;
+    out->start.z = job->start.z;
+    
+    out->end.x = job->end.x + job->start.x;
+    out->end.y = job->end.y + job->start.y;
+    out->end.z = job->end.z + job->start.z;
+}
+
+#ifdef __cplusplus
+}
+#endif
+#endif

head_hip/cuda_utility.h

+#ifndef _CUDA_UTILITY_H
+#define _CUDA_UTILITY_H
+
+#include "cuda_commen.h"
+#include "config_parameters.h"
+
+
+#ifdef __cplusplus
+extern "C"{
+#endif
+
+void * malloc_me_d(unsigned int * pitch , unsigned int size_x , unsigned size_y , unsigned size_z);
+
+
+
+#define get_Field(Field , x,y,z,...) (*(Field.ptr + (__VA_ARGS__ + x + Field.pitch *(y + (z)*ny_d) ) ) )
+#define get_Field_LAP(Field , x,y,z,...) (*(Field.ptr + (__VA_ARGS__ + x + Field.pitch *(y + (z)*ny_2lap_d) ) ) )
+
+void new_cudaField( cudaField ** pField, unsigned int size_x , unsigned int size_y , unsigned int size_z );
+void delete_cudaField(cudaField * pField);
+
+void new_cudaField_int( cudaField_int ** pField, unsigned int size_x , unsigned int size_y , unsigned int size_z );
+void delete_cudaField_int(cudaField_int * pField);
+
+
+#define get_SoA(SoA , x,y,z , var,...) (*( SoA.ptr + (__VA_ARGS__ + x + SoA.pitch*(y + ny_d*(z+ (var)*nz_d)))))
+#define get_SoA_LAP(SoA , x,y,z , var,...) (*( SoA.ptr + (__VA_ARGS__ + x + SoA.pitch*(y + ny_2lap_d*(z+ (var)*nz_2lap_d)))))
+
+void new_cudaSoA( cudaSoA ** pSoA, unsigned int size_x , unsigned int size_y , unsigned int size_z );
+void delete_cudaSoA(cudaSoA * pSoA);
+
+void new_cudaFieldPack(cudaFieldPack ** pack , unsigned int size_x , unsigned int size_y , unsigned int size_z);
+void delete_cudaFieldPack(cudaFieldPack * pField);
+
+
+void cal_grid_block_dim(dim3 * pgrid , dim3 * pblock , unsigned int threadx , unsigned int thready , unsigned int threadz ,  unsigned int size_x , unsigned int size_y , unsigned int size_z);
+
+
+
+enum {
+    H2D = 0,
+    D2H = 1
+};
+void memcpy_All(REAL *hostPtr, REAL *devPtr, unsigned int pitch, int mode, unsigned int size_x, unsigned int size_y, unsigned int size_z);
+void memcpy_All_int(int *hostPtr, int *devPtr, unsigned int pitch, int mode, unsigned int size_x, unsigned int size_y, unsigned int size_z);
+void memcpy_inner(REAL *hostPtr, REAL *devPtr, unsigned int pitch, int mode, unsigned int size_x, unsigned int size_y, unsigned int size_z);
+void memcpy_bound_x(REAL *hostPtr, REAL *devPtr, unsigned int pitch, int mode, unsigned int size_x, unsigned int size_y, unsigned int size_z);
+void memcpy_bound_y(REAL *hostPtr, REAL *devPtr, unsigned int pitch, int mode, unsigned int size_x, unsigned int size_y, unsigned int size_z);
+void memcpy_bound_z(REAL *hostPtr, REAL *devPtr, unsigned int pitch, int mode, unsigned int size_x, unsigned int size_y, unsigned int size_z);
+
+#ifdef __cplusplus
+}
+#endif
+#endif
\ No newline at end of file

head_hip/init.h

+#ifndef __INIT_H
+#define __INIT_H
+#include "parameters.h"
+
+#ifdef __cplusplus
+extern "C"{
+#endif
+
+#ifdef __HIPCC__
+void opencfd_para_init_dev();
+#endif
+#ifdef __cplusplus
+}
+#endif
+#endif
\ No newline at end of file

head_hip/io_warp.h

+#ifndef __IO_WARP_H
+#define __IO_WARP_H
+
+#include "stdio.h"
+
+#define FWRITE(ptr , size , num , stream) \
+    {   \
+        int _ENCODE_START = num*size;\
+        int _ENCODE_END = num*size;\
+        fwrite(&_ENCODE_START , sizeof(int) , 1 , stream);\
+        fwrite(ptr , size , num , stream);\
+        fwrite(&_ENCODE_END , sizeof(int) , 1 , stream);\
+    }
+    
+#define FREAD(ptr , size , num , stream) \
+    {   int tmp_buffer;\
+        fread(&tmp_buffer , sizeof(int) , 1 , stream);\
+        fread(ptr , size , num , stream);\
+        fread(&tmp_buffer , sizeof(int) , 1 , stream);\
+    }
+
+#endif
+    
\ No newline at end of file

head_hip/parameters.h

+#ifndef __PARAMETER_H
+#define __PARAMETER_H
+#include "mpi.h"
+#include "pthread.h"
+#include "config_parameters.h"
+#include "OCFD_Schemes_hybrid_auto.h"
+#include "stdarg.h"
+
+#ifdef __cplusplus
+extern "C"{
+#endif
+
+// ------For Doubleprecision  (real*8)------------------------------------------------------------------
+extern int OCFD_REAL_KIND;
+extern MPI_Datatype OCFD_DATA_TYPE;   //double precison computing
+//  ------For Single precision (real*4)-----------------------
+// typedef float REAL;
+// int OCFD_REAL_KIND=4
+// MPI_Datatype OCFD_DATA_TYPE;             //  single precision computing
+// ===========  Parameters for MPI ==========================================================----------------
+ 
+
+
+// -----constant-----
+extern REAL Re , Pr , Ama , Gamma , Ref_T , epsl_SW , PI;
+extern REAL Cv , Cp , Tsb , amu_C0;
+extern REAL split_C1 , split_C3, tmp0;
+extern REAL vis_flux_init_c;
+
+// --------MPI-------------------
+extern int my_id,npx,npy,npz;  //全局即方向id
+extern int NPX0 , NPY0 , NPZ0; // proc number on each direction
+extern int ID_XP1,ID_XM1,ID_YP1,ID_YM1,ID_ZP1,ID_ZM1; //邻居全局id
+
+extern MPI_Status status;
+extern MPI_Comm MPI_COMM_X,MPI_COMM_Y,MPI_COMM_Z,MPI_COMM_XY,MPI_COMM_XZ,MPI_COMM_YZ;
+extern MPI_Datatype TYPE_LAPX1,TYPE_LAPY1,TYPE_LAPZ1,TYPE_LAPX2,TYPE_LAPY2,TYPE_LAPZ2;
+
+extern int *i_offset,*j_offset,*k_offset,*i_nn,*j_nn,*k_nn; //某个方向的分块信息
+extern int MSG_BLOCK_SIZE;
+
+extern unsigned int nx,ny,nz; // 某方向所处理的个数
+extern unsigned int NX_GLOBAL,NY_GLOBAL,NZ_GLOBAL;
+extern unsigned int nx_lap,ny_lap,nz_lap;
+extern unsigned int nx_2lap,ny_2lap,nz_2lap;
+
+extern int Stream_MODE; //Stream 模式
+extern int TEST;
+extern pthread_t* thread_handles;
+// --------------------------------------------------------------------------------------------------------
+extern REAL dt,end_time,tt;
+extern REAL cpu_time;
+extern int Istep , end_step;
+
+// -----------Analysis and Save------------------------------------------
+
+extern int OCFD_ANA_time_average; 
+extern int OCFD_ana_flatplate,  OCFD_ana_saveplaneYZ;
+extern int OCFD_ana_saveplaneXZ, OCFD_ana_savePoints;
+extern int OCFD_ana_saveplaneXY, OCFD_ana_saveblock;
+extern int OCFD_ana_getQ;
+
+extern int Kstep_save, Kstep_show,N_ana,*K_ana,*Kstep_ana,KRK;
+
+//------------Scheme invis and vis----------------------------------------
+extern int Scheme_invis_ID;
+extern int Scheme_vis_ID;
+
+//---------------------WENO_SYMBO_Limiter------------------------------------------
+extern REAL WENO_TV_Limiter;
+extern REAL WENO_TV_MAX;
+
+// -----------Boundary Condition and Initial Condition-----------------------------
+extern int Iperiodic[3], Jacbound[3], D0_bound[6];
+extern int Non_ref[6];
+
+extern int Y_Jacbound,Y_bound0;
+
+extern int OCFD_BC_Liftbody3d;
+extern int IBC_USER;
+extern int Init_stat;
+
+extern char IF_SYMMETRY;
+extern char IF_WITHLEADING;  // 0 不含头部， 1 含头部 
+extern char IFLAG_UPPERBOUNDARY; // 0 激波外； 1 激波
+extern REAL AOA,TW,EPSL_WALL,EPSL_UPPER,WALL_DIS_BEGIN,WALL_DIS_END;
+extern REAL Sin_AOA , Cos_AOA;
+
+extern REAL *BC_rpara, (*ANA_rpara)[100]; 
+extern int *BC_npara, (*ANA_npara)[100]; 
+
+
+// ------Filter----------------------------------------------- 
+extern int FILTER_1,FILTER_2;   // 1: 5-point filter, 2: 7-point wide-band filter
+extern int NF_max;   // maximum of filtering 
+extern int Filter_Fo9p, Filter_Fopt_shock;
+extern int NFiltering,(*Filter_para)[11];      //Filtering
+extern REAL (*Filter_rpara)[3];   // s0  (0<s0<1), rth
+extern char IF_Filter_X , IF_Filter_Y , IF_Filter_Z ;
+extern int fiter_judge_X, fiter_judge_Y, fiter_judge_Z;
+// --------------------------------------------------
+
+
+
+// ----------------------------------------------------------
+// Coordinate parameters 
+extern REAL hx,hy,hz;  //参考空间网格尺寸
+extern REAL *pAxx,*pAyy,*pAzz,*pAkx,*pAky,*pAkz,*pAix,*pAiy,*pAiz,*pAsx,*pAsy,*pAsz,*pAjac;  // 度量系数矩阵
+
+
+// calculate memory
+extern REAL *pAmu; // viscous 3d [nz][nt][nx]
+extern REAL *pd,*pu,*pv,*pw,*pT,*pP; //  [nz+2*LAP][ny+2*LAP][nx+2*LAP]
+extern REAL *pf,*pfn,*pdu; // [nz][ny][nx][5]
+
+// used in filtering
+extern REAL *pf_lap; // [nz+2*LAP][ny+2*LAP][nx+2*LAP][5]
+
+// used in invis jacobian , is part of ptmpa
+extern REAL *pfp; // [5][nz+2*LAP][ny+2*LAP][nx+2*LAP]
+extern REAL *pfm; // [5][nz+2*LAP][ny+2*LAP][nx+2*LAP]
+extern REAL *pcc; // [nz+2*LAP][ny+2*LAP][nx+2*LAP]
+// used in invis jacobian , is part of ptmpb
+extern REAL *pdfp , *pdfm; // [nz][ny][nx]
+
+// used in ana
+extern REAL *pQ , *pLamda2;// [nx][ny][nx]
+extern REAL *pdm, *pum, *pvm, *pwm, *pTm;//[nx][ny][nz]
+extern int average_IO;
+extern int Istep_average;
+extern REAL tt_average;
+
+// used in vis jacobian , is part of ptmpb
+extern REAL * pEv1,*pEv2,*pEv3,*pEv4;  // [nz+2*LAP][ny+2*LAP][nx+2*LAP]
+// used in vis jacobian , is part of ptmpb
+extern REAL *puk,*pui,*pus,*pvk,*pvi,*pvs,*pwk,*pwi,*pws,*pTk,*pTi,*pTs;  //[nz][ny][nx]
+// used in mecpy
+extern REAL *pack_send_x,* pack_recv_x;
+extern REAL *pack_send_y,* pack_recv_y;
+extern REAL *pack_send_z,* pack_recv_z;
+// used in boundary_compressible_conner***********************************************************
+extern int MZMAX, MTMAX, INLET_BOUNDARY, IFLAG_WALL_NOT_NORMAL;
+extern REAL EPSL, X_DIST_BEGIN, X_DIST_END, BETA;
+extern REAL X_WALL_BEGIN, X_UP_BOUNDARY_BEGIN;
+
+// used in SCHEME_HYBRIDAUTO ********************************************************************
+extern int IFLAG_HybridAuto;
+extern int HybridA_Stage, Patch_max;
+extern int IFLAG_mem;
+extern HybridAuto_TYPE HybridAuto;
+extern int *scheme_x, *scheme_y, *scheme_z;
+
+extern int IF_CHARTERIC;
+
+void read_parameters();
+
+#ifdef __cplusplus
+}
+#endif
+#endif
+
+
+

head_hip/parameters_d.h

+#ifndef __PARAMETERS_D_H
+#define __PARAMETERS_D_H
+
+#include "config_parameters.h"
+#include "cuda_commen.h"
+
+
+// -----constant-----
+#ifdef __cplusplus
+extern "C"{
+#endif
+
+extern hipStream_t Stream[15];
+extern hipEvent_t  Event[15];
+
+extern __device__ __constant__ REAL Ama_d , Gamma_d , epsl_sw_d;
+extern __device__ __constant__ REAL Cv_d , Cp_d , Tsb_d , amu_C0_d;
+extern __device__ __constant__ REAL split_C1_d , split_C3_d;
+extern __device__ __constant__ REAL vis_flux_init_c_d;
+
+extern __device__ __constant__ unsigned int nx_d,ny_d,nz_d; // 某方向所处理的个数
+extern __device__ __constant__ unsigned int nx_lap_d,ny_lap_d,nz_lap_d; // 某方向所处理的个数
+extern __device__ __constant__ unsigned int nx_2lap_d,ny_2lap_d,nz_2lap_d; // 某方向所处理的个数
+
+extern __device__ __constant__ REAL dt_d;
+extern __device__ __constant__ REAL Sin_AOA_d , Cos_AOA_d;
+extern __device__ __constant__ REAL TW_d;
+
+//---------------------WENO_SYMBO_Limiter------------------------------------------
+extern __device__ __constant__ REAL WENO_TV_Limiter_d;
+extern __device__ __constant__ REAL WENO_TV_MAX_d;
+
+// ----------------------------------------------------------
+// Coordinate parameters 
+extern __device__ __constant__ REAL hx_d,hy_d,hz_d;  //参考空间网格尺寸
+
+extern  cudaField *pAxx_d,*pAyy_d,*pAzz_d,*pAkx_d,*pAky_d,*pAkz_d,*pAix_d,*pAiy_d,*pAiz_d,*pAsx_d,*pAsy_d,*pAsz_d,*pAjac_d;  // 度量系数矩阵
+
+
+// calculate memory
+extern cudaField *pAmu_d; // viscous 3d [nz][nt][nx]
+extern cudaField *pd_d,*pu_d,*pv_d,*pw_d,*pT_d,*pP_d; //  [nz+2*LAP][ny+2*LAP][nx+2*LAP]
+extern cudaSoA *pf_d,*pfn_d,*pdu_d; // [nz][ny][nx][5]
+
+// used in filtering
+extern cudaSoA *pf_lap_d; // [nz+2*LAP][ny+2*LAP][nx+2*LAP][5]
+
+// used in analysis
+extern cudaField *pdm_d, *pum_d, *pvm_d, *pwm_d, *pTm_d;//[nz][ny][nx]
+
+// used in invis jacobian , is part of ptmpa
+extern cudaSoA *pfp_x_d; // [5][nz+2*LAP][ny+2*LAP][nx+2*LAP]
+extern cudaSoA *pfm_x_d; // [5][nz+2*LAP][ny+2*LAP][nx+2*LAP]
+
+extern cudaSoA *pfp_y_d; // [5][nz+2*LAP][ny+2*LAP][nx+2*LAP]
+extern cudaSoA *pfm_y_d; // [5][nz+2*LAP][ny+2*LAP][nx+2*LAP]
+
+extern cudaSoA *pfp_z_d; // [5][nz+2*LAP][ny+2*LAP][nx+2*LAP]
+extern cudaSoA *pfm_z_d; // [5][nz+2*LAP][ny+2*LAP][nx+2*LAP]
+
+extern cudaField *pcc_d; // [nz+2*LAP][ny+2*LAP][nx+2*LAP]
+// used in invis jacobian , is part of ptmpb
+extern cudaField *pdfp_d , *pdfm_d; // [nz][ny][nx]
+
+
+
+// used in vis jacobian , is part of ptmpb
+extern cudaField * pEv1_d,*pEv2_d,*pEv3_d,*pEv4_d;  // [nz+2*LAP][ny+2*LAP][nx+2*LAP]
+// used in vis jacobian , is part of ptmpb
+extern cudaField *puk_d,*pui_d,*pus_d,*pvk_d,*pvi_d,*pvs_d,*pwk_d,*pwi_d,*pws_d,*pTk_d,*pTi_d,*pTs_d;  //[nz][ny][nx]
+extern cudaField *vis_u_d,*vis_v_d,*vis_w_d,*vis_T_d;  //[nz][ny][nx]
+
+extern cudaField grad_P;
+extern cudaField *pPP_d;
+
+
+#ifdef __cplusplus
+}
+#endif
+#endif
+
+
+