opencl benchmark fixes

benchmarks/opencl/bfs/CLHelper.h

@@ -233,7 +233,7 @@ free(allPlatforms);*/
   //--cambine-4: Create an OpenCL command queue
   oclHandles.queue = clCreateCommandQueue(
       oclHandles.context, oclHandles.devices[DEVICE_ID_INUSED], 0, &resultCL);
-  printf("resultCL=%d, queue=0x%x\n", resultCL, oclHandles.queue);
+  //printf("resultCL=%d, queue=0x%x\n", resultCL, oclHandles.queue);
 
   if ((resultCL != CL_SUCCESS) || (oclHandles.queue == NULL))
     throw(string("InitCL()::Creating Command Queue. (clCreateCommandQueue)"));
@@ -383,8 +383,8 @@ void _clRelease() {
         errorFlag = true;
       }
       oclHandles.kernel[nKernel] = NULL;
+      printf("clReleaseKernel()\n");
     }
-    oclHandles.kernel.clear();
   }
 
   if (oclHandles.program != NULL) {
@@ -394,6 +394,7 @@ void _clRelease() {
       errorFlag = true;
     }
     oclHandles.program = NULL;
+    printf("clReleaseProgram()\n");
   }
 
   if (oclHandles.queue != NULL) {
@@ -403,10 +404,9 @@ void _clRelease() {
       errorFlag = true;
     }
     oclHandles.queue = NULL;
+    printf("clReleaseCommandQueue()\n");
   }
 
-  free(oclHandles.devices);
-
   if (oclHandles.context != NULL) {
     cl_int resultCL = clReleaseContext(oclHandles.context);
     if (resultCL != CL_SUCCESS) {
@@ -414,6 +414,17 @@ void _clRelease() {
       errorFlag = true;
     }
     oclHandles.context = NULL;
+    printf("clReleaseContext()\n");
+  }
+
+  if (oclHandles.devices != NULL) {
+    cl_int resultCL = clReleaseDevice(oclHandles.devices[0]);
+    if (resultCL != CL_SUCCESS) {
+      cerr << "ReleaseCL()::Error: In clReleaseDevice" << endl;
+      errorFlag = true;
+    }
+    free(oclHandles.devices);
+    printf("clReleaseDevice()\n");
   }
 
   if (errorFlag)
@@ -675,7 +686,7 @@ void _clFinish() throw(string) {
 void _clInvokeKernel(int kernel_id, int work_items,
                      int work_group_size) throw(string) {
   cl_uint work_dim = WORK_DIM;
-  cl_event e[1];
+  //cl_event e[1];
   if (work_items % work_group_size != 0) // process situations that work_items
                                          // cannot be divided by work_group_size
     work_items =
@@ -684,7 +695,7 @@ void _clInvokeKernel(int kernel_id, int work_items,
   size_t global_work_size[] = {work_items, 1};
   oclHandles.cl_status = clEnqueueNDRangeKernel(
       oclHandles.queue, oclHandles.kernel[kernel_id], work_dim, 0,
-      global_work_size, local_work_size, 0, 0, &(e[0]));
+      global_work_size, local_work_size, 0, 0, NULL);
 #ifdef ERRMSG
   oclHandles.error_str = "excpetion in _clInvokeKernel() -> ";
   switch (oclHandles.cl_status) {
@@ -749,13 +760,13 @@ void _clInvokeKernel2D(int kernel_id, int range_x, int range_y, int group_x,
   cl_uint work_dim = WORK_DIM;
   size_t local_work_size[] = {group_x, group_y};
   size_t global_work_size[] = {range_x, range_y};
-  cl_event e[1];
+  //cl_event e[1];
   /*if(work_items%work_group_size != 0)	//process situations that work_items
     cannot be divided by work_group_size
     work_items = work_items + (work_group_size-(work_items%work_group_size));*/
   oclHandles.cl_status = clEnqueueNDRangeKernel(
       oclHandles.queue, oclHandles.kernel[kernel_id], work_dim, 0,
-      global_work_size, local_work_size, 0, 0, &(e[0]));
+      global_work_size, local_work_size, 0, 0, NULL);
 #ifdef ERRMSG
   oclHandles.error_str = "excpetion in _clInvokeKernel() -> ";
   switch (oclHandles.cl_status) {

benchmarks/opencl/bfs/main.cc

@@ -78,14 +78,15 @@ void run_bfs_gpu(int no_of_nodes, Node *h_graph_nodes, int edge_list_size,
   char h_over;
   cl_mem d_graph_nodes, d_graph_edges, d_graph_mask, d_updating_graph_mask,
       d_graph_visited, d_cost, d_over;
+
   try {
     //--1 transfer data from host to device
     _clInit();
+
     d_graph_nodes = _clMalloc(no_of_nodes * sizeof(Node), h_graph_nodes);
     d_graph_edges = _clMalloc(edge_list_size * sizeof(int), h_graph_edges);
     d_graph_mask = _clMallocRW(no_of_nodes * sizeof(char), h_graph_mask);
-    d_updating_graph_mask =
-        _clMallocRW(no_of_nodes * sizeof(char), h_updating_graph_mask);
+    d_updating_graph_mask = _clMallocRW(no_of_nodes * sizeof(char), h_updating_graph_mask);
     d_graph_visited = _clMallocRW(no_of_nodes * sizeof(char), h_graph_visited);
 
     d_cost = _clMallocRW(no_of_nodes * sizeof(int), h_cost);
@@ -94,8 +95,7 @@ void run_bfs_gpu(int no_of_nodes, Node *h_graph_nodes, int edge_list_size,
     _clMemcpyH2D(d_graph_nodes, no_of_nodes * sizeof(Node), h_graph_nodes);
     _clMemcpyH2D(d_graph_edges, edge_list_size * sizeof(int), h_graph_edges);
     _clMemcpyH2D(d_graph_mask, no_of_nodes * sizeof(char), h_graph_mask);
-    _clMemcpyH2D(d_updating_graph_mask, no_of_nodes * sizeof(char),
-                 h_updating_graph_mask);
+    _clMemcpyH2D(d_updating_graph_mask, no_of_nodes * sizeof(char), h_updating_graph_mask);
     _clMemcpyH2D(d_graph_visited, no_of_nodes * sizeof(char), h_graph_visited);
     _clMemcpyH2D(d_cost, no_of_nodes * sizeof(int), h_cost);
 
@@ -106,6 +106,7 @@ void run_bfs_gpu(int no_of_nodes, Node *h_graph_nodes, int edge_list_size,
     kernel_timer.reset();
     kernel_timer.start();
 #endif
+
     do {
       h_over = false;
       _clMemcpyH2D(d_over, sizeof(char), &h_over);
@@ -136,9 +137,8 @@ void run_bfs_gpu(int no_of_nodes, Node *h_graph_nodes, int edge_list_size,
       _clInvokeKernel(kernel_id, no_of_nodes, work_group_size);
 
       _clMemcpyD2H(d_over, sizeof(char), &h_over);
-    } while (h_over);
+    } while (h_over);    
 
-    _clFinish();
 #ifdef PROFILING
     kernel_timer.stop();
     kernel_time = kernel_timer.getTimeInSeconds();

benchmarks/opencl/bfs/util.h

@@ -60,10 +60,10 @@ void compare_results(const datatype *cpu_results, const datatype *gpu_results, c
       }
     }
     if (passed){
-        std::cout << "--cambine:passed:-)" << endl;
+        std::cout << "--cambine: passed: -)" << endl;
     }
     else{
-        std::cout << "--cambine: failed:-(" << endl;
+        std::cout << "--cambine: failed :-(" << endl;
     }
     return ;
 }