eic-opticks/qudarap/QMultiFilm.cu

0001
0002
0003 #include "curand_kernel.h"
0004 #include "scuda.h"
0005 #include "squad.h"
0006 #include "qmultifilm.h"
0007 //#include "qscint.h"
0008 #include "stdio.h"
0009
0010 /**
0011
0012  **/
0013
0014 __global__ void _QMultiFilm_check(unsigned width, unsigned height)
0015 {
0016     unsigned ix = blockIdx.x*blockDim.x + threadIdx.x;
0017     unsigned iy = blockIdx.y*blockDim.y + threadIdx.y;
0018     // printf("hello world!");
0019     if( ix % 10 == 0 )   printf( "//_QMultiFilm_check  ix %d iy %d width %d height %d \n", ix, iy, width, height );
0020 }
0021
0022 __global__ void _QMultiFilm_lookup(cudaTextureObject_t tex, quad4* meta, float4* lookup, unsigned num_lookup, unsigned width, unsigned height)
0023 {
0024     unsigned ix = blockIdx.x*blockDim.x + threadIdx.x;
0025     unsigned iy = blockIdx.y*blockDim.y + threadIdx.y;
0026     if (ix >= width || iy >= height) return;
0027
0028     int index = iy * width + ix;
0029     unsigned nx = meta->q0.u.x ;
0030     unsigned ny = meta->q0.u.y ;
0031
0032     float x = (float(ix)+0.5f)/float(nx) ;
0033     float y = (float(iy)+0.5f)/float(ny) ;
0034
0035     float4 v = tex2D<float4>( tex, x, y );
0036
0037     //printf( "//_QMultiFilm_lookup  ix %d iy %d width %d height %d index %d nx %d ny %d x %10.4f y %10.4f v.x %10.4f v.y %10.4f v.z %10.4f v.w %10.4f  \n", ix, iy, width, height, index, nx, ny, x, y, v.x, v.y, v.z, v.w );
0038     lookup[ index ] = v ;
0039 }
0040
0041 __global__ void _QMultiFilm_mock_lookup(qmultifilm* d_multifilm, quad2* d_input, float4* d_out, unsigned num_lookup, unsigned width, unsigned height)
0042 {
0043
0044     unsigned ix = blockIdx.x*blockDim.x + threadIdx.x;
0045     unsigned iy = blockIdx.y*blockDim.y + threadIdx.y;
0046     if (ix >= width || iy >= height) return;
0047
0048     int index = iy * width + ix;
0049     int pmtcat  = d_input[index].q0.i.x;
0050     float wv_nm = d_input[index].q0.f.y;
0051     float aoi   = d_input[index].q0.f.z;
0052
0053     float Rs = d_input[index].q1.f.x;
0054     float Ts = d_input[index].q1.f.y;
0055     float Rp = d_input[index].q1.f.z;
0056     float Tp = d_input[index].q1.f.w;
0057
0058     float4 value = d_multifilm->lookup(pmtcat, wv_nm, aoi);
0059     d_out[ index ] = value ;
0060
0061
0062     int item = 10;
0063     if(index < item || index > (num_lookup - item)){
0064         printf( "//_QMultiFilm_mock_lookup  ix %d iy %d width %d height %d index %d Rs %10.4f Ts %10.4f Rp %10.4f Tp %10.4f Rs_i %10.4f Ts_i %10.4f Rp_i %10.4f Tp_i %10.4f \n", ix, iy, width, height, index, Rs, Ts, Rp, Tp, value.x, value.y , value.z, value.w );
0065     }
0066 }
0067 extern "C" void QMultiFilm_lookup(dim3 numBlocks, dim3 threadsPerBlock, cudaTextureObject_t tex, quad4* meta, float4* lookup, unsigned num_lookup, unsigned width, unsigned height)
0068 {
0069     printf("//QMultiFilm_lookup num_lookup %d width %d height %d \n", num_lookup, width, height);
0070     _QMultiFilm_lookup<<<numBlocks,threadsPerBlock>>>( tex, meta, lookup, num_lookup, width, height );
0071 }
0072
0073
0074 extern "C" void QMultiFilm_check(dim3 numBlocks, dim3 threadsPerBlock, unsigned width, unsigned height )
0075 {
0076     printf("//QMultiFilm_check width %d height %d threadsPerBlock.x %d threadsPerBlock.y %d  numBlocks.x %d numBlocks.y %d \n", width, height,threadsPerBlock.x, threadsPerBlock.y , numBlocks.x, numBlocks.y);
0077
0078     _QMultiFilm_check<<<numBlocks,threadsPerBlock>>>( width, height  );
0079 }
0080
0081
0082 extern "C" void QMultiFilm_mock_lookup(dim3 numBlocks, dim3 threadsPerBlock, qmultifilm* d_multifilm, quad2* d_input, float4* d_out, unsigned num_lookup, unsigned width, unsigned height)
0083 {
0084     printf("//QMultiFilm_mock_lookup num_lookup %d width %d height %d \n", num_lookup, width, height);
0085     _QMultiFilm_mock_lookup<<<numBlocks,threadsPerBlock>>>( d_multifilm, d_input, d_out, num_lookup, width, height );
0086 }
0087
0088