#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 для остальной части изображения, вот так:
Если я инвертирую 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 в каждом направлении:
Комментарии:
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 вычисляет необходимое количество групп потоков, при необходимости предоставляя частичные группы потоков. Когда команда вычисления кодируется, все необходимые ссылки на параметры или ресурсы, ранее установленные в кодере, записываются как часть команды. После кодирования команды вы можете безопасно изменить состояние кодирования, чтобы настроить параметры, необходимые для кодирования других команд.