ios - iOS Metal:读取只读数据的最快方法？

Question

情况:
在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页))吗？

Answer 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