Backend Primitives#

This section describes the types related to data management backend primitives. Those primitives cannot be used directly in the user code but are used internally by oneDAL to have a convenient and efficient way to work with data.

Data Management Backend Primitives#

Primitive

Description

Multidimensional array order

An enumeration of multidimensional data orders used to store contiguous data blocks inside the table.

Multidimensional array shape

A class that represents the shape of a multidimensional array.

Multidimensional view

An implementation of a multidimensional data container that provides a view of the homogeneous data stored in an externally-managed memory block.

Multidimensional array

A class that provides a way to store and manipulate homogeneous data in a multidimensional structure.

Table to ndarray conversion functions

Functions that create ndarray objects from data tables. If possible, the data is stored in the same memory location as the data table and no data copying is performed.

Usage Example#

The following listing provides an example of ndarray API to illustrate how data management backend primitives can be used with oneDAL table types to perform computations on the data:

namespace pr = dal::backend::primitives;

// Create a 2D ndarray from input data table
auto data_nd = pr::table2ndarray<float>(queue, data_table, sycl::usm::alloc::device);
const std::int64_t row_count = data_table.get_row_count();
const std::int64_t column_count = data_table.get_column_count();

// Create a 1D ndarray to store the sum of each column
auto sum_nd = pr::ndarray<float, 1>::empty(queue,
                                           { column_count },
                                           sycl::usm::alloc::device);

// Get USM pointers to the data and sum arrays
const float * data_ptr = data_nd.get_data();
float * sum_ptr = sum_nd.get_mutable_data();

constexpr std::int64_t row_block_size    = 1024;
constexpr std::int64_t column_block_size = 512;

// Define the SYCL* range for the kernel
const auto range = sycl::nd_range<2>({ row_block_size, column_block_size }, { 1, column_block_size });

// Compute the sum of each data row
auto event = queue.submit([&](sycl::handler& cgh) {
      queue.submit([&](sycl::handler& cgh) {
      cgh.parallel_for(range, [=](auto it) {
          const std::int64_t row_shift = it.get_global_id(0);
          const std::int64_t col_shift = it.get_local_id(1);
          for (auto row_idx = row_shift; row_idx < row_count; row_idx += row_block_size) {
              const auto start = row_idx * column_count;
              const auto end = start + column_count;
              // Exclusive storage of the partial sum for execution unit
              float local_sum = 0.0f;
              for (auto idx = start + col_shift; idx < end; idx += column_block_size) {
                  local_sum += data_ptr[idx];
              }
              // Reduction over the workgroup
              sum_ptr[row_idx] = sycl::reduce_over_group(it.get_group(), local_sum, sycl::plus<float>());
          }
      });
  });

Programming interface#

Refer to API: Data Management Backend Primitives.