You can not select more than 25 topics
Topics must start with a letter or number, can include dashes ('-') and can be up to 35 characters long.
259 lines
10 KiB
259 lines
10 KiB
//
|
|
// Created by jundafeng on 3/23/22.
|
|
//
|
|
|
|
#ifndef DGTD_KERNELS_CUH
|
|
#define DGTD_KERNELS_CUH
|
|
|
|
#include "Constants.h"
|
|
#include "../dgtd-performance.hpp"
|
|
|
|
typedef struct{
|
|
// Plane Wave Excitation
|
|
int ExcitationFlag; // Scattering Planewave Waveform: (0)Time Harmonic (1)Gauss (2)Neumann (3) Modulated Gauss
|
|
fp_t_ts to; //Tdelay: Delay of excitation pulse (sec)
|
|
fp_t_ts tau; //Tau: Parameter for pulse width (sec)
|
|
fp_t_ts freq_m; //Freq: Central Modulation Frequency (MHz)
|
|
fp_t_ts Emagnitude; //magnitude of the E field
|
|
fp_t_ts theta; //angle of excitation
|
|
fp_t_ts phi; //angle of excitation
|
|
fp_t_ts Epol[3]; //Efield direction
|
|
fp_t_ts ro[3]; //position where the excitation starts
|
|
|
|
// ===============================================================================
|
|
// General Port Excitation
|
|
int TimeDistributionFlag; // Port Waveform (0)TH (1)Gauss
|
|
fp_t_ts PortDirection[3];
|
|
int portNum;
|
|
int BCNum;
|
|
|
|
// General waveport selector
|
|
// 0 = no ports
|
|
// 1 = TEM rectangular, 2 = TEM coax, 3 = TE_mn rectangular
|
|
int PortFlag;
|
|
|
|
// Medium params (if you want analytic impedances); if unused, leave defaults
|
|
fp_t_ts epr; // relative permittivity
|
|
fp_t_ts mur; // relative permeability (usually 1.0)
|
|
|
|
// (A) TEM rectangular
|
|
// E is tangential along vpath projected to the port plane.
|
|
fp_t_ts vpath[3]; // desired E-field tangential direction seed (not necessarily unit)
|
|
|
|
// (B) TEM coax (annulus) — define center and radii in the port plane
|
|
fp_t_ts r0_port[3]; // center point of coax cross-section (on the port plane)
|
|
fp_t_ts r1_port[3]; // a point on the inner conductor circle (defines inner radius a)
|
|
fp_t_ts r2_port[3]; // a point on the outer conductor circle (defines outer radius b)
|
|
|
|
// (C) TE_mn rectangular
|
|
int m, n; // TE indices
|
|
fp_t_ts rect_a; // size along local tangential axis t1
|
|
fp_t_ts rect_b; // size along local tangential axis t2
|
|
fp_t_ts uv0[3]; // reference corner where (u,v)=(0,0) in the port plane
|
|
fp_t_ts t1[3],t2[3]; // Direction vector of the port
|
|
|
|
// Port impedance you want the BC to use (kernel already sets Y=1/Zport)
|
|
// For TEM you can set Zport = eta = sqrt(mu/eps), for TE set Z_TE(ω).
|
|
fp_t_ts PortImpedance;
|
|
|
|
// Chirp bandwidth for TimeModulationInc
|
|
fp_t_ts CHIRP_BW_MHZ;
|
|
// ===============================================================================
|
|
|
|
} ExcitationProp;
|
|
|
|
// Ceiling funciton for X / Y.
|
|
__host__ __device__ static inline int ceil_div(int x, int y){return (x - 1) / y + 1;}
|
|
__host__ static inline size_t ceil_div(size_t x, size_t y){return (x - 1) / y + 1;}
|
|
__host__ static inline size_t floor_div(size_t x, size_t y){return (x - 1) / y;}
|
|
|
|
constexpr int BLOCK_SIZE_Matrix_EVAL = 128;
|
|
|
|
#if defined(DGTD_USE_CUDA)
|
|
#include <cuda.h>
|
|
#include <cuda_runtime.h>
|
|
#include <cooperative_groups.h>
|
|
#include "cuda_utils.h"
|
|
#include <cuComplex.h>
|
|
|
|
//#include <cooperative_groups/memcpy_async.h>
|
|
//#include <cuda/pipeline>
|
|
namespace cg = cooperative_groups;
|
|
#define WARP_SIZE 32
|
|
__global__ void makeNeighField(int* MappingArray, fp_t_ts* Field, fp_t_ts* auxiliar, int elements);
|
|
__global__ void checker(fp_t_ts* a, fp_t_ts* b, int length);
|
|
__global__ void addCouplingResults(fp_t_ts* Field, fp_t_ts* auxiliar, int* neighs, int* neighsOffset, int elements);
|
|
|
|
__global__ void makeNeighFieldRegular(int* MappingArray, fp_t_ts* Field, fp_t_ts* auxiliar, int elements, int polyOrder);
|
|
|
|
__global__ void makeNeighFieldRegular_custom(int* MappingArray, fp_t_ts* Field, fp_t_ts* auxiliar, int elements, int polyOrder, int numFaces);
|
|
|
|
__global__ void addCouplingResultsRegular(fp_t_ts* Field, fp_t_ts* auxiliar, int elements);
|
|
|
|
__global__ void addCouplingResultsRegular_custom(fp_t_ts* Field, fp_t_ts* auxiliar, int elements, int numFaces);
|
|
|
|
// Testing
|
|
__global__ void fill_ones_cols(fp_t_ts* __restrict__ M, int rows, int cols, int ld);
|
|
|
|
|
|
__global__ void addExcitationH_PML( int* ExcitationFacesCnt,
|
|
int* ExcitationFacesOffset,
|
|
int* ExcitationFacesNum,
|
|
int excitationCNT,
|
|
int excitationPMLCNT,
|
|
int8_t* map,
|
|
ExcitationProp exci_prop,
|
|
int PolyFlag,
|
|
fp_t_ts dt_muo, fp_t_ts t,
|
|
fp_t_ts* nd_coords_tet,
|
|
fp_t_ts* nd_coords_face,
|
|
fp_t_ts* Z_face,
|
|
fp_t_ts* InvMat,
|
|
fp_t_ts* Field);
|
|
|
|
__global__ void addExcitationE_PML( int* ExcitationFacesCnt,
|
|
int* ExcitationFacesOffset,
|
|
int* ExcitationFacesNum,
|
|
int excitationCNT,
|
|
int excitationPMLCNT,
|
|
int8_t* map,
|
|
ExcitationProp exci_prop,
|
|
int PolyFlag,
|
|
fp_t_ts dt_epsilon, fp_t_ts t,
|
|
fp_t_ts* nd_coords_tet,
|
|
fp_t_ts* nd_coords_face,
|
|
fp_t_ts* Z_face,
|
|
fp_t_ts* InvMat,
|
|
fp_t_ts* Field);
|
|
|
|
__global__ void addExcitationH( int* ExcitationFacesCnt,
|
|
int* ExcitationFacesOffset,
|
|
int* ExcitationFacesNum,
|
|
int excitationCNT,
|
|
int8_t* map,
|
|
ExcitationProp exci_prop,
|
|
int PolyFlag,
|
|
fp_t_ts dt_muo, fp_t_ts t,
|
|
fp_t_ts* nd_coords_tet,
|
|
fp_t_ts* nd_coords_face,
|
|
fp_t_ts* Z_face,
|
|
fp_t_ts* InvMat,
|
|
fp_t_ts* Field);
|
|
|
|
__global__ void addExcitationE( int* ExcitationFacesCnt,
|
|
int* ExcitationFacesOffset,
|
|
int* ExcitationFacesNum,
|
|
int excitationCNT,
|
|
int8_t* map,
|
|
ExcitationProp exci_prop,
|
|
int PolyFlag,
|
|
fp_t_ts dt_epsilon, fp_t_ts t,
|
|
fp_t_ts* nd_coords_tet,
|
|
fp_t_ts* nd_coords_face,
|
|
fp_t_ts* Z_face,
|
|
fp_t_ts* InvMat,
|
|
fp_t_ts* Field);
|
|
|
|
|
|
__global__ void CalculatePadeFreq(fp_t* a_k, fp_t* b_k, int M_global, int N_2, int* constantM, cuDoubleComplex* Hf);
|
|
|
|
|
|
|
|
__global__ void addExcitationE_port(
|
|
int* ExcitationFacesCnt,
|
|
int* ExcitationFacesOffset,
|
|
int* ExcitationFacesNum,
|
|
int excitationCNT,
|
|
int8_t* map,
|
|
ExcitationProp* exci_prop,
|
|
int* PortFacePidx_d,
|
|
int PolyFlag,
|
|
fp_t_ts dt_epsilon, fp_t_ts t,
|
|
fp_t_ts* nd_coords_tet,
|
|
fp_t_ts* nd_coords_face,
|
|
fp_t_ts* Z_face,
|
|
fp_t_ts* InvMat,
|
|
fp_t_ts* Field,
|
|
fp_t_ts* Etan_qp_d,
|
|
fp_t_ts* Htan_qp_d);
|
|
|
|
|
|
|
|
__global__ void addExcitationH_port( int* ExcitationFacesCnt,
|
|
int* ExcitationFacesOffset,
|
|
int* ExcitationFacesNum,
|
|
int excitationCNT,
|
|
int8_t* map,
|
|
ExcitationProp* exci_prop,
|
|
int* PortFacePidx_d,
|
|
int PolyFlag,
|
|
fp_t_ts dt_muo, fp_t_ts t,
|
|
fp_t_ts* nd_coords_tet,
|
|
fp_t_ts* nd_coords_face,
|
|
fp_t_ts* Z_face,
|
|
fp_t_ts* InvMat,
|
|
fp_t_ts* Field,
|
|
fp_t_ts* Etan_qp_d, fp_t_ts* Htan_qp_d);
|
|
|
|
|
|
|
|
__device__ void TimeModulationInc(int TimeDistributionFlag, fp_t_ts t, fp_t_ts to, fp_t_ts freq_m,
|
|
fp_t_ts Emag, fp_t_ts tau, fp_t_ts CHIRP_BW_MHZ, fp_t_ts &IncidExcit);
|
|
|
|
/*
|
|
__device__ void H_updateinc_PORT(
|
|
fp_t_ts t, fp_t_ts to, fp_t_ts tau, fp_t_ts freq_m,
|
|
int TimeDistributionFlag,
|
|
fp_t_ts* PortDirection, // unit propagation dir
|
|
fp_t_ts* Et_pre, // precomputed tangential E at quad point
|
|
fp_t_ts* Ht_pre, // precomputed tangential H at quad point
|
|
fp_t_ts CHIRP_BW_MHZ, // bandwidth for chirp
|
|
fp_t_ts* nxEinc, // out: n × E_inc
|
|
fp_t_ts* nxnxHinc); // out: n × (n × H_inc)
|
|
|
|
|
|
__device__ void E_updateinc_PORT(
|
|
fp_t_ts t, fp_t_ts to, fp_t_ts tau, fp_t_ts freq_m,
|
|
int TimeDistributionFlag,
|
|
fp_t_ts* PortDirection, // unit propagation dir
|
|
fp_t_ts* Et_pre, // precomputed tangential E at quad point
|
|
fp_t_ts* Ht_pre, // precomputed tangential H at quad point
|
|
fp_t_ts CHIRP_BW_MHZ, // bandwidth for chirp
|
|
fp_t_ts* nxHinc, // out: n × H_inc
|
|
fp_t_ts* nxnxEinc); // out: n × (n × E_inc)
|
|
*/
|
|
|
|
__device__ void updateinc_PORT(fp_t_ts t, const ExcitationProp& exci_prop,
|
|
const fp_t_ts r[3], const fp_t_ts *normalvtr, fp_t_ts* nxField,
|
|
fp_t_ts* nxnxField, int mode);
|
|
|
|
__global__ void addExcitationE_port( int* ExcitationFacesCnt,
|
|
int* ExcitationFacesOffset,
|
|
int* ExcitationFacesNum,
|
|
int excitationCNT,
|
|
int8_t* map,
|
|
ExcitationProp* exci_prop,
|
|
int* PortFacePidx_d,
|
|
int PolyFlag,
|
|
fp_t_ts dt_muo, fp_t_ts t,
|
|
fp_t_ts* nd_coords_tet,
|
|
fp_t_ts* nd_coords_face,
|
|
fp_t_ts* InvMat,
|
|
fp_t_ts* Field);
|
|
|
|
__global__ void addExcitationH_port( int* ExcitationFacesCnt,
|
|
int* ExcitationFacesOffset,
|
|
int* ExcitationFacesNum,
|
|
int excitationCNT,
|
|
int8_t* map,
|
|
ExcitationProp* exci_prop,
|
|
int* PortFacePidx_d,
|
|
int PolyFlag,
|
|
fp_t_ts dt_muo, fp_t_ts t,
|
|
fp_t_ts* nd_coords_tet,
|
|
fp_t_ts* nd_coords_face,
|
|
fp_t_ts* InvMat,
|
|
fp_t_ts* Field);
|
|
|
|
#endif
|
|
#endif //DGTD_KERNELS_CU
|
|
|