[Beignet] [PATCH] output warning message if the global/local_work_size is not good

[Guo Yejun]

the known issue is that utest will output lots of warning messages
since the total number of work-items is less than 64.

Signed-off-by: Guo Yejun <yejun.guo at intel.com>
---
 src/cl_api.c | 11 ++++++++++-
 1 file changed, 10 insertions(+), 1 deletion(-)

diff --git a/src/cl_api.c b/src/cl_api.c
index d8ccd7e..9c83c85 100644
--- a/src/cl_api.c
+++ b/src/cl_api.c
@@ -2963,6 +2963,7 @@ clEnqueueNDRangeKernel(cl_command_queue  command_queue,
         fixed_local_sz[1] = 1;
     } else {
       uint j, maxDimSize = 64 /* from 64? */, maxGroupSize = 256; //MAX_WORK_GROUP_SIZE may too large
+      size_t realGroupSize = 1;
       for (i = 0; i< work_dim; i++) {
         for (j = maxDimSize; j > 1; j--) {
           if (global_work_size[i] % j == 0 && j <= maxGroupSize) {
@@ -2972,7 +2973,10 @@ clEnqueueNDRangeKernel(cl_command_queue  command_queue,
             break;  //choose next work_dim
           }
         }
+        realGroupSize *= fixed_local_sz[i];
       }
+      if (realGroupSize < 8 || realGroupSize % 8 != 0)
+        printf("WARNING: unable to find good values for local_work_size[i], please provide local_work_size[] explicitly, you can find good values with trial-and-error method.\n");
     }
   }
 
@@ -2980,8 +2984,13 @@ clEnqueueNDRangeKernel(cl_command_queue  command_queue,
     fixed_global_sz[0] = (global_work_size[0]+15) / 16 * 16;
     fixed_global_sz[1] = (global_work_size[1]+15) / 16;
   } else {
-    for (i = 0; i < work_dim; ++i)
+    size_t total_work_items = 1;
+    for (i = 0; i < work_dim; ++i) {
       fixed_global_sz[i] = global_work_size[i];
+      total_work_items *= fixed_global_sz[i];
+    }
+    if (total_work_items < 64)
+      printf("WARNING: too small work-items (see values in global_work_size[]) might result in bad performance.\n");
   }
   if (global_work_offset != NULL)
     for (i = 0; i < work_dim; ++i)
-- 
1.9.1