What exactly does thread_position_in_grid mean in a metal compute kernel shader?

Question

Although this is an inappropriate use of a compute shader I was doing some experiments to determine if I could use one to produce the general UV gradient where one channel of the image goes linearly from 0-1 across the x axis and the other channel goes from 0-1 across the y axis of the image. However I became confused when I generated this image by varying the b value of a texture by the thread_position_in_grid.x value divided by the image width. I edited the pixel of the texture at the thread_position_in_grid position:

Yes it was a gradient but it certainly did not appear to be the gradient from 0-1 that I wanted. I dropped it into an image editor and sure enough it was not linear. (The part added below shows what a linear gradient from 0-1 would look like)

It would appear that I do not understand what exactly the thread_position_in_grid value means. I know it has something to do with the threads per thread-groups and thread execution width but I dont exactly understand what. I suppose my end goal is to know whether it would be possible to generate the gradient below in a compute shader however I dont understand what is going on.

For reference I was working with a 100x100 texture with the following thread settings. Really I dont know why I use these values but this is what I saw recommended somewhere so I am sticking with them. I would love to be able to generalize this problem to other texture sizes as well including rectangles.

let w = greenPipeline.threadExecutionWidth
let h = greenPipeline.maxTotalThreadsPerThreadgroup / w
let threadsPerThreadgroup = MTLSizeMake(w, h, 1)
let threadgroupsPerGrid = MTLSize(width: (texture.width + w - 1) / w,
                                  height: (texture.height + h - 1) / h,
                                  depth: 1)
encoder.dispatchThreadgroups(threadgroupsPerGrid, threadsPerThreadgroup: threadsPerThreadgroup)

And my shader looks like this:

kernel void green(texture2d<float, access::write> outputTexture [[texture(0)]],
                             uint2 position [[thread_position_in_grid]])
{
    if (position.x >= outputTexture.get_width() || position.y >= outputTexture.get_height()) {
        return;
    }
    outputTexture.write(float4(position.x / 100.0, 0.0, 0.0, 0.0), position);
}

Two things about this shader confuse me because I cannot explain them:

I am using position as the coordinate to write to on the texture so it bothers me that position doesnt work to generate the gradient.
You cannot reaplace position.x / 100.0 value with position.x / outputTexture.getWidth() even though it should also be 100. Doing so causes a black image. Yet when I made a shader that colored everything with outputTexture.getWidth() as its value it did indeed shade everything to a value equivalent to 100 (or more accurately 101 because of rounding)
It is ok to use position to check if the kernel is within bounds but not to create the UV gradient.

What is going on?

Issue 2 is because of integer arithmetic. Both position.x and outputTexture.get_width() are integer types, and the latter is larger than the former, so the division always produces zero. Try position.x / float(outputTexture.get_width()). For issue 3, it's definitely recommended to check if the position is within bounds, but I'm not sure if that's what you're asking. What do you mean by "but not to create the UV gradient". Issue 1 is probably an RGB/sRGB mismatch somewhere, but hard to know where. — Ken Thomases

Matthijs Hollemans Matthijs Hollemans · Accepted Answer · 2018-03-22T15:45:57

The thread_position_in_grid means whatever you want it to mean because you decide how large the grid is and what each thread in the grid does.

In your example, thread_position_in_grid is the pixel coordinate in the texture because your grid size is equal to the number of pixels in the texture (rounded up to a multiple of the pipeline's max thread execution width).

You can see this if you change the threadGroupsPerGrid to:

let threadgroupsPerGrid = MTLSize(width: (texture.width/2 + w - 1) / w,
                                  height: (texture.height/2 + h - 1) / h,
                                  depth: 1)

Now only the top quarter of your texture should be filled in because the grid only covers half the texture's width and height.

As to why your texture looks weird, it's probably related to the pixel format. After all, you're writing into the red color component and your texture comes out as blue.

What exactly does thread_position_in_grid mean in a metal compute kernel shader?

1 Answers