Правый край металлической текстуры имеет аномалии - PullRequest
Новая
       29

Правый край металлической текстуры имеет аномалии

0 голосов
/ 26 января 2020

Когда я запускаю этот код на интегрированном графическом процессоре Intel на Macbook Pro, у меня нет проблем. Но когда я запускаю его на iMa c с графическим процессором AMD, этот простой «Hello World» дает мне артефакты по правому краю:

enter image description here

Шейдер очень прост: 

kernel void helloworld(texture2d<float, access::write> outTexture [[texture(0)]],
                     uint2 gid [[thread_position_in_grid]])
{
    outTexture.write(float4((float)gid.x/640,
                            (float)gid.y/360,0,1),
                     gid);
}

Я пытался просмотреть содержимое текстуры двумя разными способами, и оба вызывают проблемы:

Преобразование текстуры в CIImage и просмотр это в NSImageView, или вызывая getBytes и копируя данные пикселей напрямую и вручную создавая из него PNG (полностью пропуская CIImage). В любом случае получается этот странный артефакт, поэтому он действительно находится в самой текстуре.

Есть идеи, что вызывает проблемы такого рода?

ОБНОВЛЕНИЕ:

Увлекательно, проблема возникает быть связанным с threadsPerThreadgroup, но я не уверен, почему это так.

Приведенное выше изображение было создано с 24 потоками на группу. Если я изменю это на 16, артефакты переместятся к нижнему краю.

Что я не понимаю в этом, так это то, что позиция gid должна быть исправлена ​​независимо от того, какая группа потоков на самом деле работает, не должна Это? Потому что это положение отдельных нитей на всем изображении.

Ответы [ 2 ]

1 голос
/ 26 января 2020

С dispatchThreadgroups() ядро ​​вычислений может быть вызвано для позиций сетки за пределами вашей ширины * высоты сетки. Вы должны явно ничего не делать с чем-то вроде: 

if (gid.x >= 640 || gid.y >= 360)
    return;

В противном случае вы будете пытаться писать вне границ текстуры (с цветами, некоторые из компонентов которых больше 1). Это дает неопределенные результаты.

С dispatchThreads(), Metal позаботится об этом за вас и не будет вызывать за пределами указанного вами размера сетки.

Разница в поведении между 24 и 16 потоками в расчете на группу - делится ли это равномерно на 640 и 360. То, что не делится равномерно, - это измерение, которое вызывается чрезмерно.

0 голосов
/ 26 января 2020

Используя macOS 10.13 или более позднюю версию, можно дать ОС понять некоторые из этих вещей. Я использовал: 

commandEncoder.dispatchThreadgroups(threadgroupsPerGrid, threadsPerThreadgroup: threadsPerThreadgroup)

Делая это, я должен был сам вычислить threadgroupsPerGrid, который оказался источником проблемы.

Заменив этот вызов на: 

commandEncoder.dispatchThreads(MTLSize(width: Int(width), height: Int(height), depth: 1), threadsPerThreadgroup: threadsPerThreadgroup)

Проблемы исчезли.

...