fixed filtering

ffthelp.cl

@@ -58,7 +58,8 @@ kernel void fftshift(int Cwidth, int Cheight, __global float* Celements, __globa
 	}
 }
 
-kernel void convolutionKernelInterleaved(int Cwidth, int Cheight, __global float* Celements, __global float *output, __global float* filter, int numCoeff)
+kernel void convolutionKernelInterleaved(int Cwidth, int Cheight, __global float* Celements, 
+										 __global float *output, __global float* filter, __const int numCoeff)
 {
 
 	int row = get_global_id(0);

src/prog.c

@@ -382,6 +382,7 @@ tranformFromKSpace (Matrix * const A, Matrix *result,
 
     clFFT_DestroyPlan (fftPlan);
     OCL_CHECK_ERROR (clReleaseEvent (event));
+    OCL_CHECK_ERROR (clReleaseMemObject (tempBuffer));
     OCL_CHECK_ERROR (clReleaseMemObject (d_A.elements));
     OCL_CHECK_ERROR (clReleaseMemObject (d_C.elements));
 }
@@ -400,12 +401,10 @@ filterData (int decimFactor, int isSmooth, Matrix * const input, Matrix *output,
     unsigned long wait;
     unsigned long execution;
 
-    GTimer *timer;
-    timer = g_timer_new ();
-    double wallClock = 0.0;
 
     float *filter;
     int filterlength = 0;
+    cl_int d_filterlength = filterlength;
     readFilterFromFile (isSmooth, decimFactor, filter, &filterlength);
 
     int length = input->width * input->height;
@@ -414,6 +413,10 @@ filterData (int decimFactor, int isSmooth, Matrix * const input, Matrix *output,
                                    length * sizeof(float) * 2, NULL, &errcode);
     OCL_CHECK_ERROR (errcode);
 
+    cl_mem d_filter = clCreateBuffer (context, CL_MEM_READ_ONLY,
+                                   filterlength * sizeof(float), NULL, &errcode);
+    OCL_CHECK_ERROR (errcode);
+
     GPUMatrix d_C;
     d_C.width = input->width;
     d_C.height = input->height;
@@ -422,6 +425,9 @@ filterData (int decimFactor, int isSmooth, Matrix * const input, Matrix *output,
 
     OCL_CHECK_ERROR (errcode);
 
+    output->width = d_C.width;
+    output->height = d_C.height;
+    output->elements = (clFFT_Complex *) malloc(sizeof(clFFT_Complex)*length);
 
     size_t globalWorkSize[] = { input->height, input->width };
     cl_uint i = 0;
@@ -433,27 +439,26 @@ filterData (int decimFactor, int isSmooth, Matrix * const input, Matrix *output,
                     (void*) &d_C.elements));
     OCL_CHECK_ERROR (clSetKernelArg(convKernel, i++, sizeof (tempBuffer), 
                     (void*) &tempBuffer));
-    OCL_CHECK_ERROR (clSetKernelArg(convKernel, i++, sizeof (filter), 
-                    (void*) &filter));
-    OCL_CHECK_ERROR (clSetKernelArg(convKernel, i++, sizeof (filterlength), 
-                    (void*) &filterlength));
+    OCL_CHECK_ERROR (clSetKernelArg(convKernel, i++, sizeof (d_filter), 
+                    (void*) &d_filter));
+    OCL_CHECK_ERROR (clSetKernelArg(convKernel, i++, sizeof (d_filterlength), 
+                    (void*) &d_filterlength));
 
-    g_timer_start (timer);
+   
     OCL_CHECK_ERROR (clEnqueueNDRangeKernel (queue, convKernel, 2, NULL, 
                             globalWorkSize, NULL, 
                             0, NULL, &event));
 
     clWaitForEvents (1, &event);
-    g_timer_stop (timer);
-    get_event_times (event, &wait, &execution);
-    printf("\nwaited for device %lu us:\n time to execute: %lu us\n", wait, execution);
 
-    wallClock = g_timer_elapsed (timer, NULL);
+    errcode = clEnqueueReadBuffer (queue, tempBuffer, CL_TRUE, 0,
+                                   length * sizeof(float) * 2, output->elements,
+                                   0, NULL, &event);
 
-    printf("took %e on host\n", wallClock);
+    OCL_CHECK_ERROR (errcode);
 
-    g_timer_destroy (timer);
-    OCL_CHECK_ERROR (clReleaseEvent (event));
+    OCL_CHECK_ERROR (clReleaseMemObject (tempBuffer));
+    OCL_CHECK_ERROR (clReleaseMemObject (d_C.elements));
 
 
 }
@@ -491,7 +496,7 @@ main (int argc, const char **argv)
     OCL_CHECK_ERROR (errcode);
 
     //create Matrices
-    Matrix kSpace, decimkSpace, absimg;
+    Matrix k_Space, decimk_Space, absimg;
     Matrix refImg;
     Matrix sqD;
 
@@ -502,27 +507,26 @@ main (int argc, const char **argv)
         exit (1);
     }
 
-    getDataFromSplitBinary (bin, &kSpace);
+    getDataFromSplitBinary (bin, &k_Space);
     fclose (bin);
 
-    filterData (8, 1, &kSpace, &decimkSpace, 
+    filterData (8, 1, &k_Space, &decimk_Space, 
                         context, convolutionKernelInterleaved, queues[0]);
 
-   /* tranformFromKSpace (&kSpace, &absimg, MRI_CHANNELS, context,
+   /* tranformFromKSpace (&k_Space, &absimg, MRI_CHANNELS, context,
                         fftshiftKernel, absKernel, queues[0]);*/
 
     MatrixSplit absimgSplit;
-    convertInterleavedToSplitMatrix (&absimg, &absimgSplit);
+    convertInterleavedToSplitMatrix (&k_Space, &absimgSplit);
 
-    FILE *f = fopen("imgocl", "wb");
+    FILE *f = fopen("imgocl.bin", "wb");
     if (f == NULL)
     {
         printf("Error opening file!\n");
         exit(1);
     }
 
-    writeSplitMatrixToBinaryFile (f, &kSpace);
-    writeSplitMatrixToBinaryFile (f, &decimkSpace);
+    writeSplitMatrixToBinaryFile (f, &absimgSplit);
     fclose (f);
 
     ocl_free (ocl);