Work Group Sizes

Question

For a given Kernel, why are work_groups of always the same size? I read somewhere (for the case in which we don't specify the local work size) that openCL creates 3 work groups(of 217 work-items each) for kernel with 651 work-items(divisible by 3) while it creates 653 work-groups of 1 work-item each, as 653 is a prime number.

Suppose we specify the local_work_size(i.e. no. of work-items in a work-group), let's say,5. And we have given the total work-items(global_work_size) as 9. How will the work groups be created? is this why the global_work_size have to be a multiple of local_work_size? If the data requires only 9 work-items, how do I increase it to 10(multiple of local_work_size,5)?

Why can't host allocate the memory for result array if it doesn't know how many work groups will execute the kernel?

Please help. I read all this on this: http://www.openclblog.com/2011/09/work-group-sizes.html

Goddaughter

Answer 1

OpenCL Work groups sizes don't need to be always the same size. The Global work group size is frequently related to the problem size. The Local Work Group Size is selected based on maximizing Compute Unit throughput and the number of threads that need to share Local Memory.

Let consider a couple of examples;

A) Scale a image from N by M to X by Y.

B) Sum N numbers.

For A)

The obvious Global Work Group Size is X ,Y, 1. Why? This gives 1 thread per out pixel. The Local Work Group Size should be chosen based on the number of Input pixels that need to be processed to generate an output pixel.

Eg.

A.1)Scale an image from 4K by 3.2K to 64 by 64. GWG Size [64,64,1] LWG Size 256 A.2)Scale an image from 4k by 3.2k to 800 by 600.GWG Size [800,60,1] LWG Size 256

For B)

The obvious Global Work Group Size is N/2,1,1, Why? So each thread starts by summing 2 values together. The Local Work Group should be set to the device max.

There are some caveats;

1) Global Work Group Size is constrained by the Global Memory Size and the Max Global Memory Allocation size.

2) Each device has a max Local Work group size often 256

Whaleback

Thank you Tim!!Thanks for answering. :) So, the GWG size Doesn't have to be a multiple of LWG size?? – Goddaughter

To clarify yes the LWG size must be a mukltiple or a null. From the Man pages khronos.org/registry/cl/sdk/1.1/docs/man/xhtml – Whaleback

Since the GWG size is related to the problem size (suppose in pattern matching, I take it (Length_of_text)-(Length_of_pattern)) And I take the LWG size as Length_of_pattern(only 1 word). But IF the GWG size is not a multiple of LWG,and since it must be a multiple, I cannot take the LWG to be Length_of_pattern.How should I pass the Pattern in __Local then? – Goddaughter

You can't pass in data as __local. You can pass in data as global and have some threads copy it to local. If the GWG size is not a multiple of the LWG size, make it the next highest multiple and pass in a constant that is the problem size or the number of threads Eg. pass in N as constant int and test if ( N< get_global_id(0) ) ... – Whaleback

If the pattern is immutable ( not changing) during the kernel execution you could pass it as a constant argument . Typically its faster on a GPU. – Whaleback