Christoph Bumiller <e0425...@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;
pgpFnaog5iTvB.pgp
Description: PGP signature
_______________________________________________ mesa-dev mailing list mesa-dev@lists.freedesktop.org http://lists.freedesktop.org/mailman/listinfo/mesa-dev
