--- a/TaskManager/Gpu/GpuScheduler.cc	Sat Dec 14 20:04:04 2013 +0900
+++ b/TaskManager/Gpu/GpuScheduler.cc	Mon Dec 16 20:00:13 2013 +0900
@@ -101,7 +101,7 @@
  * kernel_event, memout_event
  */
 void
-GpuScheduler::wait_for_event(cl_event* kernel_event, GpuBufferPtr memout, memaddr* reply, TaskListPtr *taskList, int cur) {
+GpuScheduler::wait_for_event(cl_event* kernel_event, GpuBufferPtr memout, memaddr* reply, TaskListPtr taskList, int cur) {
     if (kernel_event[1-cur] == NOP_REPLY) {
         if(reply[1-cur]) {
             connector->mail_write(reply[1-cur]);
@@ -141,29 +141,27 @@
         memin[1-cur].size = 0;
     }
 
-    if(reply[1-cur]) {
-        connector->mail_write(reply[1-cur]);
-        reply[1-cur]=0;
-    }
-    if (taskList[1-cur]!=NULL){
+    if (taskList!=NULL){
         cl_ulong start = 0;
         cl_ulong end   = 0;
         clGetEventProfilingInfo(kernel_event[1-cur],CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &start, NULL);
         clGetEventProfilingInfo(kernel_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;
+        if (taskList->task_start_time == 0)
+            taskList->task_start_time = start;
+        taskList->task_end_time   = end;
     }
+
 }
 
 void
-GpuScheduler::gpuTaskError(int cur, TaskListPtr *tasklist, int ret)
+GpuScheduler::gpuTaskError(int cur, TaskListPtr tasklist, int ret)
 {
     error(convert_error_status(ret));
     kernel_event[cur] = NOP_REPLY;
     kernel[cur] = 0;
     memout[cur].buf = 0;
     memin[cur].buf = 0;
-    reply[cur] = (memaddr)tasklist[cur]->waiter;
+    reply[cur] = (memaddr)tasklist->waiter;
 
     // wait kernel[1-cur] and write[1-cur]
     wait_for_event(kernel_event, memout, reply, tasklist, cur);
@@ -179,8 +177,7 @@
 GpuScheduler::run()
 {
     int cur = 0;
-    TaskListPtr tasklist[2];
-    tasklist[0]=NULL;tasklist[1]=NULL;
+    TaskListPtr tasklist = NULL;
     initGpuBuffer(&memin[0]);initGpuBuffer(&memin[1]);
     initGpuBuffer(&memout[0]);initGpuBuffer(&memout[1]);
     memset(&flag, 0, sizeof(HTask::htask_flag));
@@ -205,18 +202,20 @@
         (*connector->start_dmawait_profile)(&(connector->start_time));
         while (params_addr) {
             // since we are on the same memory space, we don't has to use dma_load here
-            tasklist[cur] = (TaskListPtr)connector->dma_load(this, params_addr,
+            tasklist = (TaskListPtr)connector->dma_load(this, params_addr,
                                                                     sizeof(TaskList), DMA_READ_TASKLIST);
             //            tasklist[cur]->task_start_time = gettime();
+            tasklist->task_start_time = 0;
             /*
              * get flip flag
              * flip : When caluculate on input data, to treat this as a output data
              */
-            if (tasklist[cur]->self) {
-                flag = tasklist[cur]->self->flag;
+            if (tasklist->self) {
+                flag = tasklist->self->flag;
             }
-            
-            for (TaskPtr nextTask = tasklist[cur]->tasks;nextTask < tasklist[cur]->last(); nextTask = nextTask->next(),params_addr = (memaddr)tasklist[cur]->next) {
+            TaskPtr nextTask = NULL;
+            nextTask[cur] = tasklist->tasks;
+            while (nextTask < taskList->last()) {
                 if(nextTask->command==ShowTime) {
                     connector->show_profile();
                     gpuTaskError(cur,tasklist,ret);
@@ -233,7 +232,7 @@
                 }
                 cl_program& program = *gpu_task_list[nextTask->command].gputask->program;
                 const char *function = gpu_task_list[nextTask->command].name;
-
+                
                 if (kernel[cur])
                     clReleaseKernel(kernel[cur]);
                 kernel[cur] = clCreateKernel(program, function, &ret);
@@ -333,10 +332,10 @@
                 }
                 
                 memout[cur].size = param - memin[cur].size;
-                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, memin[cur].size, memin[cur].event, &kernel_event[cur]);
+                tasklist->task_start_time = gettime();
+                if (tasklist->dim > 0) {
+                    ret = clEnqueueNDRangeKernel(command_queue, kernel[cur], tasklist->dim,
+                                 NULL, &tasklist->x, 0, memin[cur].size, memin[cur].event, &kernel_event[cur]);
                 } else {
                     ret = clEnqueueTask(command_queue, kernel[cur], memin[cur].size,
                                         memin[cur].event, &kernel_event[cur]);
@@ -356,22 +355,28 @@
                         continue;
                     }
                 }
-                tasklist[cur]->task_end_time = gettime();
-
+                tasklist->task_end_time = gettime();
+                
                 if (ret<0) {
                     gpuTaskError(cur,tasklist,ret);
                     continue;
                 }
 
-                reply[cur] = (memaddr)tasklist[cur]->waiter;
-                
                 // wait kernel[1-cur] and write[1-cur]
-                wait_for_event(kernel_event, memout, reply, tasklist, cur);
+                wait_for_event(kernel_event, memout, 0, tasklist, cur);
+                // pipeline    : 1-cur
+                // no pipeline : cur
+                cur = 1 - cur;
+                nextTask = nextTask->next();
             }
             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);
-            // pipeline    : 1-cur
-            // no pipeline : cur
-            cur = 1 - cur;
+            reply = (memaddr)tasklist->waiter;
+            if(reply) {
+                connector->mail_write(reply);
+                reply=0;
+            }
+
+            params_addr = (memaddr)tasklist->next;
         }
         wait_for_event(kernel_event, memout, reply, tasklist, cur);