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(); ？？？但我不知道这是什么（更不用说它是如何工作的了），一般而言，谷歌搜索结果只是在我的atm上方飞了一点

谁能提供解决此问题的示例？否则，我会赞赏任何指针。

我正在尝试通过unity / CG / hlsl中的计算着色器将纹理转换到频域，即，我正在尝试从纹理中读取像素值并输出基函数系数数组。 ...

Answer 1

首先，每当一个线程writes

到同一位置的另一个位置reads or writes