openclbank-conflict

Bank conflict in 2D kernel

Suppose our hardware has 32 banks of 4 byte width. And we have a 1D kernel of size 32, and a local 1D array of ints.

Then, ensuring that each consecutive thread accesses consecutive memory locations in the array should avoid bank conflicts.

But, suppose we have an 8 x 4 2D kernel and the same 1D array. How can I ensure that there are no bank conflicts? How do we define "consecutive thread" for a 2D array?

Solution

  • You can get the same global work-item IDs that you get in the 1D case with get_global_id(0) in the 2D case with this code:

    get_global_id(1) * get_global_size(0) + get_global_id(0);

    Just change the globals to locals if you want to get local work-item ID within a work-group.