GLSL转MSL实现FFT位反转时Metal输出异常求助

问题背景 #

我用GLSL编写FFT的位反转代码，在Windows平台GLSL环境下运行正常，但通过glslang和SpirvCross编译为MSL适配macOS Metal时，输出结果完全异常。经Metal Shader调试器排查，问题出在位反转循环的第三次迭代：reversed与number执行按位或操作后值变为0，逻辑完全错误。已在CPU上验证算法正确性，修改线程索引赋值代码时，部分场景能修复偶数线程结果，但另一种修改会导致所有结果失效，怀疑是驱动bug，寻求进一步排查思路。

相关代码 #

GLSL源码 #

#define SIZE 8
#define LOG_SIZE int(log2(SIZE))

layout (binding = 0) buffer ssbo
{
    vec2 data[SIZE];
};

layout (local_size_x = SIZE) in;

void main()
{
  uint thread = uint(gl_LocalInvocationID.x);
  uint number = thread;
  uint reversed = 0;

  for (uint i = 0; i < LOG_SIZE; i++)
  {
      reversed <<= 1;
      reversed |= (number & 1);
      number >>= 1;
  }

  data[thread] = vec2(reversed, thread);
}

生成的MSL代码 #

#include <metal_stdlib>
#include <simd/simd.h>

using namespace metal;

struct ssbo
{
    float2 data[8];
};

constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(8u, 1u, 1u);

kernel void main0(device ssbo& _49 [[buffer(0)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
{
    uint thread0 = gl_LocalInvocationID.x;
    uint number = thread0;
    uint reversed = 0u;

    for (uint i = 0u; i < 3u; i++)
    {
        reversed = reversed << uint(1);
        reversed |= (number & 1u);
        number = number >> uint(1);
    }

    _49.data[thread0] = float2(float(reversed), float(thread0));
}

输出对比 #

GLSL预期输出 #

{ 0, 0 }, { 4, 1 }, { 2, 2 }, { 6, 3 }, { 1, 4 }, { 5, 5 }, { 3, 6 }, { 7, 7 }

Metal实际输出 #

{ 0, 0 }, { 0, 1 }, { 0, 2 }, { 0, 3 }, { 2, 4 }, { 2, 5 }, { 2, 6 }, { 2, 7 }

排查思路 #

• 强制统一类型与显式操作数：将GLSL中LOG_SIZE改为uint类型（#define LOG_SIZE uint(log2(SIZE))），避免循环中int与uint比较的潜在问题；同时把移位操作的右操作数统一为无符号常量（如reversed <<= 1u），减少类型转换带来的驱动优化异常。
• 关闭编译器优化：在Metal编译选项中添加-O0禁用优化，若恢复正常则说明是编译器过度优化导致逻辑篡改，可提交问题至Apple开发者反馈渠道。
• 替换位反转实现：用预计算表替代循环逻辑，比如在CPU上生成SIZE=8的位反转表uint reverse_table[8] = {0,4,2,6,1,5,3,7}，在Shader中直接查表，验证是否是循环移位操作的驱动兼容性问题。
• 验证线程索引有效性：输出gl_LocalInvocationID.x的原始值确认线程索引传递正确；尝试用gl_GlobalInvocationID.x替代gl_LocalInvocationID.x，排除线程ID获取逻辑的异常。
• 检查编译工具链参数：确认glslang和SpirvCross的编译参数，比如指定正确的--msl-version，关闭--msl-force-active-bindings等可能影响变量处理的选项。
• 最小化测试用例：将代码简化为仅保留位反转逻辑，缩小SIZE至4等更小值，确认问题是否依然存在，进一步定位触发条件。

内容的提问来源于stack exchange，提问作者James51332