Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
208 changes: 208 additions & 0 deletions src/libraries/AMPTOOLS_AMPS/GPUTwoPiAngles_Delta_DoubleSDMEs_kernel.cu
Original file line number Diff line number Diff line change
@@ -0,0 +1,208 @@

#include <stdio.h>

#include "GPUManager/GPUCustomTypes.h"
#include "GPUManager/CUDA-Complex.cuh"


__global__ void
//
GPUTwoPiAngles_Delta_DoubleSDMEs_kernel( GPU_AMP_PROTO,
//order is like the appearing order in the paper
//Wmm
GDouble r00_33_0,
GDouble r11_33_0, GDouble r11_11_0,
GDouble r1m1_33_0, GDouble r1m1_11_0,
GDouble r10_33_0, GDouble r10_11_0,

//Wbar
GDouble r00_31_0, GDouble r00_3m1_0,
GDouble r11_31_0, GDouble r11_3m1_0,
GDouble r1m1_31_0, GDouble r1m1_3m1_0,
GDouble r10_31_0, GDouble r10_3m1_0,

//Wtil
GDouble rt1m1_31_0, GDouble rt1m1_3m1_0,
GDouble rt10_31_0, GDouble rt10_3m1_0,
// //alpha = 1 (same structure)
GDouble r00_33_1, GDouble r00_11_1,
GDouble r11_33_1, GDouble r11_11_1,
GDouble r1m1_33_1, GDouble r1m1_11_1,
GDouble r10_33_1, GDouble r10_11_1,
GDouble r00_31_1, GDouble r00_3m1_1,
GDouble r11_31_1, GDouble r11_3m1_1,
GDouble r1m1_31_1, GDouble r1m1_3m1_1,
GDouble r10_31_1, GDouble r10_3m1_1,
GDouble rt1m1_31_1, GDouble rt1m1_3m1_1,
GDouble rt10_31_1, GDouble rt10_3m1_1,
//alpha = 2 (different structure)
GDouble rt00_31_2, GDouble rt00_3m1_2,
GDouble rt11_31_2, GDouble rt11_3m1_2,
GDouble r1m1_33_2, GDouble r1m1_11_2,
GDouble r1m1_31_2, GDouble r1m1_3m1_2,
GDouble rt1m1_31_2, GDouble rt1m1_3m1_2,
GDouble r10_33_2, GDouble r10_11_2,
GDouble r10_31_2, GDouble r10_3m1_2,
GDouble rt10_31_2, GDouble rt10_3m1_2
)
{

int iEvent = GPU_THIS_EVENT;

// here we need to be careful to index the user-defined
// data with the proper integer corresponding to the
// enumeration in the C++ header file


GDouble Pgamma = GPU_UVARS(0);
GDouble cosTheta_pim = GPU_UVARS(1);
GDouble sinSqTheta_pim = GPU_UVARS(2);
GDouble sin2Theta_pim = GPU_UVARS(3);
GDouble phi_pim = GPU_UVARS(4);
GDouble cosTheta_proton = GPU_UVARS(5);
GDouble sinSqTheta_proton = GPU_UVARS(6);
GDouble sin2Theta_proton = GPU_UVARS(7);
GDouble phi_proton = GPU_UVARS(8);
GDouble bigPhi = GPU_UVARS(9);
// flag for pol (0.0 = no pol, 1.0 = pol)

GDouble sinSqTh = sinSqTheta_proton;
GDouble sin2Th = sin2Theta_proton;
GDouble cosSqTh = cosTheta_proton * cosTheta_proton;

GDouble cphi = cos(phi_proton);
GDouble sphi = sin(phi_proton);
GDouble c2phi = cos(2.0 * phi_proton);
GDouble s2phi = sin(2.0 * phi_proton);

GDouble sinSqThPi = sinSqTheta_pim;
GDouble sin2ThPi = sin2Theta_pim;
GDouble cosSqThPi = cosTheta_pim * cosTheta_pim;

GDouble cphiPi = cos(phi_pim);
GDouble sphiPi = sin(phi_pim);
GDouble c2phiPi = cos(2.0 * phi_pim);
GDouble s2phiPi = sin(2.0 * phi_pim);

GDouble r00_11_0 = 0.5 - r00_33_0 - 2.0 * r11_11_0 - 2.0 * r11_33_0;

GDouble W00_0 = r00_33_0 * sinSqTh + r00_11_0 * (1.0/3.0 + cosSqTh);
GDouble Wb00_0 = (2.0 / sqrt(3.0)) * (r00_31_0 * cphi * sin2Th + r00_3m1_0 * c2phi * sinSqTh);

GDouble W11_0 = r11_33_0 * sinSqTh + r11_11_0 * (1.0/3.0 + cosSqTh);
GDouble Wb11_0 = (2.0 / sqrt(3.0)) * (r11_31_0 * cphi * sin2Th + r11_3m1_0 * c2phi * sinSqTh);

GDouble W1m1_0 = r1m1_33_0 * sinSqTh + r1m1_11_0 * (1.0/3.0 + cosSqTh);
GDouble Wb1m1_0 = (2.0 / sqrt(3.0)) * (r1m1_31_0 * cphi * sin2Th + r1m1_3m1_0 * c2phi * sinSqTh);
GDouble Wt1m1_0 = (2.0 / sqrt(3.0)) * (rt1m1_31_0 * sphi * sin2Th + rt1m1_3m1_0 * s2phi * sinSqTh);

GDouble W10_0 = r10_33_0 * sinSqTh + r10_11_0 * (1.0/3.0 + cosSqTh);
GDouble Wb10_0 = (2.0 / sqrt(3.0)) * (r10_31_0 * cphi * sin2Th + r10_3m1_0 * c2phi * sinSqTh);
GDouble Wt10_0 = (2.0 / sqrt(3.0)) * (rt10_31_0 * sphi * sin2Th + rt10_3m1_0 * s2phi * sinSqTh);
GDouble W0 = cosSqThPi * (W00_0 - Wb00_0) + sinSqThPi * (W11_0 - Wb11_0) - sinSqThPi * (c2phiPi * (W1m1_0 - Wb1m1_0) + s2phiPi * Wt1m1_0) - sqrt(2.0) * sin2ThPi * (cphiPi * (W10_0 - Wb10_0) + sphiPi * Wt10_0);

//alpha = 1 (same structure)
GDouble W00_1 = r00_33_1 * sinSqTh + r00_11_1 * (1.0/3.0 + cosSqTh);
GDouble Wb00_1 = (2.0 / sqrt(3.0)) * (r00_31_1 * cphi * sin2Th + r00_3m1_1 * c2phi * sinSqTh);

GDouble W11_1 = r11_33_1 * sinSqTh + r11_11_1 * (1.0/3.0 + cosSqTh);
GDouble Wb11_1 = (2.0 / sqrt(3.0)) * (r11_31_1 * cphi * sin2Th + r11_3m1_1 * c2phi * sinSqTh);

GDouble W1m1_1 = r1m1_33_1 * sinSqTh + r1m1_11_1 * (1.0/3.0 + cosSqTh);
GDouble Wb1m1_1 = (2.0 / sqrt(3.0)) * (r1m1_31_1 * cphi * sin2Th + r1m1_3m1_1 * c2phi * sinSqTh);
GDouble Wt1m1_1 = (2.0 / sqrt(3.0)) * (rt1m1_31_1 * sphi * sin2Th + rt1m1_3m1_1 * s2phi * sinSqTh);

GDouble W10_1 = r10_33_1 * sinSqTh + r10_11_1 * (1.0/3.0 + cosSqTh);
GDouble Wb10_1 = (2.0 / sqrt(3.0)) * (r10_31_1 * cphi * sin2Th + r10_3m1_1 * c2phi * sinSqTh);
GDouble Wt10_1 = (2.0 / sqrt(3.0)) * (rt10_31_1 * sphi * sin2Th + rt10_3m1_1 * s2phi * sinSqTh);

GDouble W1 =cosSqThPi * (W00_1 - Wb00_1) + sinSqThPi * (W11_1 - Wb11_1) - sinSqThPi * (c2phiPi * (W1m1_1 - Wb1m1_1) + s2phiPi * Wt1m1_1) - sqrt(2.0) * sin2ThPi * (cphiPi * (W10_1 - Wb10_1) + sphiPi * Wt10_1);

// alpha = 2 (different structure)
GDouble Wt00_2 = (2.0 / sqrt(3.0)) * (rt00_31_2 * sphi * sin2Th + rt00_3m1_2 * s2phi * sinSqTh);
GDouble Wt11_2 = (2.0 / sqrt(3.0)) * (rt11_31_2 * sphi * sin2Th + rt11_3m1_2 * s2phi * sinSqTh);

GDouble W1m1_2 = r1m1_33_2 * sinSqTh + r1m1_11_2 * (1.0/3.0 + cosSqTh);
GDouble Wb1m1_2 = (2.0 / sqrt(3.0)) * (r1m1_31_2 * cphi * sin2Th + r1m1_3m1_2 * c2phi * sinSqTh);
GDouble Wt1m1_2 = (2.0 / sqrt(3.0)) * (rt1m1_31_2 * sphi * sin2Th + rt1m1_3m1_2 * s2phi * sinSqTh);

GDouble W10_2 = r10_33_2 * sinSqTh + r10_11_2 * (1.0/3.0 + cosSqTh);
GDouble Wb10_2 = (2.0 / sqrt(3.0)) * (r10_31_2 * cphi * sin2Th + r10_3m1_2 * c2phi * sinSqTh);
GDouble Wt10_2 = (2.0 / sqrt(3.0)) * (rt10_31_2 * sphi * sin2Th + rt10_3m1_2 * s2phi * sinSqTh);
GDouble W2 = cosSqThPi * Wt00_2+ sinSqThPi * Wt11_2+ sinSqThPi * (s2phiPi * (W1m1_2 - Wb1m1_2) - c2phiPi * Wt1m1_2) + sqrt(2.0) * sin2ThPi * (sphiPi * (W10_2 - Wb10_2) + cphiPi * Wt10_2);

GDouble Wpol = W0 - Pgamma * cos(2.0 * bigPhi) * W1 - Pgamma * sin(2.0 * bigPhi) * W2;

WCUComplex amp = { sqrt(fabs(Wpol)), 0.0 };
pcDevAmp[iEvent] = amp;

}


void
GPUTwoPiAngles_Delta_DoubleSDMEs_exec( dim3 dimGrid, dim3 dimBlock, GPU_AMP_PROTO,
GDouble r00_33_0,
GDouble r11_33_0, GDouble r11_11_0,
GDouble r1m1_33_0, GDouble r1m1_11_0,
GDouble r10_33_0, GDouble r10_11_0,
GDouble r00_31_0, GDouble r00_3m1_0,
GDouble r11_31_0, GDouble r11_3m1_0,
GDouble r1m1_31_0, GDouble r1m1_3m1_0,
GDouble r10_31_0, GDouble r10_3m1_0,
GDouble rt1m1_31_0, GDouble rt1m1_3m1_0,
GDouble rt10_31_0, GDouble rt10_3m1_0,
GDouble r00_33_1, GDouble r00_11_1,
GDouble r11_33_1, GDouble r11_11_1,
GDouble r1m1_33_1, GDouble r1m1_11_1,
GDouble r10_33_1, GDouble r10_11_1,
GDouble r00_31_1, GDouble r00_3m1_1,
GDouble r11_31_1, GDouble r11_3m1_1,
GDouble r1m1_31_1, GDouble r1m1_3m1_1,
GDouble r10_31_1, GDouble r10_3m1_1,
GDouble rt1m1_31_1, GDouble rt1m1_3m1_1,
GDouble rt10_31_1, GDouble rt10_3m1_1,
GDouble rt00_31_2, GDouble rt00_3m1_2,
GDouble rt11_31_2, GDouble rt11_3m1_2,
GDouble r1m1_33_2, GDouble r1m1_11_2,
GDouble r1m1_31_2, GDouble r1m1_3m1_2,
GDouble rt1m1_31_2, GDouble rt1m1_3m1_2,
GDouble r10_33_2, GDouble r10_11_2,
GDouble r10_31_2, GDouble r10_3m1_2,
GDouble rt10_31_2, GDouble rt10_3m1_2
)
{

GPUTwoPiAngles_Delta_DoubleSDMEs_kernel<<< dimGrid, dimBlock >>>
( GPU_AMP_ARGS,
r00_33_0,
r11_33_0, r11_11_0,
r1m1_33_0, r1m1_11_0,
r10_33_0, r10_11_0,
r00_31_0, r00_3m1_0,
r11_31_0, r11_3m1_0,
r1m1_31_0, r1m1_3m1_0,
r10_31_0, r10_3m1_0,
rt1m1_31_0, rt1m1_3m1_0,
rt10_31_0, rt10_3m1_0,
r00_33_1, r00_11_1,
r11_33_1, r11_11_1,
r1m1_33_1, r1m1_11_1,
r10_33_1, r10_11_1,
r00_31_1, r00_3m1_1,
r11_31_1, r11_3m1_1,
r1m1_31_1, r1m1_3m1_1,
r10_31_1, r10_3m1_1,
rt1m1_31_1, rt1m1_3m1_1,
rt10_31_1, rt10_3m1_1,
rt00_31_2, rt00_3m1_2,
rt11_31_2, rt11_3m1_2,
r1m1_33_2, r1m1_11_2,
r1m1_31_2, r1m1_3m1_2,
rt1m1_31_2, rt1m1_3m1_2,
r10_33_2, r10_11_2,
r10_31_2, r10_3m1_2,
rt10_31_2, rt10_3m1_2


);
}
Original file line number Diff line number Diff line change
@@ -0,0 +1,96 @@

#include <stdio.h>

#include "GPUManager/GPUCustomTypes.h"
#include "GPUManager/CUDA-Complex.cuh"


__global__ void
//
GPUTwoPiAngles_Delta_factorized_kernel( GPU_AMP_PROTO,
GDouble rho000,
GDouble rho100,
GDouble rho1m10,
GDouble rho001,
GDouble rho111,
GDouble rho101,
GDouble rho1m11,
GDouble rho102,
GDouble rho1m12,
GDouble delta_rho011,
GDouble delta_rho031,
GDouble delta_rho03m1
)
{

int iEvent = GPU_THIS_EVENT;

// here we need to be careful to index the user-defined
// data with the proper integer corresponding to the
// enumeration in the C++ header file


GDouble Pgamma = GPU_UVARS(0);
GDouble cosTheta_pim = GPU_UVARS(1);
GDouble sinSqTheta_pim = GPU_UVARS(2);
GDouble sin2Theta_pim = GPU_UVARS(3);
GDouble phi_pim = GPU_UVARS(4);
GDouble cosTheta_proton = GPU_UVARS(5);
GDouble sinSqTheta_proton = GPU_UVARS(6);
GDouble sin2Theta_proton = GPU_UVARS(7);
GDouble phi_proton = GPU_UVARS(8);
GDouble bigPhi = GPU_UVARS(9);
// flag for pol (0.0 = no pol, 1.0 = pol)

GDouble W0_meson = 3.*(0.5*(1. - rho000) + 0.5*(3.*rho000 - 1.)*cosTheta_pim*cosTheta_pim - sqrt(2.)*rho100*sin2Theta_pim*cos(phi_pim) - rho1m10*sinSqTheta_pim*cos(2.*phi_pim));
GDouble W1_meson= 3.*(rho111*sinSqTheta_pim + rho001*cosTheta_pim*cosTheta_pim - sqrt(2.)*rho101*sin2Theta_pim*cos(phi_pim) - rho1m11*sinSqTheta_pim*cos(2.*phi_pim));
GDouble W2_meson = 3. * (sqrt(2.)*rho102*sin2Theta_pim*sin(phi_pim) + rho1m12*sinSqTheta_pim*sin(2.*phi_pim));

GDouble W0_baryon = 3.*(0.5 - delta_rho011)*sinSqTheta_proton + delta_rho011*(1.+3.*cosTheta_proton*cosTheta_proton) - 2.*sqrt(3.)*delta_rho031*cos(phi_proton)*sin2Theta_proton - 2.*sqrt(3.)*delta_rho03m1*cos(2.*phi_proton)*sinSqTheta_proton;
// GDouble W1_baryon = (3.*delta_rho133*sinSqTheta_piplus + delta_rho111*(1.+3.*cosTheta_piplus*cosTheta_piplus) - 2.*sqrt(3.)*delta_rho131*cos(phi_piplus)*sin2Theta_piplus - 2.*sqrt(3.)*delta_rho13m1*cos(2.*phi_piplus)*sinSqTheta_piplus);
// GDouble W2_baryon = (2.*sqrt(3.)*delta_rho231*sin(phi_piplus)*sin2Theta_piplus + 2.*sqrt(3.)*delta_rho23m1*sin(2.*phi_piplus)*sinSqTheta_piplus);

GDouble Wpol = W0_meson*W0_baryon - Pgamma * cos(2.0 * bigPhi) *( W1_meson*W0_baryon) - Pgamma * sin(2.0 * bigPhi) *( W2_meson*W0_baryon);

Wpol *= 1/(4.*PI);

WCUComplex amp = { sqrt(fabs(Wpol)), 0.0 };
pcDevAmp[iEvent] = amp;

}


void
GPUTwoPiAngles_Delta_factorized_exec( dim3 dimGrid, dim3 dimBlock, GPU_AMP_PROTO,
GDouble rho000,
GDouble rho100,
GDouble rho1m10,
GDouble rho001,
GDouble rho111,
GDouble rho101,
GDouble rho1m11,
GDouble rho102,
GDouble rho1m12,
GDouble delta_rho011,
GDouble delta_rho031,
GDouble delta_rho03m1

)
{

GPUTwoPiAngles_Delta_factorized_kernel<<< dimGrid, dimBlock >>>
( GPU_AMP_ARGS,
rho000,
rho100,
rho1m10,
rho001,
rho111,
rho101,
rho1m11,
rho102,
rho1m12,
delta_rho011,
delta_rho031,
delta_rho03m1
);
}
Loading