执行方法#

这些方法用于运行 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 的相同即可。

如果描述符已使用线程运算符构建，并且 cufftdx::is_complete_fft_execution 为 true，则此方法可用。

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

警告

不保证在不同 CUDA 架构的 GPU 上执行完全相同的 FFT 会产生位相同的 results。

块执行方法#

// #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 的相同即可。

如果描述符已使用块运算符构建，并且 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)、

将产生位相同的 results。

警告

不保证在不同 CUDA 架构的 GPU 上执行完全相同的 FFT 会产生位相同的 results。

共享内存使用#

重要的是要注意，大型 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(cudaError_t& error, cudaStream_t stream = 0)

template<class FFT>
auto cufftdx::make_workspace(cudaStream_t stream = 0)

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

cufftdx::make_workspace<FFT>(cudaStream_t) 是一个辅助函数，用于创建工作区，如果工作区分配失败，它将抛出异常 std::runtime_error，并使用描述返回的 cudaError_t 代码的字符串进行初始化。这是与 cufftdx::make_workspace<FFT>(cudaError_t&, cudaStream_t) 辅助函数的唯一区别。

如果 FFT::requires_workspace 特性为 false，则用户不必创建工作区。
可以为 FFT::requires_workspace 等于 false 的 FFT 创建工作区：这样的工作区是一个空工作区，没有全局内存分配。
工作区对象仅对其创建的 FFT 有效。
工作区对象可以分配全局内存，但永远不会超过 FFT::workspace_size，并且它负责释放它。
工作区不能并发使用，因为所有副本都共享相同的底层全局内存分配。并发使用工作区将导致内存竞争。
分配的全局内存会在创建的工作区对象的最后一个副本销毁时释放。
工作区对象可以隐式转换为 FFT::workspace_type。

注意

以下大小的 FFT 不需要工作区

2 的幂，最大为 32768。
3 的幂，最大为 19683。
5 的幂，最大为 15625。
6 的幂，最大为 7776。
7 的幂，最大为 16807。
10 的幂，最大为 10000。
11 的幂，最大为 1331。
12 的幂，最大为 1728。
13 的幂，最大为 2187。
14 的幂，最大为 2744。
15 的幂，最大为 3375。
17 的幂，最大为 4913。
18 的幂，最大为 5832。
19 的幂，最大为 6859。
20 的幂，最大为 8000。
21 的幂，最大为 9261。
22 的幂，最大为 10649。
23 的幂，最大为 12167。
24 的幂，最大为 13824。
26 的幂，最大为 17576。
29 的幂，最大为 21952。
30 的幂，最大为 27000。
32 的幂，最大为 29781。
小于 512 的 4 的因子。

在 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 运算符的使用，特别是其 real_mode 枚举值。默认情况下，实数到复数的执行将实数值作为参数，但如果执行模式设置为 real_mode::folded，则库执行优化的执行，将实数输入视为大小减半的复数输入。在这种情况下，值格式是复数（与复数到复数或复数到实数的值格式相同），包含两个实数值。

类似地，复数到复数和实数到复数 FFT 的输出值格式是复数类型，但对于复数到实数 FFT，此属性取决于 RealFFTOptions 运算符的使用，特别是其 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{complex half precision value:}\;\; & (real_0, real_1, imaginary_0, imaginary_1) \\ \text{real half precision value:}\;\; & (real_0, real_1) \\\end{split}\]

数据布局#

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

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

以下各节提供了线程和块执行模式的输入和输出数据布局的详细描述。

复数到复数#

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

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

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

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

参与 FFT 的第 n 个线程（从 0 开始索引）应在其 input 值中包含 FFT 的以下元素：n + FFT::stride * i，其中 i 是 input 中的索引。结果稍后按照相同的规则存储在 input 中。请注意，在某些情况下，特别是当 FFT::requires_workspace 为 true 时，值不会在线程之间平均分配。

另请参阅 FFT::stride。

示例

stride 等于 2 的 8 点 FFT 的第 0 个线程应在其 input 中具有值 0、2、4 和 6。

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

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

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

实数到复数和复数到实数#

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

复数布局：complex_layout::natural (默认)、complex_layout::packed、complex_layout::full。
实数模式：real_mode::normal (默认)、real_mode::folded。

复数元素布局#

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

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 点 FFT，不可用。
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。

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

遵循块执行的共享内存中的复数输入/输出中描述的相同规则，但输入/输出数据存储在线程局部数组中。

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

示例

对于 16 点 FP32/FP64 C2R FFT，ElementsPerThread 设置或默认为 4（意味着 4 个线程，并且 FFT::stride 等于 4）

complex_layout::natural：输入布局（第一行）和数据（整个表）到线程局部数组的所需分区如下所示

线程 / 元素

(r₀, i₀ = 0)

(r₁, i₁)

(r₂, i₂)

(r₃, i₃)

(r₄, i₄)

(r₅, i₅)

(r₆, i₆)

(r₇, i₇),

(r₈, i₈ = 0)

0

X

X

X

1

X

X

2

X

X

3

X

X

complex_layout::packed：输入布局（第一行）和数据（整个表）到线程数组的所需分区如下所示

线程 / 元素

(r₀, r₈)

(r₁, i₁)

(r₂, i₂)

(r₃, i₃)

(r₄, i₄)

(r₅, i₅)

(r₆, i₆)

(r₇, i₇),

0

X

X

1

X

X

2

X

X

3

X

X

complex_layout::full：输入布局（第一行）和数据（整个表）到线程局部数组的所需分区如下所示

线程 / 元素

(r₀, i₀ = 0)

(r₁, i₁)

(r₂, i₂)

(r₃, i₃)

(r₄, i₄)

(r₅, i₅)

(r₆, i₆)

(r₇, i₇),

(r₈, i₈ = 0)

(r₉, i₉)

(r₁₀, i₁₀)

(r₁₁, i₁₁)

(r₁₂, i₁₂)

(r₁₃, i₁₃)

(r₁₄, i₁₄)

(r₁₅, i₁₅),

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 的输入布局如下所示：[(real₀, imag₀ = 0), (r₁, i₁), (r₂, i₂), (r₃, i₃),(r₄, i₄ = 0)]。

complex_layout::packed 输入布局对于 shared_memory_input 看起来像这样: [(real₀, real₄), (r₁, i₁), (r₂, i₂), (r₃, i₃)]。

complex_layout::full 输入布局对于 shared_memory_input 看起来像这样: [(real₀, imag₀ = 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 执行会使可用大小限制加倍。有关更多详细信息，请查阅支持的功能。

线程执行中的实数输入/输出在寄存器中#

遵循块执行的共享内存中的实数输入/输出中描述的相同规则，但输入/输出数据存储在线程局部数组中。

块执行中的实数输入/输出在寄存器中#

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

示例

对于具有 ElementsPerThread 设置或默认值为 4 的 16 点 FP32/FP64 R2C FFT

real_mode::normal 输入布局看起来像这样：[real₀, r₁, r₂, r₃, r₄, r₅, r₆, r₇]。数据到线程局部数组的所需分区如下所示

线程 / 元素

r₀

r₁

r₂

r₃

r₄

r₅

r₆

r₇

0

X

X

1

X

X

2

X

X

3

X

X

real_mode::folded 输入布局看起来像这样：[(real₀, r₁), (r₂, r₃), (r₄, r₅), (r₆, r₇)]。数据到线程局部数组的所需分区如下所示

线程 / 元素

(r₀, r₁)

(r₂, r₃)

(r₄, r₅)

(r₆, r₇)

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 遵循相同的规则，但必须考虑隐式批处理。

加载和存储数据#

为了处理根据配置更改输入和输出长度的复杂性，该库提供了各种 traits，以便于进行内存操作。Input Length Trait 和 Output Length Trait 分别描述输入和输出数组的长度。这涵盖了上面描述的实数到复数 (R2C) 和复数到实数 (C2R) 的情况。Input EPT Trait 描述了每个线程要加载的元素计数，假设每个元素的类型为 Input Type Trait。相应地，Output EPT Trait 和 Output Type Trait 描述了存储输出的相同属性。

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

示例

寄存器中具有数据的块 FFT

以下示例显示了基于块 trait 的通用加载方案，用于 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

以下示例显示了基于块 trait 的通用加载方案，用于 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;
}

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

// 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];
}