hlsl CG计算着色器竞态条件

平均盖茨比

我试图通过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();?? 但我不知道这是什么(更不用说它是如何工作的了),一般而言,谷歌的结果只是在我的atm上方飞了一点

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

Bizzarrus

首先,每当一个线程写入一个内存位置,另一个线程读取或写入同一位置时,就会发生竞争状态因此,是的,+=已经造成了竞争状况,并且没有“简单”的方法来解决此问题。(顺便说一句:+=隐式读取该值,因为您无法在不知道的情况下计算两个值的总和)

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

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

有关该算法的更详细说明,请参阅NVIDIA的Parallel Reduction白皮书(请注意,其代码在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]; }
}

本文收集自互联网,转载请注明来源。

如有侵权,请联系 [email protected] 删除。

编辑于
0

我来说两句

0 条评论
登录 后参与评论

相关文章

TOP 榜单

热门标签

归档