eic-opticks/qudarap/QScint.cu

0001
0002
0003 #include "curand_kernel.h"
0004 #include "scuda.h"
0005 #include "squad.h"
0006 #include "sphoton.h"
0007 #include "qscint.h"
0008
0009 #include "stdio.h"
0010
0011 /**
0012 Minimize the code here, as this "junction" code cannot be easily tested/mocked for use from multiple contexts.
0013 Instead implement in simple headers for flexibility of usage and testing.
0014
0015 TODO
0016
0017 * simplify this using a qscint.h helper instance that collects the resources : texObj, photon,  etc..
0018   in order to modularize
0019
0020 See::
0021
0022 * qsim::scint_wavelength_hd0
0023 * qsim::scint_wavelength_hd10
0024 * qsim::scint_wavelength_hd20
0025
0026
0027 **/
0028
0029 __global__ void _QScint_check(unsigned width, unsigned height)
0030 {
0031     unsigned ix = blockIdx.x*blockDim.x + threadIdx.x;
0032     unsigned iy = blockIdx.y*blockDim.y + threadIdx.y;
0033     printf( "//_QScint_check  ix %d iy %d width %d height %d \n", ix, iy, width, height );
0034 }
0035
0036 __global__ void _QScint_lookup(cudaTextureObject_t tex, quad4* d_meta, float* lookup, unsigned num_lookup, unsigned width, unsigned height )
0037 {
0038     unsigned ix = blockIdx.x*blockDim.x + threadIdx.x;
0039     unsigned iy = blockIdx.y*blockDim.y + threadIdx.y;
0040     if (ix >= width || iy >= height) return;
0041
0042     unsigned index = iy * width + ix ;
0043
0044     unsigned nx = d_meta ? d_meta->q0.u.x : 0 ;
0045     unsigned ny = d_meta ? d_meta->q0.u.y : 0 ;
0046
0047     float x = (float(ix)+0.5f)/float(nx) ;
0048     float y = (float(iy)+0.5f)/float(ny) ;
0049
0050     float v = tex2D<float>( tex, x, y );
0051
0052     if( ix % 100 == 0 ) printf( "//_QScint_lookup  ix %d iy %d width %d height %d index %d nx %d ny %d x %10.4f y %10.4f v %10.4f \n", ix, iy, width, height, index, nx, ny, x, y, v );
0053     lookup[index] = v ;
0054 }
0055
0056 extern "C" void QScint_lookup(dim3 numBlocks, dim3 threadsPerBlock, cudaTextureObject_t tex, quad4* meta, float* lookup, unsigned num_lookup, unsigned width, unsigned height )
0057 {
0058     printf("//QScint_lookup num_lookup %d width %d height %d \n", num_lookup, width, height );
0059     _QScint_lookup<<<numBlocks,threadsPerBlock>>>( tex, meta, lookup, num_lookup, width, height  );
0060 }
0061
0062 extern "C" void QScint_check(dim3 numBlocks, dim3 threadsPerBlock, unsigned width, unsigned height )
0063 {
0064     printf("//QScint_check width %d height %d \n", width, height );
0065     _QScint_check<<<numBlocks,threadsPerBlock>>>( width, height  );
0066 }
0067
0068
0069
0070