1

我正在尝试在数组数组上运行金属计算着色器。我认为因为我们使用指针来为着色器提供数组，所以逻辑解决方案是在处理 2D 数组时使用双指针。

我的核函数形式为：

kernel void foo(device float** array2D [[buffer(0)]], uint2 pid [[thread_position_in_grid]]) {
    ...
}

当我device float** array2D [[buffer(0)]]在 Metal 着色器函数中用作参数时，出现错误：

缓冲区指针类型“设备浮点 *”的地址空间限定无效

我假设这是因为编译器将其解释device float *为地址空间限定（无效）。我的问题是：如何告诉编译器我正在使用双指针？如果由于某种原因我不能使用双指针，那么处理二维数组有什么好的解决方法？

旁注：我知道可以将所有数组中的元素组合成一个数组，对单个数组进行计算，然后将数组分成更小的数组。但是，这确实效率低下并且需要很长时间。（我正在使用大型阵列，因此使用金属）。

4

1 回答 1

0

更新的答案

在评论（和修改后的问题）中，您澄清您不能将它们组合成一个平面数组。我没有一个我知道会起作用的答案，而是一些可以尝试的东西。

创建一个数组MTLCommandQueues。MTLCommandBuffer这个想法是在不同MTLCommandQueue的 s 上并行创建和提交单独的 s DispatchQueue.async。

我确信您可以创建的命令队列的数量有一些限制，因此您可以使用每个命令缓冲区的完成处理程序链接到为等待处理的内部数组创建另一个。

原始答案

我认为您需要将阵列放在一个平面阵列中。

假设内部数组的长度相同，即pid您要计算的值的“行”和“列”，您可以将数组长度作为参数之一传递给着色器：

kernel void foo(device float* array2D [[buffer(0)]], device uint rowLen [[buffer(1)]], uint2 pid [[thread_position_in_grid]]) {
    doSomeComputation(array2D[rowLen * pid.y + pid.x]); 
}

如果数组的长度不同，则需要将偏移量数组传入array2D所有行的开头。

kernel void foo(device float* array2D [[buffer(0)]], device uint* rowOffsets [[buffer(1)]], uint2 pid [[thread_position_in_grid]]) {
    doSomeComputation(array2D[rowOffsets[pid.y] + pid.x]); 
}
于 2021-04-09T18:09:27.730 回答