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

Yang, Rong R rong.r.yang at intel.com
Sun Dec 13 22:13:12 PST 2015


Is it good to printf it directly? Ruining have add debug print macro, only print in debug version, how about use it?

> -----Original Message-----
> From: Beignet [mailto:beignet-bounces at lists.freedesktop.org] On Behalf Of
> Guo Yejun
> Sent: Friday, December 4, 2015 5:48
> To: beignet at lists.freedesktop.org
> Cc: Guo, Yejun
> Subject: [Beignet] [PATCH] output warning message if the
> global/local_work_size is not good
> 
> 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
> 
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet


More information about the Beignet mailing list