0

我有一个内核Metal函数，它基本上看起来像这样：

struct Matrix {
    half arr[562500]; //enough to store 750x750 matrix
};

struct Output {
    half arr[12288];
};


kernel void compute_features(device Output& buffer [[ buffer(0) ]],
                             const device Matrix& mtx_0 [[ buffer(1) ]],
                             const device Matrix& mtx_1 [[ buffer(2) ]],
                             constant short2& matSize [[ buffer(3) ]],
                             constant float& offset [[ buffer(4) ]],
                             ushort2 gid [[ thread_position_in_grid ]]) {


for (int i = 0; i < 12; i++) {
    for (int j = 0; j < 12; j++) {

        int mat_id = i * matSize.x + j;

        half matrixValue_0 = mtx_0.mat[mat_id];
        half matrixValue_1 = mtx_1.mat[mat_id] - offset;     

        short someId_0 = 0;
        short someId_1 = 0;
        short someId_2 = 0;
        short someId_3 = 0;  //those ids will be calculated at the code below
        half value = 0.h;  //this value will be calculated at the code below

        //some math where `someId` and `value` are calculated with usage of `matrixValue_0` and `matrixValue_1`

        if (some_condition0) {
            buffer.arr[someId_0] += value;
        }

        if (some_condition1) {
            buffer.arr[someId_1] += value;
        }

        if (some_condition2) {
            buffer.arr[someId_2] += value;
        }

        if (some_condition3) {
            buffer.arr[someId_3] += value;
        }
     }
}

我知道这段代码有它的缺点——动态索引和大循环。但不幸的是，我试图表达的算法在那时无法以不同的方式实现。

现在，这段代码运行得非常好iPhone 7+，每次迭代都需要它200 us，我对这个数字很满意。

但是，我尝试运行完全相同的算法iPhone XR，我惊讶地发现该算法需要1.0-1.2 ms完成。

在它强大的 GPU 管道调试工具的帮助下XCode，我发现我的瓶颈是：

1)

    half matrixValue_0 = mtx_0.mat[mat_id];
    half matrixValue_1 = mtx_1.mat[mat_id] - offset;

似乎大部分处理时间都花在了Memory Load操作上。

2) 

if (some_condition0) {
    buffer[someId_0] += value;
}

if (some_condition1) {
    buffer[someId_1] += value;
}

if (some_condition2) {
    buffer[someId_2] += value;
}

if (some_condition3) {
    buffer[someId_3] += value;
}

主要的处理时间用于Memory Store操作。

对我来说，使用内存似乎iPhone XR很困难，device因为瓶颈在我使用device内存中的容器的地方。

我知道我正在使用动态索引 - 编译器无法真正预测容器中的哪个地址将在某些迭代中加载/存储。但是代码在 上效果很好iPhone 7+，但在iPhone XR.

我怀疑它可能与字节对齐有关。它可能与此有关吗？

我很想听听这方面的一些建议。提前致谢！

4

0 回答 0