[Mesa-dev] [PATCH] clover, gallium: add PIPE_COMPUTE_CAP_MAX_THREADS_PER_BLOCK

[Francisco Jerez]

Christoph Bumiller <e0425955 at student.tuwien.ac.at> writes:

> This is not necessarily the product of MAX_BLOCK_SIZE[i].

Thanks, looks good.  I've pushed it along with a fix for
clEnqueueNDRangeKernel() to check that the work-group dimensions
specified by the user are within the limit set by this cap.

> ---
>  src/gallium/docs/source/screen.rst                |    5 +++++
>  src/gallium/include/pipe/p_defines.h              |    1 +
>  src/gallium/state_trackers/clover/api/device.cpp  |    3 ++-
>  src/gallium/state_trackers/clover/core/device.cpp |    7 +++++++
>  src/gallium/state_trackers/clover/core/device.hpp |    1 +
>  5 files changed, 16 insertions(+), 1 deletions(-)
>
> diff --git a/src/gallium/docs/source/screen.rst b/src/gallium/docs/source/screen.rst
> index d912dc6..4e82d34 100644
> --- a/src/gallium/docs/source/screen.rst
> +++ b/src/gallium/docs/source/screen.rst
> @@ -203,6 +203,11 @@ pipe_screen::get_compute_param.
>    units.  Value type: ``uint64_t []``.
>  * ``PIPE_COMPUTE_CAP_MAX_BLOCK_SIZE``: Maximum block size in thread
>    units.  Value type: ``uint64_t []``.
> +* ``PIPE_COMPUTE_CAP_MAX_THREADS_PER_BLOCK``: Maximum number of threads that
> +  a single block can contain.  Value type: ``uint64_t``.
> +  This may be less than the product of the components of MAX_BLOCK_SIZE and is
> +  usually limited by the number of threads that can be resident simultaneously
> +  on a compute unit.
>  * ``PIPE_COMPUTE_CAP_MAX_GLOBAL_SIZE``: Maximum size of the GLOBAL
>    resource.  Value type: ``uint64_t``.
>  * ``PIPE_COMPUTE_CAP_MAX_LOCAL_SIZE``: Maximum size of the LOCAL
> diff --git a/src/gallium/include/pipe/p_defines.h b/src/gallium/include/pipe/p_defines.h
> index edcca23..2b6d1be 100644
> --- a/src/gallium/include/pipe/p_defines.h
> +++ b/src/gallium/include/pipe/p_defines.h
> @@ -545,6 +545,7 @@ enum pipe_compute_cap
>     PIPE_COMPUTE_CAP_GRID_DIMENSION,
>     PIPE_COMPUTE_CAP_MAX_GRID_SIZE,
>     PIPE_COMPUTE_CAP_MAX_BLOCK_SIZE,
> +   PIPE_COMPUTE_CAP_MAX_THREADS_PER_BLOCK,
>     PIPE_COMPUTE_CAP_MAX_GLOBAL_SIZE,
>     PIPE_COMPUTE_CAP_MAX_LOCAL_SIZE,
>     PIPE_COMPUTE_CAP_MAX_PRIVATE_SIZE,
> diff --git a/src/gallium/state_trackers/clover/api/device.cpp b/src/gallium/state_trackers/clover/api/device.cpp
> index 0376751..1a9127b 100644
> --- a/src/gallium/state_trackers/clover/api/device.cpp
> +++ b/src/gallium/state_trackers/clover/api/device.cpp
> @@ -87,7 +87,8 @@ clGetDeviceInfo(cl_device_id dev, cl_device_info param,
>                                       dev->max_block_size());
>  
>     case CL_DEVICE_MAX_WORK_GROUP_SIZE:
> -      return scalar_property<size_t>(buf, size, size_ret, SIZE_MAX);
> +      return scalar_property<size_t>(buf, size, size_ret,
> +                                     dev->max_threads_per_block());
>  
>     case CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR:
>        return scalar_property<cl_uint>(buf, size, size_ret, 16);
> diff --git a/src/gallium/state_trackers/clover/core/device.cpp b/src/gallium/state_trackers/clover/core/device.cpp
> index 8390f3f..661f33a 100644
> --- a/src/gallium/state_trackers/clover/core/device.cpp
> +++ b/src/gallium/state_trackers/clover/core/device.cpp
> @@ -143,6 +143,13 @@ _cl_device_id::max_block_size() const {
>     return get_compute_param<uint64_t>(pipe, PIPE_COMPUTE_CAP_MAX_BLOCK_SIZE);
>  }
>  
> +size_t
> +_cl_device_id::max_threads_per_block() const {
> +   return
> +      get_compute_param<uint64_t>(pipe,
> +                                  PIPE_COMPUTE_CAP_MAX_THREADS_PER_BLOCK)[0];
> +}
> +
>  std::string
>  _cl_device_id::device_name() const {
>     return pipe->get_name(pipe);
> diff --git a/src/gallium/state_trackers/clover/core/device.hpp b/src/gallium/state_trackers/clover/core/device.hpp
> index 8f284ba..b69e61f 100644
> --- a/src/gallium/state_trackers/clover/core/device.hpp
> +++ b/src/gallium/state_trackers/clover/core/device.hpp
> @@ -57,6 +57,7 @@ public:
>     cl_uint max_const_buffers() const;
>  
>     std::vector<size_t> max_block_size() const;
> +   size_t max_threads_per_block() const;
>     std::string device_name() const;
>     std::string vendor_name() const;
>     std::string ir_target() const;
-------------- next part --------------
A non-text attachment was scrubbed...
Name: not available
Type: application/pgp-signature
Size: 229 bytes
Desc: not available
URL: <http://lists.freedesktop.org/archives/mesa-dev/attachments/20120512/806959af/attachment.pgp>