Originally posted by: binarysplit Thanks for the feedback. I don't really know what to refer to the work items' processors as. The AMD APP documentation calls them "Stream Processors", so I just guessed SPU would be an appropriate acronym. Is there a better acronym, or should I always write out "Stream Processor" or "Work Item"?
|
Ok sorry - i've used a CELL BE and SPU is what they call their processors. The glossary in the amd docs say spu = shader processing unit, but it's never used.
I normally use 'thread', or 'workitem', although 'thread' is easy to confuse with cpu threads, and 'workitem' also refers to all work items, and not just the per-workgroup ones.
I'm using a fixed 16 work items per work group as my dataset is a tree that needs to be walked in an approximately depth-first way. It's more efficient to do this with a stack than a queue, and even if it was a queue(which would only allow breadth-first walking), I need the work items to be able to create more work items, which would require a rather inefficient multi-pass process. But that's another story. I'll try a few different methods after I've got it up and running. |
16 threads will use less than 1/4 of the processing power on a gpu. You want at least 64 threads.
If the 16 really is set in stone, you could simply run 4 lots at once, either as 32x32 tile or 64x16. Just make sure each group of 16 works on the right set of data, and that any barriers are properly hit (all threads in the group have to hit any barrier).
I'll use an implementation like this for the copy: for(int i = get_local_id(0); i < (TILE_W, TILE_H); i += get_local_size(0)) { write_imageui(colorBufImg, (int2)(tileLeft + (i % TILE_W), tileTop + (i / TILE_W), localColorBuf); This'll let the compiler unroll the loop fully and pre-calculate most indices. }
|
If you know it's going to be 16x16 and you have 16 work items, just hardcode the size:
for (int y=0;y < TILE_H) {
write_imageui(colorBufImg, (int2)(tileLeft + lx, tileTop + y, localColorBuf[lx + y * TILE_W]);
}
Originally posted by: binarysplit Oh, one more thing I was wondering about. You said you just read a 16x16 tile with a 16x16 work group size. Does that mean there's a way to separate the "process data" and "copy cached data from local back to global" phases into multiple kernels? Or do your "process data" kernels run at 16x16 as well?
|
Generally i'd use 16x16 work items as well, but it depends on the problem. Even if you're not using them all for the work, it can be more efficient to use extra work items to do the memory operations, although it might not be so important for image accesses.
Data load/data process can only be separated by an 'if' since local memory isn't persistent. Or just changing the addressing for different parts of the code.
i.e. load the data using more threads, then just if out the ones not involved in calculation
load using more threads
if (lx < 16) {
// do work
}
Although if you're doing that you dont want any more than 64 threads as the i/o work gang.
Still, if you're generating a 16x16 tile, that means each of the 16 threads has to generate 16 pixels? Even if you save some of the calculations, you're throwing away a potential for 16x parallelism there (unless you run 4 of them as above).
Also if you do run say 64 threads, split into groups of 16 for the calculation, you could re-arrange the 64 threads however you liked for the data load/store. e.g. in 4x4 lots of 8x8 rather than 16 lots of (64)x1
