Metal 中网格创建断言失败

Question

我正在尝试使用 Swift/C 的 Metal API 来调整 text 向量并行求和的示例代码，以便进行矩阵求和。 该代码块管理网格和组的大小。 矩阵 A、B 和 Result 的大小为 [rowsLength,colsLength]，然后分布在 rowsLengthcolsLengthsizeof(float) 的 1D GPU 缓冲区上。 我阅读了文档，但没有找到类似的代码。 编译它时出现错误：断言失败`组件 1: 2 必须是 <= 1 for index [[ thread_position_in_grid ]]' on the last command that dispatch the Threads and also there I have found only 1 question answered but that didn't helped my understanding the problem.

此时，我恳请您对computeEncoder 和此错误进行澄清。 谢谢

- (void)encodeAddCommand:(id<MTLComputeCommandEncoder>)computeEncoder {

    // Encode the pipeline state object and its parameters.
    [computeEncoder setComputePipelineState:_mAddFunctionPSO];
    [computeEncoder setBuffer:_mBufferA offset:0 atIndex:0];
    [computeEncoder setBuffer:_mBufferB offset:0 atIndex:1];
    [computeEncoder setBuffer:_mBufferResult offset:0 atIndex:2];

    MTLSize gridSize = MTLSizeMake(rowsLength*colsLength, 1, 1);

    // Calculate a threadgroup size.
    NSUInteger threadGroupSize = _mAddFunctionPSO.maxTotalThreadsPerThreadgroup;
    if (threadGroupSize > rowsLength)
    {
        threadGroupSize = rowsLength;
    }
    MTLSize threadgroupSize = MTLSizeMake(threadGroupSize, threadGroupSize, 1);

    // Encode the compute command.
    [computeEncoder dispatchThreads:gridSize
              threadsPerThreadgroup:threadgroupSize];
}

我尝试更改网格大小，将行放在第一个索引上，将列放在第二个索引上，并更改线程组大小。 我想获得一个包含 1 个块的 R*C 线程网格。

编辑1 我附上了一张照片，以进一步解释我想要获得的内容，这两种情况下程序都给出了相同的错误：

Answer 1

dispatchThreads:threadsPerThreadgroup:

采用网格大小和线程组大小。该断言告诉您至少一个线程组需要适合网格。

你的线程组Y轴大小是threadGroupSize，但是你的线程网格Y轴大小只有1。你不能将NxN的正方形放入Mx1的网格中，除非N等于1，并且N<= M.

我想获得一个包含 1 个块的 R*C 线程网格

如果您需要丢弃不符合您要求的线程，则需要向内核函数添加提前返回，而不是依赖 API 为您剔除线程。因此，您的内核会将您实际所需的线程数作为

[[buffer()]]

参数，然后将

[thread_position_in_grid]]

与它进行比较，如果不需要该线程，则返回。

Answer 2

OP，我在尝试扩展示例金属代码以处理两个二维数组时遇到了同样的问题。如果您解决了，请您发布解决方案。

...

经过进一步调查，当我更新内核函数声明以具有提供 2D 线程坐标的“gid”参数时，我的断言失败消失了。苹果示例内核函数声明具有 uint 索引。

即这有效：

kernel void add_arrays(device const unsigned int* inA,
                   device const unsigned int* inB,
                   device unsigned int* result,
                   uint2 gid [[thread_position_in_grid]])

这会生成这里讨论的断言。

kernel void add_arrays(device const float* inA,
                   device const float* inB,
                   device float* result,
                   uint index [[thread_position_in_grid]])