High performance
GPGPU
Rust

Идиоматичное программирование GPU на Rust: Библиотека Emu

Tutorial


Введение


Emu — это высокоуровневый язык программирования видеокарт, способный встраиваться в обычный код на системном языке программирования Rust.


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


Установка


  1. Обозреваемая библиотека нуждается во внешней зависимости OpenCL. Вам необходимо установить соответствующий вашему оборудованию драйвер.
  2. Дополните Cargo.toml приведённым ниже текстом. Это вызовет скачивание последних доступных версий (если нужна конкретная сборка, то вместо * поместите нужную версию):

    [dependencies]
    em = "*" // Поддержка языка Emu
    ocl = "*" // Обёртка над OpenCL

Синтаксис


Синтаксис Emu довольно прост, ведь данный язык предназначается лишь для написания функций-ядер, транслируемых в OpenCL при компиляции.


Типы данных


Язык Emu располагает девятью типами данных, которые аналогичны соответствующим в Rust. Ниже приведена таблица данных типов:


Название Описание
f32 Тридцатидвухбитное число с плавающей точкой
i8 Символ или восьмибитное число
i16 Знаковое шестнадцатибитное число
i32 Знаковое тридцатидвухбитное число
i64 Знаковое шестидесятичетырехбитное число
u8 Беззнаковое восьмибитное число
u16 Беззнаковое шестнадцатибитное число
u32 Беззнаковое тридцатидвухбитное число
u64 Беззнаковое шестидесятичетырёхбитное число
bool Булевое значение
[TYPE] Вектор, состоящий из переменных типа TYPE

Переменные


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


let age: i32 = 54;
let growth: f32 = 179.432;
let married: bool = true;

Конвертации


Конвертация примитивных типов данных осуществляется посредством бинарного оператора as, следующим за целевым типом. Замечу, что целевым типом также может быть единица измерения (смотреть следующую секцию):


let width: i16 = 324;
let converted_width: i64 = width as i64;

Единицы измерения


Язык Emu позволяет обращаться с числами как с единицами измерения, что призвано упростить научные вычисления. В данном примере переменная length изначально определена в метрах, но потом к ней прибавляются иные единицы измерения:


let length: f32 = 3455.345; // Метры
length += 7644.30405 as cm; // Сантиметры
length += 1687.3043 as mm; // Миллиметры

Предопределённые константы


Emu располагает набором предопределённых констант, которые удобно использовать на практике. Ниже приведена соответствующая таблица.


Название Значение
Y 10 в степени 24
Z 10 в степени 21
E 10 в степени 18
P 10 в степени 15
T 10 в степени 12
G 10 в степени 9
M 10 в степени 6
k 10 в степени 3
h 10 в степени 2
D 10 в степени 1
d 10 в степени -1
c 10 в степени -2
m 10 в степени -3
u 10 в степени -6
n 10 в степени -9
p 10 в степени -12
f 10 в степени -15
a 10 в степени -18
z 10 в степени -21
y 10 в степени -24

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


Условные операторы


Условные операторы Emu аналогичны соответствующим операторам в Rust. Ниже показан код, применяющий условные конструкции:


let number: i32 = 2634;
let satisfied: bool = false;

if (number > 0) && (number % 2 == 0) {
    satisfied = true;
}

Циклы for


Заголовок цикла For определяется как for NUM in START..END, где NUM — это переменная, принимающая значения из диапазона [START; END) через единицу.


let sum: u64 = 0;

for i in 0..215 {
    sum += i;
}

Циклы while


Заголовок цикла While определяется как while (CONDITION), где CONDITION — это условие перехода цикла к следующей итерации. Данный код аналогичен предыдущему примеру:


let sum: u64 = 0;

let idx: i32 = 0;
while (idx < 215) {
    sum += idx;
    idx += 1;
}

Бесконечные циклы


Бесконечные циклы не имеют явно заданного условия выхода и определяются ключевым словом loop. Они, однако, могут быть продолжены или прерваны посредством операторов break и continue (как и остальные два типа циклов).


let collapsed: u64 = 1;

let idx: i32 = 0;
loop {
    if idx % 2 == 0 { continue; }
    sum *= idx;

    if idx == 12 { break; }
}

Возвращение из функции


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


let result: i32 = 23446;
return result;

Другие операторы


  • Доступные операторы присваивания: =, +=, -=, *=, /=, %=, &=, ^=, <<=, >>=;
  • Оператор индекса — [IDX];
  • Оператор вызова — (ARGS);
  • Унарные операторы: * для разыменования, ! для инверсии булевых данных, - для отрицания чисел;
  • Бинарные операторы: +, -, *, /, %, &&, ||, &, |, ^, >>, <<, >, <, >=, <=, ==, !=.

Функции


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


add(left f32, right f32) f32 {
    return left + right;
}

Как вы уже могли заметить, данная функция возвращает сумму двух переданных в неё аргументов с помощью типа данных f32.


Адресные пространства


Каждый параметр функции соответствует определённому адресному пространству. По умолчанию, все параметры соответствуют пространству __private__.


Добавление префиксов global_ и local_ к идентификатору параметра явно указывает его адресное пространство.


Документация советует использовать префикс global_ ко всем векторам и не помечать префиксом ничего другое.


Встроенные функции


Emu предоставляет небольшой набор встроенных функций (взятых из OpenCL), позволяющих вам управлять данными GPU:


  • get_work_dim() — Возвращает количество измерений;
  • get_global_size() — Возвращает количество глобальных элементов для заданного измерения;
  • get_global_id() — Возвращает уникальный идентификатор элемента для заданного измерения;
  • get_global_size() — Возвращает количество глобальных элементов для заданного измерения;
  • get_local_id() — Возвращает уникальный идентификатор локального элемента внутри конкретной рабочей группы для заданного измерения;
  • get_num_groups() — Возвращает количество рабочих групп для заданного измерения;
  • get_group_id() — Возвращает уникальный идентификатор для рабочей группы.

В прикладном коде чаще всего вы встретите выражение get_global_id(0), возвращающее текущий индекс элемента вектора, ассоциированного с вызовом вашей функции-ядра.


Выполнение кода


Рассмотрим синтаксис вызова функций Emu из обычного кода на Rust. В качестве примера будем использовать функцию, перемножающую все элементы вектора на заданное число:


use em::emu;

emu! {
    multiply(global_vector [f32], scalar f32) {
        global_vector[get_global_id(0)] *= scalar;
    }
}

Чтобы транслировать данную функцию в код на OpenCL, вам необходимо поместить её сигнатуру в макрос build! следующим образом:


use em::build;

// Необходимо для макроса build! {...}
extern crate ocl;
use ocl::{flags, Platform, Device, Context, Queue, Program, Buffer, Kernel};

build! { multiply [f32] f32 }

Дальнейшие действия сводятся к вызову написанных вами функций на Emu из кода на Rust. Проще быть не может:


fn main() {
    let vector = vec![0.4445, 433.245, 87.539503, 2.0];
    let result = multiply(vector, 2.0).unwrap();
    dbg!(result);
}

Пример прикладной программы


Данная программа первым аргументом принимает скаляр, на который необходимо домножить следующие аргументы. Результирующий вектор будет напечатан в консоль:


use em::{build, emu};

// Необходимо для макроса build! {...}
extern crate ocl;
use ocl::{flags, Buffer, Context, Device, Kernel, Platform, Program, Queue};

emu! {
    multiply(global_vector [f32], scalar f32) {
        global_vector[get_global_id(0)] *= scalar;
    }
}

build! { multiply [f32] f32 }

fn main() {
    // Получить все аргументы командной строки:
    let args = std::env::args().collect::<Vec<String>>();
    if args.len() < 3 {
        panic!("Использование: cargo run -- <SCALAR> <NUMBERS>...");
    }

    // Скаляр должен быть указан первым аргументом:
    let scalar = args[1].parse::<f32>().unwrap();

    // Сконвертировать вектор строк в вектор чисел:
    let vector = args[2..]
        .into_iter()
        .map(|string| string.parse::<f32>().unwrap())
        .collect();

    // Умножить и напечатать результат:
    let result = multiply(vector, scalar).unwrap();
    dbg!(result);
}

Выполнить данный код можно командой cargo run -- 3 2.1 3.6 6.2. Полученный вывод соответствует ожиданиям:


[src/main.rs:33] result = [
    6.2999997,
    10.799999,
    18.599998,
]

Связка с OpenCL


Как было сказано ранее, Emu — это лишь абстракция над OpenCL, а следовательно, она обладает возможностью взаимодействия с крейтом ocl. Код ниже взят из примера в официальном репозитории:


use em::emu;

// Крейт "ocl" предназначен для упаковки кода в функции на Rust:
extern crate ocl;
use ocl::{flags, Platform, Device, Context, Queue, Program, Buffer, Kernel};

// Скомпилировать Emu в промежуточный код (OpenCL) и поместить
// его в глобальную переменную "EMU: &'static str":
emu! {
    // Умножает элементы буфера на скаляр:
    multiply(global_buffer [f32], coeff f32) {
        global_buffer[get_global_id(0)] *= coeff;
    }
}

fn multiply(global_buffer: Vec<f32>, coeff: f32) -> ocl::Result<Vec<f32>> {
    // Сперва нужно определить девайс и платформу для использования,
    // создать контекст, очередь, программу и измерения:
    let platform = Platform::default();
    let device = Device::first(platform)?;
    let context = Context::builder()
        .platform(platform)
        .devices(device.clone())
        .build()?;
    let program = Program::builder()
        .devices(device)
        .src(EMU)
        .build(&context)?;
    let queue = Queue::new(&context, device, None)?;
    let dims = global_buffer.len();

    // Затем необходимо создать буфер:
    let buffer = Buffer::<f32>::builder()
        .queue(queue.clone())
        .flags(flags::MEM_READ_WRITE)
        .len(dims)
        .copy_host_slice(&global_buffer)
        .build()?;

    // Третьим пунктом необходимо создать ядро с аргументами,
    // соответствующими исходному коду выше:
    let kernel = Kernel::builder()
        .program(&program)
        .name("multiply")
        .queue(queue.clone())
        .global_work_size(dims)
        .arg(&buffer)
        .arg(&coeff)
        .build()?;

    // Запустить созданное ядро (параметры по умолчанию выбраны
    // для демонстрационных целей:
    unsafe {
        kernel.cmd()
            .queue(&queue)
            .global_work_offset(kernel.default_global_work_offset())
            .global_work_size([dims, 0, 0])
            .local_work_size(kernel.default_local_work_size())
            .enq()?;
    }

    // И наконец, считать результаты из девайса в вектор с длинной
    // "dims":
    let mut vector = vec![0.0f32; dims];
    buffer.cmd()
        .queue(&queue)
        .offset(0)
        .read(&mut vector)
        .enq()?;

    Ok(vector)
}

fn main() {
    let initial_data = vec![3.7, 4.5, 9.0, 1.2, 8.9];

    // Вызвать функцию умножения, написанную на Emu, на
    // векторе "initial_data":
    let final_data = multiply(initial_data, 3.0).unwrap();

    println!("{:?}", final_data);
}

Завершение


Надеюсь, что статья вам понравилась. Быстрый ответ на возникшие вопросы вы можете получить в русскоязычном чате по языку Rust (версия для новичков).



+36
5.1k 41
Support the author
Comments 14
Top of the day