hlsl CG 计算着色器竞争条件

Question

我正在尝试通过 unity/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)

一个简化的例子可能是对所有像素值求和，目前我的做法如下。我正在尝试使用结构化缓冲区，因为我不知道我还能如何检索数据或将其存储在 gpu 上以便之后进行全局着色器访问？？

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]
}

并且不使用结构化缓冲区，但在那种情况下我不知道如何在内核分派后检索数据。如果它归结为我正在使用的 RWStructuredBuffer 的 READ 部分，但是我只能写入的等效缓冲区是什么？因为我并没有真正阅读数据。或者通用运算符“+=”无论如何都已经导致了竞争条件？

从 google 我发现一个解决方案可能正在使用 GroupMemoryBarrierWithGroupSync(); ??但我不知道这是什么（更不用说它是如何工作的了），总的来说 google 结果只是飞过我的 atm

谁能举例说明如何解决这个问题？否则我很感激任何指示。

Answer 1

首先，当一个线程将写入内存位置而另一个线程 读取时，就会发生竞争条件或者写入 from/to 相同的位置。所以，是的，+= 已经导致了竞争条件，并且没有 "easy" 解决这个问题的方法。（顺便说一句：+= 隐式读取值，因为你无法在不知道它们的情况下计算两个值的总和）

GroupMemoryBarrierWithGroupSync() 插入内存屏障，这意味着：当前线程停在该行，直到当前组中的所有线程都到达该行。如果一个线程写入内存位置，而另一个线程需要从该位置读取之后，这一点很重要。所以就其本身而言，它根本无法帮助您（但在以下算法中是必需的）。

现在，计算所有像素（或任何类似的）总和的常见解决方案是并行计算总和。这个想法是每个线程读取两个像素，并将它们的总和写入 groupshared 数组中自己的索引（注意：没有竞争条件发生，因为每个线程都有它有自己的内存要写入，没有两个线程写入同一位置）。然后，一半的线程从这个数组中读取每两个值，并将它们的和写回，依此类推，直到只剩下一个值。在这一点上，我们计算了该组覆盖区域中所有像素的总和（在您的情况下，我们计算了 8x8=64 像素的总和）。现在，该组中的一个线程（例如 SV_GroupIndex 的线程为零，这仅适用于每个组中的第一个线程）将该总和写回特定于索引的 RWStructuredBuffer那个线程组（所以，再一次，没有竞争条件发生）。然后重复此过程，直到对所有值求和。

作为对该算法的更深入的解释，请参阅 NVIDIA's Parallel Reduction Whitepaper（请注意，它们的代码是在 CUDA 中，因此虽然它在 HLSL 中的工作方式非常相似，但语法和函数名称可能不同）。

现在，这只是计算所有像素的总和。计算频域可能有点复杂，甚至需要不同的解决方案，因为每组 groupshared 内存的总大小有限（在 DX10 硬件上为 16KB）

编辑：

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]; }
}

hlsl CG 计算着色器竞争条件

hlsl CG compute shader Race Condition

shader

hlsl

unity3d

compute-shader