Online compiler

Online compiler#

The double-batched FFT library uses online (“just-in-time”) compilation due to the vast number of FFT configurations. The online compilation API is described here.

OpenCL#

cl_program bbfft::cl::build_kernel_bundle(std::string const &source, cl_context context, cl_device_id device, std::vector<std::string> const &options = {}, std::vector<std::string> const &extensions = {})#

Compile OpenCL-C code to an OpenCL program.

Compiler options are defined in the OpenCL standard: https://registry.khronos.org/OpenCL/specs/3.0-unified/html/OpenCL_API.html#compiler-options

Parameters:

  • source – Source code

  • context – OpenCL context

  • device – OpenCL device

  • options – List of compiler options

  • extensions – List of OpenCL-C extensions

Returns:

OpenCL program

cl_program bbfft::cl::build_kernel_bundle(uint8_t const *binary, std::size_t binary_size, module_format format, cl_context context, cl_device_id device)#

Build OpenCL program from native binary.

Parameters:

  • binary – Pointer to binary blob

  • binary_size – Size of binary blob

  • format – Binary format

  • context – OpenCL context

  • device – OpenCL device

Returns:

OpenCL program

cl_kernel bbfft::cl::create_kernel(cl_program prog, std::string const &name)#

Create kernel from program.

Parameters:

  • prog – OpenCL program

  • name – Kernel name

Returns:

OpenCL kernel

aot_module bbfft::cl::create_aot_module(uint8_t const *binary, std::size_t binary_size, module_format format, cl_context context, cl_device_id device)#

Build module for ahead-of-time kernel cache (aot_cache)

Parameters:

  • binary – Pointer to native device binary blob

  • binary_size – Size of native device binary blob

  • format – Binary format

  • context – OpenCL context

  • device – OpenCL device

Returns:

ahead-of-time module

Level Zero#

ze_module_handle_t bbfft::ze::build_kernel_bundle(std::string const &source, ze_context_handle_t context, ze_device_handle_t device, std::vector<std::string> const &options = {}, std::vector<std::string> const &extensions = {})#

Compile OpenCL-C code to a Level Zero module.

Compiler options are defined in the OpenCL standard: https://registry.khronos.org/OpenCL/specs/3.0-unified/html/OpenCL_API.html#compiler-options

Parameters:

  • source – Source code

  • context – Level Zero context

  • device – Level Zero device

  • options – List of compiler options

  • extensions – List of OpenCL-C extensions

Returns:

Level Zero module

ze_module_handle_t bbfft::ze::build_kernel_bundle(uint8_t const *binary, std::size_t binary_size, module_format format, ze_context_handle_t context, ze_device_handle_t device)#

Build Level Zero module from native binary.

Parameters:

  • binary – Pointer to binary blob

  • binary_size – Size of binary blob

  • format – Binary format

  • context – Level Zero context

  • device – Level Zero device

Returns:

Level Zero module

ze_kernel_handle_t bbfft::ze::create_kernel(ze_module_handle_t mod, std::string const &name)#

Create kernel from module.

Parameters:

  • mod – Level Zero module

  • name – Kernel name

Returns:

Level Zero kernel

std::vector<uint8_t> bbfft::ze::compile_to_spirv(std::string const &source, std::string const &device_type = "", std::vector<std::string> const &options = {}, std::vector<std::string> const &extensions = {})#

Takes OpenCL-C code and outputs SPIR-V.

Note: some OpenCL extensions require the device type even for SPIR-V-only compilation

Parameters:

  • source – OpenCL-C source code

  • device_type – Target device type; see ocloc compile &#8212;help for possible targets

  • options – List of compiler options

  • extensions – List of OpenCL-C extensions

Returns:

binary

std::vector<uint8_t> bbfft::ze::compile_to_native(std::string const &source, std::string const &device_type, std::vector<std::string> const &options = {}, std::vector<std::string> const &extensions = {})#

Takes OpenCL-C code and outputs the native device binary.

This function is a thin wrapper around ocloc

Parameters:

  • source – OpenCL-C source code

  • device_type – Target device type; see ocloc compile &#8212;help for possible targets

  • options – List of compiler options

  • extensions – List of OpenCL-C extensions

Returns:

binary

aot_module bbfft::ze::create_aot_module(uint8_t const *binary, std::size_t binary_size, module_format format, ze_context_handle_t context, ze_device_handle_t device)#

Build module for ahead-of-time kernel cache (aot_cache)

Parameters:

  • binary – Pointer to native device binary blob

  • binary_size – Size of native device binary blob

  • format – Binary format

  • context – Level Zero context

  • device – Level Zero device

Returns:

ahead-of-time module

SYCL#

auto bbfft::sycl::build_native_module(std::string const &source, ::sycl::context context, ::sycl::device device, std::vector<std::string> const &options = {}, std::vector<std::string> const &extensions = {}) -> module_handle_t#

Build native module of SYCL back-end.

Compiler options are defined in the OpenCL standard: https://registry.khronos.org/OpenCL/specs/3.0-unified/html/OpenCL_API.html#compiler-options

Parameters:

  • source – OpenCL-C code

  • context – context

  • device – device

  • options – List of compiler options

  • extensions – List of OpenCL-C extensions

Returns:

Handle

auto bbfft::sycl::build_native_module(uint8_t const *binary, std::size_t binary_size, module_format format, ::sycl::context context, ::sycl::device device) -> module_handle_t#

Build native module of SYCL back-end from native binary.

Parameters:

  • binary – Pointer to binary blob

  • binary_size – Size of binary blob

  • format – Binary format

  • context – context

  • device – device

Returns:

Handle

auto bbfft::sycl::make_shared_handle(module_handle_t mod, ::sycl::backend be) -> shared_handle<module_handle_t>#

Make shared handle from native handle.

Parameters:

  • mod – native handle

  • be – backend

Returns:

Shared native handle

auto bbfft::sycl::make_kernel_bundle(module_handle_t native_module, bool keep_ownership, ::sycl::context context) -> ::sycl::kernel_bundle<::sycl::bundle_state::executable>#

Create kernel bundle from native module.

Parameters:

  • native_module – Native module

  • keep_ownership – False if ownership shall be passed to SYCL kernel bundle

  • context – context

Returns:

Kernel bundle

auto bbfft::sycl::create_kernel(::sycl::kernel_bundle<::sycl::bundle_state::executable> bundle, std::string const &name) -> ::sycl::kernel#

Create kernel from bundle.

Parameters:

  • bundle – kernel bundle

  • name – Kernel name

Returns:

Kernel

aot_module bbfft::sycl::create_aot_module(uint8_t const *binary, std::size_t binary_size, module_format format, ::sycl::context context, ::sycl::device device)#

Build module for ahead-of-time kernel cache (aot_cache)

Parameters:

  • binary – Pointer to native device binary blob

  • binary_size – Size of native device binary blob

  • format – Binary format

  • context – context

  • device – device

Returns:

ahead-of-time module

Enumerations#

enum class bbfft::module_format#

Binary module format.

Values:

enumerator spirv#

SPIR-V.

enumerator native#

Native device binary.