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 部分，但是我只能写入的等效缓冲区是什么？因为我并没有真正阅读数据。或者无论如何，通用运算符“+=”是否已经导致竞争条件？

从谷歌我发现一个解决方案可能正在使用GroupMemoryBarrierWithGroupSync();？但我不知道这是什么（更不用说它是如何工作的）并且一般来说谷歌的结果只是在我的头上飞过一点

谁能提供一个如何解决这个问题的例子？否则我很感激任何指点。

Answer 1

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

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

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

作为对该算法更深入的解释，请参阅NVIDIA 的并行缩减白皮书（请注意，他们的代码位于 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]; }
}