在 Apple 芯片上运行时，跨步内存访问会影响 Metal 计算内核的性能吗？

Question

我有以下 Metal 内核，它简单地并行添加两个数组（大小 Nx*Ny）：

kernel void add_arrays(device const float* inA,
                       device const float* inB,
                       device float* result,
                       uint tid [[thread_position_in_grid]])
{
    const unsigned int Nx = 1024*5;
    const unsigned int Ny = 128*5;
        
    // Contiguous memory access in SIMD-group
    uint x = tid % Nx;
    uint y = tid / Nx;
    uint index = y * Nx + x;
    
    // Strided memory access in SIMD-group
    //uint x = tid / Ny;
    //uint y = tid % Ny;
    //uint index = x * Ny + y;
    
    result[index] = inA[index] + inB[index];
}

作为来自 CUDA 的 Metal 新手，与连续访问相比，我预计跨步内存访问的性能要差得多。然而，这两个版本在我的 M1 Max 上平均运行时间大致相同（~80 微秒）。这有意义吗？我有另一个更复杂的内核，其中跨步访问实际上比连续访问快 30%。我原以为 SIMD 组最好访问连续的内存位置。我是否遗漏了有关 SIMD 组在 Apple 芯片上的组织方式的信息？