HLSL CG 计算着色器竞争条件

Question

HLSL CG 计算着色器竞争条件

unity-game-engineshaderhlslcompute-shader

4

我正在尝试通过Unity/CG/HLSL中的计算着色器将纹理转换为频域，即我正在尝试从纹理中读取像素值并输出一组基函数系数的数组。我该如何做？我对计算着色器很陌生，所以有点迷茫。我理解竞争条件的原因以及计算着色器如何分配工作负载，但是否有任何方法来处理这个问题？总的来说，对于没有相关背景的人来说，关于缓冲区和其他内容的文档似乎有些令人失望。

我得到的错误信息： “在‘Compute.compute’着色器中发现写入共享资源的竞争条件，请考虑使此写入有条件。在Compute.compute（xxx）的testBuffer内核中（在d3d11上）”

一个简单的例子可以是将所有像素值相加，目前我的方法如下。我正在尝试使用StructuredBuffer，因为我不知道如何以其他方式检索数据或在全局着色器访问后将其存储在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]
}

不要使用StructuredBuffer，但在这种情况下，我不知道如何在内核调度后检索数据。如果是我正在使用的RWStructuredBuffer的读取部分，但是有什么等效的缓冲区可以只写入？因为我实际上不读取数据。或者一般的操作符"+="是否已经引起了竞态条件无论如何？

从Google中，我发现一个解决方案可能是使用GroupMemoryBarrierWithGroupSync(); 但我不知道这是什么（更不用说它是如何工作的），而且通常Google的结果对我来说有点难以理解。

请问是否有人能提供一个如何解决此问题的示例？否则，我很感激任何指针。

- AverageGatsby

1个回答

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- Bizzarrus · Accepted Answer

首先，每当一个线程在另一个线程读取或写入相同位置的内存时对该内存位置进行写入，就会发生竞争条件。因此，是的，+已经导致了竞争条件，没有“简单”的方法来解决这个问题。(顺便说一下：+隐式地读取值，因为在不知道它们的情况下无法计算两个值的总和) GroupMemoryBarrierWithGroupSync()插入内存屏障，这意味着：当前线程在该行停止，直到所有线程都到达该行。如果一个线程写入一个内存位置，另一个线程需要在之后从该位置读取，那么这非常重要。因此，单独使用它并不能帮助您（但在以下算法中是必需的）。

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

作为算法更深入的解释，请参见NVIDIA的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]; }
}