[Beignet] [PATCH] When local_work_size is null, try to choose a local_work_size.

Song, Ruiling ruiling.song at intel.com
Thu Nov 28 18:48:53 PST 2013


LGTM

-----Original Message-----
From: beignet-bounces at lists.freedesktop.org [mailto:beignet-bounces at lists.freedesktop.org] On Behalf Of Yang Rong
Sent: Friday, November 29, 2013 11:00 AM
To: beignet at lists.freedesktop.org
Cc: Yang, Rong R
Subject: [Beignet] [PATCH] When local_work_size is null, try to choose a local_work_size.


Signed-off-by: Yang Rong <rong.r.yang at intel.com>
---
 src/cl_api.c | 11 +++++++++--
 1 file changed, 9 insertions(+), 2 deletions(-)

diff --git a/src/cl_api.c b/src/cl_api.c index 6acc1a2..08f01d2 100644
--- a/src/cl_api.c
+++ b/src/cl_api.c
@@ -2361,7 +2361,7 @@ clEnqueueNDRangeKernel(cl_command_queue  command_queue,
   size_t fixed_global_sz[] = {1,1,1};
   size_t fixed_local_sz[] = {1,1,1};
   cl_int err = CL_SUCCESS;
-  cl_uint i;
+  cl_uint i, j;
   enqueue_data *data, no_wait_data = { 0 };
 
   CHECK_QUEUE(command_queue);
@@ -2407,9 +2407,16 @@ clEnqueueNDRangeKernel(cl_command_queue  command_queue,
   //FATAL_IF(event_wait_list != NULL, "Events are not supported");
   //FATAL_IF(event != NULL, "Events are not supported");
 
-  if (local_work_size != NULL)
+  if (local_work_size != NULL) {
     for (i = 0; i < work_dim; ++i)
       fixed_local_sz[i] = local_work_size[i];
+  } else {
+    for (i = 0; i< work_dim; i++)
+      for (j = 64; j > 1; j--) {   //check from 64?
+        if (global_work_size[i] % j == 0) //global_work_size always non null
+          fixed_local_sz[i] = j;
+      }
+  }
   if (global_work_size != NULL)
     for (i = 0; i < work_dim; ++i)
       fixed_global_sz[i] = global_work_size[i];
--
1.8.1.2

_______________________________________________
Beignet mailing list
Beignet at lists.freedesktop.org
http://lists.freedesktop.org/mailman/listinfo/beignet


More information about the Beignet mailing list