Commit 24f394c6 authored by Uldis Locans's avatar Uldis Locans

add enableRutherfordScattering option to OPALs collimatorPhysics GPU version

parent 8f00d2a5
......@@ -19,7 +19,8 @@ public:
virtual ~DKSCollimatorPhysics() { }
virtual int CollimatorPhysics(void *mem_ptr, void *par_ptr, int numpartices) = 0;
virtual int CollimatorPhysics(void *mem_ptr, void *par_ptr, int numpartices,
bool enableRutherfordScattering = true) = 0;
virtual int CollimatorPhysicsSoA(void *label_ptr, void *localID_ptr,
void *rx_ptr, void *ry_ptr, void *rz_ptr,
......
......@@ -23,6 +23,7 @@
#define X0_M 9
#define I_M 10
#define DT_M 11
#define LOWENERGY_THR 12
#define BLOCK_SIZE 128
#define NUMPAR 12
......@@ -81,7 +82,7 @@ __device__ inline void energyLoss(double &Eng, bool &pdead, curandState &state,
Eng = Eng + delta_E / 1E3;
}
pdead = ((Eng<1E-4) || (dEdx>0));
pdead = ( (Eng < par[LOWENERGY_THR]) || (dEdx > 0) );
}
......@@ -117,7 +118,9 @@ __device__ inline void Rot(double &px, double &pz, double &x, double &z, double
pz = -pxz*sin(Psixz)*sin(thetacou) + pxz*cos(Psixz)*cos(thetacou);
}
__device__ inline void coulombScat(double3 &R, double3 &P, curandState &state, double* par) {
__device__ inline void coulombScat(double3 &R, double3 &P, curandState &state,
double* par, bool enableRutherfordScattering)
{
double Eng = sqrt(dot(P, P) + 1.0) * M_P - M_P;
double gamma = (Eng + M_P) / M_P;
......@@ -145,7 +148,7 @@ __device__ inline void coulombScat(double3 &R, double3 &P, curandState &state, d
Rot(P.x, P.z, R.x, R.z, xplane, normP, thetacou, deltas, 1, par);
double P2 = curand_uniform_double(&state);//gsl_rng_uniform(rGen_m);
if(P2 < 0.0047) {
if( (P2 < 0.0047) && enableRutherfordScattering) {
double P3 = curand_uniform_double(&state);//gsl_rng_uniform(rGen_m);
double thetaru = 2.5 * sqrt(1 / P3) * sqrt(2.0) * theta0;
double P4 = curand_uniform_double(&state);//gsl_rng_uniform(rGen_m);
......@@ -171,7 +174,7 @@ __device__ inline void coulombScat(double3 &R, double3 &P, curandState &state, d
Rot(P.y,P.z,R.y,R.z, yplane, normP, thetacou, deltas, 2, par);
P2 = curand_uniform_double(&state);//gsl_rng_uniform(rGen_m);
if(P2 < 0.0047) {
if( (P2 < 0.0047) && enableRutherfordScattering) {
double P3 = curand_uniform_double(&state);//gsl_rng_uniform(rGen_m);
double thetaru = 2.5 * sqrt(1 / P3) * sqrt(2.0) * theta0;
double P4 = curand_uniform_double(&state);//gsl_rng_uniform(rGen_m);
......@@ -185,7 +188,7 @@ __device__ inline void coulombScat(double3 &R, double3 &P, curandState &state, d
template <typename T>
__global__ void kernelCollimatorPhysics(T *data, double *par, curandState *state,
int numparticles)
int numparticles, bool enableRutherfordScattering)
{
//get global id and thread id
......@@ -227,7 +230,7 @@ __global__ void kernelCollimatorPhysics(T *data, double *par, curandState *state
P.x = P.x * ptot / sq;
P.y = P.y * ptot / sq;
P.z = P.z * ptot / sq;
coulombScat(R[tid], P, s, p);
coulombScat(R[tid], P, s, p, enableRutherfordScattering);
data[idx].Pincol = P;
} else {
......@@ -250,7 +253,8 @@ __global__ void kernelCollimatorPhysics(T *data, double *par, curandState *state
}
__global__ void kernelCollimatorPhysics2(CUDA_PART2_SMALL data, double *par,
curandState *state, int numparticles)
curandState *state, int numparticles,
bool enableRutherfordScattering)
{
//get global id and thread id
......@@ -288,7 +292,7 @@ __global__ void kernelCollimatorPhysics2(CUDA_PART2_SMALL data, double *par,
P.x = P.x * ptot / sq;
P.y = P.y * ptot / sq;
P.z = P.z * ptot / sq;
coulombScat(R[tid], P, s, p);
coulombScat(R[tid], P, s, p, enableRutherfordScattering);
data.Pincol[idx] = P;
} else {
......@@ -611,7 +615,8 @@ struct less_then
}
};
int CudaCollimatorPhysics::CollimatorPhysics(void *mem_ptr, void *par_ptr, int numparticles)
int CudaCollimatorPhysics::CollimatorPhysics(void *mem_ptr, void *par_ptr, int numparticles,
bool enableRutherfordScattering)
{
int threads = BLOCK_SIZE;
......@@ -624,7 +629,8 @@ int CudaCollimatorPhysics::CollimatorPhysics(void *mem_ptr, void *par_ptr, int n
kernelCollimatorPhysics<<<blocks, threads, smem_size>>>((CUDA_PART_SMALL*)mem_ptr,
(double*)par_ptr,
m_base->cuda_getCurandStates(),
numparticles);
numparticles,
enableRutherfordScattering);
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess)
......
......@@ -110,7 +110,7 @@ public:
*
*/
int CollimatorPhysics(void *mem_ptr, void *par_ptr,
int numpartices);
int numpartices, bool enableRutherfordScattering = true);
int CollimatorPhysicsSoA(void *label_ptr, void *localID_ptr,
void *rx_ptr, void *ry_ptr, void *rz_ptr,
......
......@@ -746,11 +746,13 @@ int DKSBase::callCollimatorPhysics(void *mem_ptr, void *par_ptr,
}
int DKSBase::callCollimatorPhysics2(void *mem_ptr, void *par_ptr, int numparticles)
int DKSBase::callCollimatorPhysics2(void *mem_ptr, void *par_ptr, int numparticles,
bool enableRutherfordScattering)
{
if (apiCuda())
return CUDA_SAFECALL( ccol->CollimatorPhysics(mem_ptr, par_ptr, numparticles) );
return CUDA_SAFECALL( ccol->CollimatorPhysics(mem_ptr, par_ptr, numparticles,
enableRutherfordScattering) );
else if (apiOpenMP())
return MIC_SAFECALL( miccol->CollimatorPhysics(mem_ptr, par_ptr, numparticles) );
......
......@@ -1041,7 +1041,8 @@ public:
* For specifics check OPAL docs and CudaCollimatorPhysics class documentation.
* TODO: opencl and mic implementations.
*/
int callCollimatorPhysics2(void *mem_ptr, void *par_ptr, int numparticles);
int callCollimatorPhysics2(void *mem_ptr, void *par_ptr, int numparticles,
bool enableRutherfordScattering = true);
/**
* Monte carlo code for the degrader from OPAL classic/5.0/src/Solvers/CollimatorPhysics.cpp on device.
......
......@@ -368,7 +368,9 @@ void energyLoss(double &Eng, double &dEdx, double *par, double *randv, int ri) {
}
int MICCollimatorPhysics::CollimatorPhysics(void *mem_ptr, void *par_ptr, int numparticles) {
int MICCollimatorPhysics::CollimatorPhysics(void *mem_ptr, void *par_ptr,
int numparticles, boll enableRutherfordScattering)
{
//cast device memory pointers to appropriate types
MIC_PART_SMALL *data = (MIC_PART_SMALL*) mem_ptr;
......
......@@ -40,7 +40,8 @@ public:
~MICCollimatorPhysics() { };
int CollimatorPhysics(void *mem_ptr, void *par_ptr, int numparticles);
int CollimatorPhysics(void *mem_ptr, void *par_ptr, int numparticles,
bool enableRutherfordScattering = true);
int CollimatorPhysicsSoA(void *label_ptr, void *localID_ptr,
void *rx_ptr, void *ry_ptr, void *rz_ptr,
......
......@@ -34,7 +34,7 @@ TODO:
2. boost.compute sort for user defined structure crashes
*/
int OpenCLCollimatorPhysics::CollimatorPhysics(void *mem_ptr, void *par_ptr,
int numparticles)
int numparticles, bool enableRutherfordScattering)
{
/*
//set number of total threads, and number threads per block
......
......@@ -52,7 +52,8 @@ public:
}
/* execute degrader code on device */
int CollimatorPhysics(void *mem_ptr, void *par_ptr, int numparticles);
int CollimatorPhysics(void *mem_ptr, void *par_ptr, int numparticles,
bool enableRutherfordScattering = true);
int CollimatorPhysicsSoA(void *label_ptr, void *localID_ptr,
void *rx_ptr, void *ry_ptr, void *rz_ptr,
......
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