Запутался в thread_position_in_grid

#metal #compute-shader

#Металлические #вычислительный шейдер

Вопрос:

Я работаю над вычислительным шейдером в 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 для остальной части изображения, вот так:

2D градиент, в котором красный канал увеличивается с 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:amp;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 в каждом направлении:

Градиент, простирающийся от черного до 50% красного по горизонтали и от черного до 50% зеленого по вертикали.

Комментарии:

1. Вам нужно показать код приложения, который отправляет этот вычислительный шейдер, особенно размер группы потоков и количество потоков (групп). Кроме того, как создается текстура. (И покажите объявление outTexture в вашем вычислительном шейдере. Всегда показывайте реальный код, поскольку отредактированный код, скорее всего, не отражает вашу реальную проблему.)

2. Хорошо, я добавил их выше. Дайте мне знать, если чего-то не хватает.

3. Я бы предложил использовать threads_per_grid вместо жестко закодированных значений, а затем проверить, совпадают ли выходные данные.

4. Интересно! Это приводит к правильному выстраиванию различных плиток, поэтому я получаю плавный градиент слева направо и сверху вниз, но вместо того, чтобы получать градиент от 0 до 1 в каждом направлении, это все равно градиент от 0 до 0,5. Почему?

5. Что, если вы сделаете half4 color = half4((float)gid.x / (float)outTexture.get_width(), (float)gid.y / (float)outTexture.get_height(), 0.0, 1.0); . Кроме того, вы должны убедиться, что gid.x и gid.y никогда не превышают ширину / высоту выходной текстуры, иначе вы в конечном итоге будете записывать вне памяти текстур, и произойдут плохие вещи. (Обратите внимание, что 360 не является целым числом, кратным 16.)

Ответ №1:

Со мной происходит очень похожая вещь. Значение thread_position_in_grid , похоже, ограничено небольшим диапазоном вместо всей сетки (возможно, только размером threadgroup ). Короче говоря, это, вероятно, потому, что вы вызываете

 _commandEncoder.dispatchThreads(threadGroupCount, threadsPerThreadgroup: threadGroupSize)
  

вместо

 _commandEncoder.dispatchThreadgroups(threadGroupCount, threadsPerThreadgroup: threadGroupSize)
  

Я замечаю, что атрибут thread_position_in_grid приводит к другому значению в этих функциях. Не уверен, что это предполагаемое поведение, поскольку я не могу найти соответствующее описание в документации, и я ожидал, что этот атрибут относится к позиции во всей сетке. Кроме того, Metal определит количество threadgroup при использовании dispatchThreads() и может создать неоднородную threadgroup, которая, возможно, имеет какое-то отношение к проблеме.

dispatchThreads(_:threadsPerThreadgroup:)

Используйте этот метод, только если устройство поддерживает неоднородные размеры threadgroup. Смотрите Таблицы набора функций металла. Этот метод кодирует вызов отправки, который указывает произвольное количество потоков в сетке (threadsPerGrid). Metal вычисляет необходимое количество групп потоков, при необходимости предоставляя частичные группы потоков. Когда команда вычисления кодируется, все необходимые ссылки на параметры или ресурсы, ранее установленные в кодере, записываются как часть команды. После кодирования команды вы можете безопасно изменить состояние кодирования, чтобы настроить параметры, необходимые для кодирования других команд.