Является ли ядро ​​металла атомным? - PullRequest
Новая
       7

Является ли ядро ​​металла атомным?

1 голос
/ 23 апреля 2019

Может ли произойти переключение контекста между строками внутри функции ядра?

Поскольку я устанавливаю некоторые значения перед внесением изменений, я хочу убедиться, что значение установлено, изменение было сделано.

1 Ответ

1 голос
/ 23 апреля 2019

Короткий ответ: да, переключение контекста, скорее всего, произойдет «между строк».В этом весь смысл переключения контекста: если некоторым строкам в вашем шейдере (независимо от фрагмента, вершины или ядра) нужен какой-то ресурс, который еще не доступен (ALU, модуль специальной функции, блоки текстуры, память), GPU наверняка переключитсяконтексты.Это называется скрытие задержки , и это очень важно для производительности графических процессоров, поскольку без него ядра графического процессора будут тратить большую часть времени на остановку для различных ресурсов, упомянутых выше.И все это означает, что функции ядра Metal определенно не атомарны.

Что касается вашей проблемы, если вы хотите, чтобы что-то происходило атомарно, есть два основных способа сделать это на языке Metal Shading Language:

  1. Вы можете использовать атомарные типы и функции из заголовка metal_atomic.Это подмножество заголовка C ++ 14 atomic, которое содержит атомарные функции сохранения, загрузки, обмена, сравнения, обмена, выборки и изменения.
  2. Вы можете использовать барьеры SIMD-группы и группы потоков.Барьеры позволяют подождать, пока все потоки в группе выполнят все операции, прежде чем любому потоку будет разрешено продолжить.В зависимости от флагов, передаваемых в функцию барьера, барьер может также упорядочить доступ к памяти или просто использоваться в качестве барьера выполнения (что, вероятно, не то, что вам нужно, если вы записываете некоторые данные из своего шейдера).

Вы можете обратиться к Спецификация языка металлического шейдинга для получения дополнительной информации об этих функциях (см. Раздел 5.8.1 для получения информации о «функциях синхронизации группы нитей и SIMD-группы» и 5.13 для получения информации о «атомарных функциях»).

Обычно, если ваша функция ядра обрабатывает некоторые данные, которые вам позже понадобится уменьшить, вы должны сделать что-то вроде этого (это очень простой пример): 

kernel void
my_kernel(texture2d<half> src [[ texture(0) ]],
    texture2d<half, access::write> dst [[ texture(1) ]],
    threadgroup float *intermediate [[ threadgroup(0) ]],
    ushort2 lid [[ thread_position_in_threadgroup ]],
    ushort ti [[ thread_index_in_threadgroup ]],
    ushort2 gid [[ thread_position_in_grid ]])
{
    // Read data
    half4 clr = src.read(gid);

    // Do some work for each thread
    intermediate[ti] = intermediateResult;

    // Make sure threadhroup memory writes are visible to other threads
    threadgroup_barrier(mem_flags::mem_threadgroup);

    // One thread in the whole threadhgroup calculates some final result
    if (lid.x == 0 && lid.y == 0)
    {
        // Do some work per threadgroup
        dst.write(finalResult, gid);
    }
}

Здесь все потокив группе потоков считайте данные из src текстуры, выполните работу, сохраните промежуточный результат в памяти группы потоков, а затем вычислите и запишите окончательный результат в текстуру dst.threadgroup_barrier гарантирует, что другие потоки (включая поток с thread_position_in_threadgroup, равным (0, 0), который будет вычислять конечный результат) могут видеть записи в память.

...