标签: compute-shader

我正在探索使用计算着色器将骨骼变形应用于网格顶点,而不是使用流输出的顶点着色器.我发现计算着色器的执行速度远低于顶点着色器,但在我把它写下来之前,我想确定我没有做错.

对于300个骨骼的100,000个顶点和1,000帧动画数据的测试数据,顶点着色器运行大约0.22ms,而计算着色器在0.85ms处运行4倍.时间是通过D3D API计时器查询(而不是cpu计时器)完成的.

deform_structs.hlsl

struct Vertex {
  float3 position : POSITION;
  float3 normal : NORMAL;
  float2 texcoord : TEXCOORD;
  float3 tangent : TANGENT;
  float4 color : COLOR;
};

struct BoneWeights {
  uint index;
  float weight;
};

StructuredBuffer<matrix> g_bone_array : register(t0);
Buffer<uint> g_bone_offsets : register(t1);
Buffer<uint> g_bone_counts : register(t2);
StructuredBuffer<BoneWeights> g_bone_weights : register(t3);

bone_deform_cs.hlsl

#include "deform_structs.hlsl"

StructuredBuffer<Vertex> g_input_vertex : register(t4);
RWStructuredBuffer<Vertex> g_output_vertex : register(u0);

[numthreads(64,1,1)]
void BoneDeformCS(uint id : SV_DispatchThreadID) {
  Vertex vert = g_input_vertex[id.x];
  uint offset = g_bone_offsets[id.x];
  uint count …

所有DirectX书籍和教程都强烈建议将绘制调用之间的资源分配减少到最低限度 - 但我找不到任何可以深入了解细节的指南.回顾网上发现的大量示例代码,我得出结论,程序员对此主题有完全不同的编码原则.有些甚至设置和取消设置

VS/PS 
VS/PS ResourceViews
RasterizerStage 
DepthStencilState
PrimitiveTopology
... 

在每次绘制调用之前和之后(虽然设置保持不变),而其他人则没有.

我猜这有点过头了......

从我自己的实验,我发现,我不得不在每次绘制调用绑定的唯一资源是ShaderResourceViews(以VS和PS在我的情况).此要求可能是由于使用计算着色器引起的,因为我绑定/取消绑定UAVs到绑定到VS/ PS稍后的缓冲区.

在我发现这种重新绑定是必要的之前,我已经失去了许多小时的工作.而且我想许多程序员也不确定,并且更愿意解开并重新绑定"有点太多"而不是陷入类似的陷阱.

问题1:关于这个问题至少有一些经验法则吗？

问题2:我的ShaderResourceViews绑定是否可能VS/PS被驱动程序/ DirectX核心解除绑定,因为我UAVs 在CS调度调用之前绑定了相同的缓冲区(我没有解除绑定SRVs自己)？

问题3:VS/PS在使用计算着色器之前,我甚至没有设置为null.没有问题的工作但我总是不确定我是否正在使用这种"懒惰"的方法挖掘我的下一个陷阱.

我需要测量计算着色器的时间.但当然这不是微不足道的.从OpenGL Wiki -我得到的性能,在着色器调用之前和之后使用glFinish()是有用的.但他们也说使用它并不好.是否有可能测量着色器的时间？反正有可能测量计算着色器的时间吗？

我的代码看起来像这样:

renderloop()
{
  //(1)
  //(2)
  if(updateFunction) //this is done just one time at the beginning
  {
    //update Texture with a compute shader
    //...
    glDispatchCompute();
    glMemoryBarrier(GL_ALL_BARRIER_BITS);
  }
  //(3)
  //(1)

  //use the texture to do some marching cubes rendering
}

我想我必须插入glFinish()位置(1)并启动计时器(2)并将其停在(3).但我不确定它是否真的有效并且会产生正确的计时结果,因为在参考中他们谈论的是渲染而计算着色器不能渲染,不是吗？

还存在OpenGL Timer_Query,但我不确定它是如何工作的,也不知道它是否对我有用.这些东西对我来说是新的,我不确定我是否完全理解它.

这里的答案说,几乎不可能精确地测量代码的一部分.最好的方法是测量帧渲染时间,但我只需要帧渲染时间的计算着色器部分用于我的目的.

您认为最好的替代方案是什么？只测量整个帧渲染时间并使用它？或者您是否使用其他测量方法获得了更好的体验？

我一直在尝试编写一般的计算着色器高斯模糊实现.

它基本上可以工作,但它包含的工件即使在场景静止时也会改变每一帧.我花了几个小时试图调试这个.我已经走了尽可能确保不超出界限,展开所有循环,用常量替换制服,但工件仍然存在.

我已经在3个不同的机器/ GPU(2个nvidia,1个intel)上测试了原始代码和工件,它们都产生相同的结果.使用普通C++代码模拟执行向前和向后执行的工作组的代码执行的展开/常量版本不会产生这些错误.

通过分配[96] [96]而不是[16] [48]的共享数组,我可以消除大部分伪像.

这让我想到了我错过了一个逻辑错误,因此我设法生成了一个非常简单的着色器,它仍然会在较小的范围内产生错误,如果有人能指出原因,我会很感激.我检查了很多文档,找不到任何错误.

分配了一个16x48浮点数的共享数组,这是3072字节,大约是最小共享内存限制的10%.

着色器在16x16工作组中启动,因此每个线程将写入3个唯一位置,并从单个唯一位置读回

然后纹理作为HSV渲染,其中0-1之间的值将映射到色调0-360(红色 - 青色 - 红色),并且超出边界的值将是红色.

#version 430
//Execute in 16x16 sized thread blocks
layout(local_size_x=16,local_size_y=16) in;
uniform layout (r32f) restrict writeonly image2D _imageOut;
shared float hoz[16][48];
void main () 
{
    //Init shared memory with a big out of bounds value we can identify
    hoz[gl_LocalInvocationID.x][gl_LocalInvocationID.y] = 20000.0f;
    hoz[gl_LocalInvocationID.x][gl_LocalInvocationID.y+16] = 20000.0f;
    hoz[gl_LocalInvocationID.x][gl_LocalInvocationID.y+32] = 20000.0f;
    //Sync shared memory
    memoryBarrierShared();
    //Write the values we want to actually read back
    hoz[gl_LocalInvocationID.x][gl_LocalInvocationID.y] = 0.5f;
    hoz[gl_LocalInvocationID.x][gl_LocalInvocationID.y+16] = 0.5f; …

在尝试实现一个将影响灯光分配给集群的简单计算着色器时，我遇到了意外的行为（对于像我这样的菜鸟来说）：

我使用 glDispatchCompute(32, 32, 32); 调用此着色器 它应该为每次调用写入一个[轻计数器+8个索引]到“索引”缓冲区。但是在调试时，我发现即使我使用唯一的 clusterId，我对缓冲区的写入也会在调用之间重叠。我通过索引[outIndexStart] 的值超过 8 和视觉闪烁来检测它。

根据文档，gl_GlobalInitationID是gl_WorkGroupID * gl_WorkGroupSize + gl_LocalInitationID。但如果将所有本地大小设置为 1，写入问题就会消失。为什么 local_size 会以这种方式影响这段代码？我该如何推理在这里选择它的价值？

#version 430
layout (local_size_x = 4, local_size_y = 4, local_size_z = 4) in;

uniform int lightCount;

const unsigned int clusterSize = 32;
const unsigned int clusterSquared = clusterSize * clusterSize;

struct LightInfo {
    vec4 color;
    vec3 position;
    float radius;
};

layout(std430, binding = 0) buffer occupancyGrid {
    int exists[];
};

layout(std430, binding = 2) buffer lightInfos
{
  LightInfo lights …

我正在尝试通过 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) { …

我正在尝试使用texture2D（）从计算着色器中的sampler2d纹理读取值。在PC上运行正常，但在android移动设备（使用310 es版本）上，相同代码的编译失败，并出现以下错误：

'texture2D' : type is for Vulkan api only  

这个调用与计算着色器不兼容吗？

假设我有一个用GLSL编写的opengl计算着色器，在NVidia Geforce 970上执行。

在着色器的开始处，单个调用将写入“着色器存储缓冲区对象”（SSBO）。

然后，我发出合适的屏障，例如GLSL中的memoryBarrier（）。

然后，在每次调用时，我将从第一步编写的内存中读取。

第一次写入对当前计算操作中的所有调用都可见吗？

在https://www.khronos.org/opengl/wiki/Memory_Model#Ensuring_visibility上，Khronos说：

“如果您使用诸如barrier这样的机制在调用之间进行同步，请使用一致且适当的memoryBarrier *或groupMemoryBarrier调用。”

我很确定可以在工作组中以这种方式进行同步。但是，它是否适用于整个计算操作中每个工作组中的所有调用？

我不确定如何安排整个工作组。我希望它们可能按顺序运行，从而使我要问的那种同步成为不可能？

有人可以简单地解释一下 Metal 计算着色器中线程组的概念以及其他术语（例如 SIMD 组、threadExecutionWidth（波前））吗？我阅读了文档，但更加困惑。例如，如果我有一个 1024x1024 图像，我可以有多少个线程组，如何将线程映射到每个像素，有多少个线程可以并发运行，等等？我找不到描述计算着色器和这些概念的 WWDC 视频。

我试图了解整个 L1/L2 冲洗是如何工作的。假设我有一个像这样的计算着色器

layout(std430, set = 0, binding = 2) buffer Particles{
    Particle particles[];
};


layout(std430, set = 0, binding = 4) buffer Constraints{
    Constraint constraints[];
};


void main(){
    const uint gID = gl_GlobalInvocationID.x;
    for (int pass=0;pass<GAUSS_SEIDEL_PASSES;pass++){
        // first query the constraint, which contains particle_id_1 and particle_id_1
        const Constraint c = constraints[gID*GAUSS_SEIDEL_PASSES+pass]; 
        // read newest positions
        vec3 position1 = particles[c.particle_id_1].position; 
        vec3 position2 = particles[c.particle_id_2].position;
        // modify position1 and position2
        position1 += something;
        position2 -= something;
        // update positions
        particles[c.particle_id_1].position = …