From 01fdad29f7d9f126290cdc455c0a587d5aa51007 Mon Sep 17 00:00:00 2001
From: Thomas White <taw@physics.org>
Date: Mon, 6 Jun 2011 15:47:17 +0200
Subject: Add divergence to calculation

---
 src/diffraction-gpu.c | 27 ++++++++++++++++++++++-----
 1 file changed, 22 insertions(+), 5 deletions(-)

(limited to 'src/diffraction-gpu.c')

diff --git a/src/diffraction-gpu.c b/src/diffraction-gpu.c
index c0814e79..74e2292a 100644
--- a/src/diffraction-gpu.c
+++ b/src/diffraction-gpu.c
@@ -34,7 +34,8 @@
 
 
 #define SAMPLING (4)
-#define BWSAMPLING (10)
+#define BWSAMPLING (1)
+#define DIVSAMPLING (4)
 #define SINC_LUT_ELEMENTS (4096)
 
 
@@ -166,7 +167,10 @@ void get_diffraction_gpu(struct gpu_context *gctx, struct image *image,
 	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,6 +187,12 @@ void get_diffraction_gpu(struct gpu_context *gctx, struct image *image,
 	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;
@@ -201,6 +211,12 @@ void get_diffraction_gpu(struct gpu_context *gctx, struct image *image,
 	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;
 
 	/* Unit cell */
 	err = clSetKernelArg(gctx->kern, 8, sizeof(cl_float16), &cell);
@@ -210,8 +226,8 @@ void get_diffraction_gpu(struct gpu_context *gctx, struct image *image,
 	}
 
 	/* Local memory for reduction */
-	err = clSetKernelArg(gctx->kern, 13,
-	               BWSAMPLING*SAMPLING*SAMPLING*sizeof(cl_float), NULL);
+	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;
@@ -225,7 +241,8 @@ void get_diffraction_gpu(struct gpu_context *gctx, struct image *image,
 	for ( i=0; i<image->det->n_panels; i++ ) {
 
 		size_t dims[3];
-		size_t ldims[3] = {SAMPLING, SAMPLING, BWSAMPLING};
+		size_t ldims[3] = {SAMPLING, SAMPLING,
+		                   BWSAMPLING * DIVSAMPLING * DIVSAMPLING};
 		struct panel *p;
 		cl_mem tt;
 		size_t tt_size;
@@ -273,7 +290,7 @@ void get_diffraction_gpu(struct gpu_context *gctx, struct image *image,
 
 		dims[0] = pan_width * SAMPLING;
 		dims[1] = pan_height * SAMPLING;
-		dims[2] = BWSAMPLING;
+		dims[2] = BWSAMPLING * DIVSAMPLING * DIVSAMPLING;
 
 		err = clEnqueueNDRangeKernel(gctx->cq, gctx->kern, 3, NULL,
 		                             dims, ldims, 0, NULL, NULL);
-- 
cgit v1.2.3