eic-opticks/qudarap/QCerenkov.cu

0001
0002 #include "curand_kernel.h"
0003 #include "scuda.h"
0004 #include "squad.h"
0005 //#include "qcerenkov.h"
0006 #include "stdio.h"
0007
0008 /**
0009 Minimize the code here, as this "junction" code cannot be easily tested/mocked for use from multiple contexts.
0010 Instead implement in simple headers like qcerenkov.h for flexibility of usage and testing and for modularization.
0011
0012 * TODO: qcerenkov.h GPU side helper instance collecting resources, metadata etc.. : texObj, photon,  etc..
0013   in order to modularize  and keep qsim.h simple
0014
0015 **/
0016
0017 __global__ void _QCerenkov_check(unsigned width, unsigned height)
0018 {
0019     unsigned ix = blockIdx.x*blockDim.x + threadIdx.x;
0020     unsigned iy = blockIdx.y*blockDim.y + threadIdx.y;
0021     printf( "//_QCerenkov_check  ix %d iy %d width %d height %d \n", ix, iy, width, height );
0022 }
0023
0024 __global__ void _QCerenkov_lookup(cudaTextureObject_t tex, quad4* d_meta, float* lookup, unsigned num_lookup, unsigned width, unsigned height )
0025 {
0026     unsigned ix = blockIdx.x*blockDim.x + threadIdx.x;
0027     unsigned iy = blockIdx.y*blockDim.y + threadIdx.y;
0028     if (ix >= width || iy >= height) return;
0029
0030     unsigned index = iy * width + ix ;
0031
0032     unsigned nx = d_meta ? d_meta->q0.u.x : 0 ;
0033     unsigned ny = d_meta ? d_meta->q0.u.y : 0 ;
0034
0035     float x = (float(ix)+0.5f)/float(nx) ;
0036     float y = (float(iy)+0.5f)/float(ny) ;
0037
0038     float v = tex2D<float>( tex, x, y );
0039
0040     if( ix % 100 == 0 ) printf( "//_QCerenkov_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 );
0041     lookup[index] = v ;
0042 }
0043
0044 extern "C" void QCerenkov_lookup(dim3 numBlocks, dim3 threadsPerBlock, cudaTextureObject_t tex, quad4* meta, float* lookup, unsigned num_lookup, unsigned width, unsigned height )
0045 {
0046     printf("//QCerenkov_lookup num_lookup %d width %d height %d \n", num_lookup, width, height );
0047     _QCerenkov_lookup<<<numBlocks,threadsPerBlock>>>( tex, meta, lookup, num_lookup, width, height  );
0048 }
0049
0050 extern "C" void QCerenkov_check(dim3 numBlocks, dim3 threadsPerBlock, unsigned width, unsigned height )
0051 {
0052     printf("//QCerenkov_check width %d height %d \n", width, height );
0053     _QCerenkov_check<<<numBlocks,threadsPerBlock>>>( width, height  );
0054 }
0055
0056