diff options
| -rw-r--r-- | data/diffraction.cl | 10 | ||||
| -rw-r--r-- | src/diffraction-gpu.c | 20 |
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), ¢x); + 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), ¢y); + 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; |