aboutsummaryrefslogtreecommitdiff
path: root/src/diffraction-gpu.c
diff options
context:
space:
mode:
authorThomas White <taw@physics.org>2010-02-19 14:49:49 +0100
committerThomas White <taw@physics.org>2010-02-19 14:49:49 +0100
commit803fbcf98c0c68cc5000331d6e00f9297a1cf5f0 (patch)
treee8c3d49acc9542932af90bdf87c72490d38dd21a /src/diffraction-gpu.c
parente3d41026e188d6225c16a86accab05a6897977f6 (diff)
Honour particle size in GPU calculation
Diffstat (limited to 'src/diffraction-gpu.c')
-rw-r--r--src/diffraction-gpu.c11
1 files changed, 11 insertions, 0 deletions
diff --git a/src/diffraction-gpu.c b/src/diffraction-gpu.c
index 26aa7e34..5ffb576c 100644
--- a/src/diffraction-gpu.c
+++ b/src/diffraction-gpu.c
@@ -157,6 +157,7 @@ void get_diffraction_gpu(struct image *image, int na, int nb, int nc,
float *diff_ptr;
int i;
cl_float4 orientation;
+ cl_int4 ncells;
if ( image->molecule == NULL ) return;
@@ -251,6 +252,11 @@ void get_diffraction_gpu(struct image *image, int na, int nb, int nc,
orientation[2] = image->orientation.y;
orientation[3] = image->orientation.z;
+ ncells[0] = na;
+ ncells[1] = nb;
+ ncells[2] = nc;
+ ncells[3] = 0; /* unused */
+
err = clSetKernelArg(kern, 0, sizeof(cl_mem), &diff);
if ( err != CL_SUCCESS ) {
ERROR("Couldn't set arg 0: %s\n", clError(err));
@@ -306,6 +312,11 @@ void get_diffraction_gpu(struct image *image, int na, int nb, int nc,
ERROR("Couldn't set arg 10: %s\n", clError(err));
return;
}
+ clSetKernelArg(kern, 11, sizeof(cl_int4), &ncells);
+ if ( err != CL_SUCCESS ) {
+ ERROR("Couldn't set arg 11: %s\n", clError(err));
+ return;
+ }
err = clEnqueueNDRangeKernel(cq, kern, 2, NULL, dims, NULL,
0, NULL, &event_d);