Was ist falsch an den Metal-Instrumenten simdgroup_load oder simdgroup_store?

Betriebssystem: MacOS 12.2.1

Hardwear: MacBook Pro 2020, M1

Metall: 2.4

Xcode: 13.2.1

Hier ist mein Test-Computer-Kernel, der den Eingabepuffer mit simdgroup_load liest und den Ausgabepuffer mit simdgroup_store schreibt

kernel void fun(
const device half * Src  [[ buffer(0) ]],
constant uint4 & SrcShape [[ buffer(1) ]],
device half * Dst  [[ buffer(2) ]],
constant uint4 & DstShape [[ buffer(3) ]],
const device half * Weight [[ buffer(4) ]],
ushort3 threadgroup_position_in_grid [[ threadgroup_position_in_grid ]],
ushort3 thread_position_in_threadgroup [[ thread_position_in_threadgroup ]],
ushort3 threads_per_threadgroup [[ threads_per_threadgroup ]],
ushort3 thread_position_in_grid [[ thread_position_in_grid ]])
{
 const int SrcSlices = (int)SrcShape[0];
 const int SrcHeight = (int)SrcShape[1];
 const int SrcWidth = (int)SrcShape[2];
 const int DstSlices = (int)DstShape[0];
 const int DstHeight = (int)DstShape[1];
 const int DstWidth = (int)DstShape[2];
 const int Kernel_X = 3;
 const int KernelElemNum = 3 * 3;
 const int N_Pack = 8;
 // Test only 1 thread
 if(thread_position_in_grid.z!= 0|| thread_position_in_grid.y!= 0|| thread_position_in_grid.x * N_Pack!= 0) return;
 simdgroup_half8x8 sgMatY;
 simdgroup_load(sgMatY, Src);
 simdgroup_store(sgMatY, Dst);
}

Es ist ein einfacher Shader, aber der Ausgabepuffer speichert nur die ersten 2 Werte aus dem Eingabepuffer, die anderen 62 Werte sind ALLE NULL.

Hier ist das Ergebnis von Xcode Metal Capture

Wie man es debuggt oder repariert?

Lösung des Problems

Fehler behoben. Jeder, der diese Funktion nutzen möchte, kann sich auf die Implementierung in TF-Lite beziehen:

https://github.com/alpa-projects/tensorflow-alpa/blob/ee8f6612b515ada4509fa53491c5ba5b3ef8524a/tensorflow/lite/delegates/gpu/common/tasks/conv_metal_simd.cc