unity3d - hlsl CG 计算着色器竞争条件

Question

我正在尝试通过 unity/CG/hlsl 中的计算着色器将纹理转换到频域，即我正在尝试从纹理中读取像素值并输出一组基函数系数。我该怎么做？我真的是计算着色器的新手，所以我有点迷路了。我了解竞争条件的原因以及计算着色器如何划分工作负载，但有什么办法可以解决这个问题吗？对于没有这方面背景的人来说，关于缓冲区和其他东西的一般文档似乎有点让人难以理解。

我得到的错误:“Compute.compute”中的着色器错误:检测到写入共享资源的竞争条件，考虑使此写入有条件。在 Compute.compute(xxx)(在 d3d11 上)的内核 testBuffer 上

一个简化的例子可以是对所有像素值求和，目前我的方法如下。我正在尝试使用结构化缓冲区，因为我不知道我还能如何检索数据或将其存储在 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]
}

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

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

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

Answer 1

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

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