#line 50 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
#include "centroiding.h"

#line 253 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
__global__ void centroidingEngine(float *cx, float *cy,
				  float *flux, const float *frame,
				  float cx0, float cy0, float units,
				  int _N_SIDE_LENSLET_, int N_SOURCE,
				  int N_PX_CAMERA)
{
  
#line 398 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
int u, k, kp, i, ij, ij_inc, iLenslet,jLenslet, N_PX_CAMERA2, iSource, kSource;
float *buffer0, *buffer1, *buffer2;
u = threadIdx.x;
iLenslet = blockIdx.x;
jLenslet = blockIdx.y;
iSource  = blockIdx.z;
#line 260 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
  extern __shared__ float shared[];
  
#line 406 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
N_PX_CAMERA2 = N_PX_CAMERA*N_PX_CAMERA;
buffer0 = shared;
buffer1 = shared + N_PX_CAMERA2;
buffer2 = shared + N_PX_CAMERA2*2;

#line 263 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
  // CENTROIDING
  
#line 412 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
if (u<N_PX_CAMERA)
{
  buffer0[u] = buffer1[u] = buffer2[u] = 0.0;
  ij = lenslet2array(u, 0, N_PX_CAMERA, iLenslet, jLenslet, _N_SIDE_LENSLET_, iSource);
  for (i=0;i<N_PX_CAMERA;i++)
  {
    ij_inc = ij + i;
    buffer0[u]  += frame[ij_inc];
    buffer1[u]  += u*frame[ij_inc];
    buffer2[u]  += i*frame[ij_inc];
  }
}
#line 264 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
                          
  __syncthreads();

  if (u<1)
  {
    
#line 426 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
k = kp = blockIdx.x*_N_SIDE_LENSLET_ + blockIdx.y;
kSource = iSource*_N_SIDE_LENSLET_*_N_SIDE_LENSLET_;
k  += 2*kSource;
kp +=   kSource;
flux[kp]  = cx[k] = cy[k] = 0.0;
for (i=0;i<N_PX_CAMERA;i++)
{
  flux[kp] += buffer0[i];
  cx[k]   += buffer1[i];
  cy[k]   += buffer2[i];
}
#line 270 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
    if (flux[kp]>0) // NORMALIZATION
      {
	cx[k] /= flux[kp];
	cy[k] /= flux[kp];
      } 
    else 
      {
	cx[k] = cy[k] = 0.5*(N_PX_CAMERA-1);
      }
    cx[k] -= cx0;
    cx[k] *= units;
    cy[k] -= cy0;
    cy[k] *= units;
  }
}
#line 318 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
__global__ void centroidingEngineGT(float *cx, float *cy,
				    float *flux, const float *frame,
				    float *cx0, float *cy0, float units,
				    int _N_SIDE_LENSLET_, int N_SOURCE,
				    int N_PX_CAMERA, char *valid_lenslet)
{
  
#line 398 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
int u, k, kp, i, ij, ij_inc, iLenslet,jLenslet, N_PX_CAMERA2, iSource, kSource;
float *buffer0, *buffer1, *buffer2;
u = threadIdx.x;
iLenslet = blockIdx.x;
jLenslet = blockIdx.y;
iSource  = blockIdx.z;
#line 325 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
  int kLenslet;
  kLenslet = (iLenslet + iSource*_N_SIDE_LENSLET_)*_N_SIDE_LENSLET_ + jLenslet;
  if (valid_lenslet[kLenslet]) 
  {
    extern __shared__ float shared[];
    
#line 406 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
N_PX_CAMERA2 = N_PX_CAMERA*N_PX_CAMERA;
buffer0 = shared;
buffer1 = shared + N_PX_CAMERA2;
buffer2 = shared + N_PX_CAMERA2*2;

#line 332 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
    // CENTROIDING
    
#line 412 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
if (u<N_PX_CAMERA)
{
  buffer0[u] = buffer1[u] = buffer2[u] = 0.0;
  ij = lenslet2array(u, 0, N_PX_CAMERA, iLenslet, jLenslet, _N_SIDE_LENSLET_, iSource);
  for (i=0;i<N_PX_CAMERA;i++)
  {
    ij_inc = ij + i;
    buffer0[u]  += frame[ij_inc];
    buffer1[u]  += u*frame[ij_inc];
    buffer2[u]  += i*frame[ij_inc];
  }
}
#line 333 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
                            
    __syncthreads();

    if (u<1)
    {
      
#line 426 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
k = kp = blockIdx.x*_N_SIDE_LENSLET_ + blockIdx.y;
kSource = iSource*_N_SIDE_LENSLET_*_N_SIDE_LENSLET_;
k  += 2*kSource;
kp +=   kSource;
flux[kp]  = cx[k] = cy[k] = 0.0;
for (i=0;i<N_PX_CAMERA;i++)
{
  flux[kp] += buffer0[i];
  cx[k]   += buffer1[i];
  cy[k]   += buffer2[i];
}
#line 339 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
      if (flux[kp]>0) // NORMALIZATION
	{
	  cx[k] /= flux[kp];
	  cy[k] /= flux[kp];
	}
      else 
	{
	  cx[k] = cy[k] = 0.5*(N_PX_CAMERA-1);
	}
      cx[k] -= cx0[k];
      cx[k] *= units;
      cy[k] -= cy0[k];
      cy[k] *= units;
    }
  }
}
#line 286 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
__global__ void centroidingEngineSh(float *cx, float *cy,
				  float *flux, const float *frame,
				  float cx0, float cy0, float units,
				  int _N_SIDE_LENSLET_, int N_SOURCE,
				  int N_PX_CAMERA, float *shared)
{
  
#line 398 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
int u, k, kp, i, ij, ij_inc, iLenslet,jLenslet, N_PX_CAMERA2, iSource, kSource;
float *buffer0, *buffer1, *buffer2;
u = threadIdx.x;
iLenslet = blockIdx.x;
jLenslet = blockIdx.y;
iSource  = blockIdx.z;
#line 293 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
  
#line 406 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
N_PX_CAMERA2 = N_PX_CAMERA*N_PX_CAMERA;
buffer0 = shared;
buffer1 = shared + N_PX_CAMERA2;
buffer2 = shared + N_PX_CAMERA2*2;

#line 295 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
  // CENTROIDING
  
#line 412 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
if (u<N_PX_CAMERA)
{
  buffer0[u] = buffer1[u] = buffer2[u] = 0.0;
  ij = lenslet2array(u, 0, N_PX_CAMERA, iLenslet, jLenslet, _N_SIDE_LENSLET_, iSource);
  for (i=0;i<N_PX_CAMERA;i++)
  {
    ij_inc = ij + i;
    buffer0[u]  += frame[ij_inc];
    buffer1[u]  += u*frame[ij_inc];
    buffer2[u]  += i*frame[ij_inc];
  }
}
#line 296 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
                          
  __syncthreads();

  if (u<1)
  {
    
#line 426 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
k = kp = blockIdx.x*_N_SIDE_LENSLET_ + blockIdx.y;
kSource = iSource*_N_SIDE_LENSLET_*_N_SIDE_LENSLET_;
k  += 2*kSource;
kp +=   kSource;
flux[kp]  = cx[k] = cy[k] = 0.0;
for (i=0;i<N_PX_CAMERA;i++)
{
  flux[kp] += buffer0[i];
  cx[k]   += buffer1[i];
  cy[k]   += buffer2[i];
}
#line 302 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
    if (flux[kp]>0) // NORMALIZATION
      {
	cx[k] /= flux[kp];
	cy[k] /= flux[kp];
      }
    else 
      {
	cx[k] = cy[k] = 0.5*(N_PX_CAMERA-1);
      }
    cx[k] -= cx0;
    cx[k] *= units;
    cy[k] -= cy0;
    cy[k] *= units;
  }
}
#line 356 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
__global__ void centroidingEngineGTSh(float *cx, float *cy,
				    float *flux, const float *frame,
				    float *cx0, float *cy0, float units,
				    int _N_SIDE_LENSLET_, int N_SOURCE,
				    int N_PX_CAMERA, char *valid_lenslet,
				    float *shared)
{
  
#line 398 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
int u, k, kp, i, ij, ij_inc, iLenslet,jLenslet, N_PX_CAMERA2, iSource, kSource;
float *buffer0, *buffer1, *buffer2;
u = threadIdx.x;
iLenslet = blockIdx.x;
jLenslet = blockIdx.y;
iSource  = blockIdx.z;
#line 364 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
  int kLenslet;
  kLenslet = (iLenslet + iSource*_N_SIDE_LENSLET_)*_N_SIDE_LENSLET_ + jLenslet;
  if (valid_lenslet[kLenslet]) 
  {
    
#line 406 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
N_PX_CAMERA2 = N_PX_CAMERA*N_PX_CAMERA;
buffer0 = shared;
buffer1 = shared + N_PX_CAMERA2;
buffer2 = shared + N_PX_CAMERA2*2;

#line 370 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
    // CENTROIDING
    
#line 412 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
if (u<N_PX_CAMERA)
{
  buffer0[u] = buffer1[u] = buffer2[u] = 0.0;
  ij = lenslet2array(u, 0, N_PX_CAMERA, iLenslet, jLenslet, _N_SIDE_LENSLET_, iSource);
  for (i=0;i<N_PX_CAMERA;i++)
  {
    ij_inc = ij + i;
    buffer0[u]  += frame[ij_inc];
    buffer1[u]  += u*frame[ij_inc];
    buffer2[u]  += i*frame[ij_inc];
  }
}
#line 371 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
                            
    __syncthreads();

    if (u<1)
    {
      
#line 426 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
k = kp = blockIdx.x*_N_SIDE_LENSLET_ + blockIdx.y;
kSource = iSource*_N_SIDE_LENSLET_*_N_SIDE_LENSLET_;
k  += 2*kSource;
kp +=   kSource;
flux[kp]  = cx[k] = cy[k] = 0.0;
for (i=0;i<N_PX_CAMERA;i++)
{
  flux[kp] += buffer0[i];
  cx[k]   += buffer1[i];
  cy[k]   += buffer2[i];
}
#line 377 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
      if (flux[kp]>0) // NORMALIZATION
	{
	  cx[k] /= flux[kp];
	  cy[k] /= flux[kp];
	}
      else 
	{
	  cx[k] = cy[k] = 0.5*(N_PX_CAMERA-1);
	}
      cx[k] -= cx0[k];
      cx[k] *= units;
      cy[k] -= cy0[k];
      cy[k] *= units;
    }
  }
}

#line 96 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
void centroiding::setup(int _N_SIDE_LENSLET, int _N_SOURCE) {
  _N_SIDE_LENSLET_ = _N_SIDE_LENSLET;
  N_LENSLET = _N_SIDE_LENSLET_*_N_SIDE_LENSLET_;
  N_SOURCE = _N_SOURCE;
  HANDLE_ERROR( cudaMalloc( (void**)&d__c   , sizeof(float)*N_LENSLET*N_SOURCE*2 ) );
  HANDLE_ERROR( cudaMemset( d__c, 0, sizeof(float)*N_LENSLET*N_SOURCE*2 ) );
  d__cx = d__c;
  d__cy = d__c + N_LENSLET;
  HANDLE_ERROR( cudaMalloc( (void**)&d__mass , sizeof(float) *N_LENSLET*N_SOURCE ) );     
  HANDLE_ERROR( cudaMalloc( (void**)&lenslet_mask , sizeof(char) *N_LENSLET ) );     
  MASK_SET = 0;
  
#line 124 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
cudaDeviceProp dev_prop;
int dev_id;
HANDLE_ERROR( cudaGetDevice( &dev_id ) );
HANDLE_ERROR( cudaGetDeviceProperties( &dev_prop, dev_id ) );
DEV_SHARED_MEM = dev_prop.sharedMemPerBlock;
DEV_MAX_THREADS = dev_prop.maxThreadsPerBlock;
//printf("Maximum shared memory:  %d\n",DEV_SHARED_MEM);
//printf("Maximum threads per block: %d\n",DEV_MAX_THREADS);
#line 108 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
}

#line 112 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
void centroiding::cleanup(void) {
  INFO("@(CEO)>centroiding: freeing memory!\n");
  HANDLE_ERROR( cudaFree( d__c ) );
  HANDLE_ERROR( cudaFree( d__mass ) );
  if (lenslet_mask) {
    HANDLE_ERROR( cudaFree( lenslet_mask ) );
  }
}
#line 60 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
   
#line 158 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
void centroiding::get_data(imaging *wfs) {
  int N_PX_CAMERA, N_PX_CAMERA_BYTE;
  float cx0, cy0, units, *frame;
  cx0 = cy0 = 0.0;
  units = 1.0;\
  frame = wfs->d__frame;
  N_PX_CAMERA = wfs->N_PX_CAMERA;
  N_PX_CAMERA_BYTE = sizeof(float)*N_PX_CAMERA*N_PX_CAMERA*3;
  
#line 196 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
if (N_PX_CAMERA>DEV_MAX_THREADS) {
  fprintf(stdout,"\n\x1B[31m@(CEO)>centroiding: The required number thread (%d) is larger\n that the maximum number of threads per block (%d)\x1B[0m\n",N_PX_CAMERA,DEV_MAX_THREADS);
  exit( EXIT_FAILURE );
  }
//printf("Required shared memory:  %d\n",N_PX_CAMERA_BYTE);
dim3 blockGrid(_N_SIDE_LENSLET_,_N_SIDE_LENSLET_,N_SOURCE);
dim3 threadGrid(N_PX_CAMERA,1);
if (N_PX_CAMERA_BYTE>DEV_SHARED_MEM) {
//  fprintf(stdout,"\n\x1B[31m@(CEO)>centroiding: The required shared memory (%dbytes) is larger\n that the maximum shared memory per block (%dbytes)\x1B[0m\n",N_PX_CAMERA_BYTE,DEV_SHARED_MEM);
  float *shared;
  HANDLE_ERROR( cudaMalloc( (void**)&shared , N_PX_CAMERA_BYTE ) );    
  centroidingEngineSh <<< blockGrid , threadGrid >>> 
    (d__cx, d__cy,
     d__mass, frame,
     cx0, cy0, units,
     _N_SIDE_LENSLET_, N_SOURCE,
     N_PX_CAMERA, shared);
  HANDLE_ERROR( cudaFree( shared ));
} else {
  centroidingEngine <<< blockGrid , threadGrid , N_PX_CAMERA_BYTE >>> 
    (d__cx, d__cy,
     d__mass, frame,
     cx0, cy0, units,
     _N_SIDE_LENSLET_, N_SOURCE,
     N_PX_CAMERA);
}
#line 167 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
}
#line 171 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
void centroiding::get_data(float *frame, int N_PX_CAMERA) {
  int N_PX_CAMERA_BYTE = sizeof(float)*N_PX_CAMERA*N_PX_CAMERA*3;
  float cx0, cy0, units;
  cx0 = cy0 = 0.0;
  units = 1.0;
  
#line 196 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
if (N_PX_CAMERA>DEV_MAX_THREADS) {
  fprintf(stdout,"\n\x1B[31m@(CEO)>centroiding: The required number thread (%d) is larger\n that the maximum number of threads per block (%d)\x1B[0m\n",N_PX_CAMERA,DEV_MAX_THREADS);
  exit( EXIT_FAILURE );
  }
//printf("Required shared memory:  %d\n",N_PX_CAMERA_BYTE);
dim3 blockGrid(_N_SIDE_LENSLET_,_N_SIDE_LENSLET_,N_SOURCE);
dim3 threadGrid(N_PX_CAMERA,1);
if (N_PX_CAMERA_BYTE>DEV_SHARED_MEM) {
//  fprintf(stdout,"\n\x1B[31m@(CEO)>centroiding: The required shared memory (%dbytes) is larger\n that the maximum shared memory per block (%dbytes)\x1B[0m\n",N_PX_CAMERA_BYTE,DEV_SHARED_MEM);
  float *shared;
  HANDLE_ERROR( cudaMalloc( (void**)&shared , N_PX_CAMERA_BYTE ) );    
  centroidingEngineSh <<< blockGrid , threadGrid >>> 
    (d__cx, d__cy,
     d__mass, frame,
     cx0, cy0, units,
     _N_SIDE_LENSLET_, N_SOURCE,
     N_PX_CAMERA, shared);
  HANDLE_ERROR( cudaFree( shared ));
} else {
  centroidingEngine <<< blockGrid , threadGrid , N_PX_CAMERA_BYTE >>> 
    (d__cx, d__cy,
     d__mass, frame,
     cx0, cy0, units,
     _N_SIDE_LENSLET_, N_SOURCE,
     N_PX_CAMERA);
}
#line 177 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
}
#line 181 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
void centroiding::get_data(float *frame, int N_PX_CAMERA, float cx0, float cy0, float units) {
  int N_PX_CAMERA_BYTE = sizeof(float)*N_PX_CAMERA*N_PX_CAMERA*3;
  
#line 196 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
if (N_PX_CAMERA>DEV_MAX_THREADS) {
  fprintf(stdout,"\n\x1B[31m@(CEO)>centroiding: The required number thread (%d) is larger\n that the maximum number of threads per block (%d)\x1B[0m\n",N_PX_CAMERA,DEV_MAX_THREADS);
  exit( EXIT_FAILURE );
  }
//printf("Required shared memory:  %d\n",N_PX_CAMERA_BYTE);
dim3 blockGrid(_N_SIDE_LENSLET_,_N_SIDE_LENSLET_,N_SOURCE);
dim3 threadGrid(N_PX_CAMERA,1);
if (N_PX_CAMERA_BYTE>DEV_SHARED_MEM) {
//  fprintf(stdout,"\n\x1B[31m@(CEO)>centroiding: The required shared memory (%dbytes) is larger\n that the maximum shared memory per block (%dbytes)\x1B[0m\n",N_PX_CAMERA_BYTE,DEV_SHARED_MEM);
  float *shared;
  HANDLE_ERROR( cudaMalloc( (void**)&shared , N_PX_CAMERA_BYTE ) );    
  centroidingEngineSh <<< blockGrid , threadGrid >>> 
    (d__cx, d__cy,
     d__mass, frame,
     cx0, cy0, units,
     _N_SIDE_LENSLET_, N_SOURCE,
     N_PX_CAMERA, shared);
  HANDLE_ERROR( cudaFree( shared ));
} else {
  centroidingEngine <<< blockGrid , threadGrid , N_PX_CAMERA_BYTE >>> 
    (d__cx, d__cy,
     d__mass, frame,
     cx0, cy0, units,
     _N_SIDE_LENSLET_, N_SOURCE,
     N_PX_CAMERA);
}
#line 184 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
}
#line 188 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
void centroiding::get_data(float *frame, int N_PX_CAMERA, 
			   float *cx0, float *cy0, float units,
			   char *valid_lenslet) {
  int N_PX_CAMERA_BYTE = sizeof(float)*N_PX_CAMERA*N_PX_CAMERA*3;
  
#line 224 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
if (N_PX_CAMERA>DEV_MAX_THREADS) {
  fprintf(stdout,"\n\x1B[31m@(CEO)>centroiding: The required number thread (%d) is larger\n that the maximum number of threads per block (%d)\x1B[0m\n",N_PX_CAMERA,DEV_MAX_THREADS);
  exit( EXIT_FAILURE );
  }
//printf("Required shared memory:  %d\n",N_PX_CAMERA_BYTE);
dim3 blockGrid(_N_SIDE_LENSLET_,_N_SIDE_LENSLET_,N_SOURCE);
dim3 threadGrid(N_PX_CAMERA,1);
if (N_PX_CAMERA_BYTE>DEV_SHARED_MEM) {
//  fprintf(stdout,"\n\x1B[31m@(CEO)>centroiding: The required shared memory (%dbytes) is larger\n that the maximum shared memory per block (%dbytes)\x1B[0m\n",N_PX_CAMERA_BYTE,DEV_SHARED_MEM);
  float *shared;
  HANDLE_ERROR( cudaMalloc( (void**)&shared , N_PX_CAMERA_BYTE ) );    
  centroidingEngineGTSh <<< blockGrid , threadGrid >>> 
    (d__cx, d__cy,
     d__mass, frame,
     cx0, cy0, units,
     _N_SIDE_LENSLET_, N_SOURCE,
     N_PX_CAMERA, valid_lenslet, shared);
  HANDLE_ERROR( cudaFree( shared ));
} else {
  centroidingEngineGT <<< blockGrid , threadGrid , N_PX_CAMERA_BYTE >>> 
    (d__cx, d__cy,
     d__mass, frame,
     cx0, cy0, units,
     _N_SIDE_LENSLET_, N_SOURCE,
     N_PX_CAMERA, valid_lenslet);
}
#line 193 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
}

#line 136 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
void centroiding::fried_geometry(mask *dm, mask *pupil, int n, float threshold) {
    MASK_SET = fried_geometry_setup_vs_pupil(lenslet_mask, dm->m, 
			                     _N_SIDE_LENSLET_, n, threshold, pupil->m);
    printf("Setting filter!\n");
    dm->set_filter();
    dm->area = pupil->area;
}
#line 147 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
void centroiding::reset(void)
{
  HANDLE_ERROR( cudaMemset( d__c, 0, sizeof(float)*N_LENSLET*N_SOURCE*2 ) );
}
#line 443 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
  void centroiding::filter(float *F)
{
  int k, n_byte, N;
  float alpha, beta;
  float *c;
  n_byte = sizeof(float)*N_LENSLET*N_SOURCE*2;

  HANDLE_ERROR( cudaMalloc( (void**)&c   , n_byte ) );
  HANDLE_ERROR( cudaMemcpy( c, d__c, n_byte, cudaMemcpyDeviceToDevice) );
  cublasCreate(&handle);
  alpha = 1;
  beta = 0;
  N = N_LENSLET*2;
  for (k=0;k<N_SOURCE;k++) {
    CUBLAS_ERROR( cublasSgemv(handle, CUBLAS_OP_N,
			    N, N,
			    &alpha, 
			    F, N,
			    c+k*N, 1,
			    &beta,
			    d__c+k*N , 1));
  }
  HANDLE_ERROR( cudaFree( c ) );
  cublasDestroy(handle);

  //  HANDLE_ERROR( cudaMemset( d__c, 0, sizeof(float)*N_LENSLET*N_SOURCE*2 ) );
}
#line 477 "/home/ubuntu/projects/crseo/sys/CEO/centroiding/centroiding.nw"
void centroiding::noisify(float stddev)
{
  cublasHandle_t handle;
  curandGenerator_t gen;
  float *d__n, alpha;
  int N;
  alpha = 1.0;
  N = N_LENSLET*N_SOURCE*2;

  HANDLE_ERROR( cudaMalloc( (void**)&d__n   , sizeof(float)*N ) );

  cublasCreate(&handle);
  HANDLE_ERROR_CURAND(curandCreateGenerator(&gen,
					    CURAND_RNG_PSEUDO_DEFAULT),
		      "Failed to create pseudo-random number generator!");

  HANDLE_ERROR_CURAND(curandSetPseudoRandomGeneratorSeed(gen,
							 1234ULL),
		      "Failed to set seed!");

  HANDLE_ERROR_CURAND(curandGenerateNormal(gen, d__n, N, 0, stddev),
		      "Failed to generate variates on device!");

  CUBLAS_ERROR( cublasSaxpy(handle, N, &alpha, d__n, 1, d__c, 1) );

  HANDLE_ERROR_CURAND(curandDestroyGenerator(gen),
		      "Failed to the cleanup!");
  cublasDestroy(handle);
  HANDLE_ERROR( cudaFree( d__n ) );

}