Introduction

Emu is a high-level programming language for video cards that can be embedded in the usual code in the system programming language Rust .

This article focuses on the Emu syntax, its features, and will also show a few illustrative examples of its use in real code.

Installation

1. The overlooked library needs an external OpenCL dependency. You need to install the appropriate driver for your hardware.
2. Cargo.toml text below. This will cause the latest available versions to be downloaded (if you need a specific build, instead of * put the required version):
   
   

    [dependencies] em = "*" //   Emu ocl = "*" //   OpenCL 

Syntax

Emu's syntax is quite simple, because this language is intended only for writing kernel functions that are translated to OpenCL during compilation.

Data types

The Emu language has nine data types that are similar to those in Rust. Below is a table of these types:

Variables

Variables are declared using the let keyword, which is located after the identifier, the colon, the data type, the equal sign, the value assigned, and the semicolon.

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

Conversions

Conversion of primitive data types is carried out by means of the binary operator as , following the target type. I note that the target type can also be a unit of measurement (see the following section):

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

Units

The Emu language allows you to treat numbers as units of measure, which is designed to simplify scientific calculations. In this example, the variable length initially defined in meters, but then other units of measure are added to it:

 let length: f32 = 3455.345; //  length += 7644.30405 as cm; //  length += 1687.3043 as mm; //  

Predefined constants

Emu has a set of predefined constants that are convenient to use in practice. Below is the corresponding table.

Also defined and constants corresponding to scientific data. The table, consisting of these constants, you can find here .

Conditional statements

Emu conditional operators are similar to the corresponding operators in Rust. Below is the code that applies conditional constructions:

 let number: i32 = 2634; let satisfied: bool = false; if (number > 0) && (number % 2 == 0) { satisfied = true; } 

For loops

The header of a For loop is defined as for NUM in START..END , where NUM is a variable that takes values ​​from the [START; END) [START; END) through the unit.

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

While loops

The header of the While loop is defined as while (CONDITION) , where CONDITION is the condition for the loop to go to the next iteration. This code is similar to the previous example:

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

Infinite loops

Infinite loops do not have an explicit exit condition and are defined by the keyword loop . However, they can be continued or interrupted by means of the break and continue operators (like the other two types of cycles).

 let collapsed: u64 = 1; let idx: i32 = 0; loop { if idx % 2 == 0 { continue; } sum *= idx; if idx == 12 { break; } } 

Return from function

As in all other programming languages, the return serves as an exit from the current function. It can also return some value if the function signature (see the next sections) allows it.

 let result: i32 = 23446; return result; 

Other operators

• Available assignment operators: = , += , -= , *= , /= , %= , &= , ^= , <<= , >>= ;
• The index operator is [IDX] ;
• Call Operator - (ARGS) ;
• Unary operators: * for dereference ! to invert boolean data, to negate numbers;
• Binary operators: + , - , * , / , % , && , || , & , | , ^ , >> , << , > , < , >= , <= , == != .

Functions

In total there are three parts of functions on Emu: the identifier, the parameters and the function body, consisting of a sequence of executable instructions. Consider the addition function of two numbers:

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

As you can see, this function returns the sum of the two arguments passed to it using the f32 data f32 .

Address spaces

Each parameter of the function corresponds to a specific address space . By default, all parameters correspond to the __private__ space.

Adding the prefixes global_ and local_ to the parameter identifier explicitly indicates its address space.

The documentation recommends using the global_ prefix to all vectors and not prefixing anything else.

Built-in functions

Emu provides a small set of built-in functions (taken from OpenCL) that allow you to manage GPU data:

• get_work_dim() - Returns the number of dimensions;
• get_global_size() - Returns the number of global elements for the specified dimension;
• get_global_id() - Returns the unique identifier of the element for the specified dimension;
• get_global_size() - Returns the number of global elements for the specified dimension;
• get_local_id() - Returns the unique identifier of a local element within a specific workgroup for a given dimension;
• get_num_groups() - Returns the number of work groups for the specified dimension;
• get_group_id() - Returns a unique identifier for a workgroup.

In the application code, most often you will see the expression get_global_id(0) , which returns the current index of the element of the vector associated with the call to your kernel function.

Code execution

Consider the syntax for calling Emu functions from regular Rust code. As an example, we will use a function that multiplies all elements of a vector by a given number:

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

To translate this function into OpenCL code, you need to put its signature in the build! macro build! in the following way:

 use em::build; //    build! {...} extern crate ocl; use ocl::{flags, Platform, Device, Context, Queue, Program, Buffer, Kernel}; build! { multiply [f32] f32 } 

Further actions are reduced to calling the functions you wrote on Emu from the code in Rust. It can not be easier:

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

Sample Application Program

This program takes the first argument of a scalar, which is necessary to multiply the following arguments. The resulting vector will be printed to the console:

 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); } 

You can execute this code with the cargo run -- 3 2.1 3.6 6.2 command cargo run -- 3 2.1 3.6 6.2 . The resulting output meets expectations:

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

Completion

I hope you enjoyed the article. You can get a quick answer to your questions in the Russian-language chat using the Rust language ( beginner version ).