From e1a9cd72af84ff153edd6d230ed63e2194dcb518 Mon Sep 17 00:00:00 2001 From: Thomas White Date: Mon, 20 Feb 2012 16:03:35 +0100 Subject: Remove all bandwidth and divergence stuff, fix pattern_sim Bandwidth and divergence didn't work very well --- src/diffraction-gpu.c | 72 +++++++++----------------------------- src/diffraction.c | 97 ++++++++++++--------------------------------------- src/pattern_sim.h | 2 +- 3 files changed, 39 insertions(+), 132 deletions(-) (limited to 'src') diff --git a/src/diffraction-gpu.c b/src/diffraction-gpu.c index b3950570..464adf02 100644 --- a/src/diffraction-gpu.c +++ b/src/diffraction-gpu.c @@ -34,9 +34,6 @@ #include "pattern_sim.h" -#define SAMPLING (4) -#define BWSAMPLING (10) -#define DIVSAMPLING (1) #define SINC_LUT_ELEMENTS (4096) @@ -160,18 +157,12 @@ void get_diffraction_gpu(struct gpu_context *gctx, struct image *image, double ax, ay, az; double bx, by, bz; double cx, cy, cz; - float klow, khigh; int i; cl_float16 cell; cl_int4 ncells; - const int sampling = SAMPLING; - cl_float bwstep; int n_inf = 0; int n_neg = 0; - cl_float divxlow, divxstep; - cl_float divylow, divystep; int n_nan = 0; - int sprod; if ( gctx == NULL ) { ERROR("GPU setup failed.\n"); @@ -183,17 +174,6 @@ void get_diffraction_gpu(struct gpu_context *gctx, struct image *image, cell.s[3] = bx; cell.s[4] = by; cell.s[5] = bz; cell.s[6] = cx; cell.s[7] = cy; cell.s[8] = cz; - /* Calculate wavelength */ - klow = 1.0/(image->lambda*(1.0 + image->beam->bandwidth/2.0)); - khigh = 1.0/(image->lambda*(1.0 - image->beam->bandwidth/2.0)); - bwstep = (khigh-klow) / BWSAMPLING; - - /* Calculate divergence stuff */ - divxlow = -image->beam->divergence/2.0; - divylow = -image->beam->divergence/2.0; - divxstep = image->beam->divergence / DIVSAMPLING; - divystep = image->beam->divergence / DIVSAMPLING; - ncells.s[0] = na; ncells.s[1] = nb; ncells.s[2] = nc; @@ -204,20 +184,12 @@ void get_diffraction_gpu(struct gpu_context *gctx, struct image *image, check_sinc_lut(gctx, nb); check_sinc_lut(gctx, nc); - if ( set_arg_float(gctx, 2, klow) ) return; + if ( set_arg_float(gctx, 2, 1.0/image->lambda) ) return; if ( set_arg_mem(gctx, 9, gctx->intensities) ) return; - if ( set_arg_int(gctx, 12, sampling) ) return; - if ( set_arg_float(gctx, 14, bwstep) ) return; - if ( set_arg_mem(gctx, 15, gctx->sinc_luts[na-1]) ) return; - if ( set_arg_mem(gctx, 16, gctx->sinc_luts[nb-1]) ) return; - if ( set_arg_mem(gctx, 17, gctx->sinc_luts[nc-1]) ) return; - if ( set_arg_mem(gctx, 18, gctx->flags) ) return; - if ( set_arg_float(gctx, 23, divxlow) ) return; - if ( set_arg_float(gctx, 24, divxstep) ) return; - if ( set_arg_int(gctx, 25, DIVSAMPLING) ) return; - if ( set_arg_float(gctx, 26, divylow) ) return; - if ( set_arg_float(gctx, 27, divystep) ) return; - if ( set_arg_int(gctx, 28, DIVSAMPLING) ) return; + if ( set_arg_mem(gctx, 10, gctx->sinc_luts[na-1]) ) return; + if ( set_arg_mem(gctx, 11, gctx->sinc_luts[nb-1]) ) return; + if ( set_arg_mem(gctx, 12, gctx->sinc_luts[nc-1]) ) return; + if ( set_arg_mem(gctx, 13, gctx->flags) ) return; /* Unit cell */ err = clSetKernelArg(gctx->kern, 8, sizeof(cl_float16), &cell); @@ -226,14 +198,6 @@ void get_diffraction_gpu(struct gpu_context *gctx, struct image *image, return; } - /* Local memory for reduction */ - sprod = BWSAMPLING*SAMPLING*SAMPLING*DIVSAMPLING*DIVSAMPLING; - err = clSetKernelArg(gctx->kern, 13, sprod*sizeof(cl_float), NULL); - if ( err != CL_SUCCESS ) { - ERROR("Couldn't set local memory: %s\n", clError(err)); - return; - } - /* Allocate memory for the result */ image->data = calloc(image->width * image->height, sizeof(float)); image->twotheta = calloc(image->width * image->height, sizeof(double)); @@ -241,9 +205,8 @@ void get_diffraction_gpu(struct gpu_context *gctx, struct image *image, /* Iterate over panels */ for ( i=0; idet->n_panels; i++ ) { - size_t dims[3]; - size_t ldims[3] = {SAMPLING, SAMPLING, - BWSAMPLING * DIVSAMPLING * DIVSAMPLING}; + size_t dims[2]; + size_t ldims[2] = {1, 1}; struct panel *p; cl_mem tt; size_t tt_size; @@ -282,18 +245,15 @@ void get_diffraction_gpu(struct gpu_context *gctx, struct image *image, if ( set_arg_float(gctx, 5, p->cny) ) return; if ( set_arg_float(gctx, 6, p->res) ) return; if ( set_arg_float(gctx, 7, p->clen) ) return; - if ( set_arg_int(gctx, 10, p->min_fs) ) return; - if ( set_arg_int(gctx, 11, p->min_ss) ) return; - if ( set_arg_float(gctx, 19, p->fsx) ) return; - if ( set_arg_float(gctx, 20, p->fsy) ) return; - if ( set_arg_float(gctx, 21, p->ssx) ) return; - if ( set_arg_float(gctx, 22, p->ssy) ) return; - - dims[0] = pan_width * SAMPLING; - dims[1] = pan_height * SAMPLING; - dims[2] = BWSAMPLING * DIVSAMPLING * DIVSAMPLING; - - err = clEnqueueNDRangeKernel(gctx->cq, gctx->kern, 3, NULL, + if ( set_arg_float(gctx, 14, p->fsx) ) return; + if ( set_arg_float(gctx, 15, p->fsy) ) return; + if ( set_arg_float(gctx, 16, p->ssx) ) return; + if ( set_arg_float(gctx, 17, p->ssy) ) return; + + dims[0] = pan_width; + dims[1] = pan_height; + + err = clEnqueueNDRangeKernel(gctx->cq, gctx->kern, 2, NULL, dims, ldims, 0, NULL, NULL); if ( err != CL_SUCCESS ) { ERROR("Couldn't enqueue diffraction kernel: %s\n", diff --git a/src/diffraction.c b/src/diffraction.c index de994133..a37c2fcb 100644 --- a/src/diffraction.c +++ b/src/diffraction.c @@ -27,9 +27,6 @@ #include "pattern_sim.h" -#define SAMPLING (4) -#define BWSAMPLING (10) -#define DIVSAMPLING (1) #define SINC_LUT_ELEMENTS (4096) @@ -355,11 +352,9 @@ void get_diffraction(struct image *image, int na, int nb, int nc, double ax, ay, az; double bx, by, bz; double cx, cy, cz; - float klow, khigh, bwstep; double *lut_a; double *lut_b; double *lut_c; - double divxlow, divylow, divxstep, divystep; cell_get_cartesian(cell, &ax, &ay, &az, &bx, &by, &bz, &cx, &cy, &cz); @@ -369,15 +364,6 @@ void get_diffraction(struct image *image, int na, int nb, int nc, /* Needed later for Lorentz calculation */ image->twotheta = malloc(image->width * image->height * sizeof(double)); - klow = 1.0/(image->lambda*(1.0 + image->beam->bandwidth/2.0)); - khigh = 1.0/(image->lambda*(1.0 - image->beam->bandwidth/2.0)); - bwstep = (khigh-klow) / BWSAMPLING; - - divxlow = -image->beam->divergence/2.0; - divylow = -image->beam->divergence/2.0; - divxstep = image->beam->divergence / DIVSAMPLING; - divystep = image->beam->divergence / DIVSAMPLING; - lut_a = get_sinc_lut(na); lut_b = get_sinc_lut(nb); lut_c = get_sinc_lut(nc); @@ -385,73 +371,34 @@ void get_diffraction(struct image *image, int na, int nb, int nc, for ( fs=0; fswidth; fs++ ) { for ( ss=0; ssheight; ss++ ) { - int fs_step, ss_step, kstep; - int divxval, divyval; - int idx = fs + image->width*ss; - - for ( fs_step=0; fs_steplambda; - /* y divergence */ - q.v = qn.v*cos(ydiv) +qn.w*sin(ydiv); - q.w = -qn.v*sin(ydiv) +qn.w*cos(ydiv); + q = get_q(image, fs, ss, &twotheta, k); - f_lattice = lattice_factor(q, ax, ay, az, - bx, by, bz, - cx, cy, cz, - lut_a, lut_b, lut_c); + f_lattice = lattice_factor(q, ax, ay, az, + bx, by, bz, + cx, cy, cz, + lut_a, lut_b, lut_c); - I_molecule = molecule_factor(intensities, - phases, flags, q, - ax,ay,az,bx,by,bz,cx,cy,cz, - m, sym); + I_molecule = molecule_factor(intensities, + phases, flags, q, + ax,ay,az,bx,by,bz,cx,cy,cz, + m, sym); - I_lattice = pow(f_lattice, 2.0); - intensity = I_lattice * I_molecule; - - image->data[idx] += intensity; - - if ( fs_step + ss_step + kstep == 0 ) { - image->twotheta[idx] = twotheta; - } - - } - } - } - } - } + I_lattice = pow(f_lattice, 2.0); - image->data[idx] /= (SAMPLING*SAMPLING*BWSAMPLING - *DIVSAMPLING*DIVSAMPLING); + idx = fs + image->width*ss; + image->data[idx] = I_lattice * I_molecule; + image->twotheta[idx] = twotheta; } diff --git a/src/pattern_sim.h b/src/pattern_sim.h index fb376596..5d31dadf 100644 --- a/src/pattern_sim.h +++ b/src/pattern_sim.h @@ -20,7 +20,7 @@ /* Maxmimum index to hold values up to (can be increased if necessary) * WARNING: Altering this value constitutes an ABI change, and means you must * update data/diffraction.cl then recompile and reinstall everything. */ -#define INDMAX 200 +#define INDMAX 120 /* Array size */ #define IDIM (INDMAX*2 +1) -- cgit v1.2.3