Graphics Processing Units (GPUs) have become an essential component of modern computing systems, powering not only computer graphics but also a wide range of general-purpose applications. However, GPU software development remains challenging, since the main tools available are APIs (Application Programming Interfaces) designed for low-level languages such as C and C++.

Developing in low-level languages for accelerators introduces several challenges, such as manual memory management, device synchronization, and explicit parallelism orchestration. These complexities have motivated the search for simpler and more productive approaches to GPU programming.

To address these challenges, PolyHok was developed: a Domain-Specific Language (DSL) embedded in the Elixir programming language. PolyHok adopts several functional programming features to provide high expressiveness and ease of use, most notably by its polymorphic higher-order kernels (hence the name!). These are functions that run on the GPU and can receive other functions as parameters, allowing users to build abstractions such as GPU-accelerated algorithmic skeletons. PolyHok also provides dynamic typing, automatic device memory management via the Erlang VM (BEAM) garbage collector, and support for multiple execution backends.

Introduced in version 1.1, PolyHok supports multiple execution backends that users can choose from. Backends provide the necessary functions and Native Implemented Functions (NIFs) for Elixir to interact with GPU hardware. Each backend can use completely different APIs and execution platforms to interact with the GPU, as long as it satisfies the interface expected by the PolyHok front-end.

Currently, PolyHok offers two backends. The first uses CUDA and is recommended for NVIDIA GPUs. The second is based on OpenCL and is designed to run on accelerators from any vendors, such as AMD GPUs/iGPUs and Intel integrated GPUs.

Before version 1.1, PolyHok did not support swappable backends. Instead, the DSL was tightly coupled to each backend implementation. When the PolyHok Team developed the OpenCL version of the language, it was nicknamed “OCL-PolyHok” to distinguish it from the original CUDA-based implementation. This is the name you may find in our first papers.

Having the backends tightly coupled with the DSL implementation meant that any updates to the core language had to be manually replicated in each PolyHok version, even when those changes did not affect the backend itself. This approach was error-prone and made the development and maintenance of PolyHok more difficult.

The PolyHok Team then decided to greatly improve this architecture in PolyHok v1.1. The PolyHok front-end (the language core) is now completely decoupled from its execution backends. This means that changes to the DSL do not require modifications to the backends, and changes to a backend does not affect the language core.

All backends implement a standardized interface (using Elixir behaviours) that makes them easily interchangeable. This architecture also makes it significantly easier to add new backends to PolyHok in the future!

Below, we present a simple program to demonstrate the basic functionality of PolyHok. In this example, we create an array of numbers on the CPU, copy it to GPU memory (VRAM), add 1 to each element in parallel using a kernel, and bring the result back to the CPU.

First, we need to create a new Mix project and add PolyHok as a regular Elixir dependency.

First, create a new Elixir Mix project, like this:

$ mix new hello_polyhok
$ cd hello_polyhok

Then, add PolyHok, Nx, and a PolyHok backend of your choice as dependencies to the project's mix.exs file. For example, if we want to use the OpenCL backend we could write:

defp deps do
  [
    {:nx, "~> 0.9"},
    {:poly_hok, git: "https://github.com/lups-ufpel/poly_hok.git", sparse: "poly_hok"},
    {:opencl_backend, git: "https://github.com/lups-ufpel/poly_hok.git", sparse: "backends/opencl_backend"}
  ]
end

If you have an NVIDIA GPU and prefer CUDA, replace the OpenCL backend dependency line with:

{:cuda_backend, git: "https://github.com/lups-ufpel/poly_hok.git", sparse: "backends/cuda_backend"}

After adding the dependencies, download and configure them with:

$ mix deps.get

PolyHok's front-end is completely isolated from the backend. Therefore, we must tell PolyHok what backend should be used at runtime by configuring the project's config/runtime.exs file.

For example, when using the OpenCL backend, add this to your runtime.exs file (if the file doesn't exist, create it):

import Config

config :poly_hok, backend: OpenclBackend

For the CUDA backend, use:

import Config

config :poly_hok, backend: CudaBackend

With this configuration, PolyHok will know what backend should be called for GPU operations. Just adding it as a dependency will not configure PolyHok.

Now that the project is configured, we can write our first PolyHok program! Create a file named hello_polyhok.exs in the project root with the following code:

require PolyHok

PolyHok.defmodule SimpleTest do
  defk simple_kernel(array, size) do
    index = blockIdx.x * blockDim.x + threadIdx.x

    if (index < size) do
      array[index] = array[index] + 1.0
    end
  end
end

array_size = 100

# Create a tensor of float elements on CPU memory with values from 1 to array_size
array_cpu = Nx.tensor(Enum.to_list(1..array_size), type: {:f, 32})

IO.inspect(array_cpu, label: "CPU Array")

# Create a tensor on the GPU with the same data from the CPU tensor
array_gpu = array_cpu |> PolyHok.new_gnx()

# Spawn the kernel to run on the GPU
PolyHok.spawn(
  &SimpleTest.simple_kernel/2, # Kernel function
  {1, 1, 1},                   # Number of blocks
  {array_size, 1, 1},          # Threads per block
  [array_gpu, array_size]      # Kernel parameters
)

# Get the result back to the CPU
result = PolyHok.get_gnx(array_gpu)

IO.inspect(result, label: "Result after kernel execution")

Finally, we can run the program with:

$ mix run hello_polyhok.exs

Our program follows the classic workflow of heterogeneous computing: data preparation, transfer to the device, kernel execution, and result retrieval.

• Data Preparation: We start by creating an Nx tensor (which is basically an array) of floating-point numbers with size 100 on the CPU. We use the variable array_size and the expression Enum.to_list(1..array_size) to populate it with values from 1 to 100. Once created, we print the array to the screen so we can compare it later with the result.

• Copying to the GPU: To allow the GPU to process the data, we create a GNx (GPU Nx) from the CPU tensor using array_cpu |> PolyHok.new_gnx(). A GNx is simply an array that resides in GPU memory (VRAM) and can only be accessed and modified by kernels and device functions executed on the GPU.

• Kernel Execution: We launch the simple_kernel kernel using PolyHok.spawn/4. Notice the execution configuration: we use 1 block ({1, 1, 1}) containing 100 threads ({array_size, 1, 1}). Inside the kernel, each thread computes its own global index (index = blockIdx.x * blockDim.x + threadIdx.x) to independently determine which element of the array it should access and increment.

• Reading the Results: After execution, the modified GNx remains in GPU memory. We use PolyHok.get_gnx/1 to bring the data back to the CPU. Once retrieved, we can print the result to compare it with the original array.

$ mix run hello_polyhok.exs

CPU Array: #Nx.Tensor<
  f32[100]
  [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0, 13.0, 14.0, 15.0, 16.0, 17.0, 18.0, 19.0, 20.0, 21.0, 22.0, 23.0, 24.0, 25.0, 26.0, 27.0, 28.0, 29.0, 30.0, 31.0, 32.0, 33.0, 34.0, 35.0, 36.0, 37.0, 38.0, 39.0, 40.0, 41.0, 42.0, 43.0, 44.0, 45.0, 46.0, 47.0, 48.0, 49.0, 50.0, 51.0, 52.0, 53.0, 54.0, 55.0, 56.0, 57.0, 58.0, 59.0, 60.0, 61.0, 62.0, 63.0, 64.0, 65.0, 66.0, 67.0, 68.0, 69.0, 70.0, 71.0, 72.0, 73.0, 74.0, 75.0, 76.0, 77.0, 78.0, 79.0, 80.0, 81.0, 82.0, 83.0, 84.0, 85.0, 86.0, 87.0, 88.0, 89.0, 90.0, 91.0, 92.0, 93.0, 94.0, 95.0, 96.0, 97.0, 98.0, 99.0, 100.0]
>
Result after kernel execution: #Nx.Tensor<
  f32[100]
  [2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0, 13.0, 14.0, 15.0, 16.0, 17.0, 18.0, 19.0, 20.0, 21.0, 22.0, 23.0, 24.0, 25.0, 26.0, 27.0, 28.0, 29.0, 30.0, 31.0, 32.0, 33.0, 34.0, 35.0, 36.0, 37.0, 38.0, 39.0, 40.0, 41.0, 42.0, 43.0, 44.0, 45.0, 46.0, 47.0, 48.0, 49.0, 50.0, 51.0, 52.0, 53.0, 54.0, 55.0, 56.0, 57.0, 58.0, 59.0, 60.0, 61.0, 62.0, 63.0, 64.0, 65.0, 66.0, 67.0, 68.0, 69.0, 70.0, 71.0, 72.0, 73.0, 74.0, 75.0, 76.0, 77.0, 78.0, 79.0, 80.0, 81.0, 82.0, 83.0, 84.0, 85.0, 86.0, 87.0, 88.0, 89.0, 90.0, 91.0, 92.0, 93.0, 94.0, 95.0, 96.0, 97.0, 98.0, 99.0, 100.0, 101.0]
>