eic-opticks/qudarap/QSim_cerenkov_dev.cu

0001 #include "stdio.h"
0002 #include "qrng.h"
0003
0004 #include "scuda.h"
0005 #include "squad.h"
0006 #include "srec.h"
0007 #include "scerenkov.h"
0008 #include "sstate.h"
0009
0010 #include "qprop.h"
0011 #include "qbnd.h"
0012 #include "qsim.h"
0013 #include "qcerenkov_dev.h"
0014
0015
0016
0017
0018 /**
0019 genstep provisioning ? gensteps need to be uploaded with pointer held in qsim
0020 but for testing need to be able to manually fabricate a genstep
0021 **/
0022
0023 __global__ void _QSim_cerenkov_dev_wavelength_rejection_sampled(qsim* sim, float* wavelength, unsigned num_wavelength )
0024 {
0025     unsigned idx = blockIdx.x*blockDim.x + threadIdx.x;
0026     if (idx >= num_wavelength) return;
0027
0028     RNG rng ;  sim->rng->init(rng, 0, idx);
0029
0030     float wl = qcerenkov_dev::wavelength_rejection_sampled(sim, idx, rng);
0031
0032     if(idx % 100000 == 0) printf("//_QSim_cerenkov_dev_wavelength_rejection_sampled idx %d wl %10.4f    \n", idx, wl  );
0033     wavelength[idx] = wl ;
0034 }
0035
0036
0037 extern void QSim_cerenkov_dev_wavelength_rejection_sampled(dim3 numBlocks, dim3 threadsPerBlock, qsim* sim, float* wavelength, unsigned num_wavelength )
0038 {
0039     printf("//QSim_cerenkov_dev_wavelength_rejection_sampled num_wavelength %d \n", num_wavelength );
0040     _QSim_cerenkov_dev_wavelength_rejection_sampled<<<numBlocks,threadsPerBlock>>>( sim, wavelength, num_wavelength );
0041 }
0042
0043 __global__ void _QSim_cerenkov_dev_generate(qsim* sim, quad4* photon, unsigned num_photon )
0044 {
0045     unsigned idx = blockIdx.x*blockDim.x + threadIdx.x;
0046     if (idx >= num_photon) return;
0047
0048     RNG rng ; sim->rng->init(rng, 0, idx);
0049
0050     quad4 p ;
0051     qcerenkov_dev::generate(sim, p, idx, rng);
0052
0053     if(idx % 100000 == 0) printf("//_QSim_cerenkov_dev_generate idx %d \n", idx  );
0054     photon[idx] = p ;
0055 }
0056
0057 extern void QSim_cerenkov_dev_generate(dim3 numBlocks, dim3 threadsPerBlock, qsim* sim, quad4* photon, unsigned num_photon )
0058 {
0059     printf("//QSim_cerenkov_dev_generate num_photon %d \n", num_photon );
0060     _QSim_cerenkov_dev_generate<<<numBlocks,threadsPerBlock>>>( sim, photon, num_photon );
0061 }
0062
0063 template<typename T>
0064 __global__ void _QSim_cerenkov_dev_generate_enprop(qsim* sim, quad4* photon, unsigned num_photon )
0065 {
0066     unsigned idx = blockIdx.x*blockDim.x + threadIdx.x;
0067     if (idx >= num_photon) return;
0068
0069     RNG rng ;  sim->rng->init(rng, 0, idx);
0070
0071     quad4 p ;
0072     qcerenkov_dev::generate_enprop<T>(sim, p, idx, rng);
0073
0074     if(idx % 100000 == 0) printf("//_QSim_cerenkov_dev_generate_enprop idx %d \n", idx  );
0075     photon[idx] = p ;
0076 }
0077
0078
0079 template<typename T>
0080 extern void QSim_cerenkov_dev_generate_enprop(dim3 numBlocks, dim3 threadsPerBlock, qsim* sim, quad4* photon, unsigned num_photon)
0081 {
0082     printf("//QSim_cerenkov_dev_generate_enprop num_photon %d \n", num_photon );
0083     _QSim_cerenkov_dev_generate_enprop<T><<<numBlocks,threadsPerBlock>>>( sim, photon, num_photon );
0084 }
0085
0086
0087 template void QSim_cerenkov_dev_generate_enprop<float>(dim3, dim3, qsim*, quad4*, unsigned );
0088 template void QSim_cerenkov_dev_generate_enprop<double>(dim3, dim3, qsim*, quad4*, unsigned );
0089
0090
0091
0092
0093
0094
0095 __global__ void _QSim_cerenkov_dev_generate_expt_double(qsim* sim, quad4* photon, unsigned num_photon )
0096 {
0097     unsigned idx = blockIdx.x*blockDim.x + threadIdx.x;
0098     if (idx >= num_photon) return;
0099
0100     RNG rng ;
0101     sim->rng->init(rng, 0, idx);
0102
0103     quad4 p ;
0104     qcerenkov_dev::generate_expt_double(sim, p, idx, rng);
0105
0106     if(idx % 100000 == 0) printf("//_QSim_cerenkov_dev_generate_expt_double idx %d \n", idx  );
0107     photon[idx] = p ;
0108 }
0109
0110 extern void QSim_cerenkov_dev_generate_expt_double(dim3 numBlocks, dim3 threadsPerBlock, qsim* sim, quad4* photon, unsigned num_photon )
0111 {
0112     printf("//QSim_cerenkov_dev_generate_expt_double num_photon %d \n", num_photon );
0113     _QSim_cerenkov_dev_generate_expt_double<<<numBlocks,threadsPerBlock>>>( sim, photon, num_photon );
0114 }
0115
0116