ios - 通过索引访问GPU缓冲区

注意:我的问题与Apple的Metal API有关，但我认为该概念足够通用，也可以转换为其他GPU框架。

我的目标:在1 x N矩阵b的每一行中添加一个M x N行向量A。

我的内核归结为我遇到的麻烦:

kernel void vmadd(const device float* A [[ buffer(0) ]],
                  const device float* b [[ buffer(1) ]],
                  device float* C [[ buffer(2) ]],
                  constant ushort& aWidth [[ buffer(3) ]],
                  ushort2 gid [[ thread_position_in_grid ]]) {

    int idx = gid.y * aWidth + gid.x; // Compute absolute index in C
    C[idx] = A[idx] + b[gid.x];

}

假设:我的理解是gid是C中单个元素的位置:gid.x是列，gid.y是行。如果不是这样，请有人纠正我。

现在，如果我用8 x 8个零填充A:

A = 0 0 0 0 0 0 0 0
    0 0 0 0 0 0 0 0
    0 0 0 0 0 0 0 0
    0 0 0 0 0 0 0 0
    0 0 0 0 0 0 0 0
    0 0 0 0 0 0 0 0
    0 0 0 0 0 0 0 0
    0 0 0 0 0 0 0 0

和b像这样:

b = 1 2 3 4 5 6 7 8

然后执行后，C应该是8 x 8的矩阵，其中每一行都是1 2 3 4 5 6 7 8。

相反，我得到这个:

C = 1 2 3 4 5 6 7 8
    1 2 3 4 5 6 7 8
    1 2 3 4 5 6 7 8
    0 0 0 0 0 0 0 0
    1 2 3 4 5 6 7 8
    1 2 3 4 5 6 7 8
    1 2 3 4 5 6 7 8
    0 0 0 0 0 0 0 0

我正在根据Apple的推荐here调整线程组的大小:

let w = computePipeline.threadExecutionWidth
let h = computePipeline.maxTotalThreadsPerThreadgroup / w
let threadsPerThreadgroup = MTLSizeMake(w, h, 1)
let threadgroupsPerGrid = MTLSize(width: (cWidth + w - 1) / w,
                                  height: (cHeight + h - 1) / h,
                                  depth: 1)

在我的机器上，哪个在每个网格上生成(64, 16, 1)线程组，在每个网格上生成(1, 1, 1)线程组。

但是，如果我手动将threadsPerThreadgroup设置为(1, 1, 1)，并且将threadgroupsPerGrid设置为(8, 8, 1)，则在C中可以获得正确的结果。

问题:

我可以肯定我的问题与线程组的大小和缓冲区的管理方式有关，但是我对GPU编程还很陌生，所以我不太了解它。

为什么减小线程组的大小会产生正确的结果？甚至更普遍地讲，为什么线程组大小对此计算完全没有影响？

在我看来，如果gid始终对应于C中的索引，并且我要求b[gid.x]，则该值应该可用。我了解数据通常会被分解以适合线程组缓存-如果是这种情况，它将遵循哪些规则，我该如何处理？

最佳答案

我的理解是gid是C中单个元素的位置:gid.x是列，gid.y是行。如果不是这样，请有人纠正我。

这不是真的。 gid是网格中的位置。

因为网格恰好是64x16，所以将调用位于8x8矩阵(A和C)和8元素向量(b)之外的位置的计算功能。发生这种情况时，读取A可能会访问错误的行，甚至可能访问A的末尾。同样，对b的读取将超出其末尾。

例如，考虑gid是(8，0)的时间。 idx将为8。您将阅读A[8]，它实际上位于(0，1)。您将阅读b[8]，这已经结束了。这在技术上是不确定的，但实际上对于一个相对较短的缓冲区，它很可能为0。您将写入也是(0，1)的C[8]。这与应该在(0，1)编写的函数调用大致同时发生，并且存在一个竞争。

您的函数应该在开始时测试gid是否超出范围，如果是，则尽早返回:

if (any(gid > aWidth))
    return;

(这假定A和C始终为正方形，因此可以根据单个值检查宽度和高度。)

您可以尝试调整threadsPerThreadgroup和threadgroupsPerGrid的计算，以使网格正好与您的矩阵大小相同，但是要在所有情况下正确进行操作都可能很麻烦。也就是说，您当然可以避免threadsPerThreadgroup太大:

let w = min(computePipeline.threadExecutionWidth, cWidth)
let h = min(computePipeline.maxTotalThreadsPerThreadgroup / w, cHeight)

但是您仍然需要在计算功能中进行检查，因为总网格仍然可能太大。例如，假设computePipeline.threadExecutionWidth至少为8，而computePipeline.maxTotalThreadsPerThreadgroup为60。那么，w将为8，但h将为7。然后，threadgroupsPerGrid将为(1、2、1)，总网格大小将为8x14x1，再次比矩阵大