Я новичок в Metal и пытаюсь понять реализацию свертки в Metal в TFLite . Прочитав эту строку кода и обнаружив все случаи использования, я действительно запутался, почему значение work_group_launch_order имеет значение. Как работает этот механизм на самом деле? Связано ли это с тем, как графический процессор линеаризует 3D threadgroup?
[Tensorflow GitHub]
Для простоты позвольте мне попытаться объяснить Стратегия диспетчеризации потоков для определенного ядра (с именем ConvolutionGeneric) в TFLite.
Для заданного номера потока 3D t=<tx, ty, tz> и формы группы потоков s=<sx, sy, sz>, TFLite вычисляет количество желаемых групп потоков n=<nx, ny, nz> на nx=ceil(tx/sx), ny=ceil(tx/sy) и nz=ceil(tz/sz). Что абсолютно нормально.
Нормальным способом мы можем отправить <nx, ny, nz> группы потоков для 3 измерений и получить позицию потока в сетке в функции ядра Metal по аргументу gid с атрибутом [[thread_position_in_grid]]. Позиция потока может решить, за какую область отвечает текущий поток.
Однако TFLite выбрал странный способ, отправляет <nz, nx, ny> группы нитей для 3 измерений и получает положение нити в сетке в функции ядра ядра Metal путем вычисления из tid3d [[thread_position_in_threadgroup]] и group_id [[threadgroup_position_in_grid]] как
gid_x = group_id.y * sx + tid3d.x;
gid_y = group_id.z * sy + tid3d.y;
gid_z = group_id.x * sz + tid3d.z
Что меня удивляет, так это то, что эта стратегия действительно повышает производительность (~ 10% ускорение) .
Может кто-нибудь помочь мне объяснить основной механизм, стоящий за странным threadgroup стратегия отправки?