Pixel Coordinates, Texture Coordinates and gid within Kernel

Question

Let's say I have a texture of size 100 in width and 100 in height and dispatch it to the kernel function in {10, 10, 1} thread groups and size.

I am having trouble understanding whether thread_position_in_grid goes from 0 - 9 or 0 - 99 (as the size of my texture is 100 x 100)?

I am using float4 c = in.read(gid); to retrieve the color of the texture, but I would like to know how gid is mapped to the texture coordinates to retrieve the result which is of type float4 (would there be cases where it is between pixel (texture?) coordinates?).

I would like to understand in depth how the above works, because what I would like to achieve is to be able to retrieve the exact position as defined by the size of my texture in the kernel function; i.e.:

Dispatch size: {10, 10, 1} and texture size 100 x 100; and retrieving values in kernel function:

0, 1, 2, 3, 4, 5..., 99.

In summary, how are gid, texture coordinates and pixel coordinates related within kernel function? How are they mapped from one to another? I have been reading the documentation and similar questions for the past 4 days, but I have not been successful in obtaining a concrete and authoritative answer.

Thank you.

Answer 1

In your example, you're dispatching a grid of work like this:

MTLSize threadsPerThreadgroup = { 10, 10, 1 };
MTLSize threadgroupCount = { 10, 10, 1 };
[computeEncoder dispatchThreadgroups:threadgroupCount
               threadsPerThreadgroup:threadsPerThreadgroup];

Your gid, with attribute thread_position_in_grid, is probably of type uint2, since the domain of the grid is two-dimensional. So far, so good.

gid will range from (0, 0) to (99, 99). How you map it to other quantities, including texture coordinates is entirely up to you.

For example, say you're not operating on a 100x100 texture, but instead a 300x300 texture, and you want to perform some operation on the lower-right corner. In that case, you might add uint2(200, 200) to gid before reading/writing the texture, in order to address the region of pixels with coordinates (200, 200) to (299, 299). Other, arbitrary transformations are just as possible.

Basically, specifying the grid dimensions is a convenient way for you to describe the logical shape of the operation you want to perform, while also (through threadsPerThreadgroup) allowing you to optimize how much of that work can be executed in parallel.

As for (normalized) texture coordinates, in the simple case where your grid dimensions exactly match your texture dimensions, you can get the normalized coordinates corresponding to a kernel function invocation by dividing by a parameter that is qualified with the [[threads_per_grid]] attribute, remembering to cast first to avoid truncation:

float2 coords = gid / float2(tpg)

where tpg is another parameter to your kernel function declared as uint2 tpg [[threads_per_grid]].