HSA.jl

This is a package wrapping the Heterogeneous System Architecture (HSA) runtime libraries.

For links to official HSA Documentation, check out the wiki

Library Discovery

The package needs to know where to find the HSA runtime. The code responsible for discovering the HSA libraries and headers is located in src/discovery.jl. It expects to find the path containing the HSA libraries (libhsa-runtime64, libhsakmt [and libhsa-runtime-ext64.so]) in ENV["LD_LIBRARY_PATH"].

For a port of the vector_copy example that comes with the HSA runtime, see the sample directory.

Julia functions as Kernels

In combination with a modified Julia binary, HSA.jl can be used to execute Kernel functions written in Julia on an HSA agent.

Minimal example:

using HSA

@hsa_kernel function vector_copy(a, b)
    i = get_global_id()
    b[i] = a[i]
end

N =  1024*1024

a_in = Array(Int, N)
rand!(a_in)

b_out = Array(Int, N)
rand!(b_out)

assert(a_in != b_out)

@hsa (N) vector_copy_kernel(a_in, b_out)

assert(a_in == b_out)

More examples using this facility can be found in samples/codegen

Module overview

• HSA Container Module for all HSA functionality.
  • Wrappers for functions and constants defined by the HSA runtime.
  • Functionality to emulate execution of Julia kernels
    • On a non-HSA CPU
    • if there is no codegen support
• HSA.Builtins Functions that compile down to device intrinsics for use in Kernels
• HSA.ExtFinalization Custom Wrappers for the finalization extension to the HSA Runtime

Source Folder

• binding Contains the code that interfaces with the various parts of the HSA Runtime
  • generated Files generated from the hsa headers via gen/generate.jl
  • custom Custom wrapper code
• codegen Contains all code that relies on the modified julia code generator to work
• emulation Provides emulation for some parts of the codegen infrastructure even when no hsa or codegen are available

Source Generation

The files in src/binding/generated are generated by the script gen/generate.jl. Regenerating these is usually unnecessary unless the HSA headers have changed or you made changes to the generation script. The script has some additional dependencies (mainly Clang.jl) and needs to be able to find the HSA headers.

Known Issues

• HSA.executable_get_symbol does not find symbols
  Maybe the name string for the symbol is not being passed correctly?
  Workaround: iterate over all symbols and inspect their names