[Beignet] [PATCH 4/5] runtime: set CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE to kernel's SIMD_WIDTH.

[Zhigang Gong]

On Thu, Nov 12, 2015 at 04:47:04PM +0800, Zhigang Gong wrote:
> It makes sense to set CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE to the
> corresponding SIMD size. Then it provides a way for intel's OCL application
> to get SIMD width at runtime and make some SIMD width dependant optimization
> possible.
> 
> Signed-off-by: Zhigang Gong <zhigang.gong at intel.com>
> ---
>  src/cl_api.c                |  3 ++-
>  src/cl_command_queue_gen7.c |  2 +-
>  src/cl_device_id.c          | 11 ++++++++++-
>  src/cl_device_id.h          |  2 --
>  src/cl_gt_device.h          |  1 -
>  5 files changed, 13 insertions(+), 6 deletions(-)
> 
> diff --git a/src/cl_api.c b/src/cl_api.c
> index a18bc99..64206eb 100644
> --- a/src/cl_api.c
> +++ b/src/cl_api.c
> @@ -3001,6 +3001,7 @@ clEnqueueNDRangeKernel(cl_command_queue  command_queue,
>      err = cl_command_queue_flush(command_queue);
>    }
>  
> +error:
>    if(b_output_kernel_perf)
>    {
>      if(kernel->program->build_opts != NULL)
> @@ -3008,7 +3009,7 @@ clEnqueueNDRangeKernel(cl_command_queue  command_queue,
>      else
>        time_end(command_queue->ctx, cl_kernel_get_name(kernel), "", command_queue);
>    }
> -error:
> +

The above change is to fix a dead lock when enable kernel performance measurement and
ran into error in cl_command_queue_ND_range(). Forgot to mention it in the commit log.

Thanks,
Zhigang Gong.

>    return err;
>  }
>  
> diff --git a/src/cl_command_queue_gen7.c b/src/cl_command_queue_gen7.c
> index 2edc3be..f0ee20a 100644
> --- a/src/cl_command_queue_gen7.c
> +++ b/src/cl_command_queue_gen7.c
> @@ -329,7 +329,7 @@ cl_command_queue_ND_range_gen7(cl_command_queue queue,
>  
>    /* Compute the number of HW threads we need */
>    if(UNLIKELY(err = cl_kernel_work_group_sz(ker, local_wk_sz, 3, &local_sz) != CL_SUCCESS)) {
> -    fprintf(stderr, "Beignet: Work group size exceed Kerne's work group size.\n");
> +    fprintf(stderr, "Beignet: Work group size exceed Kernel's work group size.\n");
>      return err;
>    }
>    kernel.thread_n = thread_n = (local_sz + simd_sz - 1) / simd_sz;
> diff --git a/src/cl_device_id.c b/src/cl_device_id.c
> index 4551aa8..8186ac8 100644
> --- a/src/cl_device_id.c
> +++ b/src/cl_device_id.c
> @@ -966,7 +966,16 @@ cl_get_kernel_workgroup_info(cl_kernel kernel,
>          return CL_SUCCESS;
>        }
>      }
> -    DECL_FIELD(PREFERRED_WORK_GROUP_SIZE_MULTIPLE, device->preferred_wg_sz_mul)
> +    case CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE:
> +    {
> +      if (param_value && param_value_size < sizeof(size_t))
> +        return CL_INVALID_VALUE;
> +      if (param_value_size_ret != NULL)
> +        *param_value_size_ret = sizeof(size_t);
> +      if (param_value)
> +        *(size_t*)param_value = interp_kernel_get_simd_width(kernel->opaque);
> +      return CL_SUCCESS;
> +    }
>      case CL_KERNEL_LOCAL_MEM_SIZE:
>      {
>        size_t local_mem_sz =  interp_kernel_get_slm_size(kernel->opaque) + kernel->local_mem_sz;
> diff --git a/src/cl_device_id.h b/src/cl_device_id.h
> index 4a923ef..c5f9e57 100644
> --- a/src/cl_device_id.h
> +++ b/src/cl_device_id.h
> @@ -108,8 +108,6 @@ struct _cl_device_id {
>    size_t driver_version_sz;
>    size_t spir_versions_sz;
>    size_t built_in_kernels_sz;
> -  /* Kernel specific info that we're assigning statically */
> -  size_t preferred_wg_sz_mul;
>    /* SubDevice specific info */
>    cl_device_id parent_device;
>    cl_uint      partition_max_sub_device;
> diff --git a/src/cl_gt_device.h b/src/cl_gt_device.h
> index de7a636..12987b7 100644
> --- a/src/cl_gt_device.h
> +++ b/src/cl_gt_device.h
> @@ -39,7 +39,6 @@
>  .native_vector_width_float = 4,
>  .native_vector_width_double = 2,
>  .native_vector_width_half = 8,
> -.preferred_wg_sz_mul = 16,
>  .address_bits = 32,
>  .max_mem_alloc_size = 512 * 1024 * 1024,
>  .image_support = CL_TRUE,
> -- 
> 1.9.1
> 
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet