Skip to content

Commit

Permalink
call params.data in the host
Browse files Browse the repository at this point in the history
  • Loading branch information
marco garofalo committed Apr 13, 2021
1 parent 3e6a514 commit 12abc04
Showing 1 changed file with 20 additions and 16 deletions.
36 changes: 20 additions & 16 deletions modules/DFT.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -17,18 +17,22 @@
void compute_FT(const Viewphi phi, cluster::IO_params params , int iconf, Viewphi::HostMirror &h_phip){
int T=params.data.L[0];
size_t Vs=params.data.V/T;

double norm0=sqrt(2*params.data.kappa0);
double norm1=sqrt(2*params.data.kappa1);
int L1=params.data.L[1], L2=params.data.L[2], L3=params.data.L[3];

Viewphi phip("phip",2,params.data.L[0]*Vp);

typedef Kokkos::TeamPolicy<> team_policy;//team_policy ( number of teams , team size)
typedef Kokkos::TeamPolicy<>::member_type member_type;


Kokkos::parallel_for( "FT_loop", team_policy( T*Vp*2, Kokkos::AUTO), KOKKOS_LAMBDA ( const member_type &teamMember ) {
const int ii = teamMember.league_rank();
//ii = comp+ 2*myt
//myt= t +T*(reim+ p*2)
//p=px+py*4+pz*16
double norm[2]={sqrt(2.*params.data.kappa0),sqrt(2.*params.data.kappa1)};// need to be inside the loop for cuda<10
double norm[2]={norm0,norm1};// need to be inside the loop for cuda<10
const int p=ii/(4*T);
int res=ii-p*4*T;
const int reim=res/(2*T);
Expand All @@ -54,16 +58,16 @@ void compute_FT(const Viewphi phi, cluster::IO_params params , int iconf, Viewp
Kokkos::parallel_reduce( Kokkos::TeamThreadRange( teamMember, Vs ), [&] ( const size_t x, double &inner) {

size_t i0= x+t*Vs;
int ix=x%params.data.L[1];
int iz=x /(params.data.L[1]*params.data.L[2]);
int iy=(x- iz*params.data.L[1]*params.data.L[2])/params.data.L[1];
int ix=x%L1;
int iz=x /(L1*L2);
int iy=(x- iz*L1*L2)/L1;
#ifdef DEBUG
if (x!= ix+ iy*params.data.L[1]+iz*params.data.L[1]*params.data.L[2]){
printf("error %ld = %d + %d *%d+ %d*%d*%d\n",x,ix,iy,params.data.L[1],iz,params.data.L[1],params.data.L[2]);
if (x!= ix+ iy*L1+iz*L1*L2){
printf("error %ld = %d + %d *%d+ %d*%d*%d\n",x,ix,iy,L1,iz,L1,L2);
// exit(1);
}
#endif
double wr=6.28318530718 *( px*ix/(double (params.data.L[1])) + py*iy/(double (params.data.L[2])) +pz*iz/(double (params.data.L[3])) );
double wr=6.28318530718 *( px*ix/(double (L1)) + py*iy/(double (L2)) +pz*iz/(double (L3)) );
wr=cos(wr);

// phip(comp,xp)+=phi(comp,i0)*wr; }
Expand All @@ -74,16 +78,16 @@ void compute_FT(const Viewphi phi, cluster::IO_params params , int iconf, Viewp
else if(reim==1){
Kokkos::parallel_reduce( Kokkos::TeamThreadRange( teamMember, Vs ), [&] ( const size_t x, double &inner) {
size_t i0= x+t*Vs;
int ix=x%params.data.L[1];
int iz=x /(params.data.L[1]*params.data.L[2]);
int iy=(x- iz*params.data.L[1]*params.data.L[2])/params.data.L[1];
int ix=x%L1;
int iz=x /(L1*L2);
int iy=(x- iz*L1*L2)/L1;
#ifdef DEBUG
if (x!= ix+ iy*params.data.L[1]+iz*params.data.L[1]*params.data.L[2]){
printf("error %ld = %d + %d *%d+ %d*%d*%d\n",x,ix,iy,params.data.L[1],iz,params.data.L[1],params.data.L[2]);
// exit(1);
}
#endif
double wr=6.28318530718 *( px*ix/(double (params.data.L[1])) + py*iy/(double (params.data.L[2])) +pz*iz/(double (params.data.L[3])) );
double wr=6.28318530718 *( px*ix/(double (L1)) + py*iy/(double (L2)) +pz*iz/(double (L3)) );
wr=sin(wr);

inner+=phi(comp,i0)*wr;// /((double) Vs *norm[comp]);
Expand Down Expand Up @@ -259,8 +263,8 @@ void compute_cuFFT(const Viewphi phi, cluster::IO_params params , int iconf, Vi

cudaMalloc((void**)&idata, sizeof(cufftReal)*Vs);
cudaMalloc((void**)&odata, sizeof(cufftComplex)*Vs);
int n[3] = {params.data.L[1], params.data.L[2],params.data.L[3]};
cufftPlanMany(&plan, 3, n,
int L[3] = {params.data.L[1], params.data.L[2],params.data.L[3]};
cufftPlanMany(&plan, 3, L,
NULL, 1, Vs , // *inembed, istride, idist
NULL, 1, Vs, // *onembed, ostride, odist
CUFFT_R2C, 2*T);
Expand Down Expand Up @@ -291,7 +295,7 @@ void compute_cuFFT(const Viewphi phi, cluster::IO_params params , int iconf, Vi
const int px=p%Lp;
const int pz=p /(Lp*Lp);
const int py= (p- pz*Lp*Lp)/Lp;
int pcuff=(px+py*(params.data.L[1]/2 +1)+pz* (params.data.L[1]/2+1)*(params.data.L[2]));
int pcuff=(px+py*(L[1]/2 +1)+pz* (L[1]/2+1)*(L[2]));
int ip=t+pp*T;
if(reim==0)
Kphi(comp,ip)=odata[pcuff].x/(Vs*sqrt(2*params.data.kappa0));
Expand All @@ -317,7 +321,7 @@ void compute_cuFFT(const Viewphi phi, cluster::IO_params params , int iconf, Vi
const int px=p%Lp;
const int pz=p /(Lp*Lp);
const int py= (p- pz*Lp*Lp)/Lp;
int pcuff=(px+py*(params.data.L[1]/2 +1)+pz* (params.data.L[1]/2+1)*(params.data.L[2]));
int pcuff=(px+py*(L[1]/2 +1)+pz* (L[1]/2+1)*(L[2]));
int ip=t+pp*T;
if (fabs(Kphi(0,ip)-phip(0,ip))>1e-6 )
printf("p=%d= (%d,%d,%d) reim=%d pp=%ld ip=%d t=%d pcuff=%d cuFFT =%g DFT =%g\n", p,px,py,pz,reim, pp, ip,t,pcuff,Kphi(comp,ip) ,phip(comp,ip));
Expand Down

0 comments on commit 12abc04

Please sign in to comment.