# How is the dimensionality of an OpenCL problem specified?

When writing an OpenCL kernel, calls to get_global_id() are used to determine the "position" within the problem that a particular instance of the kernel is working on. However, I'm confused about how the dimensionality of a problem is specified or determined to begin with. I can write simple kernels that do the right thing, but I don't know why they're doing the right thing! When moving to more complex kernels to do real work, I'm going to need to understand what's actually happening.

Am I right in saying that if a kernel takes $n$ arguments that are qualified as __global, then I can call get_global_id(0), …, get_global_id( $n-1$ )? Am I further correct in saying that each __global argument is treated as a one-dimensional array and that Mathematica runs one kernel (i.e., one OpenCL thread) for each element of the longest such array (but that I can specify the number of threads as the last argument to the Mathematica function created by OpenCLFunctionLoad if I want to control the dimensionality of the problem)? Assuming this, and $k$ threads, presumably values returned by get_global_id(0) will range from 0 to $k-1$; but if so, what values could get_global_id(1), get_global_id(2), etc. take? How can this be specified when running a kernel via the OpenCLLink Mathematica functions?

Alternatively, does Mathematica look at the dimensionality of the largest __global argument, somehow pass this to the underlying OpenCL, such that get_global_id(0) return values that index the first dimension of this argument, get_global_id(1) return values that index the second dimension of this same argument, and so on? If this is the case, how do I specify the dimensionality of the problem when running a Mathematica function created by OpenCLFunctionLoad?

So, in short, how is the dimensionality of a problem specified and what values can be taken by get_global_id($i$)?

The maximum dimensionality you can have is 3 and so the maximum number you can pass to get_global_id() is 2. I'd only ever used this in 1Dand 2D before and understand what is going on. However, there appears to be a problem with Mathematica with 3D kernels.

Consider this source code for three kernels. They take an input array and output the values of get_global_id(0), get_global_id(1) and get_global_id(2) respectively.

Needs["OpenCLLink"];
srcXindex =
"__kernel void myKernel( __global mint * data, mint width, mint \
height,mint depth) {
int xIndex = get_global_id(0);
int yIndex = get_global_id(1);
int zIndex = get_global_id(2);
int index = xIndex + yIndex*width +zIndex*width*height;
if (xIndex < width && yIndex < height && zIndex<depth) {
data[index] = xIndex;
}
}";
srcYindex =
"__kernel void myKernel( __global mint * data, mint width, mint \
height,mint depth) {
int xIndex = get_global_id(0);
int yIndex = get_global_id(1);
int zIndex = get_global_id(2);
int index = xIndex + yIndex*width +zIndex*width*height;
if (xIndex < width && yIndex < height && zIndex<depth) {
data[index] = yIndex;
}
}";

srcZindex =
"__kernel void myKernel( __global mint * data, mint width, mint \
height,mint depth) {
int xIndex = get_global_id(0);
int yIndex = get_global_id(1);
int zIndex = get_global_id(2);
int index = xIndex + yIndex*width +zIndex*width*height;
if (xIndex < width && yIndex < height && zIndex<depth) {
data[index] = zIndex;
}
}";


Let's create OpenCL functions from this source

funXindex1D =
"myKernel", {{_Integer}, _Integer, _Integer, _Integer}, 8];
funYindex1D =
"myKernel", {{_Integer}, _Integer, _Integer, _Integer}, 8];
funZindex1D =
"myKernel", {{_Integer}, _Integer, _Integer, _Integer}, 8];


Now, lets give them a 1D input array

width = 5;
height = 4;
depth = 2;
array1D = ConstantArray[0, {width}];

In:= funXindex1D[array1D, width, height, depth]

Out= {{0, 1, 2, 3, 4}}

In:= funYindex1D[array1D, width, height, depth]

Out= {{0, 0, 0, 0, 0}}

In:= funZindex1D[array1D, width, height, depth]

Out= {{0, 0, 0, 0, 0}}


The xIndex (and hence the value of get_global_id(0)) ranges from 0 to 4, as we'd expect while the value of yIndex (get_global_id(1)) and zIndex (get_global_id(2)) stay at 0..again as we'd expect since this is a 1D calculation.

Moving to 2D

Let's create a 2D input array and just hope that our 1D function works

In:= array2D = ConstantArray[0, {height, width}];

In:= funXindex1D[array2D, width, height, depth]

During evaluation of In:= OpenCLFunction::invgrdblk: OpenCLLink was called with invalid grid dimensions. The kernel cannot be launched with {4,5} set as the grid dimensions, while {8} set as the block dimensions. >>

Out= OpenCLFunction["<>",
"myKernel", {{_Integer}, "Integer64", "Integer64",
"Integer64"}][{{0, 0, 0, 0, 0}, {0, 0, 0, 0, 0}, {0, 0, 0, 0,
0}, {0, 0, 0, 0, 0}}, 5, 4, 2]


No such luck. We've explicitly set our block dimensions to be {8} and yet Mathematica has detected that the grid dimensions are {4,5}...i.e. that we have a 2D problem. So, we create some 2D functions from the same kernel source code simply by changing our specification of the block dimensions.

funXindex2D =
"myKernel", {{_Integer}, _Integer, _Integer, _Integer}, {8, 8}];
funYindex2D =
"myKernel", {{_Integer}, _Integer, _Integer, _Integer}, {8, 8}];
funZindex2D =
"myKernel", {{_Integer}, _Integer, _Integer, _Integer}, {8, 8}];


Running these on our 2D array is instructive.

In:= funXindex2D[array2D, width, height, depth] (*Values given by get_global_id(0)*)

Out= {{{0, 1, 2, 3, 4}, {0, 1, 2, 3, 4}, {0, 1, 2, 3, 4}, {0, 1,
2, 3, 4}}}

In:= funYindex2D[array2D, width, height, depth] (*values given by get_global_id(1)*)

Out= {{{0, 0, 0, 0, 0}, {1, 1, 1, 1, 1}, {2, 2, 2, 2, 2}, {3, 3,
3, 3, 3}}}

In:= funZindex2D[array2D, width, height, depth] (*values given by get_global_id(2)*)

Out= {{{0, 0, 0, 0, 0}, {0, 0, 0, 0, 0}, {0, 0, 0, 0, 0}, {0, 0,
0, 0, 0}}}


Moving to 3D

At this point, you may feel like you understand what's going on and that we are almost ready to see the zIndex change by using a 3D array. Let's build one

width = 5;
height = 4;
depth = 2;
array3D = ConstantArray[0, {depth, height, width}];


Recall that when we switched from 1D->2D, we had to change from {8} to {8,8} in the call to OpenCLFunctionLoad. Otherwise, we received an error. So, I expect the following not to work:

In:= funXindex2D[array3D, width, height, depth]

Out= {{{{0, 1, 2, 3, 4}, {0, 1, 2, 3, 4}, {0, 1, 2, 3, 4}, {0, 1,
2, 3, 4}}, {{0, 0, 0, 0, 0}, {0, 0, 0, 0, 0}, {0, 0, 0, 0, 0}, {0,
0, 0, 0, 0}}}}


The 2D function worked on a 3D array. Well, it worked in the sense that there was no error but it appears to be ignoring the second half of my 3D array. Same thing happens when we look at yIndex and zIndex:

In:= funYindex2D[array3D, width, height, depth]

Out= {{{{0, 0, 0, 0, 0}, {1, 1, 1, 1, 1}, {2, 2, 2, 2, 2}, {3, 3,
3, 3, 3}}, {{0, 0, 0, 0, 0}, {0, 0, 0, 0, 0}, {0, 0, 0, 0,
0}, {0, 0, 0, 0, 0}}}}

In:= funZindex2D[array3D, width, height, depth]

Out= {{{{0, 0, 0, 0, 0}, {0, 0, 0, 0, 0}, {0, 0, 0, 0, 0}, {0, 0,
0, 0, 0}}, {{0, 0, 0, 0, 0}, {0, 0, 0, 0, 0}, {0, 0, 0, 0,
0}, {0, 0, 0, 0, 0}}}}


zIndex is clearly set to 0 the whole time. Behavior so far is as if Mathematica hasn' t noticed that we are passing a 3 D array. Let' s be explicit

In:= funZindex3D =
"myKernel", {{_Integer, 3}, _Integer, _Integer, _Integer}, {8, 8, 8}];

In:= funZindex3D[array3D, width, height, depth]

During evaluation of In:= OpenCLFunction::invgrdblk:
OpenCLLink was called with invalid grid dimensions.
The kernel cannot be launched with {2,4} set as the grid dimensions,
while {8,8,8} set as the block dimensions. >>

Out=
OpenCLFunction["<>",
"myKernel", {{_Integer, 3}, "Integer64", "Integer64",
"Integer64"}][{{{0, 0, 0, 0, 0}, {0, 0, 0, 0, 0}, {0, 0, 0, 0,
0}, {0, 0, 0, 0, 0}}, {{0, 0, 0, 0, 0}, {0, 0, 0, 0, 0}, {0, 0, 0,
0, 0}, {0, 0, 0, 0, 0}}}, 5, 4, 2]
`

Here, I believe, we seen the root of the problem. Mathematica thinks the grid dimensions are {2, 4} rather than {2, 4, 5}. I can't see how to get around this and wonder if it's a bug. I'm going to contact Wolfram Support to see if they can shed any light on the issue.

Regardless of the 3D case, I hope this sheds some light into 1D and 2D uses of get_global_id()

• Thanks so much. I had run pretty much this exact same experiment (and saw the exact same results), but was so confused about the programming model that I assumed my experiment was flawed! Out of interest, which GPU are you using? I'm using a MacBookPro11,1 with an Intel Iris GPU. Finally, out of interest, if you force your kernel to run on the CPU, do you see the same behavior? – Chris Feb 11 '14 at 17:39
• For the above answer, the GPU was an NVIDIA quadro 2000 and I was running Mathematica 9 on Windows. It was the only OpenCL capable device on that machine. I get identical results on my MacBook Air. Now, Mathematica on the MBA repoers 2 open – WalkingRandomly Feb 11 '14 at 17:52
• Now, Mathematica on the MBA reports 2 OpenCL devices: Intel i5-4250U and HD Graphics 5000. Seems I was using the HD-5000. When I force to use the CPU, I still get the same results. That's assuming that it really is running on the CPU. Never figured out how to check – WalkingRandomly Feb 11 '14 at 18:00
• Thanks. I guess that's good evidence that it's not just one specific GPU. – Chris Feb 12 '14 at 9:59