Форма параллелизма, известная как одна инструкция для нескольких данных (SIMD), относится к способности большинства современных микропроцессоров выполнять математические операции с несколькими элементами данных параллельно, используя одну машинную команду. Комбинация SIMD и многопоточности формирует параллелизм, известный как многопоточность с одной инструкцией (SIMT), основу всех современных графических процессоров.

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

kernel void (device float* r [[buffer (0)]],
 constant float* a [[buffer (1)]],
 constant float* b [[buffer (2)]],
 uint pid [[thread_position_in_grid]]){
 
 r[pid] = a[pid] * b[pid];
 }

GPU против CPU

Кто-то может спросить, почему между GPU и CPU существует фундаментальная разница в производительности параллелизма. Ответ кроется в философии дизайна между двумя типами процессоров.

Графический процессор должен иметь возможность перемещать очень большие объемы данных в свою основную DRAM и из нее из-за требований к буферу графического кадра, это называется пропускной способностью. Хотя ЦП должен удовлетворять требованиям устаревшей ОС, приложения и операции ввода-вывода затрудняют увеличение пропускной способности памяти, поэтому они предназначены для минимизации задержки выполнения одного потока.

Важным наблюдением является то, что уменьшение задержки намного дороже увеличения пропускной способности с точки зрения мощности и площади чипа. Следовательно, преобладающим решением является оптимизация производительности большого числа потоков.

Можно сделать вывод, что идеальная установка для приложения должна быть спроектирована с низкой задержкой и высокой пропускной способностью. Где низкая задержка позволяет ЦП выполнять команды последовательно, при этом система не работает медленно или не отвечает; а высокая пропускная способность позволяет графическому процессору обрабатывать данные параллельно.

Кодирование

Наконец, самое интересное - кодирование. Существует множество API-интерфейсов для параллельного программирования, таких как CUDA от NVIDIA, openCL от Kronos group, DirectX от Microsoft и, что наиболее важно, от Apple от Metal. Как обычно на этом сайте, мы сосредоточимся на Apple Metal API.

Мы начнем с изучения основных задач программирования Metal для целей GPGPU, помните, что графическое программирование выходит за рамки этой статьи. Как всегда, вы можете скачать код из моего репозитория Github (пожалуйста, не забудьте подписаться на меня на Github 😁).

Металлическая инициализация

Для связи с графическим процессором требуется следующая строка:

var device = MTLCreateSystemDefaultDevice()!

Ответственность за устройство будет заключаться в создании прямо или косвенно объектов, которые можно использовать только с этим объектом устройства. Приложения, использующие несколько графических процессоров, будут использовать несколько объектов устройства и создавать аналогичную иерархию металлических объектов для каждого из них.

  1. CommandQueue: этот объект отвечает за создание и организацию металлических буферов для отправки и планирования задач на GPU.
  2. Библиотека: похожа на commandQueue, она создается устройством для хранения указателя на металлические функции, вершину, фрагмент и ядро.
  3. Конвейер: конвейер определяет шаги, которые графический процессор выполняет для выполнения определенной задачи; Преобразуя функцию в исполняемый код на GPU. Поскольку мы фокусируемся на GPGPU, мы используем функцию ядра.
  4. Буфер: содержит данные и команды для графического процессора.
  5. Кодировщик: commandEncoder используется для записи команд в буферный объект.

Подводя итог, вы начнете с MTLDevice для объекта устройства, который, используя этот объект, создает один объект MTLCommandQueue и один объект MTLLibrary в вашем приложении. У вас будет хотя бы один объект MTLComputePipeline и хотя бы один объект MTLBuffer.

Потоки

Массивный параллелизм требует массивных потоков! В Metal потоки организованы в сетки 1D, 2D и 3D. Металл делит сетки на резьбовые группы до 3-х размеров. Потоки в группе потоков используют одно и то же пространство памяти.

Подсчет потоков является частью настройки конвейера.

Код

Код на стороне процессора:

import Foundation
import MetalKit
let row : uint = 30000
var column : uint = 4000
var array  = Array(repeating: Array<Float>(repeating: 0, count: Int(column)), count: Int(row))
let start = DispatchTime.now() // <<<<<<<<<< Start time
//1
var device = MTLCreateSystemDefaultDevice()!
var commandQueue = device.makeCommandQueue()!
var library = device.makeDefaultLibrary()
//2
let commandBuffer = commandQueue.makeCommandBuffer()
let computeEncoder = commandBuffer?.makeComputeCommandEncoder()
//3
var computeFunction = library?.makeFunction(name: "kernel_main")!
var computePipelineState = try! device.makeComputePipelineState(function: computeFunction!)
//4
var matrixBuffer = device.makeBuffer(bytes: &array, length: Int(row*column) * MemoryLayout<Float>.stride, options: [])
//5
computeEncoder?.pushDebugGroup("settingup")
computeEncoder?.setComputePipelineState(computePipelineState)
computeEncoder?.setBuffer(matrixBuffer, offset: 0, index: 0)
computeEncoder?.setBytes(&column, length: MemoryLayout<uint>.stride, index: 1)
let threadsPerThreadGrid = MTLSizeMake(Int(row * column), 1, 1)
computeEncoder?.dispatchThreadgroups(threadsPerThreadGrid, threadsPerThreadgroup: MTLSizeMake(1, 1, 1))
//6
computeEncoder?.endEncoding()
computeEncoder?.popDebugGroup()
commandBuffer?.commit()
commandBuffer?.waitUntilCompleted()
let end = DispatchTime.now()   // <<<<<<<<<<   end time
let nanoTime = end.uptimeNanoseconds - start.uptimeNanoseconds // <<<<< Difference in nano seconds (UInt64)
let timeInterval = Double(nanoTime) / 1_000_000_000 // Technically could overflow for long running tests
print("Time to execute: \(timeInterval) seconds")
let contents = matrixBuffer?.contents()
let pointer = contents?.bindMemory(to: Float.self, capacity: Int(row*column)

Просмотр кода

  1. Создание объекта устройства, который напрямую создает объект commandQueue и библиотеку. Эти объекты будут одинаковыми на протяжении всего жизненного цикла приложения.
  2. Создание commandBuffer для хранения данных и команд, которые ЦП настроит для объекта commandQueue.
  3. Создание объекта, содержащего указатель на файл, в котором определена функция. Затем создается объект конвейера, который будет выполнять функцию только для текущего буферного объекта.
  4. Объект matrixBuffer - это макет памяти графического процессора, который будет инициализирован переменной массива.
  5. Объект кодировщика задает конвейер, все объекты, которые необходимо отправить в память графического процессора, и количество потоков, которые графическому процессору необходимо создать и выполнить.
  6. Наконец, мы завершим кодирование и передадим буфер в GPU для выполнения. Для расчета времени дождемся GPU.

На следующем рисунке показан приведенный выше код:

Код на GPU

#include <metal_stdlib>
using namespace metal;
kernel void kernel_main(device float* factors [[buffer(0)]],
                        constant uint& column [[buffer(1)]],
                        uint pid [[thread_position_in_grid]]){
    factors[pid] = (pid / column) * (pid % column);
}

Пройдемся по параметрам нашей функции:

  1. Расположение памяти устройства позволяет выполнять операции чтения и записи, массив факторов будет хранить результат умножения. Буфер атрибутов находится в позиции 0, где мы устанавливаем его на стороне процессора.
  2. Переменная столбца находится в постоянной памяти, поскольку это будет только операция чтения. Буфер атрибутов находится в ячейке 1, где он был установлен на стороне ЦП.
  3. Последний параметр - это используемые потоки, здесь мы устанавливаем каждый поток в своей рабочей группе, потому что между какими-либо потоками не было взаимодействия. Каждый поток независимо вычислял значение своего местоположения в массиве.

Тело функции представляет собой одну строку, наш двумерный массив на стороне процессора был передан, имеет одномерный массив. Вот почему нам нужно, чтобы количество столбцов было разбито на следующую строку.

Металл низкий уровень, очень низкий, вся память относится к типу C. Обратите внимание, что все функции ядра в Metal являются функциями void.

Сравнение

Из нашего предыдущего блога, используя CPU с GCD и задавая 30000 строк с 4000 столбцами, мы получаем следующее:

Time to execute: 31.385106194 seconds
Program ended with exit code: 0

По сравнению с использованием графического процессора вывод приведенного выше кода:

Time to execute: 0.488734338 seconds
Program ended with exit code: 0

Заключение

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

Надеюсь, вы нашли этот пост полезным, пожалуйста, хлопайте и подписывайтесь, чтобы побудить меня написать больше о Metal и Swift. Я был бы очень признателен за посещение моего блога здесь и следование за мной. Порекомендовать другим - это всегда отличная идея, поверьте мне 😀.

Спасибо и до следующего раза, удачного кодирования!