How to make metal shaders not clamp texture writes to (0-1) range?

Question

The input to this metal shader is a bgra image from iOS camera (in 0-1 range) and output is supposed to be rgba scaled to 0-255 range (note the x256 in kernel). But I don't understand why the output of the kernel appears to be getting clamped back to 0-1 range. Does Metal bake a clamp into their texture write? Or am I missing something?

kernel void bgraScaleKernel(
        texture2d<float, access::read> inputImage [[texture(0)]],
        texture2d<float, access::write> outputImage [[texture(1)]],
        const device int *starting [[ buffer(0) ]],
        uint2 gid [[thread_position_in_grid]])
{
    uint i = gid.x+starting[0]; 
    uint j = gid.y+starting[1]; 
    float4 pixel_new = (inputImage.read(uint2 (i, j)).bgra)*256;
    outputImage.write(pixel_new, (uint2)(i,j));
};

I have a bunch of other filter layers that are applied after this and they assume rgba (0-255) float values (its ported from OpenCL and changing it will require changing Androi, Windows and Mac versions too). So I need the output of the metal layers to be in 0-255 range.

What is the pixel format of the texture that you're binding to texture(1)? — Ken Thomases
Yeah, Ken. That was the problem. I was using MTLPixelFormatRGBA8Unorm. MTLPixelFormatRGBA32Float doesn't do that clamping. I don't understand why they clamp the write values. Seems a little controlling. Shouldn't that be up to the programmer? I mean I had to switch everything to RGBA32Float, except the last layer - because of course the texture to image function only takes unorms. So I had pass a isFinalLayer flag all the way into the texture construction part of the code. Just seems unnecessary to me, but maybe because I'm more used to OCL. — Hashman
Since you are writing a custom filter I guess you already have found sucess with what I am stuck at. I am trying to apply a simple vignette filter to a live camera feed using metal. The results are pretty slow and laggy, please check this if you can tell me what is missing:stackoverflow.com/q/53898780/1364053 — nr5

Ken Thomases Ken Thomases · Accepted Answer · 2018-04-08T16:06:38

An 8Unorm pixel format represents values in memory using integers in the range [0, 255], but in shaders those values are presented as floating point numbers in the range [0.0, 1.0].

To achieve the result you apparently desired, you simply had to assign the floating point value you read from inputImage directly to outputImage without multiplying by 256 (or 255). Metal does that for you. By doing it yourself, you made most values greater than 1.0 and thus Metal clamped them to 1.0 before doing its own conversion to 8Unorm.

How to make metal shaders not clamp texture writes to (0-1) range?

1 Answers