eic-opticks/qudarap/QBnd.cu

0001 #include <stdio.h>
0002 #include "curand_kernel.h"
0003 #include "scuda.h"
0004 #include "qgs.h"
0005
0006 __global__ void _QBnd_lookup_0(cudaTextureObject_t tex, quad4* meta, quad* lookup, int num_lookup, int width, int height )
0007 {
0008     int ix = blockIdx.x * blockDim.x + threadIdx.x;
0009     int iy = blockIdx.y * blockDim.y + threadIdx.y;
0010     int index = iy * width + ix ;
0011     if (ix >= width | iy >= height ) return;
0012
0013     // Excluding hangover threads for a 2d launch based on 1d index is a bug
0014     // as the hangovers would overwrite into the output.
0015     // Must exclude based on both ix and iy.
0016
0017     int nx = meta->q0.u.x  ;
0018     int ny = meta->q0.u.y  ;
0019
0020     // HMM: this is assuming normalizedCoordinates:true thats not the normal way for boundary_tex
0021     float x = (float(ix)+0.5f)/float(nx) ;
0022     float y = (float(iy)+0.5f)/float(ny) ;
0023
0024     quad q ;
0025     q.f = tex2D<float4>( tex, x, y );
0026
0027 #ifdef DEBUG
0028     // debug launch config by returning coordinates
0029     printf(" ix %d iy %d index %d nx %d ny %d x %10.3f y %10.3f \n", ix, iy, index, nx, ny, x, y );
0030     q.u.x = ix ;
0031     q.u.y = iy ;
0032     q.u.z = index ;
0033     q.u.w = nx ;
0034 #endif
0035
0036     lookup[index] = q ;
0037 }
0038
0039 extern "C" void QBnd_lookup_0(
0040     dim3 numBlocks,
0041     dim3 threadsPerBlock,
0042     cudaTextureObject_t tex,
0043     quad4* meta,
0044     quad* lookup,
0045     int num_lookup,
0046     int width,
0047     int height  )
0048 {
0049     _QBnd_lookup_0<<<numBlocks,threadsPerBlock>>>( tex, meta, lookup, num_lookup, width, height );
0050 }
0051
0052