Two-Dimensional NDRange

These concepts can be extended to a two-dimensional NDRanges. This type of NDRange works well with two-dimensional data such as matrices. Consider the following matrix adder kernel:

__kernel __attribute__ ((reqd_work_group_size(2, 2, 1)))
void madd(__global int* a, __global int* b, __global int* output) {
 int index = get_global_id(1)*get_global_size(0) + get_global_id(0);

 output[index] = a[index] + b[index];
}

This kernel defines a local work size of 2x2, specified as a required size of (2, 2, 1). The calls to get_global_id() provide the index in the global work size, while get_global_size() provides the total range value (e.g., 64 for a 64x64 matrix). Alternatively, the kernel could also index the local work indices and sizes using get_local_id() and get_local_size().

Figure: Two-Dimensional NDRange

The preceding figure illustrates how this two-dimensional space is defined and indexed. While the ND range is 64x64x1, the local work size is 2x2x1. Similar to the one-dimensional work size, this enables simpler coding as well as concurrent implementation across the vast resources of the FPGA.