I know this is a very common question but I cannot seem to figure this out.
These are the properties of my GPU and the kernel:
max work item dimensions: 3
max work item sizes:
value[0]: 1024
value[1]: 1024
value[2]: 64
kernel max work group size: 256
kernel preferred work group size multiple: 32
From what I read online “max work group size” is the maximum number of work items in one work group and “shape” of the group is given by dimension count (3) and max work item sizes (1024, 1024, 64).
What does that mean when kernel max work group size has only one value (one dimension) equal to 256? Is that the max number of work items per work group in all dimensions or in only one dimension (so the actual size is 256x256x256)?
Why are work item sizes all larger than the max work group size? I don’t understand why those are specified only to be limited by max work group size?
Working example: I have a batch of RGBA images (1600 x 600 x 4 bytes), say a 1000 of those (batch size = 1000). I’d like to write a simple kernel that takes each color component (RGBA) and puts it together with same components (so 1600x600 RGBA becomes 1600x600 R + 1600x600 G + …). Input and output size would be the same just the data layout would be different.
Given GPU and kernel specs above what would be a good way to split that data into work groups and work items? What would be the global work size vs local work size? I’d appreciate a thorough explanation, I’m trying to get better at OpenCL.
This is the kernel code I came up with, seems to work OK although I’m not sure if this is the optimal way to do this:
// split RGBA color components into individual color bands
__kernel void splitComponents(__global unsigned char *input, __global unsigned char *output) {
const uint wRgba = get_global_size(0);
const uint w = wRgba / 4;
const uint h = get_global_size(1);
const uint x = get_global_id(0);
const uint c = x % 4;
const uint y = get_global_id(1);
const uint imageIndex = get_global_id(2);
const uint imageOffset = wRgba * h * imageIndex;
const uint cOffset = w * h * c;
const uint src = imageOffset + y * wRgba + x;
const uint dst = imageOffset + cOffset + y * w + x / 4;
output[dst] = input[src];
}
In my test I used the following values (input dim = 3):
wRgba = 1600 * 4 (dim(0)), h = 600 (dim(1)) and image count = 1 (dim(2))
I did not use any local storage in this kernel function so I set local dim to NULL. I still don’t know why kernel group size is lower than the one supported by my GPU but I was not able to find anything on that.
I tested my code by splitting and merging color components back, for 1000 images that took 806.29 milliseconds (not counting time spent on data transfer from GPU to host, 3662 MB).
In terms of efficiency, you’re doing an awful lot of maths in the kernel just to load and store one byte at a time. It might be OK on some hardware, but it smells expensive.
Assuming data rows in the image are in linear layout and tightly packed without padding, my initial stab would be to make this a 1D problem and then update 4 color components from a single work item (passing in precomputed channel offsets).
You might get faster by loading 4 texels per work item, permuting in shader code, and then storing a vec4 of R, G, B and A, but not sure how well all GPUs will handle the permute.
Total size (x * y * z) must be less than the max workgroup size.
It’s a weird implementation quirk, I agree. However, nothing breaks if you follow all the spec requirements (i.e. in your case you would be limited to 256, even if the underlying feature reports 1024 for a single axis).
In my experience, individual axis limits are often underlying native hardware limits. The overall workgroup limit often has additional contraints imposed by the driver software, which means it can be lower than the harware limit.
Thank you for your reply. It makes sense that driver could add its own restrictions, that’s a good point.
I modified my code based on your example, I could not make it exactly the same because my example is a little different (bands are grouped within each image and not globally) but I did make use of char4 vector type and vload function. Average execution time changed from 0.6 ms to 0.52 ms, so that’s about 13% faster =)
I’ve never used vector types before so I’ll be reading up on those. Thanks to your comments I found a tutorial on how to use those along with local kernel memory for further optimization, I’ll see if I can make sense of it.