Идиоматичное программирование GPU на Rust: Библиотека Emu04.06.2019 09:32

y_lbgimanxfupq85ezcf_pn_nk4.jpeg

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

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


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

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


Типы данных

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


Переменные

Переменные объявляются с помощью ключевого слова 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 располагает набором предопределённых констант, которые удобно использовать на практике. Ниже приведена соответствующая таблица.

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


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

Условные операторы 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::>();
    if args.len() < 3 {
        panic!("Использование: cargo run --  ...");
    }

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

    // Сконвертировать вектор строк в вектор чисел:
    let vector = args[2..]
        .into_iter()
        .map(|string| string.parse::().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,
]

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


© Habrahabr.ru