hlsl CG compute shader Условие гонки

Я пытаюсь преобразовать текстуру в частотную область с помощью вычислительного шейдера в единстве / CG / hlsl, т.е. я пытаюсь прочитать значения пикселей из текстуры и вывести массив коэффициентов базовой функции. как бы я это сделал? Я действительно новичок в вычислении шейдеров, поэтому я немного потерялся. Я понимаю причину состояния гонки и то, как вычислительные шейдеры разделяют рабочую нагрузку, но есть ли способ с этим справиться? В целом документация по буферам и другим вещам кажется немного разочаровывающей для тех, кто не разбирается в этом вопросе.

ошибка, которую я получаю: Shader error in 'Compute.compute': race condition writing to shared resource detected, consider making this write conditional. at kernel testBuffer at Compute.compute(xxx) (on d3d11)

упрощенный пример может заключаться в суммировании всех значений пикселей, в настоящее время мой подход будет следующим. Я пытаюсь использовать структурированные буферы, так как я не знаю, как еще я смогу впоследствии получить данные или сохранить их на графическом процессоре для доступа к глобальному шейдеру ??

struct valueStruct{
float4 values[someSize];
}

RWStructuredBuffer<valueStruct> valueBuffer;

// same behaviour if using RWStructuredBuffer<float3> valueBuffer;
// if using 'StructuredBuffer<float3> valueBuffer;' i get the error:
// Shader error in 'Compute.compute': l-value specifies const object at kernel testBuffer at Compute.compute(xxx) (on d3d11)

Texture2D<float4> Source;

[numthreads(8, 8, 1)]
void testBuffer(uint3 id : SV_DispatchThreadID) {

      valueBuffer[0].values[0] +=  Source[id.xy];  // in theory the vaules 
      valueBuffer[0].values[1] +=  Source[id.xy];  // would be different
      valueBuffer[0].values[2] +=  Source[id.xy];  // but it doesn't really 
      valueBuffer[0].values[3] +=  Source[id.xy];  // matter for this, so 
      valueBuffer[0].values[4] +=  Source[id.xy];  // they are just Source[id.xy]
      //.....

}

Все это не вызывает ошибку состояния гонки, если я разворачиваю буфер в отдельные значения, такие как

    float3 value0;
    float3 value1;
    float3 value2;
    float3 value3;
    float3 value4;
    float3 value5;
    float3 value6;
    float3 value7;
    float3 value8;


[numthreads(8, 8, 1)]
void testBuffer(uint3 id : SV_DispatchThreadID) {

      value0 +=  Source[id.xy];  // in theory the vaules 
      value1 +=  Source[id.xy];  // would be different
      value1 +=  Source[id.xy];  // but it doesn't really 
      value1 +=  Source[id.xy];  // matter for this, so 
      value1 +=  Source[id.xy];  // they are just Source[id.xy]
}


и не использую структурированный буфер, но в этом случае я не знаю, как получить данные после отправки ядра. Если дело доходит до части READ RWStructuredBuffer, которую я использую, но какой будет эквивалентный буфер, в который я могу только писать? Поскольку я действительно не читаю данные. Или общий оператор «+ =» уже вызывает состояние гонки, несмотря ни на что?

из Google я обнаружил, что решение может использовать GroupMemoryBarrierWithGroupSync(); ?? но я понятия не имею, что это такое (не говоря уже о том, как это работает), и в целом результаты Google просто немного летают над моей головой в банкомате

может ли кто-нибудь привести пример того, как решить эту проблему? В противном случае я оцениваю любые указатели.


person AverageGatsby    schedule 15.12.2019    source источник


Ответы (1)


Прежде всего, состояние гонки возникает всякий раз, когда один поток записывает в ячейку памяти , в то время как другой поток читает ИЛИ записывает из / в это же место. Итак, да, += уже вызывает состояние гонки, и нет "простого" способа исправить это. (Кстати: += неявно считывает значение, потому что вы не можете вычислить сумму двух значений, не зная их)

GroupMemoryBarrierWithGroupSync() вставляет барьер памяти, что означает: текущий поток останавливается на этой строке, пока все потоки в текущей группе не достигнут этой строки. Это важно, если один поток записывает в область памяти, а другой поток должен читать из этого места впоследствии. Так что сам по себе он вам совсем не поможет (но он требуется в следующем алгоритме).

Теперь обычным решением для вычисления суммы всех пикселей (или чего-либо подобного) является параллельное вычисление суммы. Идея состоит в том, что каждый поток считывает два пикселя и записывает их сумму в собственный индекс в массиве groupshared (Примечание: не возникает условий гонки, поскольку каждый поток имеет собственную память для записи, никакие два потока не пишут в одно и то же место). Затем половина потоков считывает каждые два значения из этого массива и записывает свою сумму обратно, и так далее, пока не останется только одно значение. На этом этапе мы вычислили сумму всех пикселей в области, которую охватывает эта группа (в вашем случае мы суммировали 8x8 = 64 пикселя). Теперь один поток в этой группе (например, тот, что SV_GroupIndex равен нулю, что верно только для первого потока в каждой группе) записывает эту сумму обратно в RWStructuredBuffer по индексу, специфичному для этой группы потоков (так что, опять же, состояние гонки не происходит). Затем этот процесс повторяется до тех пор, пока не будут суммированы все значения.

Для более подробного объяснения алгоритма см. Технический документ NVIDIA по параллельному сокращению (обратите внимание, что их код находится на CUDA, поэтому, хотя он работает очень похоже на HLSL, синтаксис и имена функций могут отличаться).

Теперь это просто вычисление суммы всех пикселей. Вычисление частотной области может быть немного сложнее или даже потребовать совсем другого решения, поскольку общий размер groupshared памяти на группу ограничен (16 КБ на оборудовании DX10)

Изменить:

Небольшой пример кода в HLSL (предполагается, что изображение уже было загружено в линейный StructuredBuffer), вычисляющее сумму 128 последовательных пикселей:

StructuredBuffer<float> Source : register(t0);
RWStructuredBuffer<float> Destination : register(u0);
groupshared float TotalSum[64];

[numthreads(64,1,1)]
void mainCS(uint3 groupID : SV_GroupID, uint3 dispatchID : SV_DispatchThreadID, uint groupIndex : SV_GroupIndex)
{
    uint p = dispatchID.x * 2;
    float l = Source.Load(p);
    float r = Source.Load(p + 1);
    TotalSum[groupIndex] = l + r;
    GroupMemoryBarrierWithGroupSync();
    for(uint k = 32; k > 0; k >>= 1)
    {
        if(groupIndex < k)
        {
            TotalSum[groupIndex] += TotalSum[groupIndex + k];
        }
        GroupMemoryBarrierWithGroupSync();
    }
    if(groupIndex == 0) { Destination[groupID.x] = TotalSum[0]; }
}
person Bizzarrus    schedule 16.12.2019