% -*- mode: Noweb; noweb-code-mode: c-mode -*- @ \index{shackHartmann} \section{The files} \label{sec:files} \subsection{Header} \label{sec:header} <>= #ifndef __SHACKHARTMANN__ #define __SHACKHARTMANN__ #ifndef __UTILITIES_H__ #include "utilities.h" #endif #ifndef __SOURCE_H__ #include "source.h" #endif #ifndef __IMAGING_H__ #include "imaging.h" #endif #ifndef __CENTROIDING_H__ #include "centroiding.h" #endif struct shackHartmann { <> void setup(int N_SIDE_LENSLET, int N_PX_LENSLET, float d, int DFT_osf_, int N_PX_IMAGE, int BIN_IMAGE, int N_GS); <> float* get_frame_dev_ptr(void); void update_lenslet(float *filter); }; struct geometricShackHartmann { <> <> void setup(int N_SIDE_LENSLET, float d, int N_GS); <> void folded_slopes(float *d__valid_slopes, mask *lenslet_mask); void reset(void); }; struct tt7 { <> void setup(int N_SIDE_LENSLET, int N_PX_LENSLET, float d, int DFT_osf_, int N_PX_IMAGE, int BIN_IMAGE, int N_GS); <> float* get_frame_dev_ptr(void); void update_lenslet(float *filter); }; #endif // __SHACKHARTMANN__ @ with <>= void cleanup(void); void identify_valid_lenslet(source *src, float threshold); void set_reference_slopes(source *src); void calibrate(source *src, float threshold); void propagate(source *gs); void propagate(source *gs, int *maks); void process(void); void analyze(source *gs); void get_valid_reference_slopes(float *d__valid_slopes); void get_valid_slopes(float *d__valid_slopes); void masked_slopes(float *d__valid_slopes, mask *lenslet_mask); void get_valid_slopes_norm(float *d__valid_slopes_norm); @ \subsection{Source} \label{sec:source} <>= #include "shackHartmann.h" <> <> <> <> <> <> <> <> <> <> <> <> <> <> <> <> <> <> <> @ \section{Parameters} \label{sec:parameters} \index{shackHartmann!shackHartmann} \index{shackHartmann!tt7} \index{shackHartmann!geometricShackHartmann} The parameters of the [[shackHartmann]] structure are: \begin{itemize} \item the number of wavefront sensor: <>= int N_WFS; @ \item the number of lenslets (linear and total): <>= int N_SIDE_LENSLET, N_LENSLET; @ \item the total number of actuators in the Fried geometry: <>= int N_ACTUATOR; @ \item the total number of slopes: <>= int N_SLOPE; @ \item the reference slopes: <>= float *d__c0, *d__cx0, *d__cy0; @ \item the instantaneous geometric centroids: <>= float *_d__c_, *_d__cx_, *_d__cy_; @ \item the geometric centroids accumulation factor: <>= int N_FRAME; @ \item the valid lenslet mask: <>= mask valid_lenslet; @ \item the valid actuator mask for a Fried geometry: <>= mask valid_actuator; @ \item the imager which combines a lenslet array with a camera: <>= imaging camera; @ \item the data processing algorithm: <>= centroiding data_proc; @ \item the oversampling factor of the DFT [[DFT_osf]]: <>= int DFT_osf; @ \item the lenslet pitch: <>= float lenslet_pitch; @ \item the pixel scale: <>= float pixel_scale; @ \item the lenslet intensity threshold: <>= float intensity_threshold; @ \item the slopes gain: <>= float slopes_gain; @ \item CUBLAS handle: <>= cublasHandle_t handle; @ \end{itemize} \section{Functions} \label{sec:functions} \subsection{Setup \& Cleanup} \label{sec:setup--cleanup} The [[shackHartmann]] structure is initialized with: \begin{itemize} \item the linear size of the lenslet array [[N_SIDE_LENSLET]], \item the number of pixels per lenslet [[N_PX_LENSLET]], \item the oversampling factor of the DFT [[DFT_osf]]$\geq 2$, \item the pixel size of the image plane [[N_PX_IMAGE]], \item the binning factor of the image plane [[BIN_IMAGE]]$\geq 1$, \item the number of guide stars [[N_GS]]. \end{itemize} \index{shackHartmann!shackHartmann!setup} <>= void shackHartmann::setup(int N_SIDE_LENSLET_, int N_PX_LENSLET, float d, int DFT_osf_, int N_PX_IMAGE, int BIN_IMAGE, int N_GS) { N_WFS = N_GS; <> } @ with <>= lenslet_pitch = d; DFT_osf = DFT_osf_; N_SIDE_LENSLET = N_SIDE_LENSLET_; N_LENSLET = N_SIDE_LENSLET*N_SIDE_LENSLET; N_ACTUATOR = (N_SIDE_LENSLET+1)*(N_SIDE_LENSLET+1); N_SLOPE = N_LENSLET*N_WFS*2; slopes_gain = 1.0; HANDLE_ERROR( cudaMalloc( (void**)&d__c0 , sizeof(float) *N_SLOPE ) ); HANDLE_ERROR( cudaMemset( d__c0, 0, sizeof(float)*N_SLOPE ) ); d__cx0 = d__c0; d__cy0 = d__c0 + N_LENSLET; valid_lenslet.setup(N_LENSLET*N_WFS);//,(N_SIDE_LENSLET-1)*d); valid_actuator.setup(N_ACTUATOR*N_WFS);//,N_SIDE_LENSLET*d); camera.setup(N_PX_LENSLET, N_SIDE_LENSLET, DFT_osf, N_PX_IMAGE, BIN_IMAGE, N_WFS); data_proc.setup(N_SIDE_LENSLET, N_WFS); <> @ The [[lenslet_mask]] in [[data_proc]] now points to the mask in the [[valid_lenslet]] [[mask]] structure <>= HANDLE_ERROR( cudaFree( data_proc.lenslet_mask) ); data_proc.lenslet_mask = valid_lenslet.m; @ The GMT TT7 wavefront sensor is a special kind of Shack--Hartmann wavefront sensor with one imaging system per GMT pupil segment all using the same single guide star. \index{shackHartmann!tt7!setup} <>= void tt7::setup(int N_SIDE_LENSLET_, int N_PX_LENSLET, float d, int DFT_osf_, int N_PX_IMAGE, int BIN_IMAGE, int N_GS) { N_WFS = 7; <> } @ The default [[shackHartmann]] structure uses a diffractive model to propagate the wavefront onto the detector and then computes the centroids. The [[geometricShackHartmann]] structure does not propagate the wavefront instead the centroids are computed directly from the wavefront phase. \index{shackHartmann!geometricShackHartmann!setup} <>= void geometricShackHartmann::setup(int N_SIDE_LENSLET_, float d, int N_GS) { lenslet_pitch = d; N_SIDE_LENSLET = N_SIDE_LENSLET_; N_LENSLET = N_SIDE_LENSLET*N_SIDE_LENSLET; N_ACTUATOR = (N_SIDE_LENSLET+1)*(N_SIDE_LENSLET+1); N_WFS = N_GS; N_SLOPE = N_LENSLET*N_GS*2; slopes_gain = 1.0; HANDLE_ERROR( cudaMalloc( (void**)&d__c0 , sizeof(float) *N_SLOPE ) ); HANDLE_ERROR( cudaMemset( d__c0, 0, sizeof(float)*N_SLOPE ) ); d__cx0 = d__c0; d__cy0 = d__c0 + N_LENSLET; HANDLE_ERROR( cudaMalloc( (void**)&_d__c_ , sizeof(float) *N_SLOPE ) ); HANDLE_ERROR( cudaMemset( _d__c_, 0, sizeof(float)*N_SLOPE ) ); _d__cx_ = _d__c_; _d__cy_ = _d__c_ + N_LENSLET; valid_lenslet.setup(N_LENSLET*N_WFS);//,(N_SIDE_LENSLET-1)*d); valid_actuator.setup(N_ACTUATOR*N_WFS);//,N_SIDE_LENSLET*d); data_proc.setup(N_SIDE_LENSLET, N_GS); <> cublasCreate(&handle); N_FRAME = 0; } @ Allocated memory is freed with: \index{shackHartmann!shackHartmann!cleanup} <>= void shackHartmann::cleanup(void) { <> } @ \index{shackHartmann!tt7!cleanup} <>= void tt7::cleanup(void) { <> } @ with <>= <> INFO(" |-"); camera.cleanup(); @ or with \index{shackHartmann!geometricShackHartmann!cleanup} <>= void geometricShackHartmann::cleanup(void) { <> HANDLE_ERROR( cudaFree( _d__c_) ); cublasDestroy(handle); } @ where <>= INFO("@(CEO)>shackHartmann: freeing memory!\n"); HANDLE_ERROR( cudaFree( d__c0) ); INFO(" |-"); valid_lenslet.cleanup(); data_proc.lenslet_mask = 0; INFO(" |-"); valid_actuator.cleanup(); INFO(" |-"); data_proc.cleanup(); @ \subsection{Calibration} \label{sec:calibration} The WFS calibration consists in selecting the valid lenslets based on their illumination using an intensity threshold parameter [[threshold]] and setting the reference slopes. The valid lenslets are identified with: \index{shackHartmann!shackHartmann!identify\_valid\_lenslet} <>= void shackHartmann::identify_valid_lenslet(source *src, float threshold) { <> data_proc.MASK_SET = 1; } void tt7::identify_valid_lenslet(source *src, float threshold) { } @ The reference slopes are set with: \index{shackHartmann!shackHartmann!set\_reference\_slopes} <>= void shackHartmann::set_reference_slopes(source *src) { data_proc.get_data(camera.d__frame,camera.N_PX_CAMERA); <> } void tt7::set_reference_slopes(source *src) { } @ The WFS full calibration routine is given by: \index{shackHartmann!shackHartmann!calibrate} <>= void shackHartmann::calibrate(source *src, float threshold) { <> } @ \index{shackHartmann!shackHartmann!calibrate} <>= void tt7::calibrate(source *src, float threshold) { <> } @ with <>= <> <> <> data_proc.MASK_SET = 1; camera.reset(); @ The wavefront sensor need first to be calibrated with a reference calibration source. <>= //camera.pixel_scale = src->wavelength()/(lenslet_pitch*DFT_osf); //pixel_scale = camera.pixel_scale*camera.BIN_IMAGE*slopes_gain; propagate(src); data_proc.get_data(camera.d__frame,camera.N_PX_CAMERA); @ The slopes derived from the calibration wavefront will be set a the reference slopes. <>= HANDLE_ERROR( cudaMemcpy( d__c0, data_proc.d__c, N_SLOPE*sizeof(float), cudaMemcpyDeviceToDevice ) ); HANDLE_ERROR( cudaMemset( data_proc.d__c, 0, sizeof(float)*N_SLOPE ) ); @ The total intensity per lenslet is used to compute the valid lenslet. The valid lenslet and actuator masks are first set to 0. <>= HANDLE_ERROR( cudaMemset(valid_lenslet.m, 0, sizeof(char)*N_LENSLET*N_WFS ) ); HANDLE_ERROR( cudaMemset(valid_actuator.m, 0, sizeof(char)*N_ACTUATOR*N_WFS ) ); @ The maximum lenslet intensity is sought: <>= cublasHandle_t handle; cublasCreate(&handle); int idx; CUBLAS_ERROR( cublasIsamax(handle, N_LENSLET*N_WFS, data_proc.d__mass, 1, &idx) ); cublasDestroy(handle); idx--; @ The masks are computed with the [[set_valid_lenslet]] kernel: <>= int NL = camera.N_SIDE_LENSLET; dim3 blockDim(16,16); dim3 gridDim(NL/16+1,NL/16+1,N_WFS); // The 3rd argument sets the WFS which intensity is used for calibration set_valid_lenslet LLL gridDim,blockDim RRR (valid_lenslet.m, valid_actuator.m, NL, data_proc.d__mass, idx, threshold); valid_lenslet.set_filter(); valid_lenslet.set_index(); valid_actuator.set_filter(); @ with the kernel: <>= __global__ void set_valid_lenslet(char *lenslet_mask, char *actuator_mask, int NL, float *lenslet_intensity_map, int max_idx, float threshold) { int iL, jL, kL, kA, iSource; float max_intensity_threshold; // lenslet index iL = blockIdx.x * blockDim.x + threadIdx.x; jL = blockIdx.y * blockDim.y + threadIdx.y; iSource = blockIdx.z; if ( (iL=max_intensity_threshold) { lenslet_mask[kL] = 1; kA = lenset2actuator(iL,jL,NL, 1, 1); kA += iSource*(1+NL)*(1+NL); actuator_mask[kA] = 1; kA = lenset2actuator(iL,jL,NL,-1, 1); kA += iSource*(1+NL)*(1+NL); actuator_mask[kA] = 1; kA = lenset2actuator(iL,jL,NL,-1,-1); kA += iSource*(1+NL)*(1+NL); actuator_mask[kA] = 1; kA = lenset2actuator(iL,jL,NL, 1,-1); kA += iSource*(1+NL)*(1+NL); actuator_mask[kA] = 1; } else lenslet_intensity_map[kL] = 0.0; } } @ The valid lenslet mask can be modified with <>= void shackHartmann::update_lenslet(float *filter) { valid_lenslet.alter(filter); valid_lenslet.set_filter(); valid_lenslet.set_index(); valid_actuator.set_filter(); } void tt7::update_lenslet(float *filter) { } @ \subsubsection{Geometric model} \label{sec:geometric-model} The valid lenslets are identified with: \index{shackHartmann!geometricShackHartmann!identify\_valid\_lenslet} <>= void geometricShackHartmann::identify_valid_lenslet(source *src, float threshold) { int n; n = (int) sqrt(src->wavefront.N_PX/src->wavefront.N); n -= 1; n/= N_SIDE_LENSLET; HANDLE_ERROR( cudaMemset(valid_actuator.m, 0, sizeof(char)*N_ACTUATOR*N_WFS ) ); dim3 blockDim(16,16); dim3 gridDim(N_SIDE_LENSLET/16+1,N_SIDE_LENSLET/16+1,N_WFS); set_geometric_valid_lenslet LLL gridDim,blockDim RRR (valid_lenslet.f, valid_lenslet.m, src->wavefront.M->f, valid_actuator.m, threshold, N_SIDE_LENSLET, n); valid_lenslet.set_filter(); valid_lenslet.set_index(); valid_actuator.set_filter(); } @ The reference slopes are set with: \index{shackHartmann!geometricShackHartmann!set\_reference\_slopes} <>= void geometricShackHartmann::set_reference_slopes(source *src) { src->wavefront.finite_difference(data_proc.d__cx, data_proc.d__cy, N_SIDE_LENSLET, lenslet_pitch, &valid_lenslet); <> } @ The WFS full calibration routine is given by: \index{shackHartmann!geometricShackHartmann!calibrate} <>= void geometricShackHartmann::calibrate(source *src, float threshold) { identify_valid_lenslet(src, threshold); set_reference_slopes(src); } @ The valid lenslet are identified by computing the normalized sum of the wavefront mask [[M]] in each lenslet: <>= __global__ void set_geometric_valid_lenslet(float *valid_lenslet_f, char *valid_lenslet_m, float *wavefront_mask_f, char *actuator_mask, float threshold, const int NL, const int n) { int iL, jL, i, j, kL, kA, u, v, w0, w, NP, iSource, n2; iL = blockIdx.x * blockDim.x + threadIdx.x; jL = blockIdx.y * blockDim.y + threadIdx.y; iSource = blockIdx.z; kL = iL*NL+jL; kL += iSource*NL*NL; if ( (iL>= void geometricShackHartmann::reset() { N_FRAME = 0; HANDLE_ERROR( cudaMemset( data_proc.d__c, 0, sizeof(float)*N_SLOPE ) ); } @ \subsection{Wavefront processing} \label{sec:wavefront-processing} The guide star [[gs]] is propagated through the lenslet array to the detector: \index{shackHartmann!shackHartmann!propagate} <>= void shackHartmann::propagate(source *gs) { <> camera.propagate(gs); } @ \index{shackHartmann!tt7!propagate} <>= void tt7::propagate(source *gs) { <> camera.propagateTT7(gs); } void tt7::propagate(source *gs, int *mask) { <> camera.propagateTT7(gs, mask); } @ with <>= /* printf("slopes gain:%.2E\n",slopes_gain); */ camera.pixel_scale = gs->wavelength()/(lenslet_pitch*DFT_osf); /* printf("pixel scale:%.2E\n",pixel_scale); */ pixel_scale = camera.pixel_scale*camera.BIN_IMAGE; @ and the resulting frame is processed to extract the centroids \index{shackHartmann!shackHartmann!process} <>= void shackHartmann::process(void) { <> } @ \index{shackHartmann!tt7!process} <>= void tt7::process(void) { <> } @ with <>= data_proc.get_data(camera.d__frame,camera.N_PX_CAMERA, d__cx0,d__cy0,pixel_scale*slopes_gain, valid_lenslet.m); @ Noiseless propagation and centroiding can be performed, for the diffractive and geometric model, respectively, with a single call to: \index{shackHartmann!shackHartmann!analyze} <>= void shackHartmann::analyze(source *gs) { <> } @ \index{shackHartmann!tt7!analyze} <>= void tt7::analyze(source *gs) { <> } @ with <>= propagate(gs); process(); @ \index{shackHartmann!geometricShackHartmann!propagate} <>= void geometricShackHartmann::propagate(source *gs) { gs->wavefront.finite_difference(_d__cx_, _d__cy_, N_SIDE_LENSLET, lenslet_pitch, &valid_lenslet); float alpha = 1.0; CUBLAS_ERROR( cublasSaxpy(handle, N_SLOPE, &alpha, _d__c_, 1, data_proc.d__c, 1) ); ++N_FRAME; } @ \index{shackHartmann!geometricShackHartmann!process} <>= void geometricShackHartmann::process() { float alpha; alpha= 1.0/N_FRAME; CUBLAS_ERROR( cublasSscal(handle, N_SLOPE, &alpha, data_proc.d__c, 1) ); alpha = -1.0; CUBLAS_ERROR( cublasSaxpy(handle, N_SLOPE, &alpha, d__c0, 1, data_proc.d__c, 1) ); } @ \index{shackHartmann!geometricShackHartmann!analyze} <>= void geometricShackHartmann::analyze(source *gs) { propagate(gs); process(); } @ A pointer to the detector frame is obtained with <>= float* shackHartmann::get_frame_dev_ptr(void) { return camera.d__frame; } @ The slopes in the array [[data_proc.d__c]] contains first all the slopes along the X--axis direction and then all the slopes along the Y--axis direction. This pattern repeats itself for multiple guide starts i.e. [[data_proc.d__c]]$=[c^1_x,c^1_y,c^2_x,c^2_y,\dots]$. The slopes in the valid lenslet are extracted with \index{shackHartmann!shackHartmann!get\_valid\_slopes} <>= void shackHartmann::get_valid_slopes(float *d__valid_slopes) { <> } @ \index{shackHartmann!tt7!get\_valid\_slopes} <>= void tt7::get_valid_slopes(float *d__valid_slopes) { <> } @ \index{shackHartmann!geometricShackHartmann!get\_valid\_slopes} <>= void geometricShackHartmann::get_valid_slopes(float *d__valid_slopes) { <> } @ where <>= dim3 blockDim(16,16); dim3 gridDim(valid_lenslet.nnz/256+1,1); get_valid_slopes_kern LLL gridDim, blockDim RRR (d__valid_slopes, valid_lenslet.idx, valid_lenslet.nnz, data_proc.d__c, N_LENSLET); @ with <>= __global__ void get_valid_slopes_kern(float *d__valid_slopes, int *idx, const int NNZ, float *d__c, const int N_LENSLET) { int i, j, k, iSource, _idx_; i = blockIdx.x * blockDim.x + threadIdx.x; j = blockIdx.y * blockDim.y + threadIdx.y; k = j * gridDim.x * blockDim.x + i; if (k>= void shackHartmann::get_valid_reference_slopes(float *d__valid_slopes) { <> } void tt7::get_valid_reference_slopes(float *d__valid_slopes) { } @ \index{shackHartmann!geometricShackHartmann!get\_valid\_reference\_slopes} <>= void geometricShackHartmann::get_valid_reference_slopes(float *d__valid_slopes) { <> } @ where <>= dim3 blockDim(16,16); dim3 gridDim(valid_lenslet.nnz/256+1,1); get_valid_slopes_kern LLL gridDim, blockDim RRR (d__valid_slopes, valid_lenslet.idx, valid_lenslet.nnz, d__c0, N_LENSLET); @ The valid slopes norm $\bar s$ is defined as $\bar s^2 = (s_x+s_{0,x})^2 + (s_y+s_{0,y})^2$. It is computed with: \index{shackHartmann!shackHartmann!get\_valid\_slopes\_norm} <>= void shackHartmann::get_valid_slopes_norm(float *d__valid_slopes_norm) { <> } void tt7::get_valid_slopes_norm(float *d__valid_slopes_norm) { } @ \index{shackHartmann!geometricShackHartmann!get\_valid\_slopes\_norm} <>= void geometricShackHartmann::get_valid_slopes_norm(float *d__valid_slopes_norm) { <> } @ where <>= dim3 blockDim(16,16); dim3 gridDim(valid_lenslet.nnz/256+1,1); get_valid_slopes_norm_kern LLL gridDim, blockDim RRR (d__valid_slopes_norm, valid_lenslet.idx, valid_lenslet.nnz, data_proc.d__c, d__c0, N_LENSLET); @ with <>= __global__ void get_valid_slopes_norm_kern(float *d__valid_slopes_norm, int *idx, const int NNZ, float *d__c, float *d__c0, const int N_LENSLET) { int i, j, k, iSource, _idx_; i = blockIdx.x * blockDim.x + threadIdx.x; j = blockIdx.y * blockDim.y + threadIdx.y; k = j * gridDim.x * blockDim.x + i; if (k>= void shackHartmann::masked_slopes(float *d__valid_slopes, mask *lenslet_mask) { <> } @ \index{shackHartmann!tt7!masked\_slopes} <>= void tt7::masked_slopes(float *d__valid_slopes, mask *lenslet_mask) { <> } @ \index{shackHartmann!geometricShackHartmann!masked\_slopes} <>= void geometricShackHartmann::masked_slopes(float *d__valid_slopes, mask *lenslet_mask) { <> } @ where <>= dim3 blockDim(16,16); dim3 gridDim(lenslet_mask->nnz/256+1,1); get_valid_slopes_kern LLL gridDim, blockDim RRR (d__valid_slopes, lenslet_mask->idx, lenslet_mask->nnz, data_proc.d__c, N_LENSLET); @ The masked slopes can be averaged on the WFSs: \index{shackHartmann!geometricShackHartmann!folded\_slopes} <>= void geometricShackHartmann::folded_slopes(float *d__valid_slopes, mask *lenslet_mask) { dim3 blockDim(16,16); dim3 gridDim(lenslet_mask->nnz/256+1,1); folded_slopes_kern LLL gridDim, blockDim RRR (d__valid_slopes, lenslet_mask->idx, lenslet_mask->nnz, data_proc.d__c, N_LENSLET,N_WFS); } @ with <>= __global__ void folded_slopes_kern(float *d__valid_slopes, int *idx, const int NNZ, float *d__c, const int N_LENSLET, const int N_WFS) { int i, j, k, iSource, _idx_; i = blockIdx.x * blockDim.x + threadIdx.x; j = blockIdx.y * blockDim.y + threadIdx.y; k = j * gridDim.x * blockDim.x + i; if (k>= #ifndef __CEO_H__ #include "h" #endif <>= <> @ \subsection{Test 1} \label{sec:test-1} In this section, the generation of the imagelets and slopes of a Shack--Hartmann WFS are described. @ The main function is: <>= <
> <> int main(int argc,char *argv[]) { printf(" ##### TEST 1 #####\n"); if (argc != 4) { printf("Usage: %s N_SIDE_LENSLET N_PX_LENSLET N_GS\n",argv[0]); return 1; } cudaSetDevice(0); <> wfs.calibrate(&gs,0.5); atm.get_phase_screen(&gs,N_GS,p,N_PX,p,N_PX,0); gs.wavefront.masked(); wfs.analyze(&gs); gs.wavefront.show_amplitude("Shack-Hartmann/pupil"); gs.wavefront.show_phase("Shack-Hartmann/wavefront"); wfs.camera.show_frame("Shack-Hartmann/frame"); wfs.data_proc.show_centroids("Shack-Hartmann/centroids"); wfs.data_proc.show_flux("Shack-Hartmann/flux"); lmmse.estimation(&wfs.data_proc); atm.get_phase_screen(&gs_dm,N_GS,d,N_ACTUATOR,d,N_ACTUATOR,0); gs_dm.wavefront.masked(&DM_mask); gs_dm.wavefront.show_phase("Shack-Hartmann/wavefront low res."); phase_screen.M = &DM_mask; phase_screen.add_phase(1,lmmse.d__phase_est); //phase_screen.add_phase(-1,gs_dm.wavefront.phase); phase_screen.show_phase("Shack-Hartmann/recon. wavefront"); /* atm.get_phase_screen_gradient(&wfs.data_proc,N_SIDE_LENSLET,d,&gs,N_GS,0); wfs.data_proc.show_centroids("Shack-Hartmann/geometric centroids"); lmmse.estimation(&wfs.data_proc); phase_screen.reset(); phase_screen.add_phase(1,lmmse.d__phase_est); //phase_screen.add_phase(-1,gs_dm.wavefront.phase); phase_screen.show_phase("Shack-Hartmann/geom. recon. wavefront"); */ <> } @ The input parameters are \begin{itemize} \item the size of the lenslet array [[N_SIDE_LENSLET]]$\times$[[N_SIDE_LENSLET]], <>= int N_SIDE_LENSLET, N_ACTUATOR, N_ACTUATOR2; N_SIDE_LENSLET = atoi( argv[1] ); N_ACTUATOR = N_SIDE_LENSLET + 1; N_ACTUATOR2 = N_ACTUATOR*N_ACTUATOR; printf("__ SHACK-HARTMANN __\n"); printf(" . %dX%d lenslets: \n", N_SIDE_LENSLET, N_SIDE_LENSLET); @ \item the number of pixel per lenslet [[N_PX_LENSLET]]$\times$[[N_PX_LENSLET]] <>= int N_PX_LENSLET; N_PX_LENSLET = atoi( argv[2] ); printf(" . %dX%d pixels per lenslet: \n", N_PX_LENSLET, N_PX_LENSLET); @ \item the number of sources, <>= int N_GS; N_GS = atoi( argv[3] ); printf(" . source #: %d\n", N_GS); @ \end{itemize} The total number of pixel is given by <>= int N_PX, N_PX2; N_PX = N_SIDE_LENSLET*N_PX_LENSLET; N_PX2 = N_PX*N_PX; float p, d; p = 25.0/N_PX; d = 25.0/N_SIDE_LENSLET; @ The guide star is defined with <>= float zenith[] = {ARCSEC(30),ARCSEC(30),ARCSEC(30),ARCSEC(30),ARCSEC(30),ARCSEC(30)}, azimuth[] = {0,2.*PI/6.,4.*PI/6.,6.*PI/6.,8.*PI/6.,10.*PI/6.,12.*PI/6.}; source gs; gs.setup("K", zenith, azimuth, INFINITY, N_GS, N_PX2); <>= gs.cleanup(); @ The telescope pupil is defined with the [[mask]] structure: <>= mask telescope; telescope.setup_circular(N_PX); gs.wavefront.masked(&telescope); <>= telescope.cleanup(); @ The wavefront sampled on the actuators location is from the following guide star: <>= source gs_dm; gs_dm.setup("K", zenith, azimuth, INFINITY, N_GS, N_ACTUATOR2); <>= gs_dm.cleanup(); @ The wavefront sensor is defined with: <>= shackHartmann wfs; wfs.setup(N_SIDE_LENSLET,N_PX_LENSLET,d,2,N_PX_LENSLET,1,N_GS); <>= wfs.cleanup(); @ The timing is done with the [[stopwatch]] structure: <>= stopwatch tid; @ The square pupil is made of tiles [[N_PX_LENSLET]]$\times$[[N_PX_LENSLET]] pixels with intensity equal to lenslet index <>= dim3 blockDim(16,16); dim3 gridDim(1+N_PX_LENSLET*N_SIDE_LENSLET/16,1+N_PX_LENSLET*N_SIDE_LENSLET/16,N_GS); squareTiledPupil LLL gridDim,blockDim RRR (gs.wavefront.amplitude, N_PX_LENSLET, N_SIDE_LENSLET, N_GS); gs.wavefront.show_amplitude("Shack-Hartmann/amplitude",N_PX,N_PX*N_GS); <>= __global__ void squareTiledPupil(float* amplitude, const int N_PX_LENSLET, const int N_SIDE_LENSLET, const int N_GS) { int i, j, iLenslet, jLenslet, ij_PUPIL, iSource, kLenslet; iSource = blockIdx.z; if ( threads2lenslet(threadIdx, blockIdx, blockDim, &i, &j, N_PX_LENSLET, &iLenslet, &jLenslet, N_SIDE_LENSLET) ) { ij_PUPIL = lenslet2array(i,j,N_PX_LENSLET,iLenslet,jLenslet,N_SIDE_LENSLET,iSource); kLenslet = iLenslet*N_SIDE_LENSLET + jLenslet; amplitude[ij_PUPIL] = ((float) (kLenslet + 1)) + iSource*0.5; } } @ The atmosphere model is created with: <>= atmosphere atm; atm.gmt_setup(20e-2,30); <>= atm.cleanup(); @ The reconstructed wavefront is saved in: <>= complex_amplitude phase_screen; phase_screen.setup((N_SIDE_LENSLET+1)*(N_SIDE_LENSLET+1)); <>= phase_screen.cleanup(); @ The DM mask: <>= mask DM_mask; DM_mask.setup( (N_SIDE_LENSLET+1)*(N_SIDE_LENSLET+1) ); <>= DM_mask.cleanup(); @ The Fried geometry for a circular pupil with the intensity [[threshold]] is enforced:: <>= float threshold = 0.5; wfs.data_proc.MASK_SET = fried_geometry_setup(wfs.data_proc.lenslet_mask, DM_mask.m, N_SIDE_LENSLET, 16, threshold); @ The filtering properties associated with the pupil are set with: <>= DM_mask.set_filter(); @ The wavefront is reconstructed from the centroids: <>= LMMSE lmmse; lmmse.setup(&atm,&gs,&gs,d,N_SIDE_LENSLET,&DM_mask,"MINRES"); <>= lmmse.cleanup(); @ \subsection{Test 2} \label{sec:test-2} The input parameters are \begin{itemize} \item the size of the lenslet array [[N_SIDE_LENSLET]]$\times$[[N_SIDE_LENSLET]], <>= int N_SIDE_LENSLET, N_ACTUATOR, N_ACTUATOR2; N_SIDE_LENSLET = atoi( argv[1] ); int _N_LENSLET_ = (N_SIDE_LENSLET*N_SIDE_LENSLET); N_ACTUATOR = N_SIDE_LENSLET + 1; N_ACTUATOR2 = N_ACTUATOR*N_ACTUATOR; printf("__ SHACK-HARTMANN __\n"); printf(" . %dX%d lenslets: \n", N_SIDE_LENSLET, N_SIDE_LENSLET); @ \item the number of pixel per lenslet [[N_PX_LENSLET]]$\times$[[N_PX_LENSLET]] <>= int N_PX_LENSLET; N_PX_LENSLET = atoi( argv[2] ); printf(" . %dX%d pixels per lenslet: \n", N_PX_LENSLET, N_PX_LENSLET); @ \item the number of sources, <>= int N_GS; N_GS = atoi( argv[3] ); printf(" . source #: %d\n", N_GS); @ \end{itemize} The total number of pixel is given by <>= int N_PX, N_PX2; N_PX = N_SIDE_LENSLET*N_PX_LENSLET; N_PX2 = N_PX*N_PX; float p, d, D; D = 25.0; p = D/N_PX; d = D/N_SIDE_LENSLET; @ The components common to all the systems are defined first: \begin{itemize} \item the science source, <>= source src; src.setup("J",ARCSEC(0) , 0, INFINITY,(N_SIDE_LENSLET+1)*(N_SIDE_LENSLET+1), "SRC"); <>= src.cleanup(); @ \item the atmosphere, <>= atmosphere atm; //atm.setup(20e-2,30,10e3,10,0); atm.gmt_setup(15e-2,60); /* float altitude[] = {0, 10e3}, xi0[] = {0.5, 0.5}, wind_speed[] = {10, 10}, wind_direction[] = {0, 0}; atm.setup(20e-2,30,altitude, xi0, wind_speed, wind_direction); */ <>= atm.cleanup(); @ \item the pupil mask. <>= //mask pupil_mask; //pupil_mask.setup( (N_SIDE_LENSLET+1)*(N_SIDE_LENSLET+1) ); <>= //pupil_mask.cleanup(); @ \item the wavefront sensor centroid container, <>= //centroiding cog; //cog.setup(N_SIDE_LENSLET,1); <>= //cog.cleanup(); @ \item the diffraction limited science imager, <>= imaging imager; imager.setup(N_SIDE_LENSLET+1,1,4,1,1); <>= imager.cleanup(); @ \item the turbulence limited science imager, <>= imaging imager_turb; imager_turb.setup(N_SIDE_LENSLET+1,1,4,1,1); <>= imager_turb.cleanup(); @ \item the science imager for NGSAO LMMSE, <>= imaging imager_ngsao_lmmse; imager_ngsao_lmmse.setup(N_SIDE_LENSLET+1,1,4,1,1); <>= imager_ngsao_lmmse.cleanup(); @ \item the statistical tool. <>= stats S; S.setup(); <>= S.cleanup(); @ \end{itemize} The natural guide star (NGS) is given by: <>= source ngs; ngs.setup("R",ARCSEC(0) , 0, INFINITY, (N_SIDE_LENSLET+1)*(N_SIDE_LENSLET+1), "NGS"); <>= ngs.cleanup(); @ The wavefront sensor is defined with: <>= shackHartmann wfs; wfs.setup(N_SIDE_LENSLET,N_PX_LENSLET,d,2,N_PX_LENSLET,1,N_GS); wfs.calibrate(&ngs,0.5); <>= wfs.cleanup(); @ The Fried geometry for a circular pupil with the intensity [[threshold]] is enforced:: <>= //float threshold = 0.5; //cog.MASK_SET = fried_geometry_setup(cog.lenslet_mask, pupil_mask.m, // N_SIDE_LENSLET, 16, threshold); @ The filtering properties associated with the pupil are set with: <>= wfs.valid_actuator.set_filter(); //pupil_mask.set_filter(); @ The science is propagated through the [[pupil_mask]] to the focal plane of the [[imager]]: <>= src.wavefront.masked(&wfs.valid_actuator); imager.propagate(&src); //imager.frame2file("refFrame.bin"); char plotly_name[64], plotly_dir[64]; sprintf(plotly_dir,"Shack-Hartmann+LMMSE/D=%.1fm, N=%d (%d samples)/",D,N_SIDE_LENSLET,N_SAMPLE); SET_PLOTLY_NAME(plotly_name,plotly_dir,"frames/diffraction limited"); imager.show_frame(plotly_name); @ A few useful variables are defined here: <>= int NP, NP2, k_SAMPLE; float tau=0.; NP = N_SIDE_LENSLET+1; NP2 = NP*NP; k_SAMPLE = 0; @ The science wavefront is propagated through the atmosphere from [[src]] to [[pupil_mask]]. <>= atm.get_phase_screen(&src,d,NP,d,NP,tau); src.wavefront.masked(); if (k_SAMPLE==(N_SAMPLE-1)) { SET_PLOTLY_NAME(plotly_name,plotly_dir,"phases/science phase screen"); src.wavefront.show_phase(plotly_name); } <>= float science_wf_rms=0.; <>= science_wf_rms += S.var(src.wavefront.phase, &wfs.valid_actuator, NP2); @ and then propagated to the focal plane of the imager: <>= imager_turb.propagate(&src); if (k_SAMPLE==(N_SAMPLE-1)) { SET_PLOTLY_NAME(plotly_name,plotly_dir,"frames/turbulence limited"); imager_turb.show_frame(plotly_name); } @ The turbulence wavefront is saved apart for later use: <>= complex_amplitude phase_screen; phase_screen.setup((N_SIDE_LENSLET+1)*(N_SIDE_LENSLET+1)); phase_screen.add_phase(1,src.wavefront.phase); @ All CEO programs must include the following headers which also contains the headers for all the CEO library modules. <
>= #ifndef __CEO_H__ #include "h" #endif @ The number of atmosphere sample <
>= #define N_SAMPLE 1 #define PLOTLY_LIM (N_SAMPLE-1) // (N_SAMPLE-1) for plotting, higher for disabling plotting @ The main function is: <>= <
> void SET_PLOTLY_NAME(char *name,char *dir,char *path) { strcpy(name, dir); strcat(name,path); } int main(int argc,char *argv[]) { printf(" ##### TEST 2 #####\n"); if (argc != 4) { printf("Usage: %s N_SIDE_LENSLET N_PX_LENSLET N_GS\n",argv[0]); return 1; } cudaSetDevice(0); <> INFO(" Samples: %d:\n",N_SAMPLE); for (k_SAMPLE=0;k_SAMPLE> <> } INFO("\n"); <> <> } @ The NGS wavefront gradients are computed with: <>= atm.get_phase_screen_gradient(&wfs.data_proc,N_SIDE_LENSLET,d,&ngs,tau); @ The wavefront is reconstructed from the NGS centroids: <>= LMMSE ngs_lmmse; ngs_lmmse.setup(&atm,&ngs,&src,&wfs,"MINRES"); int cvgce_iteration=0; float elapsed_time=0.; <>= ngs_lmmse.estimation(&wfs.data_proc); elapsed_time += ngs_lmmse.elapsed_time; cvgce_iteration += ngs_lmmse.iSolve.cvgce_iteration; @ The LMMMSE NGS wavefront estimate is subtracted from the science wavefront: <>= src.wavefront.reset(phase_screen); src.wavefront.add_phase(-1,ngs_lmmse.d__phase_est); if (k_SAMPLE==PLOTLY_LIM) { SET_PLOTLY_NAME(plotly_name,plotly_dir,"phases/LMMSE NGS AO"); src.wavefront.show_phase(plotly_name); } <>= float ngs_wfe_rms_lmmse=0.; <>= ngs_wfe_rms_lmmse += S.var(src.wavefront.phase, &wfs.valid_actuator, NP2); //ngs_lmmse.toFile("phaseEstNgsItp.bin"); <>= ngs_lmmse.cleanup(); @ and the residual wavefront corresponding image is computed. <>= imager_ngsao_lmmse.propagate(&src); //imager.frame2file("ngsLmmseFrame.bin"); if (k_SAMPLE==PLOTLY_LIM) { SET_PLOTLY_NAME(plotly_name,plotly_dir,"frames/LMMSE NGS AO"); imager_ngsao_lmmse.show_frame(plotly_name,&imager); } @ \subsubsection{Results} \label{sec:results} <>= science_wf_rms = 1E9*sqrtf(science_wf_rms/N_SAMPLE); ngs_wfe_rms_lmmse = 1E9*sqrtf(ngs_wfe_rms_lmmse/N_SAMPLE); elapsed_time /= N_SAMPLE; cvgce_iteration /= N_SAMPLE; printf("------------------------------\n"); printf("___ TURBULENCE WAVEFRONT ___\n"); printf("\n NGS WF RMS: %7.2fnm\n", science_wf_rms); printf("\n___ ON-AXIS WAVEFRONT ESTIMATE FROM NGS ___\n"); printf("\n NGS MMSE WFE RMS: %8.3fnm in %.2fms with %d iterations\n", ngs_wfe_rms_lmmse,elapsed_time,cvgce_iteration); printf("------------------------------\n"); @ \subsection{Python} \label{sec:python-1} <>= #!/usr/bin/env python import sys sys.path.append("/home/rconan/CEO/python") from pylab import * from import * nLenslet = 20 D = 10.0 n = 7 nPx = n*nLenslet nn = array([nPx,nPx],dtype=int32) nSrc = 1 zen = zeros( nSrc, dtype=float32) azi = zeros( nSrc, dtype=float32) src = Source("K",zen,azi,float("inf"),nn,fwhm=2) tel = Telescope(n*nLenslet) src.masked(tel) wfs = ShackHartmann(nLenslet, n, D/nLenslet, 2, n, 1, nSrc) wfs.calibrate(src,0.5) wfs.analyze(src) frame0 = zeros( (nPx*nSrc,nPx), order="c", dtype=float32) wfs.getframe(frame0) figure(1) ax = imshow(frame0,interpolation='None') colorbar(ax) draw() atm = GmtAtmosphere(15e-2,30) p = D/nPx atm.get_phase_screen(src, nSrc, p, nPx, p, nPx, 0.0) src.masked(tel) ps = zeros( (nPx*nSrc,nPx), order="c", dtype=float32) src.getphase(ps) figure(2) subplot(1, 3, 1) ax = imshow(ps,interpolation='None') colorbar(ax) wfs.reset() wfs.analyze(src) wfs.getframe(frame0) subplot(1, 3, 2) ax = imshow(frame0,interpolation='None') colorbar(ax) c = zeros( nSrc*2*nLenslet**2, order="c", dtype=float32 ) wfs.getc(c) subplot(1, 3, 3) ax = imshow(reshape(c,[nSrc*nLenslet*2,nLenslet]),interpolation='None') colorbar(ax) show()