c++ - Metal 防护设备地址模式

Question

我正在创建一个使用Metal渲染所有内容的图形应用程序。当我在所有绘画调用的流水线统计信息下进行帧调试时，都会出现!!名为“Prevented Device Address Mode Load”的优先级警报，详细信息:

Indexing using unsigned int for offset prevents addressing calculation in device. To prevent this extra ALU operation use int for offset.

因此，对于涉及此的最简单的平局，这就是正在进行的事情。索引缓冲区后面紧跟着大量的顶点数据。索引缓冲区是在开始时创建并填充的，此后一直保持不变。顶点数据一直在不断变化。

我有以下几种类型:

struct Vertex {
    float3 data;
};
typedef int32_t indexType; 

然后下面的画图调用

[encoder drawIndexedPrimitives:MTLPrimitiveTypeTriangle indexCount:/*int here*/ indexType:MTLIndexTypeUInt32 indexBuffer:indexBuffer indexBufferOffset:0];

转到以下顶点函数

vertex VertexOutTC vertex_fun(constant Vertex * vertexBuffer [[ buffer(0) ]],
                                indexType vid [[ vertex_id ]],
                                 constant matrix_float3x3* matrix [[buffer(1)]]) {

    const float2 coords[] = {float2(-1, -1), float2(-1, 1), float2(1, -1), float2(1, 1)};
    CircleVertex vert = vertexBuffer[vid];
    VertexOutTC out;
    out.position = float4((*matrix * float3(vert.data.x, vert.data.y, 1.0)).xy, ((float)((int)vid/4))/10000.0, 1.0);
    out.color = HSVtoRGB(vert.data.z, 1.0, 1.0);
    out.tc = coords[vid % 4];
    return out;
}

我很困惑我到底在做什么错。该错误似乎表明我不应该为我猜为索引缓冲区的偏移量使用无符号类型。

最终，对于索引缓冲区来说，只有 MTLIndexTypeUInt32和 MTLIndexTypeUInt16都是未签名的。此外，如果我尝试使用原始 int作为着色器无法编译的类型。这里发生了什么？

Answer 1

在“金属着色语言规范”的表5.1中，他们将vertex_id的“对应数据类型”列出为ushort或uint。 (该文档中所有其他类型都有类似的表，我的示例将使用相同的thread_position_in_grid)。
同时，硬件更喜欢使用签名类型进行寻址。所以如果你这样做

kernel void test(uint position [[thread_position_in_grid]], device float *test) {
    test[position] = position;
    test[position + 1] = position;
    test[position + 2] = position;
}

我们用一个无符号整数索引 test。调试此着色器，我们可以看到它涉及23条指令，并显示“Prevented Device Mode Store”警告:

如果我们改为转换为 int，则仅使用18条指令:

kernel void test(uint position [[thread_position_in_grid]], device float *test) {
    test[(int)position] = position;
    test[(int)position + 1] = position;
    test[(int)position + 2] = position;
}

但是，并不是所有的 uint都适合 int，因此此优化仅适用于uint范围的一半。仍然有很多用例。
那 ushort呢？好，

kernel void test(ushort position [[thread_position_in_grid]], device float *test) {
    test[position] = position;
    test[position + 1] = position;
    test[position + 2] = position;
}

这个版本只有17条指令。我们也“警告”在这里使用无符号索引，即使它比上面的有符号版本更快。这对我来说，警告不是特别精心设计的警告，需要大量的解释。

kernel void test(ushort position [[thread_position_in_grid]], device float *test) {
    short p = position;
    test[p] = position;
    test[p + 1] = position;
    test[p + 2] = position;
}

这是short的签名版本，可以修复警告，但也包含17条指令。因此，它使Xcode更加快乐，但我不确定它实际上是否更好。
最后，这是我的情况。我的 position范围在 signed short以上，但在 unsigned short以下。将 short提升为 int进行索引是否有意义？

kernel void test(ushort position [[thread_position_in_grid]], device float *test) {
    int p = position;
    test[p] = position;
    test[p + 1] = position;
    test[p + 2] = position;
}

这也是17条指令，并生成设备存储警告。我相信编译器证明 ushort适合 int，并且忽略了转换。然后，即使我确实这样做了，这种“无符号”算法也会产生一个警告，告诉我使用int。
总之，这些警告有些天真，应该通过设备上的测试予以确认或驳斥。