aboutsummaryrefslogtreecommitdiff
path: root/src/diffraction-gpu.c
diff options
context:
space:
mode:
authorThomas White <taw@physics.org>2010-02-18 12:10:15 +0100
committerThomas White <taw@physics.org>2010-02-18 12:10:15 +0100
commit1cbd38bd817f4b73403e348fa715c6c5356be98a (patch)
treef5c5b7a3602eedbf7c76ae7efdaa50e7595e01eb /src/diffraction-gpu.c
parentb04ea39a8bf04b12765b73da3ff9eda7f2143bde (diff)
CL fixes
Diffstat (limited to 'src/diffraction-gpu.c')
-rw-r--r--src/diffraction-gpu.c53
1 files changed, 36 insertions, 17 deletions
diff --git a/src/diffraction-gpu.c b/src/diffraction-gpu.c
index a6db1980..89655601 100644
--- a/src/diffraction-gpu.c
+++ b/src/diffraction-gpu.c
@@ -29,6 +29,33 @@
#define BANDWIDTH (1.0 / 100.0)
+static const char *clstrerr(cl_int err)
+{
+ switch ( err ) {
+ case CL_SUCCESS : return "no error";
+ case CL_INVALID_PLATFORM : return "invalid platform";
+ case CL_INVALID_KERNEL : return "invalid kernel";
+ case CL_INVALID_ARG_INDEX : return "invalid argument index";
+ case CL_INVALID_ARG_VALUE : return "invalid argument value";
+ case CL_INVALID_MEM_OBJECT : return "invalid memory object";
+ case CL_INVALID_SAMPLER : return "invalid sampler";
+ case CL_INVALID_ARG_SIZE : return "invalid argument size";
+ case CL_INVALID_COMMAND_QUEUE : return "invalid command queue";
+ case CL_INVALID_CONTEXT : return "invalid context";
+ case CL_INVALID_VALUE : return "invalid value";
+ case CL_INVALID_EVENT_WAIT_LIST : return "invalid wait list";
+ case CL_MAP_FAILURE : return "map failure";
+ case CL_MEM_OBJECT_ALLOCATION_FAILURE : return "object allocation failure";
+ case CL_OUT_OF_HOST_MEMORY : return "out of host memory";
+ case CL_OUT_OF_RESOURCES : return "out of resources";
+ case CL_INVALID_KERNEL_NAME : return "invalid kernel name";
+ default :
+ ERROR("Error code: %i\n", err);
+ return "unknown error";
+ }
+}
+
+
static cl_device_id get_first_dev(cl_context ctx)
{
cl_device_id dev;
@@ -86,7 +113,7 @@ static cl_program load_program(const char *filename, cl_context ctx,
r = clBuildProgram(prog, 0, NULL, "-Werror", NULL, NULL);
if ( r != CL_SUCCESS ) {
- ERROR("Couldn't build program\n");
+ ERROR("Couldn't build program '%s'\n", filename);
show_build_log(prog, dev);
*err = r;
return 0;
@@ -168,14 +195,6 @@ void get_diffraction_gpu(struct image *image, int na, int nb, int nc,
ctx = clCreateContextFromType(prop, CL_DEVICE_TYPE_GPU, NULL, NULL, &err);
if ( err != CL_SUCCESS ) {
ERROR("Couldn't create OpenCL context: %i\n", err);
- switch ( err ) {
- case CL_INVALID_PLATFORM :
- ERROR("Invalid platform\n");
- break;
- case CL_OUT_OF_HOST_MEMORY :
- ERROR("Out of memory\n");
- break;
- }
return;
}
@@ -189,7 +208,7 @@ void get_diffraction_gpu(struct image *image, int na, int nb, int nc,
/* Create buffer for the picture */
diff_size = image->width*image->height*sizeof(cl_float)*2; /* complex */
- diff = clCreateBuffer(ctx, CL_MEM_READ_WRITE, diff_size, NULL, &err);
+ diff = clCreateBuffer(ctx, CL_MEM_WRITE_ONLY, diff_size, NULL, &err);
if ( err != CL_SUCCESS ) {
ERROR("Couldn't allocate diffraction memory\n");
return;
@@ -202,7 +221,7 @@ void get_diffraction_gpu(struct image *image, int na, int nb, int nc,
sfac_ptr[2*i+0] = creal(image->molecule->reflections[i]);
sfac_ptr[2*i+1] = cimag(image->molecule->reflections[i]);
}
- sfacs = clCreateBuffer(ctx, CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR,
+ sfacs = clCreateBuffer(ctx, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
sfac_size, sfac_ptr, &err);
if ( err != CL_SUCCESS ) {
ERROR("Couldn't allocate sfac memory\n");
@@ -210,7 +229,7 @@ void get_diffraction_gpu(struct image *image, int na, int nb, int nc,
}
tt_size = image->width*image->height*sizeof(cl_float);
- tt = clCreateBuffer(ctx, CL_MEM_READ_WRITE, tt_size, NULL, &err);
+ tt = clCreateBuffer(ctx, CL_MEM_WRITE_ONLY, tt_size, NULL, &err);
if ( err != CL_SUCCESS ) {
ERROR("Couldn't allocate twotheta memory\n");
return;
@@ -249,20 +268,20 @@ void get_diffraction_gpu(struct image *image, int na, int nb, int nc,
clSetKernelArg(kern, 10, sizeof(cl_float4), &orientation);
err = clEnqueueNDRangeKernel(cq, kern, 2, NULL, dims, NULL,
- 0, NULL, NULL);
+ 1, &event_q, &event_d);
if ( err != CL_SUCCESS ) {
- ERROR("Couldn't enqueue kernel\n");
+ ERROR("Couldn't enqueue diffraction kernel\n");
return;
}
diff_ptr = clEnqueueMapBuffer(cq, diff, CL_TRUE, CL_MAP_READ, 0,
- diff_size, 0, NULL, NULL, &err);
+ diff_size, 1, &event_d, NULL, &err);
if ( err != CL_SUCCESS ) {
- ERROR("Couldn't map sfac buffer\n");
+ ERROR("Couldn't map diffraction buffer: %s\n", clstrerr(err));
return;
}
tt_ptr = clEnqueueMapBuffer(cq, tt, CL_TRUE, CL_MAP_READ, 0,
- tt_size, 0, NULL, NULL, &err);
+ tt_size, 1, &event_d, NULL, &err);
if ( err != CL_SUCCESS ) {
ERROR("Couldn't map tt buffer\n");
return;