Loading...
3D FFT extends Fourier analysis to volumetric data, crucial for molecular dynamics, MRI reconstruction, and turbulence simulation. Memory requirements grow as O(n³), making large transforms challenging. cuFFT provides optimized 3D transforms with options for distributed multi-GPU execution.
Half memory for real volumetric data.
Distribute large transforms across GPUs.
Decompose along one axis for parallel execution.
Simple 3D FFT without optimization.
void fft3d_naive(cufftComplex* d_data, int nx, int ny, int nz) {
cufftHandle plan;
cufftPlan3d(&plan, nz, ny, nx, CUFFT_C2C);
cufftExecC2C(plan, d_data, d_data, CUFFT_FORWARD);
cufftDestroy(plan);
}Optimized 3D R2C FFT with manual workspace and multi-GPU support.
class FFT3D {
cufftHandle plan_r2c, plan_c2r;
int nx, ny, nz;
size_t work_size;
public:
void init(int nx_, int ny_, int nz_) {
nx = nx_; ny = ny_; nz = nz_;
// Real-to-complex 3D plan
int n[3] = {nz, ny, nx};
int inembed[3] = {nz, ny, nx};
int onembed[3] = {nz, ny, nx/2 + 1};
cufftCreate(&plan_r2c);
cufftSetAutoAllocation(plan_r2c, 0); // Manual memory management
cufftMakePlanMany(plan_r2c, 3, n,
inembed, 1, nx * ny * nz,
onembed, 1, (nx/2 + 1) * ny * nz,
CUFFT_R2C, 1, &work_size);
cufftCreate(&plan_c2r);
cufftSetAutoAllocation(plan_c2r, 0);
cufftMakePlanMany(plan_c2r, 3, n,
onembed, 1, (nx/2 + 1) * ny * nz,
inembed, 1, nx * ny * nz,
CUFFT_C2R, 1, &work_size);
// Allocate workspace
void* d_work;
cudaMalloc(&d_work, work_size);
cufftSetWorkArea(plan_r2c, d_work);
cufftSetWorkArea(plan_c2r, d_work);
}
void forward(float* d_real, cufftComplex* d_complex) {
cufftExecR2C(plan_r2c, d_real, d_complex);
}
void inverse(cufftComplex* d_complex, float* d_real) {
cufftExecC2R(plan_c2r, d_complex, d_real);
size_t N = (size_t)nx * ny * nz;
normalize<<<(N+255)/256, 256>>>(d_real, N, 1.0f / N);
}
size_t complex_size() { return (size_t)(nx/2 + 1) * ny * nz * sizeof(cufftComplex); }
size_t real_size() { return (size_t)nx * ny * nz * sizeof(float); }
};
// Multi-GPU 3D FFT
void fft3d_multigpu(float* d_data[], int nx, int ny, int nz, int ngpus) {
cufftHandle plan;
cufftCreate(&plan);
int n[3] = {nz, ny, nx};
size_t work_size[ngpus];
cufftXtMakePlanMany(plan, 3, n, NULL, 1, 0, CUDA_R_32F,
NULL, 1, 0, CUDA_C_32F, 1, work_size, CUDA_C_32F);
cudaLibXtDesc* desc;
cufftXtMalloc(plan, &desc, CUFFT_XT_FORMAT_INPLACE);
// Copy data to distributed format
cufftXtMemcpy(plan, desc, d_data[0], CUFFT_COPY_HOST_TO_DEVICE);
// Execute distributed FFT
cufftXtExecDescriptorR2C(plan, desc, desc);
}| Metric | Naive | Optimized | Improvement |
|---|---|---|---|
| 512³ R2C single GPU | 180ms | 95ms | 1.9x faster |
| 1024³ 4-GPU vs 1-GPU | 8.2s (1 GPU) | 2.4s (4 GPU) | 3.4x faster |
| Memory 512³ R2C vs C2C | 1GB (C2C) | 537MB (R2C) | 1.9x less |
Depends on GPU memory. 512³ C2C = 1GB, 1024³ = 8GB. For larger: use R2C to halve memory, use multi-GPU, or out-of-core methods. NVIDIA A100 80GB can do ~2048³ C2C.
cuFFT uses slab decomposition: volume is split along one axis (typically Z). Each GPU handles nz/ngpus slabs. All-to-all transpose needed between 1D FFT stages. cuFFTXt handles this automatically.
Ready to optimize your CUDA code? Download RightNow AI and get real-time performance analysis for your kernels.