Максимизация производительности: запуск кода .metal на современных графических процессорах

Введение

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

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

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

Графический процессор против процессора

Так в чем же разница между GPU (графическим процессором) и CPU (центральным процессором)?

ЦП имеют несколько мощных ядер, оптимизированных для последовательной обработки. Они предназначены для вычислений общего назначения и могут решать широкий спектр задач со сложными потоками управления и разветвлениями. Ядра ЦП отлично справляются с однопоточной логикой и имеют более высокую тактовую частоту, чем графический процессор, но поскольку их гораздо меньше (около 8–12 против 1000–2500+ для графического процессора), они плохо работают с параллелизмом.
Графические процессоры имеют тысячи меньших и простых ядер, оптимизированных для параллельной обработки. Они предназначены для задач, связанных с параллельными данными и интенсивными вычислениями, что делает их высокоэффективными для задач, которые можно распараллелить. Однако их тактовая частота на ядро ниже, а одно ядро графического процессора гораздо менее мощное, чем ядро ЦП.

Хорошей аналогией является представление о процессоре как о автомобиле, а о графическом процессоре как о реактивном самолете. Автомобиль универсален и может использоваться для доставки во множество мест — продуктовый магазин, спортзал, школу и т. д. Самолет — на с другой стороны — может отлично перевезти вас через всю страну (и намного быстрее, чем на машине), но, вероятно, не подойдет для того, чтобы отвезти вас по дороге за молоком.

Как работает написание кода для графического процессора?

Итак, как мы можем выполнить некоторый выбор (высокопараллельный код) на графическом процессоре? Ну, есть несколько форматов файлов, используемых для запуска кода на графическом процессоре. .cuработает на графических процессорах NVIDIA CUDA, .clнаписан для использования на различных архитектурах графических процессоров с использованием спецификации OpenCL, а .metal предназначен для iOS. framework, который сочетает в себе некоторые функции OpenGL и OpenCL. В этой статье мы будем писать код с использованием .metal.

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

Нити, блоки и сетки

Это три наиболее важных строительных блока для выполнения графических процессоров.

Поток. Поток представляет собой наименьшую единицу работы в ядре графического процессора. Каждый поток обычно выполняет один и тот же код, но может обрабатывать разные данные или следовать разным путям внутри кода. Потоки организованы в сетки и блоки.
Блок потока. Блок потока — это группа потоков, которые могут взаимодействовать и синхронизироваться друг с другом. Потоки внутри одного блока могут обмениваться данными через общую память, что делает их пригодными для совместных задач. Размер блока определяется во время запуска ядра и часто выбирается исходя из характеристик проблемы и оборудования. Обычные размеры блоков составляют 128, 256 или 512 потоков, но оптимальный размер блока зависит от конкретного графического процессора и проблемы.
Сетка. Сетка представляет собой набор блоков потоков. Потоки внутри разных блоков одной и той же сетки не могут напрямую взаимодействовать или синхронизироваться друг с другом. Размер сетки также указывается во время запуска ядра и определяет количество блоков потоков, необходимых для выполнения ядра. Организация потоков в блоки и сетки обеспечивает эффективное параллельное выполнение на всем графическом процессоре.

Код

Установить код

import MetalKit

let count: Int = 50000000

let array1: [Float] = getRandomArray()
let array2: [Float] = getRandomArray()

print("REGULAR:")
regular(arr1: array1, arr2: array2)
print("COMPUTE:")
compute(arr1: array1, arr2: array2)

Этот код сначала запустит добавление массива, используя базовый однопоточный цикл for, а затем воспользуется методом вычислений на графическом процессоре для распараллеливания вычислений.

func getRandomArray() ->[Float] {
    var result = [Float].init(repeating: 0.0, count: count)
    for i in 0..<count {
        result[i] = Float(arc4random_uniform(10))
    }
    return result
}

Мы используем этот код для создания большого списка (независимо от того, какое значение мы устанавливаем) случайных чисел от 0 до 10.

Обычный цикл For

func regular(arr1: [Float], arr2: [Float]) {
    let startTime = CFAbsoluteTimeGetCurrent()
    var result: [Float] = [Float].init(repeating: 0.0, count: count)
    
    for i in 0..<count {
        result[i] = arr1[i] + arr2[i]
    }
    
    let timeElapsed = CFAbsoluteTimeGetCurrent() - startTime
    print("Time Taken: \(String(format:"%.05f", timeElapsed)) seconds")
}

Здесь нет ничего особенного, мы просто перебираем наши два массива и суммируем их в один массив результатов.

Time Taken: 44.99706 seconds

Вычисления на графическом процессоре

#include <metal_stdlib>
using namespace metal;

kernel void addition_function(constant float *arr1 [[buffer(0)]], constant float *arr2 [[buffer(1)]], device float *resultArr [[buffer(2)]],
                              uint index [[thread_position_in_grid]]) {
    resultArr[index] = arr1[index] + arr2[index];
}

Это наш код Compute.metal, который будет выполняться на графическом процессоре. Код графического процессора указывается функцией ядра, и мы даем ему имя add_function. Для этого потребуется три буфера массива с плавающей запятой (два для суммируемых массивов и один для буфера результатов), мы также принимаем значение позиции потока, который мы используем для выполнения этой конкретной функции (это также позволит нам чтобы получить соответствующую запись в двух массивах, которые нам нужно сложить).

func compute(arr1: [Float], arr2: [Float]) {
    let startTime = CFAbsoluteTimeGetCurrent()
    
    //1. set up GPU and the metal function
    let gpuDevice = MTLCreateSystemDefaultDevice()
    let commandQueue = gpuDevice?.makeCommandQueue()
    let gpuFunctions = gpuDevice?.makeDefaultLibrary()
    let addFunction = gpuFunctions?.makeFunction(name: "addition_function")
    
    //2. set up pipeline for the function
    var additionPipelineState: MTLComputePipelineState!
    do {
        additionPipelineState = try gpuDevice?.makeComputePipelineState(function: addFunction!)
    } catch {
        print(error)
    }
    
    //3. create the shared buffers
    let arr1Buf = gpuDevice?.makeBuffer(bytes: arr1, length: MemoryLayout<Float>.size * count,options: .storageModeShared)
    let arr2Buf = gpuDevice?.makeBuffer(bytes: arr2, length: MemoryLayout<Float>.size * count,options: .storageModeShared)
    let resultBuf = gpuDevice?.makeBuffer(length: MemoryLayout<Float>.size * count,options: .storageModeShared)
    
    //4. create command encoder
    let commandBuffer = commandQueue?.makeCommandBuffer()
    let commandEncoder = commandBuffer?.makeComputeCommandEncoder()
    commandEncoder?.setComputePipelineState(additionPipelineState)
    commandEncoder?.setBuffer(arr1Buf, offset: 0, index: 0)
    commandEncoder?.setBuffer(arr2Buf, offset: 0, index: 1)
    commandEncoder?.setBuffer(resultBuf, offset: 0, index: 2)
    
    //5. set up threads
    let threadsPerGrid = MTLSize(width: count, height: 1, depth: 1)
    let maxThreadsPerThreadGroup = additionPipelineState.maxTotalThreadsPerThreadgroup
    let threadsPerThreadGroup = MTLSize(width: maxThreadsPerThreadGroup, height: 1, depth: 1)
    
    //6. execute
    commandEncoder?.dispatchThreads(threadsPerGrid, threadsPerThreadgroup: threadsPerThreadGroup)
    commandEncoder?.endEncoding()
    commandBuffer?.commit()
    commandBuffer?.waitUntilCompleted()
    
    //7. get results and print first 3
    var resultBufferPointer = resultBuf?.contents().bindMemory(to: Float.self, capacity: MemoryLayout<Float>.size * count)
    for i in 0..<3 {
        print("\(arr1[i]) + \(arr2[i]) = \(Float(resultBufferPointer!.pointee) as Any)")
        resultBufferPointer = resultBufferPointer?.advanced(by: 1)
    }
    
    let timeElapsed = CFAbsoluteTimeGetCurrent() - startTime
    print("Time Taken: \(String(format:"%.05f", timeElapsed)) seconds")
}

Чтобы выполнить это параллельное добавление, нам нужно гораздо больше накладных расходов.

Первое, что мы делаем, это получаем наше графическое устройство по умолчанию и создаем функцию add_function, объявленную в нашем коде .metal.

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

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

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

Затем мы создаем нашу сетку потоков, у нас есть счетчик количества потоков в нашей сетке, и мы позволяем конвейеру сообщать нам, сколько потоков у нас будет на группу потоков (то есть, сколько потоков на блок потоков).

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

После завершения мы извлекаем буфер результатов из хранилища общих буферов, теперь в нем есть значения результатов, поэтому мы напечатаем первые 3.

0.0 + 4.0 = 4.0
9.0 + 7.0 = 16.0
8.0 + 7.0 = 15.0
Time Taken: 0.58681 seconds

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

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

REGULAR:
Time Taken: 0.00905 seconds
COMPUTE:
Time Taken: 0.09852 seconds

Заключение

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