Get the optimal local / global workgroup size in OpenCL?

Question

I use the following function to get the best local and workgroup size for my OpenCL application.

//maxWGSize == CL_KERNEL_WORK_GROUP_SIZE
//wgMultiple == CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE
//compUnits == CL_DEVICE_MAX_COMPUTE_UNITS
//rems == max required work items

void MyOpenCL::getBestWGSize(cl_uint maxWGSize, cl_uint wgMultiple, cl_uint compUnits, cl_uint rems, size_t *gsize, size_t *lsize) const
{
    cl_uint cu = 1;
    if(wgMultiple <= rems)
    {
        bool flag = true;
        while(flag)
        {
            if(cu < compUnits)
            {
                cu++;
                if((wgMultiple * cu) > rems)
                {
                    cu--;
                    flag = false;
                    break;
                }
            }
            else if(wgMultiple < maxWGSize)
            {
                wgMultiple *= 2;
                if((wgMultiple * cu) > rems)
                {
                    wgMultiple /= 2;
                    flag = false;
                    break;
                }
            }
            else
            {
                cu++;
                if(((wgMultiple * cu) > rems) || (cu > 2 * compUnits))
                {
                    cu--;
                    flag = false;
                    break;
                }
            }
        }
    }
    else
    {
        bool flag = true;
        wgMultiple = 2;
        while(flag)
        {
            if(cu < compUnits)
            {
                cu++;
                if((wgMultiple * cu) > rems)
                {
                    cu--;
                    flag = false;
                    break;
                }
            }
            else
            {
                wgMultiple *= 2;
                if((wgMultiple * cu) > rems)
                {
                    wgMultiple /= 2;
                    flag = false;
                    break;
                }
                else
                {
                    cl_int temp = rems - (wgMultiple * cu);
                    if((temp == 0) || (temp == 1))
                    {
                       flag = false;
                       break;
                    }
                }
            }
        }
    }

    *gsize = wgMultiple * cu;
    *lsize = wgMultiple;
    if(rems < *gsize)
    {
        *gsize = rems;
        *lsize = rems;
    }
    if(cu != compUnits)
    {
        while((cu * 2) <= compUnits)
        {
            cu *= 2;
            if(*lsize % 2 == 0)
                *lsize /= 2;
        }
    }
}

Algorithm:

• Determine how much workgroup is required if local size == CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE
• If more work units are required, multiply the local size by 2 until it reaches CL_KERNEL_WORK_GROUP_SIZE

Any suggestions for improving the algorithm?

Some results that I get:

for GPU if max required work items == 99
maxWGSize    256 
wgMultiple   64 
compUnits    6 
rems     99 
*gsize   64 
*lsize   16 


for GPU if max required work items == 35
maxWGSize    256 
wgMultiple   4 
compUnits    6 
rems     35 
*gsize   24 
*lsize   4 

for GPU if max required work items == 57
maxWGSize    256 
wgMultiple   8 
compUnits    6 
rems     57 
*gsize   48 
*lsize   8 

for CPU if max required work items == 99
maxWGSize    1024 
wgMultiple   16 
compUnits    4 
rems     99 
*gsize   64 
*lsize   16 

for CPU if max required work items == 35
maxWGSize    1024 
wgMultiple   8 
compUnits    4 
rems     35 
*gsize   32 
*lsize   8

for CPU if max required work items == 57
maxWGSize    1024 
wgMultiple   8 
compUnits    4 
rems     57 
*gsize   32 
*lsize   8 

Answer 1

, ( ) , , : , .

(, DarkZeros : .., null , OpenCL .

. OpenCL "" . : ( ), OpenCL 1...

, . , , , . :

__kernel void reduceURatios(
    __global myreal *coef, 
    __global myreal *row, 
    myreal ratio,
    int sizeOfArrays)  // Add this parameter
{
    size_t gid = get_global_id(0);
    if (gid >= sizeOfArrays)
    {
        return; // Don't access invalid elements
    }

    myreal pCoef = coef[gid];
    myreal pRow = row[gid];

    pCoef = pCoef - (pRow * ratio);
    coef[gid] = pCoef;
}

. CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE, . , CL_DEVICE_MAX_WORK_GROUP_SIZE ( , , , , - ).

, CL_DEVICE_MAX_WORK_GROUP_SIZE :

// As queried with CL_DEVICE_MAX_WORK_GROUP_SIZE
int maxWorkGroupSize = ...
int numWorkGroups = (n-1) / maxWorkGroupSize + 1;
int globalSizePadded = numWorkGroups * maxWorkGroupSize;

​​ () . if, , , . ​​ null, CL_DEVICE_MAX_WORK_GROUP_SIZE (, , ).

, , , ...