Configuring data layouts
This page describes how to configure a descriptor object for a specific data layout. When non-native, all the relevant types and enumerations mentioned below belong to the oneapi::mkl::dft namespace and are declared in oneapi/mkl/dft.hpp (file to be included). The usage of prepended namespace specifiers oneapi::mkl::dft is omitted below for conciseness.
The DPC++ interface provides the configuration parameter config_param::FWD_STRIDES (resp. config_param::BWD_STRIDES) to define the data layout locating entries (or parts thereof) of relevant data sequences in the forward (resp. backward) domain. In case of batched transforms, i.e., if the configuration value for config_param::NUMBER_OF_TRANSFORMS is set to an integer \(M\) larger than \(1\), the value set for configuration parameter config_param::FWD_DISTANCE (resp. config_param::BWD_DISTANCE) completes the description of the data layout by specifying the distances between successive data sequences in the forward (resp. backward) domain.
This topic leverages the general notations from the introduction, and uses the superscript \(\text{fwd}\) (resp. \(\text{bwd}\)) for data sequences in forward (resp. backward) domain. A placeholder label \(\text{v}\) is also used to capture a possible distinction between the real (if \(\text{v}\) is \(\text{r}\)) and imaginary (if \(\text{v}\) is \(\text{i}\)) parts of a complex data entry; naturally, that placeholder label \(\text{v}\) is relevant only for data layouts that distinguish real and imaginary parts of complex data entries.
A non-redundant entry \(\left(\cdot\right)^{m}_{k_{1}, k_{2}, \ldots, k_{d}}\) (or its real or imaginary part, if relevant) is stored at index \(J^{\text{xwd}}\left(k_{1}, \ldots, k_{d}, m; \text{v}\right)\) of the appropriate data container (sycl::buffer object or device-accessible USM allocation) provided to a compute function, the base data type of which is (possibly implicitly re-interpreted as) documented in the table below. That index value is defined as
\[J^{\text{xwd}}\left(k_{1}, \ldots, k_{d}, m; \text{v}\right) = s^{\text{xwd}}_{0} + k_{1}\ s^{\text{xwd}}_{1} + k_{2}\ s^{\text{xwd}}_{2} + \ldots + k_{d-1}\ s^{\text{xwd}}_{d-1} + p^{\text{xwd}}_{d}\left(k_{d};\text{v}\right)\ s^{\text{xwd}}_{d} + m\ l^{\text{xwd}},\]
wherein
\(s^{\text{xwd}}_{j}\), \(\forall j \in \lbrace 0, \ldots, d\rbrace\) represents the offset and generalized strides defining the locations of relevant values within each \(d\)-dimensional data sequence in the forward (resp. backward) domain if \(\text{x} = \text{f}\) (resp. if \(\text{x} = \text{b}\)), counted in number of elements of the relevant implicitly-assumed elementary data type;
\(l^{\text{xwd}}\) represents the distance between successive \(d\)-dimensional data sequences in the forward (resp. backward) domain if \(\text{x} = \text{f}\) (resp. if \(\text{x} = \text{b}\)), counted in number of elements of the relevant implicitly-assumed elementary data type;
the relation \(p^{\text{xwd}}_{d}\left(k_{d};\text{v}\right)\) simplifies into the identity \(p^{\text{xwd}}_{d}\left(k_{d};\text{v}\right) = k_{d}\) in all recommended use cases or if \(\text{x} = \text{f}\), i.e., \(\text{v}\) is either irrelevant or unused in such cases.
In this page, it is assumed that only non-redundant data sequence entries are of interest, i.e., that \(0\leq m < M\), \(0\leq k_{j} < n_{j}, \forall j \in \lbrace 1, 2, \ldots, d-1\rbrace\) and that \(0 \leq k_{d} \leq \lfloor \frac{n_{d}}{2}\rfloor\) (resp. \(0 \leq k_{d} < n_{d}\)) for entries that do (resp. do not) belong the backward domain of a real DFT.
Note that all elements accessed as a value stored at index \(J^{\text{xwd}}\left(k_{1}, \ldots, k_{d}, m; \text{v}\right)\) of a given user-provided data container must belong to the same block allocation.
Implicitly-assumed elementary data type
When reading or writing an element at index \(J^{\text{xwd}}\left(k_{1}, \ldots, k_{d}, m; \text{v}\right)\) of any user-provided data container used at compute time, a descriptor object may first re-interpret the base data type of that data container into an implicitly-assumed elementary data type. That implicitly-assumed data type depends on the object type; that is, on the specialization values used for the template parameters when instantiating the descriptor class, and on other configuration value(s). The table below lists the implicitly-assumed data type in either domain (last 2 columns) based on the object type and its configuration value(s).
Type of descriptor and relevant configuration values |
Implicitly-assumed elementary data type in forward domain |
Implicitly-assumed elementary data type in backward domain |
|---|---|---|
Complex descriptor with config_value::COMPLEX_COMPLEX set for config_param::COMPLEX_STORAGE |
std::complex<fp_type> |
std::complex<fp_type> |
Complex descriptor with config_value::REAL_REAL set for config_param::COMPLEX_STORAGE |
fp_type |
fp_type |
Real descriptor |
fp_type |
std::complex<fp_type>` |
Descriptors that implicitly assume an elementary data type of float or double (resp. std::complex<float> or std::complex<double>) in a domain are referred to as “descriptors expecting real (resp. complex) data” in that domain.
Configuring strides in forward and backward domains
The values \(s^{\text{xwd}}_0, s^{\text{xwd}}_1, \dots, s^{\text{xwd}}_d\) defining \(J^{\text{xwd}}\left(k_{1}, \ldots, k_{d}, m; \text{v}\right)\) are to be communicated as elements (in that order) of a std::vector<std::int64_t> object of size \((d+1)\), passed as the configuration value for config_param::FWD_STRIDES if \(\text{x} = \text{f}\) (resp. config_param::BWD_STRIDES if \(\text{x} = \text{b}\)) using the relevant configuration-setting member function. The element \(s^{\text{xwd}}_0\) represents an absolute offset (or “displacement”) in the data sets while the subsequent elements \(s^{\text{xwd}}_j\ (j > 0)\) are generalized strides to be considered along dimensions \(j \in \lbrace 1, \ldots, d\rbrace\).
When created, the descriptors are default-configured for unbatched, in-place transforms using a unit stride along the last dimension, no offset and the default configuration settings documented in the above table. For real descriptors, minimal padding is used in forward domain, aligning with the data layout requirements for in-place transforms.
In other words, the default stride values are \(s^{\text{fwd}}_0 = s^{\text{bwd}}_0 = 0\), \(s^{\text{fwd}}_d = s^{\text{bwd}}_d = 1\) and, for \(d\)-dimensional transforms with \(d > 1\),
\(s^{\text{fwd}}_{d-1} = s^{\text{bwd}}_{d-1} = n_{d}\) for complex descriptors;
\(s^{\text{bwd}}_{d-1} = \lfloor \frac{n_{d}}{2} \rfloor + 1\), and \(s^{\text{fwd}}_{d-1} = 2 s^{\text{bwd}}_{d-1}\) for real descriptors;
if \(d > 2\), \(s^{\text{xwd}}_k = n_{k+1} s^{\text{xwd}}_{k+1}\) for \(k \in \lbrace 1, \ldots, d - 2\rbrace\) (for \(\text{x} = \text{f}\) and \(\text{x} = \text{b}\)).
The usage of these default strides for unbatched, in-place transforms is illustrated in the usage examples.
Configuring batched transforms
The value \(l^{\text{xwd}}\) completing the definition of \(J^{\text{xwd}}\left(k_{1}, \ldots, k_{d}, m; \text{v}\right)\) is to be set as an std::int64_t configuration value for config_param::FWD_DISTANCE if \(\text{x} = \text{f}\) (resp. config_param::BWD_DISTANCE if \(\text{x} = \text{b}\)) using the relevant configuration-setting member function. This value is irrelevant for unbatched transforms, i.e., for descriptors set to handle a number of transforms \(M\) equal to \(1\) (default behavior).
In case of batched transforms, the desired number \(M\) of DFTs must be set explicitly as an std::int64_t configuration value config_param::NUMBER_OF_TRANSFORMS using the relevant configuration-setting member function. In that case, config_param::FWD_DISTANCE and config_param::BWD_DISTANCEmust also be set explicitly since their default configuration values of \(0\) would break the data layout requirements for any \(M > 1\).
The configuration of batched transforms is illustrated in the usage examples.
Data layout requirements
In general, the distances and strides must be set so that
values of \(J^{\text{xwd}}\left(k_{1}, \ldots, k_{d}, m; \text{v}\right)\) are non-negative for all \((d+1)\)-tuples \((k_{1}, k_{2}, \dots, p^{\text{xwd}}_{d}\left(k_{d};\text{v}\right), m)\) within relevant ranges;
every value of \(J^{\text{xwd}}\left(k_{1}, \ldots, k_{d}, m; \text{v}\right)\) corresponds to a unique value relevant to the data sequences under consideration. In other words, there must not be one value of \(J^{\text{xwd}}\left(k_{1}, \ldots, k_{d}, m; \text{v}\right)\) corresponding to two different \((d+1)\)-tuples \((k_{1}, k_{2}, \dots, p^{\text{xwd}}_{d}\left(k_{d};\text{v}\right), m)\) that would both be within relevant ranges.
Additionally, for in-place transforms (configuration value config_value::INPLACE set for config_param::PLACEMENT), the following “consistency requirements” apply:
descriptors expecting the same data type in either domain (e.g., complex descriptors) must use the same offset, stride(s), and distance values in forward and backward domains;
for real descriptors expecting complex data in backward domain (default behavior for real descriptors), the memory address(es) of leading entry(ies) along the last dimension must be identical in forward and backward domains. Specifically, that requirement translates into the conditions \(s^{\text{fwd}}_{j} = 2 s^{\text{bwd}}_{j}, \ \forall j \in \lbrace 0, \ldots, d - 1\rbrace\) as well as, if \(M > 1\), \(l^{\text{fwd}} = 2 l^{\text{bwd}}\). Note that this requirement leads to some data padding to be used in forward domain if unit strides are used along dimension \(d\) in forward and backward domains (recommended usage, as set by default).
- Support for negative strides with a sufficiently large (positive) offset index guaranteeing non-negativeness of all \(J^{\text{xwd}}\left(k_{1}, \ldots, k_{d}, m; \text{v}\right)\) is not enabled yet (unimplemented);
Supported layouts on GPU devices
On GPU devices, oneMKL requires
the rank \(d\) of the transform to be no greater than \(3\);
the offset values \(s^{\text{fwd}}_{0}\) and \(s^{\text{bwd}}_{0}\) to be \(0\);
either \(l^{\text{xwd}} = n_{1} s^{\text{xwd}}_{1}\) or \(s^{\text{xwd}}_{1} = M l^{\text{xwd}}\) for batched, two-dimensional real transforms (for \(\text{x} = \text{f}\)and\(\text{x} = \text{b}\));
either \(s^{\text{xwd}}_{1} = n_{2} s^{\text{xwd}}_{2}\) (along with \(l^{\text{xwd}} = n_{1} s^{\text{xwd}}_{1}\) if \(M > 1\)) or \(s^{\text{xwd}}_{2} = n_{1} s^{\text{xwd}}_{1}\) (along with \(s^{\text{xwd}}_{1} = M l^{\text{xwd}}\) if \(M > 1\)) for three-dimensional real transforms (for \(\text{x} = \text{f}\)and\(\text{x} = \text{b}\));
real descriptors require real data in the forward domain and complex data(non-redundant entries) in the backward domain.