iOS Metal：读取只读数据的最快方法？答案

【问题标题】：iOS Metal: The fastest way to read read-only data?iOS Metal：读取只读数据的最快方法？
【发布时间】：2016-09-08 23:03:43
【问题描述】：

情况： 在 Metal 内核函数中，线程组中的每个线程一次读取完全相同的值。内核伪代码：

kernel void foo(device   int2*   ranges,  
                constant float3& readonlyBuffer,  
                device   float*  results,  
                uint lno [[ threadgroup_position_in_grid ]])  
{  
  float acc = 0.0;  

  for(int i=ranges[lno].x; i<ranges[lno].y; i++) {  
    // each thread in threadgroup processes the same value from the buffer  
    acc += process( readonlyBuffer[i] );  
  }  

  results[...] = acc;  
}

问题：为了优化缓冲区读取，我将readonlyBuffer的地址空间限定符从device改为constant。尽管Apple documentation 表示不同，但这对内核性能的影响为零：

常量地址空间针对执行图形或内核函数的多个实例访问缓冲区中的同一位置进行了优化。

问题：

如何提高常量缓冲区的内存读取时间？
我可以将缓冲区（或至少一部分）移动到片上缓存（类似于Constant Buffer Preloading（第 24 页））吗？

【问题讨论】：

标签： ios ios10 metal compute-shader

【解决方案1】：

在您的示例代码中，对readonlyBuffer 的索引会产生编译器错误。

假设readonlyBuffer被声明为一个指针，那么编译器并不静态知道大小，也无法将数据移动到常量内存空间中。

如果readonlyBuffer 很小（您只有 4KB 的常量内存可以使用），请将其放入结构体中，如下所示：

struct ReadonlyBuffer {
    float3 values[MAX_BUFFER_SIZE];
};

然后做：

kernel void foo(device   int2*   ranges,  
                constant ReadonlyBuffer& readonlyBuffer,  
                device   float*  results,  
                uint lno [[ threadgroup_position_in_grid ]])

最后，运行 GPU 跟踪（“捕获 GPU 帧”）并确保您不会收到以下错误：

编译器无法预加载您的缓冲区。内核函数，缓冲指数：1。

有关缓冲区预加载的更多信息，请参阅：https://developer.apple.com/videos/play/wwdc2016/606/?time=408

【讨论】：