Mercurial > hg > Game > Cerium

changeset 1549:68200bc3ab6b draft

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
change clfinish to wait for event
author Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
date Fri, 15 Feb 2013 11:30:11 +0900
parents 614a3f62c881
children ebaaaac6751a
files TaskManager/Gpu/GpuScheduler.cc TaskManager/Gpu/GpuScheduler.h
diffstat 2 files changed, 30 insertions(+), 25 deletions(-) [+]
line wrap: on
line diff
--- a/TaskManager/Gpu/GpuScheduler.cc	Fri Feb 15 07:37:04 2013 +0900
+++ b/TaskManager/Gpu/GpuScheduler.cc	Fri Feb 15 11:30:11 2013 +0900
@@ -45,15 +45,12 @@
         exit(EXIT_FAILURE);
     }
     context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret);
-    command_queue = new cl_command_queue[2];
-    command_queue[0] = clCreateCommandQueue(context, device_id, 0, &ret);
-    command_queue[1] = clCreateCommandQueue(context, device_id, 0, &ret);
+    command_queue = clCreateCommandQueue(context, device_id, 0, &ret);
 }
 
 GpuScheduler::~GpuScheduler()
 {
-    clReleaseCommandQueue(command_queue[0]);
-    clReleaseCommandQueue(command_queue[1]);
+    clReleaseCommandQueue(command_queue);
     clReleaseContext(context);
 }
 
@@ -71,6 +68,9 @@
     int cur = 0;
     memaddr reply[2];
     cl_kernel *kernel = new cl_kernel[2];
+    cl_event event[2];
+    event[0] = NULL;
+    event[1] = NULL;
     cl_mem *memin[2];
     cl_mem *memout[2];
     HTask::htask_flag flag;
@@ -82,8 +82,7 @@
         // read task list mail from DmaManager
 
         if ((memaddr)params_addr == (memaddr)MY_SPE_COMMAND_EXIT) {
-            clFinish(command_queue[0]);
-            clFinish(command_queue[1]);
+            clFinish(command_queue);
             return ;
         }
 
@@ -136,7 +135,7 @@
                 // set arg count
                 cl_mem memparam = clCreateBuffer(context, CL_MEM_READ_ONLY,
                                                  sizeof(memaddr)*nextTask->param_count, NULL, NULL);
-                ret = clEnqueueWriteBuffer(command_queue[cur], memparam, CL_TRUE, 0,
+                ret = clEnqueueWriteBuffer(command_queue, memparam, CL_TRUE, 0,
                                            sizeof(memaddr)*nextTask->param_count,nextTask->param(0), 0, NULL, NULL);
                 if (ret<0) {
                     const char *msg=convert_error_status(ret);
@@ -159,7 +158,7 @@
                     for(;i<nextTask->inData_count;i++) {
                         ListElement *input_buf = nextTask->inData(i);
                         memin[cur][i] = clCreateBuffer(context, mem_flag, input_buf->size, NULL, NULL);
-                        ret = clEnqueueWriteBuffer(command_queue[cur], memin[cur][i], CL_TRUE, 0,
+                        ret = clEnqueueWriteBuffer(command_queue, memin[cur][i], CL_TRUE, 0,
                                                    input_buf->size, input_buf->addr, 0, NULL, NULL);
                         if (ret<0) {
                             const char *msg=convert_error_status(ret);
@@ -195,7 +194,7 @@
                     if (flag.flip) { // use output buffer as input buffer
                         ListElement *input_buf = nextTask->inData(i);
 
-                        ret = clEnqueueWriteBuffer(command_queue[cur], memout[cur][i], CL_TRUE, 0,
+                        ret = clEnqueueWriteBuffer(command_queue, memout[cur][i], CL_TRUE, 0,
                                                    input_buf->size, input_buf->addr, 0, NULL, NULL);
                         if (ret<0) {
                             const char *msg=convert_error_status(ret);
@@ -210,25 +209,23 @@
                     param++;
                 }
 
-                cl_event ev = NULL;
                 if (flag.nd_range){
-                    ret = clEnqueueNDRangeKernel(command_queue[cur],kernel[cur],dimension,NULL,gws,lws,0,NULL,NULL);
+                    ret = clEnqueueNDRangeKernel(command_queue,kernel[cur],dimension,NULL,gws,lws,0,NULL,&event[cur]);
                 } else {
-                    ret = clEnqueueTask(command_queue[cur], kernel[cur], 0, NULL, &ev);    
+                    ret = clEnqueueTask(command_queue, kernel[cur], 0, NULL, &event[cur]);
                 }
-
                 if (ret<0) {
                     const char *msg=convert_error_status(ret);
                     error(msg);
                 }
                 // ndrange flagが0ならdim,global_work_size[0],local_work_size[0] = 1で固定に
                 // clEnqueueNDRange
-                // (command_queue[cur], kernel[cur], dim, NULL,global_work_size[0],local_work_size[0],NULL&ev);
+                // (command_queue, kernel[cur], dim, NULL,global_work_size[0],local_work_size[0],NULL&ev);
 
                 for(int i=0;i<nextTask->outData_count;i++) { // read output data
                     ListElement *output_buf = flag.flip? nextTask->inData(i) :nextTask->outData(i);
-                    ret = clEnqueueReadBuffer(command_queue[cur], memout[cur][i], CL_TRUE, 0,
-                                              output_buf->size, output_buf->addr, 1, &ev, NULL);
+                    ret = clEnqueueReadBuffer(command_queue, memout[cur][i], CL_TRUE, 0,
+                                              output_buf->size, output_buf->addr, 0, NULL, NULL);
                     if (ret<0) {
                         const char *msg=convert_error_status(ret);
                         error(msg);
@@ -237,10 +234,18 @@
             }
 
             reply[cur] = (memaddr)tasklist->waiter;
-            clFlush(command_queue[1-cur]); // waiting for queued task
+            //clFlush(command_queue); // waiting for queued task
+            if (event[1-cur] != NULL) {
+                ret=clWaitForEvents(1,&event[1-cur]);
+                clReleaseKernel(kernel[1-cur]);
+            }
+            if (ret<0) {
+                const char *msg=convert_error_status(ret);
+                error(msg);
+            }
+            // clFlush(command_queue);
             // pipeline    : 1-cur
             // no pipeline : cur
-            clReleaseKernel(kernel[1-cur]);
             /* should be released
              *  clReleaseMemObject(memin[1-cur]);
              *  clReleaseMemObject(memout[1-cur]);
@@ -253,11 +258,11 @@
             cur = 1 - cur;
         }
 
-          clFlush(command_queue[1-cur]); // waiting for queued task
-          connector->mail_write(reply[1-cur]);
-
-
-          connector->mail_write((memaddr)MY_SPE_STATUS_READY);
+        //clFlush(command_queue); // waiting for queued task
+        ret=clWaitForEvents(1,&event[1-cur]);
+        connector->mail_write(reply[1-cur]);
+        
+        connector->mail_write((memaddr)MY_SPE_STATUS_READY);
     }
     // TaskArrayの処理
 }
--- a/TaskManager/Gpu/GpuScheduler.h	Fri Feb 15 07:37:04 2013 +0900
+++ b/TaskManager/Gpu/GpuScheduler.h	Fri Feb 15 11:30:11 2013 +0900
@@ -37,7 +37,7 @@
     cl_uint ret_num_platforms;
     cl_uint ret_num_devices;
     cl_context context;
-    cl_command_queue *command_queue;
+    cl_command_queue command_queue;
     cl_int ret;
 private:
     FifoDmaManager *fifoDmaManager;