执行方法

这些方法用于运行FFT操作。

代码示例：

#include <cufftdx.hpp>

using FFT = decltype( cufftdx::Size<128>() + cufftdx::Type<fft_type::c2c>()
                    + cufftdx::Direction<fft_direction::forward>()
                    + cufftdx::Precision<float>() + cufftdx::Block() );

using complex_type = typename FFT::value_type;

__global__ kernel(... /* arguments */) {

  // Shared memory pointer
  extern __shared__ __align__(alignof(float4)) complex_type shared_mem[];

  // Register data
  complex_type thread_data[FFT::storage_size];

  // Load data into registers (thread_data)
  // ...

  FFT().execute(thread_data, shared_mem);

  // Store results (thread_data) into global memory
}

线程执行方法

void FFT().execute<typename T>(T* input)

运行由FFT描述符定义的FFT操作。T可以是任何类型（例如float2或double2），只要其对齐方式和元素大小与FFT::value_type相同即可。

如果描述符是使用Thread Operator构建的，并且cufftdx::is_complete_fft_execution为true，则此方法可用。

input数组应存放在寄存器中。input必须能容纳FFT::storage_size个FFT::value_type类型的元素。

警告

不能保证在不同CUDA架构的GPU上执行完全相同的FFT会产生完全一致的比特结果。

块执行方法

// #1
void FFT().execute<typename T>(T* input, void* shared_memory, FFT::workspace_type& workspace)

// #2: Version of #1 for FFTs which don't require workspace
void FFT().execute<typename T>(T* input, void* shared_memory)

// #3: Execute with input data in shared memory
void FFT().execute<typename T>(T* shared_memory_input, FFT::workspace_type& workspace)

// #4: Version of #3 for FFTs which don't require workspace
void FFT().execute<typename T>(T* shared_memory_input)

运行由FFT描述符定义的FFT操作。T可以是任何类型（例如float2或double2），只要其对齐方式和元素大小与FFT::value_type相同即可。

如果描述符是通过Block Operator构建的，并且cufftdx::is_complete_fft_execution为true，则此方法可用。

当FFT::requires_workspace为false时，可以使用重载版本#2和#4。否则，用户必须使用方法#1或#3并传入工作区引用。

警告

库代码假设shared_memory和shared_memory_input都按128位对齐以实现最佳内存操作。 这可以通过使用__align__或alignas编译器指令来实现。示例simple_fft_block_shared中展示了正确的对齐方式。 有关内存对齐的更多详细信息，请参阅CUDA C++编程指南。虽然非必需，但用户也可以考虑以相同方式对齐线程本地数组，以减少内核资源使用。

示例

// Examples of how to make sure shared memory pointer is aligned to 128 bits (16 bytes):
using value_type = typename FFT::value_type;
extern __shared__ alignas(float4) unsigned char shared_mem[];                               // 1
extern __shared__ __align__(16) value_type shared_mem[];                                    // 2
extern __shared__ __align__(alignof(float4)) value_type shared_mem[];                       // 3
extern __shared__ float4 shared_mem[];                                                      // 4

// Warning: std::aligned_storage became deprecated in C++23
extern __shared__ std::aligned_storage_t<sizeof(float4), alignof(float4)> shared_mem[];     // 5

在方法#1和#2中，input位于线程本地数组中，shared_memory是指向大小为FFT::shared_memory_size字节的共享内存的指针。该操作是原地进行的，意味着结果存储在input中。input必须容纳FFT::storage_size个FFT::value_type类型的元素。

注意

方法#1和#2不假设共享内存(shared_memory)可以在没有块同步的情况下安全修改或访问， 并在首次使用前执行必要的同步(__syncthreads())。此外，方法#1和#2在完成对共享内存的最后操作后， 不会同步块内的任何线程。如果该共享内存稍后将被重用，则必须先执行同步。

在方法#3和#4中，输入数据通过共享内存(shared_memory_input)传递。该操作是原地进行的，意味着结果会被存回shared_memory_input。这些方法不需要额外传递shared_memory指针，因为shared_memory_input将用于线程间所需的通信。因此，shared_memory_input必须容纳所有输入和输出值，且不能小于FFT::shared_memory_size字节（即共享内存大小以字节计最多为FFT::shared_memory_size、FFT::ffts_per_block * <FFT_input_size_in_bytes>和FFT::ffts_per_block * <FFT_output_size_in_bytes>)字节）。

注意

方法#3和#4通过共享内存获取输入，它们假设已经执行了同步操作，可以安全地访问数据。这些方法在完成对共享内存的最后操作后，不会同步块内的任何线程。在从共享内存读取或写入之前，必须先执行同步操作。

警告

不能保证执行相同的FFT（大小、方向、类型、精度）但具有不同

• 每个线程处理的元素数量 (ElementsPerThread),

• 每个CUDA块计算的FFT数量 (FFTsPerBlock), 或

• 块维度 (BlockDim),

将产生完全一致的结果。

警告

不能保证在不同CUDA架构的GPU上执行完全相同的FFT会产生完全一致的比特结果。

共享内存使用

需要注意的是，大型FFT运算可能每个CUDA块需要超过48 KB的共享内存。因此，如CUDA编程指南 (#1, #2, #3)所述，执行此类FFT运算的内核必须使用动态共享内存，而非静态大小的共享内存数组。此外，这些内核需要通过显式调用cudaFuncSetAttribute()来设置cudaFuncAttributeMaxDynamicSharedMemorySize参数。具体实现可参考以下示例代码及入门示例。

#include <cufftdx.hpp>
using namespace cufftdx;

using FFT = decltype(Size<16384>() + Precision<float>() + Type<fft_type::c2c>()
                     + Direction<fft_direction::forward>() + SM<800>() + Block());

__global__ void block_fft_kernel(FFT::value_type* data) {
  // dynamic shared memory
  extern __shared__ __align__(alignof(float4)) FFT::value_type shared_mem[];

  (...)
}

void example() {
  (...)

  // Increases the max dynamic shared memory size to match FFT requirements
  cudaFuncSetAttribute(block_fft_kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, FFT::shared_memory_size)
  // Invokes kernel with FFT::block_dim threads in CUDA block
  block_fft_kernel<FFT><<<1, FFT::block_dim, FFT::shared_memory_size>>>(data, workspace);

  (...)
}

创建工作区函数

template<class FFT>
auto cufftdx::make_workspace<FFT>(cudaError_t& error, cudaStream_t stream = 0)

cufftdx::make_workspace<FFT>(cudaError_t&, cudaStream_t) 是一个辅助函数，用于创建块execute(...)方法所需的工作空间， 当FFT::requires_workspace为true时。FFT是FFT描述符的类型。如果未传递流参数， 则使用默认的0流进行执行。如果在调用函数后error不是cudaSuccess， 则表示工作空间未正确创建且无效。

• 如果 FFT::requires_workspace 特性为 false，用户无需创建工作区。

• 可以为FFT创建工作区，当FFT::requires_workspace等于false时：此类工作区是一个没有全局内存分配的空工作区。

• 工作区对象仅对创建它的FFT有效。

• Workspace对象可以分配全局内存，但总量不超过FFT::workspace_size，并且负责释放这些内存。

• 工作空间无法并发使用，因为所有副本共享相同的底层全局内存分配。并发使用工作空间会导致内存竞争。

• 分配的全局内存在最后一个创建的workspace对象副本销毁时释放。

• Workspace对象可以隐式转换为FFT::workspace_type。

注意

以下尺寸的FFT不需要工作区：

• 2的幂次方，最大到32768

• 3的幂次方，最高至19683

• 5的幂次方，最高至15625

• 6的幂次方，最高至1296

• 7的幂次方直到2401

• 10的幂次方，最高到10000

• 11的幂次方直到1331

• 12的幂次方直到1728

In the future versions of cuFFTDx:

• 其他配置可能不需要工作区要求。

• 不需要工作空间的FFT配置将继续保持这一特性。

警告

FFT::workspace_type 对象不跟踪底层内存的生命周期，仅在转换来源的工作区对象生命周期内有效。

警告

由cufftdx::make_workspace<FFT>(cudaError_t&, cudaStream_t)返回的类型可能因不同的FFT描述而异，并且与FFT::workspace_type不同。用户在创建工作区对象时应使用auto。示例：

// Kernel
template<class FFT>
__launch_bounds__(FFT::max_threads_per_block)
__global__ void block_fft_kernel(typename FFT::value_type* data, typename FFT::workspace_type workspace) {
    // ...

    // Execute FFT
    FFT().execute(thread_data, shared_mem, workspace);
}

// Create workspace
cudaError_t error = 0;
auto workspace = cufftdx::make_workspace<FFT>(error, stream);

// ...

// Run kernel with FFT
block_fft_kernel<FFT><<<1, FFT::block_dim, FFT::shared_memory_size, stream>>>(data, workspace);

数值格式与数据布局

要正确执行FFT运算，数据需要按照数值格式章节描述的特定格式传递给库，并按照数据布局章节详述的方式在线程间进行分区。用户需要确保这两个因素正确无误，但该库提供了惯用法、特征和示例，使得API的使用变得直观易懂。

数值格式

对于单精度和双精度的复数，复数中的第一个值是实部，第二个值是虚部。

对于复数到复数以及复数到实数的FFT运算，输入值的格式采用前面提到的复数类型；但对于实数到复数的FFT运算，该属性取决于是否使用RealFFTOptions Operator及其枚举值real_mode。默认情况下，实数到复数运算将实数作为参数输入，但如果执行模式设置为real_mode::folded，库会执行优化运算，将实数输入视为尺寸减半的复数输入。此时数值格式为复数（与复数到复数或复数到实数的情况相同），包含两个实数值。

类似地，复数到复数以及实数到复数FFT的输出值格式为复数类型，但对于复数到实数FFT，此属性取决于是否使用RealFFTOptions Operator及其特定的real_mode枚举值。默认情况下，复数到实数执行会输出实数值作为结果，但如果执行模式设置为real_mode::folded，库将执行优化处理，将实数输出视为大小减半的复数输出。在这种情况下，值格式为复数（与复数到复数或复数到实数相同），否则将由实数值组成。

半精度隐式批处理

在cuFFTDx中处理半精度(fp16)FFT时采用隐式批处理方式，即单个计算过程会同时处理两个FFT批次。cuFFTDx要求半精度复数数据按以下顺序排列：2个实部后跟2个虚部（即real₀, real₁, imaginary₀, imaginary₁）。半精度的实数值（用于R2C和C2R FFT）遵循相同逻辑，每个值应包含两个实数。另请参阅FFT::implicit_type_batching特性。

\[\begin{split}\text{复数半精度值:}\;\; & (real_0, real_1, imaginary_0, imaginary_1) \\ \text{实数半精度值:}\;\; & (real_0, real_1) \\\end{split}\]

数据布局

FFT的输入和输出数据布局严格取决于其配置和所选的变换类型。数据布局描述了execute方法的输入或输出数据的大小和模式，无论输入/输出是在线程之间分区为线程本地数组还是连续存储在共享内存中。

如数值格式部分所述，输入/输出格式可以是实数或复数。另一个重要的执行属性是序列长度。 在复数到复数变换中，该属性等于变换尺寸，但对于实数到复数和复数到实数变换，该库提供了几个选项可供选择。 这一选择可以通过RealFFTOptions Operator运算符来实现。

以下部分详细描述了thread和block两种执行模式下输入输出数据的布局。

复数到复数

对于复数到复数的变换，输入和输出数据都必须是相应精度的复数数组。输入和输出数组的长度始终相同，等于FFT的大小。

线程执行中的寄存器输入/输出

FFT的输入值应按自然顺序存储在input中。 结果将按照相同规则存储在input中。

块执行中的寄存器输入/输出

n号线程(从0开始索引)参与FFT计算时，其input值应包含以下FFT元素：n + FFT::stride * i，其中i是input中的索引。 计算结果随后会按照相同规则存入input。 请注意在某些情况下，特别是当FFT::requires_workspace为true时， 各线程间的数值分配可能不均等。

另请参阅 FFT::stride。

示例

8点FFT的第0个线程，当FFT::stride等于2时，其input中应包含值0、2、4和6。

元素并不总是能在线程间完美划分：一个7点FFT的第0个线程，其FFT::stride等于2时，其input中应包含值0、2、4和6，而第1个线程应包含值1、3和5。

块执行中的共享内存输入/输出

FFT的输入值应按自然顺序存储在shared_memory_input中。 结果将按照相同规则存储在shared_memory_input中。

实数转复数与复数转实数

对于实数到复数(R2C)和复数到实数(C2R)的FFT变换，输入和输出数据布局取决于RealFFTOptions Operator。 默认情况下，RealFFTOptions设置为complex_layout::natural和real_mode::normal。

• 复杂布局：complex_layout::natural（默认），complex_layout::packed，complex_layout::full。

• 实模式: real_mode::normal (默认), real_mode::folded.

复杂元素布局

复杂元素布局是为复数到实数输入和实数到复数输出定义的。它取决于传递的complex_layout值，该值来自RealFFTOptions运算符。

• complex_layout::natural 用于偶数长度 ((r₀, i₀ = 0), (r₁, i₁), … , (r _N/2, i _N/2 = 0))，仅包含非冗余的 N/2 + 1 个元素，其中第一个和最后一个元素仅包含实部。

  • 由于复数到实数FFT输入信号的数学特性，imag₀和imag_N/2都被假定为0。

• complex_layout::natural 对于奇数长度 ((r₀, i₀ = 0), (r₁, i₁), … , (r _⌊N/2⌋, i _⌊N/2⌋))，仅包含非冗余的 ⌊N/2⌋ + 1 个元素，其中第一个元素仅包含实部。

  • 由于复数到实数FFT输入信号的数学特性，假设imag₀为0。

• complex_layout::packed ((r₀, i₀), (r₁, i₁), … , (r _{N/2 - 1}, i _{N/2 - 1}))，将最后一个实数元素打包到第一个元素的虚部中。

  • 打包意味着第一个元素x₀包含(real₀, real_⌊N/2⌋)。

  • 仅允许偶数尺寸，例如16点FFT的长度为8，15点则不可用。

• complex_layout::full 用于偶数长度 ((r₀, i₀ = 0), (r₁, i₁), … , (r_N/2, i_N/2 = 0), … , (r_{N - 1}, i_{N - 1}))，包含所有元素（包括冗余元素），因为输出是厄米特对称的。

  • 例如，对于16点FFT，其长度为16

  • 由于复数到实数FFT输入信号的数学特性，imag₀和imag_N/2都被假定为0。

• complex_layout::full 用于奇数长度 ((r₀, i₀ = 0), (r₁, i₁ ), … , (r_{N - 1}, i_{N - 1}))，包含所有元素（包括冗余元素），因为输出是厄米特矩阵。

  • 例如，对于15点FFT，其长度为15

  • 由于复数到实数FFT输入信号的数学特性，假设imag₀为0。

线程执行中寄存器的复杂输入/输出

遵循与块执行中共享内存的复杂输入/输出相同的规则，但输入/输出数据存储在线程本地数组中。

块执行中寄存器的复杂输入/输出

n号线程(从0开始索引)参与FFT计算时，其input值应包含以下FFT元素：n + FFT::stride * i，其中i是input中的索引。计算结果随后会按照相同规则存入input。 请注意在某些情况下，特别是当FFT::requires_workspace为true时， 各线程之间的数值分配可能不均等。

示例

对于16点FP32/FP64 C2R FFT，当ElementsPerThread设置为或默认为4时 (意味着4个线程，且FFT::stride等于4)：

complex_layout::natural: 下方展示了输入布局（第一行）以及将数据（整个表格）分区为线程本地数组的要求：

线程/元素	(r0, i0 = 0)	(r1, i1)	(r2, i2)	(r3, i3)	(r4, i4)	(r5, i5)	(r6, i6)	(r7, i7),	(r8, i8 = 0)
0	X				X				X
1		X				X
2			X				X
3				X				X

complex_layout::packed: 下方展示了输入布局（首行）以及将数据（整个表格）分区到线程数组所需的方式：

线程/元素	(r0, r8)	(r1, i1)	(r2, i2)	(r3, i3)	(r4, i4)	(r5, i5)	(r6, i6)	(r7, i7),
0	X				X
1		X				X
2			X				X
3				X				X

complex_layout::full: 下方展示了输入布局（首行）以及将数据（整个表格）分区为线程本地数组的要求：

线程/元素	(r0, i0 = 0)	(r1, i1)	(r2, i2)	(r3, i3)	(r4, i4)	(r5, i5)	(r6, i6)	(r7, i7),	(r8, i8 = 0)	(r9, i9)	(r10, i10)	(r11, i11)	(r12, i12)	(r13, i13)	(r14, i14)	(r15, i15),
0	X				X				X				X
1		X				X				X				X
2			X				X				X				X
3				X				X				X				X

R2C输出的布局和所需分区看起来类似。

半精度FFT遵循相同的规则，但需要考虑隐式批处理。

块执行中共享内存的复杂输入/输出

FFT的输入值应按自然顺序存储在shared_memory_input中。 结果将按照相同规则存储在shared_memory_input中。

示例

对于8点FP32/FP64 C2R FFT：

• complex_layout::natural 输入布局对于 shared_memory_input 看起来像这样: [(实部₀, 虚部₀ = 0), (r₁, i₁), (r₂, i₂), (r₃, i₃),(r₄, i₄ = 0)].

• complex_layout::packed 输入布局对于 shared_memory_input 看起来像这样: [(实部₀, 实部₄), (实部₁, 虚部₁), (实部₂, 虚部₂), (实部₃, 虚部₃)].

• complex_layout::full 输入布局对于 shared_memory_input 看起来像这样: [(实部₀, 虚部₀ = 0), (r₁, i₁), (r₂, i₂), (r₃, i₃),(r₄, i₄ = 0), (r₅, i₅), (r₆, i₆), (r₇, i₇)].

R2C输出的布局看起来类似。

半精度FFT遵循相同的规则，但需要考虑隐式批处理。

真实元素布局

实数元素布局是为复数到实数输出和实数到复数输入定义的。它取决于RealFFTOptions运算符的real_mode参数：

• real_mode::normal (x₀, x₁, … , x_{N - 1}), 长度为变换大小的实数元素数组。

• real_mode::folded (x₀, x₁, … , x _{N/2 - 1}), 长度为变换尺寸一半的复数元素数组。

注意

共享内存中元素的物理布局与real_mode::normal情况相同，但逻辑布局发生了变化：它不再是实数数组，而是一个复数数组。

注意

real_mode::folded 执行模式依赖于特定的FFT特性，目前仅在块级执行中可用，且仅支持大小为2^N的情况（其中N为能使变换适配可用尺寸范围的任意指数）。对于线程级执行，此优化支持所有大小为2*N的情况（其中N为能使变换适配可用尺寸范围的任意乘数）。使用folded执行模式可将可用尺寸限制扩大一倍。更多详情请参阅Supported Functionality。

线程执行时寄存器中的实际输入/输出

遵循与共享内存中块执行的实时输入/输出相同的规则，但输入/输出数据存储在线程本地数组中。

块执行中寄存器的实际输入/输出

参与FFT的第n个线程（从0开始索引）应包含以下Input Type Trait类型的值：n + FFT::stride * i，其中i是input中的索引。 结果随后会按照相同规则存储在input中。 请注意，在某些情况下（特别是当FFT::requires_workspace为true时）， 这些值不会在线程间均匀分配。 同样重要的是要记住，每个线程中该类型的值数量将是Input EPT Trait而非Elements Per Thread Trait。

示例

对于16点FP32/FP64 R2C FFT，当ElementsPerThread设置为或默认为4时：

real_mode::normal 输入布局如下所示：[real₀, r₁, r₂, r₃, r₄, r₅, r₆, r₇]。 下面展示了将数据分区到线程本地数组所需的方式：

线程/元素	r0	r1	r2	r3	r4	r5	r6	r7
0	X				X
1		X				X
2			X				X
3				X				X

real_mode::folded 输入布局如下所示: [(real₀, r₁), (r₂, r₃), (r₄, r₅), (r₆, r₇)]。 下面展示了将数据分区到线程本地数组所需的方式:

线程/元素	(r0, r1)	(r2, r3)	(r4, r5)	(r6, r7)
0	X
1		X
2			X
3				X

C2R输出的布局和所需分区看起来类似。

半精度FFT遵循相同的规则，但需要考虑隐式批处理。

块执行中共享内存的实际输入/输出

FFT的输入值应按自然顺序存储在shared_memory_input中。结果也遵循相同规则存储在shared_memory_input中。

示例

对于8点FP32/FP64 R2C FFT：

• real_mode::normal 输入布局对于 shared_memory_input 看起来像这样: [real₀, r₁, r₂, r₃, r₄, r₅, r₆, r₇]

• real_mode::folded 输入布局对于 shared_memory_input 看起来像这样: [(real₀, r₁), (r₂, r₃), (r₄, r₅), (r₆, r₇)]

C2R输出的布局看起来类似。

半精度FFT遵循相同的规则，但需要考虑隐式批处理。

加载与存储数据

为了处理因配置变化导致输入输出长度复杂度的问题，该库提供了多种特性以简化内存操作。Input Length Trait和Output Length Trait分别描述了输入和输出数组的长度特性，这涵盖了上文所述的实数转复数(R2C)和复数转实数(C2R)场景。Input EPT Trait定义了每个线程需要加载的元素数量（假设每个元素类型为Input Type Trait），相应地Output EPT Trait和Output Type Trait则描述了输出存储时的相同属性。

详细的惯用IO操作展示在示例中，但从全局内存加载数据的一般方法如下：

示例

使用寄存器数据的块FFT

以下示例展示了基于块特征的通用加载方案，适用于cuFFTDx寄存器API执行模式。在此模式下，一个线程组协作执行FFT运算，因此数据需要分布在所有参与者之间。以下代码片段已考虑这种分区方式，同时也支持不同的值格式（如值格式所述）和数据布局（如实数元素布局和复数元素布局所述）。

// Which FFT in this block is this thread performing
const auto local_fft_id = threadIdx.y;
// Which FFT in this grid is this thread performing
const auto global_fft_id = FFT::ffts_per_block * blockIdx.x + local_fft_id;
// Memory offset for accessing the first element of the global_fft
const auto global_offset = global_fft_id * FFT::input_length;

// Cast registers to type required as input to FFT execution
using input_t = typename FFT::input_type;
auto thread_fft_data = reinterpret_cast<input_t*>(thread_data);
auto fft_data = reinterpret_cast<const input_t*>(input);

auto index = threadIdx.x;
for (unsigned int i = 0; i < FFT::input_ept; ++i) {
  if (index < FFT::input_length) {
      thread_fft_data[i] = fft_data[global_offset + index];
      index += FFT::stride;
  }

使用共享内存中的数据进行块FFT

以下示例展示了基于块特征的cuFFTDx共享内存API执行模式的通用加载方案。在此方案中，整个CUDA块协作执行ffts-per-block-trait次FFT运算，因此共享内存中的数据需要包含所有必要的批次。为了实现合并式内存访问，所有批次的数据加载由整个线程块共同完成。以下代码片段考虑了不同的值格式（如值格式所述）以及数据布局（如实数元素布局和复数元素布局所述）。

// The index of first FFT being performed by threads of this block
const auto block_fft_id = blockIdx.x * FFT::ffts_per_block;
// Offset in memory to the first element accessed by threads in this block
const auto block_offset = block_fft_id * FFT::input_length;
// Combined length of all FFTs performed by threads of this block
constexpr auto block_input_length = FFT::ffts_per_block * FFT::input_length;

// Cast registers to type required as input to FFT execution
using input_t = typename FFT::input_type;
auto shared_fft_data = reinterpret_cast<input_t*>(shared_memory);
auto fft_data = reinterpret_cast<const input_t*>(input);

// The entire block loads all required batches in a coalesced manner,
// threads will load elements from different batches than they will later
// execute on, and this is on purpose.
const auto stride = blockDim.x * blockDim.y;

auto index  = threadIdx.y * blockDim.x + threadIdx.x;
for (int i = 0; i < FFT::input_ept; ++i) {
    if (index < block_input_length) {
        shared_fft_data[index] = fft_data[block_offset + index];
    }
    index += stride;
}

以下示例展示了基于线程特性的cuFFTDx通用加载方案。在此方案中，单个线程执行完整的FFT运算，因此无需对数据进行分区，它将加载整个序列。该方案已考虑不同的值格式（如值格式所述）和数据布局（如实数元素布局和复数元素布局所述）。

// This example assumes a 2-dimensional block and 1-dimensional grid

// Which FFT in block is this thread performing
const auto local_fft_id = threadIdx.y * blockDim.x + threadIdx.x;
// Which FFT in grid is this thread performing
const auto global_fft_id = (blockDim.x * blockDim.y) * blockIdx.x + local_fft_id;
// Memory offset for accessing first element of this FFT
const auto global_offset = global_fft_id * FFT::input_length;

// Cast registers to type required as input to FFT execution
using input_t = typename FFT::input_type;
auto thread_fft_data = reinterpret_cast<input_t*>(thread_data);
auto fft_data = reinterpret_cast<const input_t*>(input);

for (unsigned int i = 0; i < FFT::input_length; ++i) {
  thread_fft_data[i] = fft_data[global_offset + i];
}