Металлическая производительность на iPhone XR
У меня есть ядро Metal
функция, которая в основном выглядит так:
struct Matrix {
half arr[562500]; //enough to store 750x750 matrix
};
struct Output {
half arr[12288];
};
kernel void compute_features(device Output& buffer [[ buffer(0) ]],
const device Matrix& mtx_0 [[ buffer(1) ]],
const device Matrix& mtx_1 [[ buffer(2) ]],
constant short2& matSize [[ buffer(3) ]],
constant float& offset [[ buffer(4) ]],
ushort2 gid [[ thread_position_in_grid ]]) {
for (int i = 0; i < 12; i++) {
for (int j = 0; j < 12; j++) {
int mat_id = i * matSize.x + j;
half matrixValue_0 = mtx_0.mat[mat_id];
half matrixValue_1 = mtx_1.mat[mat_id] - offset;
short someId_0 = 0;
short someId_1 = 0;
short someId_2 = 0;
short someId_3 = 0; //those ids will be calculated at the code below
half value = 0.h; //this value will be calculated at the code below
//some math where `someId` and `value` are calculated with usage of `matrixValue_0` and `matrixValue_1`
if (some_condition0) {
buffer.arr[someId_0] += value;
}
if (some_condition1) {
buffer.arr[someId_1] += value;
}
if (some_condition2) {
buffer.arr[someId_2] += value;
}
if (some_condition3) {
buffer.arr[someId_3] += value;
}
}
}
Я понимаю, что у этого кода есть свои минусы - динамическое индексирование и большой цикл. Но, к сожалению, алгоритм, который я пытаюсь выразить, не может быть реализован иначе на этом этапе.
Теперь этот код отлично работает на iPhone 7+
, это занимает около 200 us
за итерацию, и я очень доволен этим числом.
НО, я попытался запустить тот же алгоритм на iPhone XR
и я был удивлен, увидев, что этот алгоритм 1.0-1.2 ms
завершить.
С помощью XCode
и это великолепный инструмент для отладки конвейера GPU. Я обнаружил, что мои узкие места:
1)
half matrixValue_0 = mtx_0.mat[mat_id];
half matrixValue_1 = mtx_1.mat[mat_id] - offset;
Кажется, что значительная часть времени обработки тратится на Memory Load
операция.
2)
if (some_condition0) {
buffer[someId_0] += value;
}
if (some_condition1) {
buffer[someId_1] += value;
}
if (some_condition2) {
buffer[someId_2] += value;
}
if (some_condition3) {
buffer[someId_3] += value;
}
Основное время обработки уходит на Memory Store
операция.
Для меня это похоже iPhone XR
довольно сложно работать с device
память, потому что узкие места находятся в местах, где я работаю с контейнерами из device
объем памяти.
Я понимаю, что использую динамическое индексирование - компилятор не может реально предсказать, какой адрес в контейнере будет загружен / сохранен в определенной итерации. Но код очень хорошо работает наiPhone 7+
, но не на iPhone XR
.
Я подозреваю, что это может иметь какое-то отношение к выравниванию байтов. Может это как-то связано с этим?
Я хотел бы услышать несколько предложений по этому поводу. Заранее спасибо!