Attachment 131090 Details for Bug 100639 – patch for clFFT to identify wrong execution of OpenCL 1.2 on Haswell

patch for clFFT to identify wrong execution of OpenCL 1.2 on Haswell

patch_clfft_haswell.patch (text/plain), 8.04 KB, created by Michal on 2017-04-27 11:52:36 UTC

(hide)

Description:

Filename:

MIME Type:

Creator: Michal

Created: 2017-04-27 11:52:36 UTC

Size: 8.04 KB

patch

obsolete

>diff -ru clFFT_orig/src/examples/fft1d.c clFFT_custom/src/examples/fft1d.c
>--- clFFT_orig/src/examples/fft1d.c	2017-04-27 11:55:47.722280825 +0200
>+++ clFFT_custom/src/examples/fft1d.c	2017-04-27 11:46:44.134309954 +0200
>@@ -20,6 +20,9 @@
> /* No need to explicitely include the OpenCL headers */
> #include <clFFT.h>
> 
>+#include <math.h>
>+#include <sys/time.h>
>+
> int main( void )
> {
>     cl_int err;
>@@ -30,9 +33,11 @@
>     cl_command_queue queue = 0;
>     cl_mem bufX;
>     float *X;
>+    float *X_result;
>     cl_event event = NULL;
>     int ret = 0;
>-    size_t N = 16;
>+    //size_t N = 4*1024*1024; // 16
>+    size_t N = 8*1024; // 16
>     char platform_name[128];
>     char device_name[128];
> 
>@@ -43,35 +48,48 @@
> 
>     /* Setup OpenCL environment. */
>     err = clGetPlatformIDs( 1, &platform, NULL );
>+    fprintf( stderr, "(err: %d)after clGetPlatformIDs\n", err );
> 
>     size_t ret_param_size = 0;
>     err = clGetPlatformInfo(platform, CL_PLATFORM_NAME,
>             sizeof(platform_name), platform_name,
>             &ret_param_size);
>-    printf("Platform found: %s\n", platform_name);
>+    fprintf( stderr, "(err: %d) Platform found: %s\n", err, platform_name);
> 
>     err = clGetDeviceIDs( platform, CL_DEVICE_TYPE_DEFAULT, 1, &device, NULL );
>+    fprintf( stderr, "(err: %d)after clGetDeviceIDs\n", err );
> 
>     err = clGetDeviceInfo(device, CL_DEVICE_NAME,
>             sizeof(device_name), device_name,
>             &ret_param_size);
>-    printf("Device found on the above platform: %s\n", device_name);
>+    fprintf( stderr, "(err: %d)Device found on the above platform: %s\n", err, device_name);
>+
>+	
>+    cl_uint c1;
>+    err = clGetDeviceInfo(device, CL_DEVICE_ADDRESS_BITS,
>+            sizeof(c1), &c1,
>+            &ret_param_size);
> 
>     props[1] = (cl_context_properties)platform;
>     ctx = clCreateContext( props, 1, &device, NULL, NULL, &err );
>+    fprintf( stderr, "(err: %d) after clCreateContext\n", err );
>     queue = clCreateCommandQueue( ctx, device, 0, &err );
>+    fprintf( stderr, "(err: %d)after clCreateCommandQueue\n", err );
> 
>     /* Setup clFFT. */
>     clfftSetupData fftSetup;
>     err = clfftInitSetupData(&fftSetup);
>+    fprintf( stderr, "(err: %d)after clfftInitSetupData\n", err );
>     err = clfftSetup(&fftSetup);
>+    fprintf( stderr, "(err: %d)after clfftSetup\n", err );
> 
>     /* Allocate host & initialize data. */
>     /* Only allocation shown for simplicity. */
>     X = (float *)malloc(N * 2 * sizeof(*X));
>+    X_result = (float *)malloc(N * 2 * sizeof(*X_result));
> 
>     /* print input array */
>-    printf("\nPerforming fft on an one dimensional array of size N = %lu\n", (unsigned long)N);
>+    fprintf( stderr, "\nPerforming fft on an one dimensional array of size N = %lu\n", (unsigned long)N);
>     int print_iter = 0;
>     while(print_iter<N) {
>         float x = (float)print_iter;
>@@ -85,42 +103,54 @@
> 
>     /* Prepare OpenCL memory objects and place data inside them. */
>     bufX = clCreateBuffer( ctx, CL_MEM_READ_WRITE, N * 2 * sizeof(*X), NULL, &err );
>+    fprintf( stderr, "(err: %d)after clCreateBuffer\n", err );
> 
>-    err = clEnqueueWriteBuffer( queue, bufX, CL_TRUE, 0,
>-            N * 2 * sizeof( *X ), X, 0, NULL, NULL );
> 
>     /* Create a default plan for a complex FFT. */
>     err = clfftCreateDefaultPlan(&planHandle, ctx, dim, clLengths);
>+    fprintf( stderr, "(err: %d)after clfftCreateDefaultPlan\n", err );
> 
>     /* Set plan parameters. */
>     err = clfftSetPlanPrecision(planHandle, CLFFT_SINGLE);
>+    fprintf( stderr, "(err: %d)after clfftSetPlanPrecision\n", err );
>     err = clfftSetLayout(planHandle, CLFFT_COMPLEX_INTERLEAVED, CLFFT_COMPLEX_INTERLEAVED);
>+    fprintf( stderr, "(err: %d)after clfftSetLayout\n", err );
>     err = clfftSetResultLocation(planHandle, CLFFT_INPLACE);
>+    fprintf( stderr, "(err: %d)after clfftSetResultLocation\n", err );
> 
>     /* Bake the plan. */
>     err = clfftBakePlan(planHandle, 1, &queue, NULL, NULL);
>+    fprintf( stderr, "(err: %d)after clfftBakePlan\n", err );
>+
>+  struct timeval tv1; 
>+struct timeval timer1;
>+gettimeofday( &timer1, NULL );
>+
>+    err = clEnqueueWriteBuffer( queue, bufX, CL_TRUE, 0,
>+            N * 2 * sizeof( *X ), X, 0, NULL, NULL );
>+    fprintf( stderr, "(err: %d)after clEnqueueWriteBuffer\n", err );
> 
>     /* Execute the plan. */
>     err = clfftEnqueueTransform(planHandle, CLFFT_FORWARD, 1, &queue, 0, NULL, NULL, &bufX, NULL, NULL);
>+    fprintf( stderr, "(err: %d)after clfftEnqueueTransform\n", err );
>+    if( err == CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST ) 
>+    	fprintf( stderr, "(err: %d)after clfftEnqueueTransform = CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST\n", err );
> 
>     /* Wait for calculations to be finished. */
>     err = clFinish(queue);
>+    fprintf( stderr, "(err: %d)after clFinish\n", err );
> 
>     /* Fetch results of calculations. */
>-    err = clEnqueueReadBuffer( queue, bufX, CL_TRUE, 0, N * 2 * sizeof( *X ), X, 0, NULL, NULL );
>+    err = clEnqueueReadBuffer( queue, bufX, CL_TRUE, 0, N * 2 * sizeof( *X_result ), X_result, 0, NULL, NULL );
>+    fprintf( stderr, "(err: %d)after clEnqueueReadBuffer\n", err );
>+
>+  gettimeofday( &tv1, NULL );
>+fprintf(stderr, "Time took 1: %f\r\n",  (float)(tv1.tv_sec*1000000ULL+tv1.tv_usec - (timer1.tv_sec*1000000ULL+timer1.tv_usec))/(float)1000000ULL );
> 
>-    /* print output array */
>-    print_iter = 0;
>-    while(print_iter<N) {
>-        printf("(%f, %f) ", X[2*print_iter], X[2*print_iter+1]);
>-        print_iter++;
>-    }
>-    printf("\n");
> 
>     /* Release OpenCL memory objects. */
>     clReleaseMemObject( bufX );
> 
>-    free(X);
> 
>     /* Release the plan. */
>     err = clfftDestroyPlan( &planHandle );
>@@ -132,5 +162,26 @@
>     clReleaseCommandQueue( queue );
>     clReleaseContext( ctx );
> 
>+    int i1;
>+    int same_count = 0;
>+    for( i1 = 0; i1 < 20 ; i1++  )
>+    {
>+      if( X[i1] == X_result[i1] )
>+        same_count++;
>+    }
>+
>+    if( same_count > 5 )
>+      fprintf( stderr, "\r\n\r\nERROR: found same numbers !!! ... FFT was not executed !!! \r\n\r\n" );
>+
>+    /* print output array */
>+    print_iter = 0;
>+    while(print_iter<N) {
>+        printf("%d (%f, %f)\n", print_iter, X_result[2*print_iter], X_result[2*print_iter+1]);
>+        print_iter++;
>+    }
>+    printf("\n");
>+
>+    free(X);
>+    free(X_result);
>     return ret;
> }
>diff -ru clFFT_orig/src/library/enqueue.cpp clFFT_custom/src/library/enqueue.cpp
>--- clFFT_orig/src/library/enqueue.cpp	2017-04-27 11:55:47.722280825 +0200
>+++ clFFT_custom/src/library/enqueue.cpp	2017-04-27 11:37:11.602340635 +0200
>@@ -657,6 +657,21 @@
>     cl_int call_status = clEnqueueNDRangeKernel( *commQueues, kern, static_cast< cl_uint >( gWorkSize.size( ) ),
>                                             NULL, &gWorkSize[ 0 ],  &lWorkSize[ 0 ], numWaitEvents, waitEvents, outEvents );
>     OPENCL_V( call_status, _T( "clEnqueueNDRangeKernel failed" ) );
>+cl_int statusxw;
>+assert( call_status == CL_SUCCESS ); // All the parameters are valid
>+call_status = clWaitForEvents( 1, outEvents );
>+if( call_status == CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST ) // 1 of the events in the queue failed
>+{
>+fprintf( stderr, "clEnqueueNDRangeKernel failed second... %d params: %d %d %d %d\r\n", call_status, static_cast< cl_uint >( gWorkSize.size( ) ), gWorkSize[ 0 ], lWorkSize[ 0 ], numWaitEvents );
>+   // Optional: Check if it's the kernel
>+   call_status = clGetEventInfo( *outEvents , CL_EVENT_COMMAND_EXECUTION_STATUS, sizeof(statusxw),&statusxw, NULL);
>+   assert( call_status == CL_SUCCESS);
>+
>+   if(statusxw != CL_COMPLETE )
>+      fprintf(stderr, "ERROR executing the kernel %d\n", statusxw);
>+}
>+
>+
> 
>     if( fftRepo.pStatTimer )
>     {
>diff -ru clFFT_orig/src/library/private.h clFFT_custom/src/library/private.h
>--- clFFT_orig/src/library/private.h	2017-04-27 11:55:47.726280824 +0200
>+++ clFFT_custom/src/library/private.h	2017-04-09 12:21:19.000000000 +0200
>@@ -274,6 +274,8 @@
> 			return _T( "CLFFT_VERSION_MISMATCH" );
> 		case CLFFT_INVALID_PLAN:
> 			return _T( "CLFFT_INVALID_PLAN" );
>+		case CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST: 
>+			return _T( "CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST" );
> 		default:
> 			return _T( "Error code not defined" );
> 		break;

Actions: View

Attachments on bug 100639: 130772 | 131090

Use of freedesktop.org services, including Bugzilla, is subject to our Code of Conduct. How we collect and use information is described in our Privacy Policy.