aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorThomas White <taw@physics.org>2010-02-19 18:07:44 +0100
committerThomas White <taw@physics.org>2010-02-19 18:07:44 +0100
commit3542ac97a332f2657c7a69b93bf1133ca81967cb (patch)
treeb88c84c54df1b0f38681a2d0169186b6378a44a5
parentab3573ec308bd8d4e4790ab2c7e54ad4b6b81710 (diff)
Simplify sampling a bit - preventing some indexing issues
-rw-r--r--data/diffraction.cl10
-rw-r--r--src/diffraction-gpu.c20
2 files changed, 15 insertions, 15 deletions
diff --git a/data/diffraction.cl b/data/diffraction.cl
index babfc6b7..138af028 100644
--- a/data/diffraction.cl
+++ b/data/diffraction.cl
@@ -52,8 +52,8 @@ float4 get_q(int x, int y, float cx, float cy, float res, float clen, float k,
float ttx, tty, tt;
float4 q;
- rx = ((float)x - cx)/res;
- ry = ((float)y - cy)/res;
+ rx = ((float)x - sampling*cx)/(res*sampling);
+ ry = ((float)y - sampling*cy)/(res*sampling);
r = sqrt(pow(rx, 2.0) + pow(ry, 2.0));
@@ -134,8 +134,8 @@ kernel void diffraction(global float2 *diff, global float *tt, float k,
int xmin, int ymin, int sampling, local float2 *tmp)
{
float ttv;
- const int x = get_global_id(0) + xmin;
- const int y = get_global_id(1) + ymin;
+ const int x = get_global_id(0) + (xmin*sampling);
+ const int y = get_global_id(1) + (ymin*sampling);
float f_lattice;
float2 f_molecule;
float4 q;
@@ -165,7 +165,7 @@ kernel void diffraction(global float2 *diff, global float *tt, float k,
diff[ax+w*ay] = sum / (sampling*sampling);
- /* Leader thread also records 2theta value.
+ /* Leader thread also records the 2theta value.
* This should really be averaged across all pixels, but
* I strongly suspect this would be a waste of time. */
tt[ax+w*ay] = ttv;
diff --git a/src/diffraction-gpu.c b/src/diffraction-gpu.c
index 9057d3cc..330a676e 100644
--- a/src/diffraction-gpu.c
+++ b/src/diffraction-gpu.c
@@ -240,11 +240,6 @@ void get_diffraction_gpu(struct gpu_context *gctx, struct image *image,
size_t dims[2];
size_t ldims[2] = {SAMPLING, SAMPLING};
- cl_float res = image->det.panels[p].res * SAMPLING;
- int offsx = image->det.panels[p].min_x * SAMPLING;
- int offsy = image->det.panels[p].min_y * SAMPLING;
- cl_float centx = image->det.panels[p].cx * SAMPLING;
- cl_float centy = image->det.panels[p].cy * SAMPLING;
/* In a future version of OpenCL, this could be done
* with a global work offset. But not yet... */
@@ -253,17 +248,20 @@ void get_diffraction_gpu(struct gpu_context *gctx, struct image *image,
dims[0] *= SAMPLING;
dims[1] *= SAMPLING;
- clSetKernelArg(gctx->kern, 4, sizeof(cl_float), &centx);
+ clSetKernelArg(gctx->kern, 4, sizeof(cl_float),
+ &image->det.panels[p].cx);
if ( err != CL_SUCCESS ) {
ERROR("Couldn't set arg 4: %s\n", clError(err));
return;
}
- clSetKernelArg(gctx->kern, 5, sizeof(cl_float), &centy);
+ clSetKernelArg(gctx->kern, 5, sizeof(cl_float),
+ &image->det.panels[p].cy);
if ( err != CL_SUCCESS ) {
ERROR("Couldn't set arg 5: %s\n", clError(err));
return;
}
- clSetKernelArg(gctx->kern, 6, sizeof(cl_float), &res);
+ clSetKernelArg(gctx->kern, 6, sizeof(cl_float),
+ &image->det.panels[p].res);
if ( err != CL_SUCCESS ) {
ERROR("Couldn't set arg 6: %s\n", clError(err));
return;
@@ -274,12 +272,14 @@ void get_diffraction_gpu(struct gpu_context *gctx, struct image *image,
ERROR("Couldn't set arg 7: %s\n", clError(err));
return;
}
- clSetKernelArg(gctx->kern, 12, sizeof(cl_int), &offsx);
+ clSetKernelArg(gctx->kern, 12, sizeof(cl_int),
+ &image->det.panels[p].min_x);
if ( err != CL_SUCCESS ) {
ERROR("Couldn't set arg 12: %s\n", clError(err));
return;
}
- clSetKernelArg(gctx->kern, 13, sizeof(cl_int), &offsy);
+ clSetKernelArg(gctx->kern, 13, sizeof(cl_int),
+ &image->det.panels[p].min_y);
if ( err != CL_SUCCESS ) {
ERROR("Couldn't set arg 13: %s\n", clError(err));
return;