Mercurial > hg > Game > Cerium

changeset 1803:fa15a19d27ef draft

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
add event flag for GpuScheduler
author Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
date Mon, 09 Dec 2013 18:43:56 +0900
parents 9a1ba9cb9557
children 1febe61a935a
files TaskManager/Gpu/GpuScheduler.cc TaskManager/Gpu/GpuScheduler.h
diffstat 2 files changed, 42 insertions(+), 48 deletions(-) [+]
line wrap: on
line diff
--- a/TaskManager/Gpu/GpuScheduler.cc	Fri Dec 06 05:31:30 2013 +0900
+++ b/TaskManager/Gpu/GpuScheduler.cc	Mon Dec 09 18:43:56 2013 +0900
@@ -65,7 +65,7 @@
 #define NOP_REPLY NULL
 
 void
-GpuScheduler::wait_for_event(cl_event* event,memaddr* reply,int cur) {
+GpuScheduler::wait_for_event(cl_event* event,memaddr* reply,TaskListPtr taskList, int cur) {
     if (event[1-cur] == NOP_REPLY) {
         if(reply[1-cur]) {
             connector->mail_write(reply[1-cur]);
@@ -75,14 +75,21 @@
     if (event[1-cur] != NULL) {
         int ret=clWaitForEvents(1,&event[1-cur]);
         if (ret<0) {
-            const char *msg=convert_error_status(ret);
-            error(msg);
+            error(convert_error_status(ret));
         }
         if(reply[1-cur]) {
             connector->mail_write(reply[1-cur]);
             reply[1-cur]=0;
         }
     }
+    if (tasklist[1-cur]!=NULL) {
+        cl_ulong start = 0;
+        cl_ulong end   = 0;
+        clGetEventProfilingInfo(event[1-cur],CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &start, NULL);
+        clGetEventProfilingInfo(event[1-cur],CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &end, NULL);
+        tasklist[1-cur]->task_start_time = start;
+        tasklist[1-cur]->task_end_time   = end;
+    }
 }
 
 
@@ -96,16 +103,8 @@
 GpuScheduler::run()
 {
     int cur = 0;
-    memaddr reply[2]={0,0};
-    cl_kernel kernel[2]={0,0};
-    cl_event event[2];
-    event[0]=NULL;event[1]=NULL;
-
-    cl_mem *memin[2];
-    cl_mem *memout[2];
     TaskListPtr tasklist[2];
     tasklist[0]=NULL;tasklist[1]=NULL;
-    HTask::htask_flag flag;
     memset(&flag, 0, sizeof(HTask::htask_flag));
 
     for (;;) {
@@ -174,7 +173,7 @@
                 }
 
                 ret = clEnqueueWriteBuffer(command_queue, memparam, CL_TRUE, 0,sizeof(memaddr)*nextTask->param_count,
-                                           nextTask->param(param), 0, NULL, NULL);
+                                           nextTask->param(param), 0, NULL, &param_event[1-cur]); // set event flag
                 param=0;
                 if (ret<0) {
                     const char *msg=convert_error_status(ret);
@@ -192,7 +191,7 @@
                 param++;
 
                 cl_mem_flags mem_flag = CL_MEM_READ_ONLY;
-                if ( memout[cur] )  clReleaseMemObject(*memout[cur]);
+                if ( memout[cur] )  clReleaseMemObject(*memout[cur]); // using preallocated memorry
                 memin[cur] = new cl_mem[nextTask->inData_count];
 
                 if (!flag.flip) { // set input data when not flip
@@ -205,7 +204,8 @@
                             error(msg);
                         }
                         ret = clEnqueueWriteBuffer(command_queue, memin[cur][i], CL_TRUE, 0,
-                                                   input_buf->size, input_buf->addr, 0, NULL, NULL);
+                                                   input_buf->size, input_buf->addr, 1, 
+                                                   param_event[1-cur], memin_event[1-cur][i]);
                         if (ret<0) {
                             const char *msg=convert_error_status(ret);
                             error(msg);
@@ -244,7 +244,8 @@
                         ListElement *input_buf = nextTask->inData(i);
 
                         ret = clEnqueueWriteBuffer(command_queue, memout[cur][i], CL_TRUE, 0,
-                                                   input_buf->size, input_buf->addr, 0, NULL, NULL);
+                                                   input_buf->size, input_buf->addr, 
+                                                   memparam_size[1-cur],param_event[1-cur], memin_event[1-cur][i]);
                         if (ret<0) {
                             const char *msg=convert_error_status(ret);
                             error(msg);
@@ -261,9 +262,10 @@
                 tasklist[cur]->task_start_time = gettime();
                 if (tasklist[cur]->dim > 0) {
                     ret = clEnqueueNDRangeKernel(command_queue, kernel[cur], tasklist[cur]->dim,
-                                                 NULL, &tasklist[cur]->x, 0, 0, NULL, &event[cur]);
+                                 NULL, &tasklist[cur]->x, 0, memin_size[1-cur], memin_event[1-cur], &kernel_event[cur]);
                 } else {
-                    ret = clEnqueueTask(command_queue, kernel[cur], 0, NULL, &event[cur]);
+                    ret = clEnqueueTask(command_queue, kernel[cur], memin_size[1-cur],
+                                        memin_event[1-cur], &kernel_event[cur]);
                 }
                 if (ret<0) {
                     const char *msg=convert_error_status(ret);
@@ -275,7 +277,7 @@
                     ListElement *output_buf = flag.flip? nextTask->inData(i) :nextTask->outData(i);
                     if (output_buf->size==0) break;
                     ret = clEnqueueReadBuffer(command_queue, memout[cur][i], CL_FALSE, 0,
-                                              output_buf->size, output_buf->addr, 0, NULL, NULL);
+                                    output_buf->size, output_buf->addr, 1, kernel_event[1-cur], memout_event[1-cur][i]);
                     if (ret<0) {
                         const char *msg=convert_error_status(ret);
                         error(msg);
@@ -295,16 +297,9 @@
                 }
 
                 reply[cur] = (memaddr)tasklist[cur]->waiter;
-
-                wait_for_event(event,reply,cur);
-                if (tasklist[1-cur]!=NULL) {
-                    cl_ulong start = 0;
-                    cl_ulong end   = 0;
-                    clGetEventProfilingInfo(event[1-cur],CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &start, NULL);
-                    clGetEventProfilingInfo(event[1-cur],CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &end, NULL);
-                    tasklist[1-cur]->task_start_time = start;
-                    tasklist[1-cur]->task_end_time   = end;
-                }
+                
+                // wait kernel[1-cur] and write[1-cur]
+                wait_for_event(event,reply,tasklist,cur);
                 event[1-cur] = NULL;
                 // pipeline    : 1-cur
                 // no pipeline : cur
@@ -313,30 +308,13 @@
             printf("GPU %d %s\t%lld\n",tasklist[cur]->self->cpu_type,(char*)(gpu_task_list[tasklist[cur]->tasks[0].command].name),tasklist[cur]->task_end_time-tasklist[cur]->task_start_time);
             cur = 1 - cur;
         }
-        wait_for_event(event,reply,cur);
-        if (tasklist[cur]!=NULL) {
-            cl_ulong start = 0;
-            cl_ulong end   = 0;
-            clGetEventProfilingInfo(event[cur],CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &start, NULL);
-            clGetEventProfilingInfo(event[cur],CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &end, NULL);
-            tasklist[cur]->task_start_time = start;
-            tasklist[cur]->task_end_time   = end;
-            event[cur] = NULL;
-        }
+        wait_for_event(memout_event,reply,tasklist,cur);
+
         unsigned long long wait = 0;
         (*connector->end_dmawait_profile)(&wait, &(connector->start_time), &(connector->stop_time));
         connector->mail_write((memaddr)MY_SPE_STATUS_READY);
     }
-    // TaskArrayの処理
-    if ( memout[cur-1] )  clReleaseMemObject(*memout[cur-1]);
-    if ( memout[cur] )  clReleaseMemObject(*memout[cur]);
-    if ( memin[cur-1] )  clReleaseMemObject(*memin[cur-1]);
-    if ( memin[cur] )  clReleaseMemObject(*memin[cur]);
-    if ( kernel[cur]) clReleaseKernel(kernel[cur]);
-    if ( kernel[cur-1]) clReleaseKernel(kernel[cur-1]);
-    if (event[0] && event[0]!=NOP_REPLY) clReleaseEvent(event[0]);
-    if (event[1] && event[0]!=NOP_REPLY) clReleaseEvent(event[1]);
-
+    /* NOT REACHED */
 }
 
 int
--- a/TaskManager/Gpu/GpuScheduler.h	Fri Dec 06 05:31:30 2013 +0900
+++ b/TaskManager/Gpu/GpuScheduler.h	Mon Dec 09 18:43:56 2013 +0900
@@ -43,6 +43,22 @@
     cl_context context;
     cl_command_queue command_queue;
     cl_int ret;
+    memaddr reply[2];
+    cl_kernel kernel[2];
+    cl_event kernel_event[2];
+    cl_mem memparam[2]; // clCreateBuffer
+    cl_int memparam_size[2];
+    cl_event param_event[2];
+    cl_event *param_wait[2]; // wait list for param input
+    cl_mem *memin[2]; // clCreateBuffer
+    cl_int *memin_size[2];
+    cl_event *memin_event[2];
+    cl_event *memin_wait[2]; // wait list for param input
+    cl_mem *memout[2]; // clCreateBuffer
+    cl_int *memout_size[2];
+    cl_event *memout_event[2];
+    cl_event *memout_wait[2]; // wait list for param input
+    HTask::htask_flag flag;
 private:
     FifoDmaManager *fifoDmaManager;
     int load_kernel(int cmd);