Эффективно подсчитывайте количество прозрачных пикселей в UIImage/CIImage с помощью Metal.

То, что вы хотите выполнить, — это операция сокращения, которая не обязательно хорошо подходит для графического процессора из-за его массивно-параллельной природы. Я бы рекомендовал не писать операцию сокращения для графического процессора самостоятельно, а использовать некоторые высокооптимизированные встроенные API, предоставляемые Apple (например, CIAreaAverage или соответствующие шейдеры производительности Metal).

Наиболее эффективный способ немного зависит от вашего варианта использования, в частности, откуда берется изображение (загружается через UIImage/CGImage или результат конвейера Core Image?) и где вам понадобится результирующий счетчик (на стороне ЦП/Swift). или как вход для другого фильтра Core Image?).
Это также зависит от того, могут ли пиксели быть полупрозрачными (альфа не 0.0 или 1.0).

Если изображение находится на графическом процессоре и/или счетчик должен использоваться на графическом процессоре, я бы рекомендовал использовать CIAreaAverage. Альфа-значение результата должно отражать процент прозрачных пикселей. Обратите внимание, что это работает только в том случае, если теперь есть полупрозрачные пиксели.

Следующим лучшим решением, вероятно, является просто повторение данных пикселей на ЦП. Это может быть несколько миллионов пикселей, но сама операция выполняется очень быстро, так что это почти не займет времени. Вы даже можете использовать многопоточность, разделив изображение на куски и используя concurrentPerform(...) из DispatchQueue.

Последним, но, вероятно, излишним решением будет использование Accelerate (это сделает @FlexMonkey счастливым): загрузите пиксельные данные изображения в буфер vDSP и используйте методы sum или average для вычисления процента с использованием векторных единиц ЦП.

Пояснение

Когда я говорил, что операция редукции не обязательно хорошо подходит для GPU, я имел в виду, что ее довольно сложно реализовать эффективным способом и далеко не так просто, как последовательный алгоритм.

Конечно, проверка прозрачности пикселя может выполняться параллельно, но результаты должны быть собраны в одно значение, что требует чтения нескольких ядер графического процессора и запись значений в ту же память. Обычно это требует некоторой синхронизации (и тем самым препятствует параллельному выполнению) и влечет за собой затраты на задержку из-за доступа к общей или глобальной памяти. Вот почему эффективные алгоритмы сбора данных для графического процессора обычно следуют многоэтапному древовидному подходу. Я настоятельно рекомендую прочитать публикации NVIDIA по этой теме (например, здесь и здесь). Именно поэтому я рекомендовал использовать встроенные API, когда это возможно, поскольку команда Apple Metal знает, как лучше всего оптимизировать эти алгоритмы для своего оборудования.

Также есть пример реализации сокращения в Спецификации языка затенения металлов Apple (стр. 158), который использует встроенные функции simd_shuffle для эффективной передачи промежуточных значений вниз по дереву. Однако общий принцип такой же, как описано в публикациях NVIDIA, ссылки на которые приведены выше.

Чтобы ответить на ваш вопрос, как сделать металл, вы бы использовали device atomic_int.

По сути, вы создаете Int MTLBuffer и передаете его своему ядру и увеличиваете его на atomic_fetch_add_explicit.

Создать буфер один раз:

var bristleCounter = 0
counterBuffer = device.makeBuffer(bytes: &bristleCounter, length: MemoryLayout<Int>.size, options: [.storageModeShared])

Сбросить счетчик на 0 и привязать буфер счетчика:

var z = 0
counterBuffer.contents().copyMemory(from: &z, byteCount: MemoryLayout<Int>.size)
kernelEncoder.setBuffer(counterBuffer, offset: 0, index: 0)

Ядро:

kernel void myKernel (device atomic_int *counter [[buffer(0)]]) {}

Увеличьте счетчик в ядре (и получите значение):

int newCounterValue = atomic_fetch_add_explicit(counter, 1, memory_order_relaxed);

Получить счетчик на стороне процессора:

kernelEncoder.endEncoding()
kernelBuffer.commit()
kernelBuffer.waitUntilCompleted()
    
//Counter from kernel now in counterBuffer
let bufPointer = counterBuffer.contents().load(as: Int.self)
print("Counter: \(bufPointer)")