Name

    ARM_non_uniform_work_group_size

Name Strings

    cl_arm_non_uniform_work_group_size

Contributors

    Scott Moyers, ARM
    Robert Elliott, ARM

Contact

    Robert Elliott, ARM (robert.elliott 'at' ARM.com)

IP Status

    No claims or disclosures are known to exist.

Version

    Revision: #1, Jun 5th, 2015

Number

    OpenCL Extension #39

Status

    Complete.

Extension Type

    OpenCL device extension

Dependencies

    Requires OpenCL version 1.0 or later.

Overview

    This extension provides a way to enqueue kernels with local work-group
    sizes that are not integer factors of the global work-group size in
    OpenCL C 1.x languages.

    Such work-groups are referred to in the OpenCL 2.0 specification as
    non-uniform work-groups.

    To enable this extension the option -cl-arm-non-uniform-work-group-size
    must be provided in the options string when building a program from
    source using clBuildProgram. Kernels created from such a program will
    be able to be enqueued via clEnqueueNDRangeKernel with a non-uniform
    local work-group size.

    This feature is enabled by default in OpenCL C 2.0. See section 5.10
    of the OpenCL 2.0 API specification. This section also details how kernels
    that are enqueued with non-uniform work-group sizes are divided into work
    groups.

    The built in function get_local_size() for kernels that have been built
    with this extension will take on the OpenCL 2.0 behaviour. See
    section 6.13.1 of the OpenCL 2.0 C specification for details.

Header File

    cl_ext.h

Glossary

    No new terminology is introduced by this extension.

New Types

    None

New Procedures and Functions

    None

New Tokens

    None

Interactions with other extensions

    None

Sample Code

        const char *source[] = { "kernel void example( void ){}\n" };

        cl_program program = clCreateProgramWithSource( context, 1, source,
                                                        NULL, NULL );

        clBuildProgram( program, 0, NULL, "-cl-arm-non-uniform-work-group-size",
                        NULL, NULL );

        cl_kernel kernel = clCreateKernel( program, "example", NULL );

        size_t global_work_size = 10;
        size_t local_work_size = 9;
        err = clEnqueueNDRangeKernel( queue, kernel, 1, NULL,
                                      &global_work_size, &local_work_size,
                                      0, NULL, NULL );
        assert( CL_SUCCESS == err );

Conformance Tests

    None

Revision History

    Revision: #1, Jun 5th, 2015 - Initial revision