Mercurial > hg > Game > Cerium
changeset 1550:ebaaaac6751a draft
add WaitMarker
| author | Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp> |
|---|---|
| date | Fri, 15 Feb 2013 16:02:43 +0900 |
| parents | 68200bc3ab6b |
| children | 57317332f6ef |
| files | TaskManager/Gpu/GpuScheduler.cc TaskManager/kernel/ppe/HTask.h |
| diffstat | 2 files changed, 41 insertions(+), 42 deletions(-) [+] |
line wrap: on
line diff
--- a/TaskManager/Gpu/GpuScheduler.cc Fri Feb 15 11:30:11 2013 +0900 +++ b/TaskManager/Gpu/GpuScheduler.cc Fri Feb 15 16:02:43 2013 +0900 @@ -91,17 +91,16 @@ TaskListPtr tasklist = (TaskListPtr)connector->dma_load(this, params_addr, sizeof(TaskList), DMA_READ_TASKLIST); + /* + * get flip flag + * flip : When caluculate on input data, to treat this as a output data + */ if (tasklist->self) { - /* - * get flip flag - * flip : When caluculate on input data, to treat this as a output data - */ flag = tasklist->self->flag; } - - for (TaskPtr nextTask = tasklist->tasks; - nextTask < tasklist->last(); nextTask = nextTask->next()) { - + + for (TaskPtr nextTask = tasklist->tasks;nextTask < tasklist->last(); nextTask = nextTask->next()) { + load_kernel(nextTask->command); cl_program& program = *task_list[nextTask->command].gputask->program; const char *function = task_list[nextTask->command].name; @@ -208,11 +207,11 @@ } param++; } - + if (flag.nd_range){ - ret = clEnqueueNDRangeKernel(command_queue,kernel[cur],dimension,NULL,gws,lws,0,NULL,&event[cur]); + ret = clEnqueueNDRangeKernel(command_queue,kernel[cur],dimension,NULL,gws,lws,0,NULL, NULL); } else { - ret = clEnqueueTask(command_queue, kernel[cur], 0, NULL, &event[cur]); + ret = clEnqueueTask(command_queue, kernel[cur], 0, NULL, NULL); } if (ret<0) { const char *msg=convert_error_status(ret); @@ -231,35 +230,38 @@ error(msg); } } - } - - reply[cur] = (memaddr)tasklist->waiter; - //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); + + ret = clEnqueueMarkerWithWaitList(command_queue,0,NULL,&event[cur]); + if (ret<0) { + const char *msg=convert_error_status(ret); + error(msg); + } + reply[cur] = (memaddr)tasklist->waiter; + //clFlush(command_queue); // waiting for queued task + if (event[1-cur] != NULL) { + ret=clWaitForEvents(1,&event[1-cur]); + } + if (ret<0) { + const char *msg=convert_error_status(ret); + error(msg); + } + // clFlush(command_queue); + // pipeline : 1-cur + // no pipeline : cur + /* should be released + * clReleaseMemObject(memin[1-cur]); + * clReleaseMemObject(memout[1-cur]); + */ + if(reply[1-cur]) { + connector->mail_write(reply[1-cur]); + } + + params_addr = (memaddr)tasklist->next; + cur = 1 - cur; } - // clFlush(command_queue); - // pipeline : 1-cur - // no pipeline : cur - /* should be released - * clReleaseMemObject(memin[1-cur]); - * clReleaseMemObject(memout[1-cur]); - */ - if(reply[1-cur]) { - connector->mail_write(reply[1-cur]); - } - - params_addr = (memaddr)tasklist->next; - cur = 1 - cur; } - - //clFlush(command_queue); // waiting for queued task - ret=clWaitForEvents(1,&event[1-cur]); + 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);
--- a/TaskManager/kernel/ppe/HTask.h Fri Feb 15 11:30:11 2013 +0900 +++ b/TaskManager/kernel/ppe/HTask.h Fri Feb 15 16:02:43 2013 +0900 @@ -128,12 +128,9 @@ void no_flip() { flag.flip = 0; } - void NDrange() { + void nd_range() { flag.nd_range = 1; } - void no_NDrange() { - flag.nd_range = 0; - } htask_flag get_flag(){ return flag;