Запутался по поводу thread_position_in_grid

Question

Я работаю над вычислительным шейдером в Metal для macOS.Я пытаюсь сделать некоторые очень простые вещи, чтобы узнать, как они работают.Я вижу некоторые результаты, которые я не понимаю.Я думал, что начну с попытки создать простой 2D-градиент.Красный канал увеличится с 0 до 1 по ширине, а зеленый канал увеличится с 0 до 1 по высоте.Итак, я написал это ядро:

kernel void myKernel(texture2d<half, access::write> outTexture [[ texture(MBKT_OutputTexture) ]],
                     uint2  gid  [[thread_position_in_grid]])
{
    half4  color = half4((float)gid.x / 480.0, (float)gid.y / 360.0, 0.0, 1.0);

    outTexture.write(color, gid);
}

И я получаю увеличение от 0 до 0,5 на полпути и твердое 0,5 для остальной части изображения, например:

Если я инвертирую 2 значения, чтобы ядро ​​рассчитало это:

half4  color = half4(1.0 - (float)gid.x / 480.0, 1.0 - (float)gid.y / 360.0, 0.0, 1.0);

результаты будут еще более странными.Я бы ожидал, что он будет равен 1,0 слева и снизу и снизится до 0,5 в середине, но вместо этого я получу это:

Чтоздесь происходит?В первом случае это как все, что выше средней точки, имеет значение 0,5.Во втором случае это похоже на левый / нижний край, равный 0,5, а середина - на 1,0, затем переворачивается на 0,0 через пиксель.

Странно, если я использую thread_position_in_grid, чтобы извлечь значения из буферов,это работает правильно.Например, я могу вычислить множество Мандельброта, и результаты верны.Но меня смущает то, что происходит с простым ядром выше.Кто-нибудь может мне это объяснить?

Вот мой код установки вычислительного ядра в MTKViewDelegate.Это основано на примере кода "Hello Compute" от Apple:

    _metalView = metalView;
    _device = metalView.device;
    _commandQueue = [_device newCommandQueue];

    _metalView.colorPixelFormat = MTLPixelFormatBGRA8Unorm_sRGB;

    // Load all the shader files with a .metal file extension in the project
    id<MTLLibrary> defaultLibrary = [_device newDefaultLibrary];

    // Load the kernel function from the library
    id<MTLFunction> kernelFunction = [defaultLibrary newFunctionWithName:@"myKernel"];

    // Create a compute pipeline state
    NSError*    error   = nil;
    _computePipelineState = [_device newComputePipelineStateWithFunction:kernelFunction
                                                                   error:&error];

    if(!_computePipelineState)
    {
        NSLog(@"Failed to create compute pipeline state, error %@", error);
        return nil;
    }

А вот код, в котором я создаю текстуру вывода и группы потоков:

MTLTextureDescriptor*   outputTextureDescriptor = [MTLTextureDescriptor texture2DDescriptorWithPixelFormat:MTLPixelFormatBGRA8Unorm_sRGB
                                                                                                     width:_viewportSize.x
                                                                                                    height:_viewportSize.y
                                                                                                 mipmapped:NO];
_outputTexture = [_device newTextureWithDescriptor:outputTextureDescriptor];

// Set the compute kernel's threadgroup size of 16x16
_threadgroupSize = MTLSizeMake(16, 16, 1);

// Calculate the number of rows and columns of threadgroups given the width of the input image
// Ensure that you cover the entire image (or more) so you process every pixel
_threadgroupCount.width  = (_viewportSize.x + _threadgroupSize.width - 1) / _threadgroupSize.width;
_threadgroupCount.height = (_viewportSize.y + _threadgroupSize.height - 1) / _threadgroupSize.height;

// Since we're only dealing with a 2D data set, set depth to 1
_threadgroupCount.depth = 1;

В моих тестах_viewportSize равно 480 x 360.

Я сделал дополнительный тест, предложенный @Egor_Shkorov в комментариях.Вместо жесткого кодирования 480 и 360, я использовал переменную threads_per_grid:

kernel void myKernel(
                             texture2d<half, access::write> outTexture [[ texture(MBKT_OutputTexture) ]],
                             uint2  gid  [[thread_position_in_grid]],
                             uint2 tpg [[threads_per_grid]])
{

    half4  color = half4((float)gid.x / tpg.x, (float)gid.y / tpg.y, 0.0, 1.0);

    outTexture.write(color, gid);
}

, которая улучшает положение вещей, заставляя градиент растягиваться полностью в каждом направлении, но он все еще идет только от 0 до 0,5вместо 1 в каждом направлении: