时间戳

首页 > 解决方案 > 金属防护设备地址模式

c++ - 金属防护设备地址模式

问题描述

我正在创建一个使用 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两者都是无符号的。此外，如果我尝试使用 rawint作为着色器不会编译的类型。这里发生了什么？

标签： c++objective-cgpumetal

解决方案

在 Metal Shading Language Specification 的表 5.1 中，他们列出了vertex_idasushort或的“对应数据类型” 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，即使这正是我所做的。

总而言之，这些警告有点幼稚，应该通过设备上的测试来确认或反驳。

推荐阅读