データレイアウトの設定#
ここでは、特定のデータレイアウトの descriptor オブジェクトを構成する方法について説明します。非ネイティブの場合、以下で説明する関連するすべてのタイプと列挙は oneapi::mkl::dft 名前空間に属し、oneapi/mkl/dft.hpp (インクルード・ファイル) で宣言されます。簡潔にするため、先頭に付加された名前空間指定子 oneapi::mkl::dft は以下では省略されています。
DPC++ インターフェイスは、順方向 (または逆方向) ドメイン内の関連するデータシーケンスのエントリー (またはその一部) を配置するデータレイアウトを定義する構成パラメーターconfig_param::FWD_STRIDES (または config_param::BWD_STRIDES) を提供します。バッチ変換の場合、つまり、config_param::NUMBER_OF_TRANSFORMS の構成値が \(1\) より大きい整数 \(M\) に設定されている場合、構成パラメーター config_param::FWD_DISTANCE (または config_param::BWD_DISTANCE) に設定された値は、順方向 (または逆方向) ドメイン内の連続するデータシーケンス間の距離を指定することによって、データレイアウトの記述を完了します。
ここでは、概要の一般的な表記法を活用し、順方向 (または逆方向) 領域のデータシーケンスに上付き文字 \(\text{fwd}\) (または \(\text{bwd}\)) を使用します。プレースホルダー・ラベル \(\text{v}\) は、複素数データエントリーの実部 (\(\text{v}\) が \(\text{r}\) の場合) と虚部 (\(\text{v}\) が \(\text{i}\)) の場合) を区別するのにも使用されます。当然、プレースホルダー・ラベル \(\text{v}\) は、複素数データエントリーの実部と虚部を区別するデータレイアウトにのみ関連します。
非冗長エントリー \(\left(\cdot\right)^{m}_{k_{1}, k_{2}, \ldots, k_{d}}\) (または、該当する場合はその実数部または虚数部) は、計算関数に提供された適切なデータコンテナー (sycl::buffer オブジェクトまたはデバイスアクセス可能な USM 割り当て) のインデックス \(J^{\text{xwd}}\left(k_{1}, \ldots, k_{d}, m; \text{v}\right)\) に格納されます。その基本データタイプは、次の表に示されています (暗黙的に再解釈される可能性があります)。インデックスは次のように定義されます
ここで、
\(s^{\text{xwd}}_{j}\)、\(\forall j \in \lbrace 0, \ldots, d\rbrace\) は、\(\text{x} = \text{f}\) の場合 (または \(\text{x} = \text{b}\) の場合) の順方向 (または逆方向) ドメインの各 \(d\) 次元データシーケンス内の関連する値の位置を定義するオフセットと一般化ストライドを表します。これは、関連する暗黙的に想定される基本データタイプの要素数でカウントされます;
\(L^{\text{xwd}}\) は、\(\text{x} = \text{f}\) の場合 (または \(\text{x} = \text{b}\) の場合) の順方向 (または逆方向) ドメインの各\(d\) 次元データシーケンス間の距離を、暗黙的に想定される基本データタイプの要素数でカウントしたものを表します;
関係 \(p^{\text{xwd}}_{d}\left(k_{d};\text{v}\right)\) は、すべての推奨される使用ケースで恒等式 \(p^{\text{xwd}}_{d}\left(k_{d};\text{v}\right) = k_{d}\) に簡略化されます。または、\(\text{x} = \text{f}\) の場合、つまり、そのような場合には \(\text{v}\) は無関係であるか使用されません。ただし、非推奨の構成を使用する一部の 1 次元実数記述子では、逆方向ドメインのエントリー \(\left(\cdot\right)^{m}_{k_{1}}\) の実部と虚部は互いに別々に考慮され、対応するインデックスは \(J^{\text{bwd}}\left(k_{1}, m;\text{r}\right) = s^{\text{bwd}}_{0} + p^{\text{bwd}}_{1}\left(k_{1};\text{r}\right)\ s^{\text{bwd}}_{1} + m\ l^{\text{bwd}}\) および \(J^{\text{bwd}}\left(k_{1}, m;\text{i}\right) = s^{\text{bwd}}_{0} + p^{\text{bwd}}_{1}\left(k_{1};\text{i}\right)\ s^{\text{bwd}}_{1} + m\ l^{\text{bwd}}\) となります。
ここでは、非冗長データシーケンスのエントリーのみが対象であると仮定します。つまり、\(0\leq m < M\)、\(0\leq k_{j} < n_{j}, \forall j \in \lbrace 1, 2, \ldots, d-1\rbrace\) であり、実 DFT の逆方向領域に属するエントリーについては \(0 \leq k_{d} \leq \lfloor \frac{n_{d}}{2}\rfloor\) (または \(0 \leq k_{d} < n_{d}\)) となります。
特定のユーザー提供のデータコンテナーのインデックス \(J^{\text{xwd}}\left(k_{1}, \ldots, k_{d}, m; \text{v}\right)\) に格納された値としてアクセスされるすべての要素は、同じブロック割り当てに属している必要があることに注意してください。
暗黙的に想定される基本データタイプ#
計算時に使用されるユーザー提供のデータコンテナーのインデックス \(J^{\text{xwd}}\left(k_{1}, \ldots, k_{d}, m; \text{v}\right)\) の要素を読み取ったり書き込んだりする場合、descriptor オブジェクトは最初にそのデータコンテナーの基本データタイプを暗黙的に想定される基本データタイプに再解釈することがあります。暗黙的に想定されるデータタイプは、オブジェクト・タイプ、つまり descriptor クラスをインスタンス化するテンプレート・パラメーターとして使用される特殊化値とその他の構成値によって異なります。以下の表は、オブジェクト・タイプとその構成値に基づいて、いずれかのドメインで暗黙的に想定されるデータタイプ (最後の 2 列) を示しています。
記述子のタイプと関連する構成値 |
順方向ドメインで暗黙的に想定される基本データタイプ |
逆方向ドメインで暗黙的に想定される基本データタイプ |
|---|---|---|
config_param::COMPLEX_STORAGE に config_value::COMPLEX_COMPLEX が設定された複素数記述子(デフォルトの動作)
|
|
|
config_param::COMPLEX_STORAGE に config_value::REAL_REAL が設定された複合記述子(oneMKL の DPC++ インターフェイス経由では実装されていません)
|
|
|
config_param::CONJUGATE_EVEN_STORAGE に config_value::COMPLEX_COMPLEX が設定された実数記述子(デフォルトの動作)
|
|
|
config_param::CONJUGATE_EVEN_STORAGE に config_value::COMPLEX_REAL が設定された実数記述子(CPU 上の 1D DFT のみサポート、非推奨)
|
|
|
ドメイン内で float または double (それぞれ std::complex<float> と std::complex<double>) の基本データタイプを暗黙的に想定する記述子は、そのドメイン内で「実数 (複素数) データを想定する記述子」と呼ばれます。
順方向ドメインと逆方向ドメインでのストライドの設定#
\(J^{\text{xwd}}\left(k_{1}, \ldots, k_{d}, m; \text{v}\right)\) を定義する値 \(s^{\text{xwd}}_0, s^{\text{xwd}}_1, \dots, s^{\text{xwd}}_d\) は、サイズ \((d+1)\) の std::vector<std::int64_t> オブジェクトの要素として (この順序で) 通信され、関連する構成設定メンバー関数を使用して、\(\text{x} = \text{f}\) の場合は config_param::FWD_STRIDES、 (\(\text{x} = \text{b}\) の場合は config_param::BWD_STRIDES) の構成値として渡されます。要素 \(s^{\text{xwd}}_0\) はデータセット内の絶対オフセット (または「変位」) を表し、後続の要素 \(s^{\text{xwd}}_j\ (j > 0)\) は次元 \(j \in \lbrace 1, \ldots, d\rbrace\) に沿って考慮される一般化されたストライドです。
記述子は作成されると、最後の次元に沿った単位ストライド、オフセットなし、および上記の表に記載されているデフォルトの構成設定を使用して、バッチ処理されていないインプレース変換用にデフォルトで構成されます。実数記述子の場合、インプレース変換のデータレイアウト要件に合わせて、順方向ドメインで最小限のパディングが使用されます。
言い換えれば、デフォルトのストライド値は \(s^{\text{fwd}}_0 = s^{\text{bwd}}_0 = 0\)、\(s^{\text{fwd}}_d = s^{\text{bwd}}_d = 1\) であり、\(d > 1\) の \(d\) 次元変換の場合、
複雑な記述子では \(s^{\text{fwd}}_{d-1} = s^{\text{bwd}}_{d-1} = n_{d}\) です;
実数記述子では \(s^{\text{bwd}}_{d-1} = \lfloor \frac{n_{d}}{2} \rfloor + 1\)、と \(s^{\text{fwd}}_{d-1} = 2 s^{\text{bwd}}_{d-1}\) となります;
\(d > 2\) の場合、\(k \in \lbrace 1, \ldots, d - 2\rbrace\) に対して \(s^{\text{xwd}}_k = n_{k+1} s^{\text{xwd}}_{k+1}\) となります (\(\text{x} = \text{f}\) と \(\text{x} = \text{b}\) の場合)。
バッチ処理されていないインプレース変換でのこれらのデフォルトのストライドの使用方法は、使用例に示されています。
バッチ変換の構成#
\(J^{\text{xwd}}\left(k_{1}, \ldots, k_{d}, m; \text{v}\right)\) の定義を完了する値 \(l^{\text{xwd}}\) は、関連する構成設定メンバー関数を使用して、\(\text{x} = \text{f}\) の場合は config_param::FWD_DISTANCE、 (\(\text{x} = \text{b}\) の場合は config_param::BWD_DISTANCE) の構成値として設定されます。この値は、バッチ処理されていない変換、つまり、\(M\) が \(1\) に等しい変換の数を処理するように設定された記述子では無関係です (デフォルトの動作)。
バッチ変換の場合、関連する構成設定メンバー関数を使用して、必要な DFT の数 \(M\) を std::int64_t 構成値 config_param::NUMBER_OF_TRANSFORMS として明示的に設定する必要があります。その場合、config_param::FWD_DISTANCE と config_param::BWD_DISTANCE も明示的に設定する必要があります。これは、デフォルトの構成値 \(0\) が、\(M > 1\) のデータレイアウト要件に違反するためです。
バッチ変換の構成は使用例に示されています。
1 次元実数変換の逆方向ドメインにおける非推奨のレイアウト#
すべての複素数記述子と、逆方向領域で複素数データを期待する実数記述子は、単純な恒等関係 \(p^{\text{bwd}}_{d}\left(k_{d}; \text{v}\right) = k_{d}\) を使用します。つまり、その場合 \(\text{v}\) は無関係です。すべてのデフォルトの動作と推奨される使用法はこのカテゴリーに該当します。結果のレイアウトとデフォルトの (または推奨される) ストライドと距離に関する詳細と図については、使用例を参照してください。
逆方向領域で実データを期待する実数記述子の場合 (非推奨の使用法、CPU 上の 1D 実数 DFT でのみサポートされます)、関係 \(p^{\text{bwd}}_{d}\left(k_{d}; \text{v}\right)\) はより複雑な形式になります。逆方向ドメインでは、このような記述子は、データシーケンスのエントリーの実数部と虚数部が必ずしもメモリー内に連続して格納されるわけではない (または格納されない) ため、実数データを期待します。\(p^{\text{bwd}}_{d}\left(k_{d}; \text{v}\right)\) の具体的な形式は、config_param::PACKED_FORMAT に設定された値によって異なります。逆方向ドメインで実データを期待する実記述子の場合、その構成パラメーターには 3 つの異なる値 (以下に説明) が可能です: config_value::CCS_FORMAT、config_value::PACK_FORMAT、および config_value::PERM_FORMAT です。CPU 上の 1D 変換のサポートが限られているため、このセクションの残りでは説明を簡略化するために \(d = 1\) が使用されます。バッチ処理を行わない場合、つまり \(M = 1\) の場合の図も示されているため、このセクションの表では、その場合の余分なバッチ・インデックス \(m = 0\) も省略されています。
config_param::PACKED_FORMAT に config_value::CCS_FORMAT 値を設定#
設定値 config_value::CCS_FORMAT が使用されている場合、
\(p^{\text{bwd}}_{1}\left(k_{1}; \text{r}\right) = 2k_{1}\);
\(p^{\text{bwd}}_{1}\left(k_{1}; \text{i}\right) = 2k_{1} + 1\)。
逆方向領域のすべての非冗長エントリーが \(0\leq k_{1} \leq \lfloor \frac{n_{1}}{2}\rfloor\) によってキャプチャーされることを考慮すると、この場合、\(n_{1}\) が偶数 (または奇数) の場合、\(p^{\text{bwd}}_{1}\left(k_{1}; \text{v}\right)\) の関連値の範囲は \(0 \leq p^{\text{bwd}}_{1}\left(k_{1}; \text{v}\right) < n_{1} + 2\) (または \(0 \leq p^{\text{bwd}}_{1}\left(k_{1}; \text{v}\right) < n_{1} + 1\)) となります。
この形式は、\(M = 1\)、\(s^{\text{bwd}}_{0} = 0\)、および \(s^{\text{bwd}}_{1} = 1\) の場合を以下の表に示します。
\(J^{\text{bwd}}\left(k_{1}; \text{v}\right)\) |
\(0\) |
\(1\) |
\(2\) |
\(3\) |
\(\ldots\) |
\(2\lfloor n_{1}/2\rfloor\) |
\(2\lfloor n_{1}/2\rfloor + 1\) |
|---|---|---|---|---|---|---|---|
\(\left(k_{1}; \text{v}\right)\) |
\(\left(0; \text{r}\right)\) |
\(\left(0; \text{i}\right)\) |
\(\left(1; \text{r}\right)\) |
\(\left(1; \text{i}\right)\) |
\(\ldots\) |
\(\left(\lfloor n_{1}/2\rfloor; \text{r}\right)\) |
\(\left(\lfloor n_{1}/2\rfloor; \text{i}\right)\) |
保存された値 |
\(\Re\left(\left(\cdot\right)_{0}\right)\) |
\(0\) |
\(\Re\left(\left(\cdot\right)_{1}\right)\) |
\(\Im\left(\left(\cdot\right)_{1}\right)\) |
\(\ldots\) |
\(\Re\left(\left(\cdot\right)_{\lfloor n_{1}/2\rfloor}\right)\) |
\(\Im\left(\left(\cdot\right)_{\lfloor n_{1}/2\rfloor}\right)\) |
config_param::PACKED_FORMAT に config_value::PACK_FORMAT 値を設定#
設定値 config_value::PACK_FORMAT が使用されている場合、
\(p^{\text{bwd}}_{1}\left(0; \text{r}\right) = 0\);
\\(p^{\text{bwd}}_{1}\left(0; \text{i}\right)\) は存在しない (\(0\) 値の虚数部は明示的に格納されない)；
\(p^{\text{bwd}}_{1}\left(k_{1}; \text{r}\right) = 2k_{1} - 1\) の任意の \(0 < k_{1} \leq \lfloor n_{1}/2\rfloor\);
\(p^{\text{bwd}}_{1}\left(k_{1}; \text{i}\right) = 2k_{1}\) の任意の \(0 < k_{1} < \lfloor n_{1}/2 \rfloor\)。\(n_{1}\) が奇数の場合、\(k_{1} = \lfloor n_{1}/2 \rfloor\) に対して \(p^{\text{bwd}}_{1}\left(k_{1}; \text{i}\right) = 2k_{1}\) が成立します。また、\(n_{1}\) が偶数の場合、\(p^{\text{bwd}}_{1}\left(\lfloor n_{1}/2 \rfloor; \text{i}\right)\) は存在しません (\(0\) 値の虚数部は明示的に格納されません)。
逆方向領域のすべての非冗長エントリーが \(0\leq k_{1} \leq \lfloor \frac{n_{1}}{2}\rfloor\) によってキャプチャーされることを考慮すると、この場合 \(p^{\text{bwd}}_{1}\left(k_{1}; \text{v}\right)\) の関連値の範囲は \(0 \leq p^{\text{bwd}}_{1}\left(k_{1}; \text{v}\right) < n_{1}\) です (\(n_{1}\) が偶数か奇数かに関係なく)。
この形式は、\(M = 1\)、\(s^{\text{bwd}}_{0} = 0\)、および \(s^{\text{bwd}}_{1} = 1\) の場合を以下の表に示します。
\(J^{\text{bwd}}\left(k_{1}; \text{v}\right)\) |
\(0\) |
\(1\) |
\(2\) |
\(3\) |
\(\ldots\) |
\(n_{1} - 2\) |
\(n_{1} - 1\) |
|---|---|---|---|---|---|---|---|
\(\left(k_{1}; \text{v}\right)\) |
\(\left(0; \text{r}\right)\) |
\(\left(1; \text{r}\right)\) |
\(\left(1; \text{i}\right)\) |
\(\left(2; \text{r}\right)\) |
\(\ldots\) |
\(\left(n_{1}/2 - 1; \text{i}\right)\) |
\(\left(n_{1}/2; \text{r}\right)\) |
保存された値 |
\(\Re\left(\left(\cdot\right)_{0}\right)\) |
\(\Re\left(\left(\cdot\right)_{1}\right)\) |
\(\Im\left(\left(\cdot\right)_{1}\right)\) |
\(\Re\left(\left(\cdot\right)_{2}\right)\) |
\(\ldots\) |
\(\Im\left(\left(\cdot\right)_{n_{1}/2 - 1}\right)\) |
\(\Re\left(\left(\cdot\right)_{n_{1}/2}\right)\) |
\(J^{\text{bwd}}\left(k_{1}; \text{v}\right)\) |
\(0\) |
\(1\) |
\(2\) |
\(3\) |
\(\ldots\) |
\(n_{1} - 2\) |
\(n_{1} - 1\) |
|---|---|---|---|---|---|---|---|
\(\left(k_{1}; \text{v}\right)\) |
\(\left(0; \text{r}\right)\) |
\(\left(1; \text{r}\right)\) |
\(\left(1; \text{i}\right)\) |
\(\left(2; \text{r}\right)\) |
\(\ldots\) |
\(\left(\lfloor n_{1}/2\rfloor; \text{r}\right)\) |
\(\left(\lfloor n_{1}/2\rfloor; \text{i}\right)\) |
保存された値 |
\(\Re\left(\left(\cdot\right)_{0}\right)\) |
\(\Re\left(\left(\cdot\right)_{1}\right)\) |
\(\Im\left(\left(\cdot\right)_{1}\right)\) |
\(\Re\left(\left(\cdot\right)_{2}\right)\) |
\(\ldots\) |
\(\Re\left(\left(\cdot\right)_{\lfloor n_{1}/2\rfloor}\right)\) |
\(\Im\left(\left(\cdot\right)_{\lfloor n_{1}/2\rfloor}\right)\) |
config_param::PACKED_FORMAT に config_value::PERM_FORMAT 値を設定#
設定値 config_value::PERM_FORMAT が使用される場合、\(p^{\text{bwd}}_{1}\left(k_{1}; \text{v}\right)\) 関係は、\(n_{1}\) が偶数か奇数かによって異なります。
\(n_{1}\) が偶数の場合、
\(p^{\text{bwd}}_{1}\left(0; \text{r}\right) = 0\) および \(p^{\text{bwd}}_{1}\left(0; \text{i}\right)\) は存在しません (\(0\) 値の虚数部は明示的に格納されません)。
\(p^{\text{bwd}}_{1}\left(n_{1}/2; \text{r}\right) = 1\) および \(p^{\text{bwd}}_{1}\left(n_{1}/2; \text{i}\right)\) は存在しません (\(0\) 値の虚数部は明示的に格納されません)。
\(p^{\text{bwd}}_{1}\left(k_{1}; \text{r}\right) = 2k_{1}\) の任意の \(0 < k_{1} < n_{1}/2\);
\(p^{\text{bwd}}_{1}\left(k_{1}; \text{i}\right) = 2k_{1} + 1\) の任意の \(0 < k_{1} < n_{1}/2\)。
\(n_{1}\) が奇数の場合 (この形式は \(n_{1}\) が奇数の場合の config_value::PACK_FORMAT と同等です)
\(p^{\text{bwd}}_{1}\left(0; \text{r}\right) = 0\) および \(p^{\text{bwd}}_{1}\left(0; \text{i}\right)\) は存在しません (\(0\) 値の虚数部は明示的に格納されません)。
\(p^{\text{bwd}}_{1}\left(k_{1}; \text{r}\right) = 2k_{1} - 1\) の任意の \(0 < k_{1} \leq \lfloor n_{1}/2\rfloor\);
\(p^{\text{bwd}}_{1}\left(k_{1}; \text{i}\right) = 2k_{1}\) の任意の \(0 < k_{1} \leq \lfloor n_{1}/2\rfloor\)。
逆方向領域のすべての非冗長エントリーが \(0\leq k_{1} \leq \lfloor \frac{n_{1}}{2}\rfloor\) によってキャプチャーされることを考慮すると、この場合 \(p^{\text{bwd}}_{1}\left(k_{1}; \text{v}\right)\) の関連値の範囲は \(0 \leq p^{\text{bwd}}_{1}\left(k_{1}; \text{v}\right) < n_{1}\) です (\(n_{1}\) が偶数か奇数かに関係なく)。
この形式は、\(M = 1\)、\(s^{\text{bwd}}_{0} = 0\)、および \(s^{\text{bwd}}_{1} = 1\) の場合を以下の表に示します。
\(J^{\text{bwd}}\left(k_{1}; \text{v}\right)\) |
\(0\) |
\(1\) |
\(2\) |
\(3\) |
\(\ldots\) |
\(n_{1} - 2\) |
\(n_{1} - 1\) |
|---|---|---|---|---|---|---|---|
\(\left(k_{1}; \text{v}\right)\) |
\(\left(0; \text{r}\right)\) |
\(\left(n_{1}/2; \text{r}\right)\) |
\(\left(1; \text{r}\right)\) |
\(\left(1; \text{i}\right)\) |
\(\ldots\) |
\(\left(n_{1}/2 - 1; \text{r}\right)\) |
\(\left(n_{1}/2 - 1; \text{i}\right)\) |
保存された値 |
\(\Re\left(\left(\cdot\right)_{0}\right)\) |
\(\Re\left(\left(\cdot\right)_{n_{1}/2}\right)\) |
\(\Re\left(\left(\cdot\right)_{1}\right)\) |
\(\Im\left(\left(\cdot\right)_{1}\right)\) |
\(\ldots\) |
\(\Re\left(\left(\cdot\right)_{n_{1}/2 - 1}\right)\) |
\(\Im\left(\left(\cdot\right)_{n_{1}/2 - 1}\right)\) |
\(J^{\text{bwd}}\left(k_{1}; \text{v}\right)\) |
\(0\) |
\(1\) |
\(2\) |
\(3\) |
\(\ldots\) |
\(n_{1} - 2\) |
\(n_{1} - 1\) |
|---|---|---|---|---|---|---|---|
\(\left(k_{1}; \text{v}\right)\) |
\(\left(0; \text{r}\right)\) |
\(\left(1; \text{r}\right)\) |
\(\left(1; \text{i}\right)\) |
\(\left(2; \text{r}\right)\) |
\(\ldots\) |
\(\left(\lfloor n_{1}/2\rfloor; \text{r}\right)\) |
\(\left(\lfloor n_{1}/2\rfloor; \text{i}\right)\) |
保存された値 |
\(\Re\left(\left(\cdot\right)_{0}\right)\) |
\(\Re\left(\left(\cdot\right)_{1}\right)\) |
\(\Im\left(\left(\cdot\right)_{1}\right)\) |
\(\Re\left(\left(\cdot\right)_{2}\right)\) |
\(\ldots\) |
\(\Re\left(\left(\cdot\right)_{\lfloor n_{1}/2\rfloor}\right)\) |
\(\Im\left(\left(\cdot\right)_{\lfloor n_{1}/2\rfloor}\right)\) |
注
逆方向ドメインで複素データを期待する実記述子は、config_param::PACKED_FORMAT に対して config_value::CCE_FORMAT 以外の値をサポートしません。したがって、config_value::CCE_FORMAT は config_param::PACKED_FORMAT のデフォルト値として使用されます (config_value::COMPLEX_COMPLEX が config_param::CONJUGATE_EVEN_STORAGE のデフォルト値セットであることと一致します)。
config_param::PACKED_FORMAT に設定される値は、逆方向ドメインで実データを期待する実際の記述子の場合は明示的に (config_value::CCS_FORMAT、config_value::PACK_FORMAT、または config_value::PERM_FORMAT のいずれかに) 設定する必要があります。これは、その記述子の動作をさらに指定するためです (上記の説明を参照)。逆方向領域で実数データを期待する実数記述子は、CPU 上の 1D 実数 DFT でのみサポートされます。それらのサポートは非推奨です。
データレイアウトの要件#
一般的に、距離とストライドは次のように設定する必要があります
\(J^{\text{xwd}}\left(k_{1}, \ldots, k_{d}, m; \text{v}\right)\) の値は、関連する範囲内のすべての \((d+1)\) タプル \((k_{1}, k_{2}, \dots, p^{\text{xwd}}_{d}\left(k_{d};\text{v}\right), m)\) に対して非負です。
\(J^{\text{xwd}}\left(k_{1}, \ldots, k_{d}, m; \text{v}\right)\) のすべての値は、検討中のデータシーケンスに関連する一意の値に対応します。言い換えれば、2 つの異なる \((d+1)\) タプル \((k_{1}, k_{2}, \dots, p^{\text{xwd}}_{d}\left(k_{d};\text{v}\right), m)\) に対応する \(J^{\text{xwd}}\left(k_{1}, \ldots, k_{d}, m; \text{v}\right)\) の値が、両方とも関連範囲内にあるということはあり得ません。
さらに、インプレース変換 (config_param::PLACEMENT に設定された構成値 config_value::INPLACE) の場合、次の一貫性の要件が適用されます。
どちらのドメインでも同じデータタイプを期待する記述子 (複素数記述子など) は、順方向ドメインと逆方向ドメインで同じオフセット、ストライド、および距離の値を使用する必要があります;
逆方向領域で複素数データを期待する実数記述子の場合 (実数記述子のデフォルトの動作)、最後の次元に沿った先頭エントリーのメモリーアドレスは、順方向領域と逆方向領域で同一である必要があります。具体的には、その要件は条件 \(s^{\text{fwd}}_{j} = 2 s^{\text{bwd}}_{j}, \ \forall j \in \lbrace 0, \ldots, d - 1\rbrace\) に変換され、\(M > 1\) の場合は \(l^{\text{fwd}} = 2 l^{\text{bwd}}\) にも変換されます。この要件により、順方向ドメインと逆方向ドメインの次元 \(d\) に沿って単位ストライドが使用される場合、順方向ドメインでデータパディングが使用されることに注意してください (デフォルトで設定されている推奨される使用法)。
注
すべての \(J^{\text{xwd}}\left(k_{1}, \ldots, k_{d}, m; \text{v}\right)\) が非負であることを保証する十分に大きい (正の) オフセット・インデックスを持つ負のストライドはまだサポートされていません (未実装);
逆方向領域の実数データを期待し、
config_param::PACKED_FORMATの構成値config_value::CCS_FORMATを使用する 1 次元実数記述子では、\(l^{\text{bwd}} \geq n_{1} + 2\) も必要です。
入力データと出力データのストライドの設定 [非推奨]#
ドメインごとにストライドを指定する代わりに、入力および出力データシーケンスのストライドを指定することもできます。\(\text{x} = \text{i}\) (または \(\text{x} = \text{o}\)) の場合、\(s^{\text{x}}_{j}, \ j \in \lbrace 0, 1, \ldots, d\rbrace\) を入力 (または出力) データシーケンスのストライド値とします。このような \(s^{\text{xwd}}_0, s^{\text{xwd}}_1, \dots, s^{\text{xwd}}_d\) 値は、サイズ \((d+1)\) の std::vector<std::int64_t> オブジェクトの要素として (この順序で) 通信され、関連する構成設定メンバー関数を使用して、\(\text{x} = \text{f}\) の場合は config_param::INPUT_STRIDES、 (\(\text{x} = \text{b}\) の場合は config_param::OUTPUT_STRIDES) の構成値として渡されます。
\(s^{\text{i}}_{j}\) と \(s^{\text{o}}_{j}\) の値は、\(s^{\text{fwd}}_{j} = s^{\text{bwd}}_{j} = 0, \forall j \in \lbrace 0, 1, \ldots, d\rbrace\) の場合にのみ、oneMKL によって使用されます。config_param::INPUT_STRIDES と config_param::OUTPUT_STRIDES が設定されていて、config_param::FWD_STRIDES と config_param::BWD_STRIDES が設定されていない場合、これは自動的に行われます (以下の注記を参照)。このような場合、descriptor オブジェクトは、2 つの計算方向に対応するデータレイアウトを個別に考慮する必要があります。前述のように、関連するデータシーケンスのエントリーは、計算関数に提供されるデータコンテナー (sycl::buffer オブジェクトまたはデバイスアクセス可能な USM 割り当て) の要素としてアクセスされます。その基本データタイプは、上記の表に記載されているとおりです (暗黙的に再解釈される可能性があります)。入力ストライドと出力ストライドを使用する場合、順方向ドメインでデータシーケンスのエントリー \(\left(\cdot\right)^{m}_{k_{1}, k_{2}, \ldots, k_{d}}\) またはその一部にアクセスする際に使用するインデックスは以下です
ここで、順方向 (または逆方向) DFT の場合 \(\text{x} = \text{i}\) (または \(\text{x} = \text{o}\))。同様に、逆方向ドメインでデータシーケンスのエントリー (またはその一部) にアクセスする際に使用されるインデックスは、
ここで、順方向 (または逆方向) DFT の場合 \(\text{x} = \text{o}\) (または \(\text{x} = \text{i}\))。
その結果、これらの非推奨の構成パラメーターを使用して descriptor オブジェクトを構成すると、順方向ドメインと逆方向ドメインで異なるストライド値が使用されると、構成の方向が依存するようになります。意図された計算方向は descriptor オブジェクトをコミットするときには不明であるため、順方向ドメインと逆方向ドメインで正当なデータレイアウトとなるすべての方向は、正常にコミットされた descriptor オブジェクトによってサポートされる必要があります。
注
これらの非推奨の構成パラメーターによって構成されたストライドを持つ descriptor オブジェクトの場合、データレイアウト要件は 2 つの計算方向のうちの 1 つ、つまり順方向または逆方向の DFT のどちらかに対してのみ満たされる可能性があります。データレイアウト要件と一致しない計算方向にそのオブジェクトを使用する場合、oneMKL の動作は未定義になります。
config_param::INPUT_STRIDES または config_param::OUTPUT_STRIDES のいずれかを設定すると、config_param::FWD_STRIDES および config_param::BWD_STRIDES の任意の (デフォルトまたは以前に設定された) 値が \(0\) 値のベクトルにリセットされ、その逆も同様です。この暗黙的な動作により、config_param::INPUT_STRIDES または config_param::OUTPUT_STRIDESconfig_param::OUTPUT_STRIDES のいずれかと config_param::FWD_STRIDES または config_param::BWD_STRIDES のいずれかが混在することが防止されますが、これは oneMKL ではサポートされていません。このような構成を試みると、暗黙的にリセットされたストライド値によって非自明な DFT のデータレイアウト要件が無効になるため、コミット時に無効な構成の例外がスローされます。
これらの非推奨の構成パラメーターを使用してデータレイアウトのストライドを指定し、ストライドが順方向ドメインと逆方向ドメインで異なる場合は、以下に示すように、逆方向で DFT を計算するため記述子を再構成して再コミットする必要があります。
// ...
desc.set_value(config_param::INPUT_STRIDES, fwd_domain_strides);
desc.set_value(config_param::OUTPUT_STRIDES, bwd_domain_strides);
desc.commit(queue);
compute_forward(desc, ...);
// ...
desc.set_value(config_param::INPUT_STRIDES, bwd_domain_strides);
desc.set_value(config_param::OUTPUT_STRIDES, fwd_domain_strides);
desc.commit(queue);
compute_backward(desc, ...);config_param::INPUT_STRIDES および config_param::OUTPUT_STRIDES パラメーターは、oneMKL 2024.1 以降では非推奨となっています。これらの構成パラメーターを使用するすべてのアプリケーションに対して、config_param::FWD_STRIDES および config_param::BWD_STRIDES の使用方法を更新するようユーザーにアドバイスするコンパイル時の非推奨警告が示されます。
GPU デバイスでサポートされているレイアウト#
GPU デバイスでは、oneMKL は以下が必要です
変換の階数 \(d\) が \(3\) 以下であること;
オフセット値 \(s^{\text{fwd}}_{0}\) と \(s^{\text{bwd}}_{0}\) を \(0\) であること;
バッチ処理された 2 次元実変換の場合、\(l^{\text{xwd}} = n_{1} s^{\text{xwd}}_{1}\) または \(s^{\text{xwd}}_{1} = M l^{\text{xwd}}\) のいずれか (\(\text{x} = \text{f}\) と \(\text{x} = \text{b}\) の場合) であること;
3次元実変換の場合 (\(\text{x} = \text{f}\) および \(\text{x} = \text{b}\))、\(s^{\text{xwd}}_{1} = n_{2} s^{\text{xwd}}_{2}\) (\(M > 1\)) の場合は \(l^{\text{xwd}} = n_{1} s^{\text{xwd}}_{1}\) もしくは、\(s^{\text{xwd}}_{2} = n_{1} s^{\text{xwd}}_{1}\) (\(s^{\text{xwd}}_{1} = M l^{\text{xwd}}\) とともに) も成り立つこと;
実際の記述子では、
config_param::CONJUGATE_EVEN_STORAGEの場合はconfig_value::COMPLEX_COMPLEXを使用し、config_param:PACKED_FORMATの場合はconfig_value::CCE_FORMATを使用します。
注
出力結果に使用されるブロック割り当ての一部が計算された DFT に無関係であることが判明した場合、GPU コミット 記述子descriptor によってそ上書きされる可能性があります (たとえば、連続するデータ シーケンス間のパディング要素など)。