问题描述
我正在尝试在一组数组上运行 Metal 计算着色器。我认为因为我们使用指针为着色器提供数组，所以逻辑解决方案是在处理 2D 数组时使用双指针。
我的核函数是这样的：
kernel void foo(device float** array2D [[buffer(0)]],uint2 pid [[thread_position_in_grid]]) {
...
}
当我在 Metal 着色器函数中使用 device float** array2D [[buffer(0)]] 作为参数时，出现错误：
缓冲区指针类型“设备浮点*”的无效地址空间限定
我假设这是因为编译器将 device float * 解释为地址空间限定（这是无效的）。
我的问题是：我如何告诉编译器我正在使用双指针？如果由于某种原因我不能使用双指针，有什么好的解决方法可以解决二维数组？
旁注：我知道可以将所有数组中的元素组合成一个数组，对单个数组进行计算，然后将该数组分成更小的数组。然而，这确实是低效的并且需要很长时间。 （我正在处理大数组，因此使用 Metal）。
解决方法
更新答案
在评论（和修改后的问题）中，您澄清您不能将它们组合成一个平面数组。我没有一个我知道行得通的答案，而是要尝试一些东西。
创建一个 MTLCommandQueues 数组。这个想法是使用 MTLCommandBuffer 在不同的 MTLCommandQueue 上并行创建和提交单独的 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]);
}