将值写入 iOS Metal 中的 float4x4 矩阵 - 替换现有的浮点值答案

【问题标题】：Writing values to a float4x4 matrix in iOS Metal - existing float values replaced将值写入 iOS Metal 中的 float4x4 矩阵 - 替换现有的浮点值
【发布时间】：2021-09-13 12:44:46
【问题描述】：

我正在尝试制作一个缓冲区来计算均匀间隔网格的某些单元格中的粒子/点（每个粒子将根据其位置添加到单元格中）。然后将每个粒子的索引存储在与每个单元格对应的 float4x4 矩阵中，以供以后查找。我正在使用原子计数器为每个单元添加一个粒子数，使粒子数不超过 16，并使用此计数器将粒子索引顺序添加到矩阵中。

当行数或列数>4时，即下一列/行，位置[0]处的索引写入0.0。我不确定这是否与在 float4x4 矩阵中覆盖列的方式有关，或者是否与分配内存的方式/float4x4 的指针有关。基本上，我只想按顺序添加和更改单个值，而不影响矩阵中的现有值。

非常感谢任何帮助，非常感谢！

这是内核：

kernel void findCellandCount(  device Particle *particles [[ buffer(0) ]],
                     volatile device atomic_uint *cellCountArray [[buffer(1)]],
                     device float4x4* cellIndicesBuffer [[ buffer(2) ]],
                     uint id [[ thread_position_in_grid ]]) {
  
    uint particleIndex = id;
    
    Particle particle = particles[particleIndex];

    const float cellSize = Params().cellWidth;

*// GET CELL INDEX:*

    int2 cellIndex = int2(fast::floor(particle.position.x / cellSize), fast::floor(particle.position.y / cellSize)); 

    uint flatCellIndex = GetFlatCellIndex(cellIndex, numberGridCells); // this is in the range 0-15 (16 cells)
    
    // This is a counter to store particle count in each cell // this is reset to zero each frame:
    int cellCounter = atomic_fetch_add_explicit(&cellCountArray[flatCellIndex], 1, memory_order_relaxed);
        
    if (cellCounter < 16) { // i.e. count is less than 4x4 float matrix 

        uint a = cellCounter % 4;
        uint b = cellCounter / 4;

    //m[index][column][row]

        cellIndicesBuffer[flatCellIndex][a][b] = id; // this writes the particle index to the float4x4
 
    }
  
}

这是单元格的输出：

Cell:  3 cell indices:  simd_float4x4([[23.0, 0.0, 0.0, 0.0], [25.0, 0.0, 0.0, 0.0], [44.0, 0.0, 0.0, 0.0], [61.0, 0.0, 0.0, 0.0]])

^ 这个输出符合预期，前 4 行存储了 4 个索引

Cell:  8 cell indices:  simd_float4x4([[0.0, 38.0, 0.0, 0.0], [0.0, 39.0, 0.0, 0.0], [0.0, 42.0, 0.0, 0.0], [0.0, 63.0, 0.0, 0.0]])

^ 这里存储了 8 个索引，但位置 [0] 处的值被 0.0 覆盖/替换。

Cell:  9 cell indices:  simd_float4x4([[0.0, 35.0, 0.0, 0.0], [0.0, 45.0, 0.0, 0.0], [13.0, 0.0, 0.0, 0.0], [28.0, 0.0, 0.0, 0.0]])

^ 这里存储了 6 个索引，但第一个位置的值再次被覆盖。

【问题讨论】：

您在哪里定义了在调用GetFlatCellIndex 时使用的cellIndex？它不在参数列表中，也没有定义为局部变量。
嗨芯片，我已经添加了单元格索引代码； getFlatCellIndex 函数只返回一个以该位置为基础的散列值：int n = p1 * cellIndex.x ^ p2*cellIndex.y; n %= partitionBucketCount; 其中 p1 和 p2 是大素数
我可能是错的，但我认为您的a 和b 矩阵索引计算可能会倒退。自从我需要使用 SIMD 矩阵类型已经有一段时间了，但我记得它们是主要的列，这意味着不仅在使用下标运算符时首先指定列索引，而且内存布局是所有一列的元素是连续的，后跟下一列的所有元素，等等...因此，如果连续的cellCounter 值在移动到下一列之前向下寻址每一列，则a 应该是cellCounter / 4 和b应该是cellCounter % 4。

标签： swift matrix metal particles

【解决方案1】：

我认为我知道发生了什么。在您拥有的代码中

cellIndicesBuffer[flatCellIndex][a][b] = id;

我认为那条线的效果本质上是这样的：

float4x4 temp = cellIndicesBuffer[flatCellIndex];
temp[a][b] = id;
cellIndicesBuffer[flatCellIndex] = temp;

虽然您使用atomic_uint 来防止cellCountArray 数组元素上的数据竞争，但您并没有使用cellIndicesBuffer 元素来防范它，这样做无论如何都会对性能产生问题。

我认为，问题源于 float4x4 是 SIMD 类型 - 本质上是 struct。我建议它将整个内容读入 gpu-thread 本地内存（可能是寄存器），更新元素，然后将整个内容写回数组，覆盖 float4x4 可能已被其他人写入的元素在中间的时间线程。为避免这种情况，您只需处理要更新的元素，而无需通过float4x4，您可以通过将cellIndicesBuffer 重铸为float* 来做到这一点。

auto i = flatCellIndex * 16 + a * 4 + b;
reinterpret_cast<device float*>(cellIndicesBuffer)[i] = id;

如果您打算在继续下一列之前将连续的 cellCounter 值填满每一列，您可以消除 a 和 b：

auto i = flatCellIndex * 16 + cellCounter;
reinterpret_cast<device float*>(cellIndicesBuffer)[i] = id;

我应该提到，负责任的 C++ 程序员可能对我在这里使用 reinterpret_cast 感到畏缩，这是正确的。如果您不需要在着色器函数中的其他任何地方使用cellIndicesBuffer 专门作为device float4x4*，而您没有在提供的代码中使用，最好将参数类型更改为device float*。那么您不必执行reinterpret_cast。您不需要对 Swift 代码进行任何更改。

【讨论】：