Inheritance diagram for q3::gpu::HIPGPUFFTGrid:

Public Member Functions
	HIPGPUFFTGrid (std::shared_ptr< execution::gpu::device::GPUDeviceBackend > device_backend)

	~HIPGPUFFTGrid () override

void	execute_forward_on_buffers () const override

void	execute_inverse_raw_on_buffers () const override

void	forward (Complex *data) const override

void	forward_single (ComplexSingle *data) const override

void	forward_single_batch (ComplexSingle *data, int batch_size) const override

void	forward_transform (Complex *data) const override

void	forward_transform_batch (Complex *data, int batch_size) const override

int	g_vector_to_fft_grid_index (int miller_1, int miller_2, int miller_3) const override

int	grid_dimension_1 () const override

int	grid_dimension_2 () const override

int	grid_dimension_3 () const override

void	init (int dimension_1, int dimension_2, int dimension_3, bool use_measured_planning=false) override

Complex *	input_buffer () const override

void	inverse (Complex *data) const override

void	inverse_raw (Complex *data) const override

void	inverse_raw_single (ComplexSingle *data) const override

void	inverse_raw_single_batch (ComplexSingle *data, int batch_size) const override

void	inverse_transform (Complex *data) const override

void	inverse_transform_batch (Complex *data, int batch_size) const override

int	local_grid_points () const override

int	local_offset () const override

Complex *	output_buffer () const override

void	prepare_forward_from_output () const override

void	reduce_g_space (Complex *data, int number_of_elements) const override

int	total_grid_points () const override

Public Member Functions inherited from q3::gpu::GPUFFTGrid
	GPUFFTGrid (std::shared_ptr< execution::gpu::device::GPUDeviceBackend > device_backend)

	~GPUFFTGrid () override=default

execution::gpu::device::GPUDeviceBackend &	device_backend ()

const execution::gpu::device::GPUDeviceBackend &	device_backend () const

Public Member Functions inherited from q3::FFTGrid
virtual	~FFTGrid ()=default

virtual void	assign_sticks (const std::vector< int > &)

virtual bool	co_locates_mirrors () const

virtual int	concurrent_transform_slots () const

virtual void	execute_forward_on_slot (int) const
	Forward transform on slot `slot`, the slot-aware `execute_forward_on_buffers()`.

virtual void	execute_inverse_raw_on_slot (int) const
	Inverse transform on slot `slot`, the slot-aware `execute_inverse_raw_on_buffers()`.

virtual void	forward_batched (Complex *data, int batch_size) const

virtual void	gather_local_g_space (const Complex full_g_space, Complex local_g_space) const

virtual int	global_real_space_index (int local_index) const

virtual void	initialize_batched_plans (int, bool=false) const

virtual int	initialized_batched_plan_size () const

virtual void	inverse_raw_batched (Complex *data, int batch_size) const

virtual void	inverse_raw_on_slot (Complex *data, int) const

virtual int	local_g_space_points () const

virtual bool	owns_stick (int, int) const

virtual void	prepare_forward_from_output_on_slot (int) const
	Output-to-input copy on slot `slot`, the slot-aware `prepare_forward_from_output()`.

virtual void	scatter_local_g_space (const Complex local_g_space, Complex full_g_space) const
	The reverse: place this rank's reciprocal-space points into a full array.

virtual int	slab_dimension_1 () const

virtual Complex *	slot_input_buffer (int) const

virtual Complex *	slot_output_buffer (int) const

Additional Inherited Members
Protected Attributes inherited from q3::gpu::GPUFFTGrid
std::shared_ptr< execution::gpu::device::GPUDeviceBackend >	device_backend_

Detailed Description

Definition at line 42 of file hip.cpp.

Constructor & Destructor Documentation

◆ HIPGPUFFTGrid()

q3::gpu::HIPGPUFFTGrid::HIPGPUFFTGrid ( std::shared_ptr< execution::gpu::device::GPUDeviceBackend > device_backend )

inlineexplicit

Construct a HIP FFT grid backed by the given device.

hipFFT plans are not created here — they are created in init() after the grid dimensions are known.

Definition at line 50 of file hip.cpp.

◆ ~HIPGPUFFTGrid()

q3::gpu::HIPGPUFFTGrid::~HIPGPUFFTGrid ( )

inlineoverride

Definition at line 53 of file hip.cpp.

Member Function Documentation

◆ execute_forward_on_buffers()

void q3::gpu::HIPGPUFFTGrid::execute_forward_on_buffers ( ) const

inlineoverridevirtual

Execute a forward FFT on the internal buffers (no copy).

Caller must have populated input_buffer() before calling. Result is available in output_buffer() after this call.

Implements q3::FFTGrid.

Definition at line 159 of file hip.cpp.

References q3::SerialFFTGrid::execute_forward_on_buffers().

◆ execute_inverse_raw_on_buffers()

void q3::gpu::HIPGPUFFTGrid::execute_inverse_raw_on_buffers ( ) const

inlineoverridevirtual

Execute a raw inverse FFT on the internal buffers (no copy, no normalization).

Caller must have populated input_buffer() before calling. Result is available in output_buffer() after this call.

Implements q3::FFTGrid.

Definition at line 156 of file hip.cpp.

References q3::SerialFFTGrid::execute_inverse_raw_on_buffers().

◆ forward()

void q3::gpu::HIPGPUFFTGrid::forward ( Complex * data ) const

inlineoverridevirtual

Forward FFT: real-space -> reciprocal-space (unnormalized).

Implements q3::FFTGrid.

Definition at line 130 of file hip.cpp.

References q3::SerialFFTGrid::forward().

◆ forward_single()

void q3::gpu::HIPGPUFFTGrid::forward_single ( ComplexSingle * data ) const

inlineoverridevirtual

Forward FP32 FFT via hipfftExecC2C with HIPFFT_FORWARD.

Uses a persistent device buffer allocated in init() to avoid per-call hipMalloc/hipFree overhead.

Implements q3::gpu::GPUFFTGrid.

Definition at line 174 of file hip.cpp.

◆ forward_single_batch()

void q3::gpu::HIPGPUFFTGrid::forward_single_batch	(	ComplexSingle *	data,
		int	batch_size
	)		const

inlineoverridevirtual

── Batched FP32 FFT (hipfftPlanMany) ─────────────────────────────────

Transforms multiple contiguous FFT grids in a single kernel launch. The batch plan and device buffer are created lazily on first use and resized if the batch size grows. Batched forward FP32 FFT via hipfftPlanMany + hipfftExecC2C.

All batch_size grids are transformed in a single hipFFT kernel launch, eliminating per-band dispatch overhead.

Reimplemented from q3::gpu::GPUFFTGrid.

Definition at line 256 of file hip.cpp.

◆ forward_transform()

void q3::gpu::HIPGPUFFTGrid::forward_transform ( Complex * data ) const

inlineoverridevirtual

Forward FP64 FFT via hipfftExecZ2Z with HIPFFT_FORWARD.

Uses a persistent device buffer allocated in init() to avoid per-call hipMalloc/hipFree overhead.

Reimplemented from q3::gpu::GPUFFTGrid.

Definition at line 212 of file hip.cpp.

◆ forward_transform_batch()

void q3::gpu::HIPGPUFFTGrid::forward_transform_batch	(	Complex *	data,
		int	batch_size
	)		const

inlineoverridevirtual

Batched forward FP64 FFT via hipfftPlanMany + hipfftExecZ2Z.

All batch_size grids are transformed in a single hipFFT kernel launch.

Reimplemented from q3::gpu::GPUFFTGrid.

Definition at line 301 of file hip.cpp.

◆ g_vector_to_fft_grid_index()

int q3::gpu::HIPGPUFFTGrid::g_vector_to_fft_grid_index	(	int	miller_1,
		int	miller_2,
		int	miller_3
	)		const

inlineoverridevirtual

Map a G-vector (integer Miller indices) to FFT array index. Handles negative indices via wrapping.

Implements q3::FFTGrid.

Definition at line 141 of file hip.cpp.

References q3::SerialFFTGrid::g_vector_to_fft_grid_index().

◆ grid_dimension_1()

int q3::gpu::HIPGPUFFTGrid::grid_dimension_1 ( ) const

inlineoverridevirtual

Grid size along the first dimension.

Implements q3::FFTGrid.

Definition at line 111 of file hip.cpp.

References q3::SerialFFTGrid::grid_dimension_1().

◆ grid_dimension_2()

int q3::gpu::HIPGPUFFTGrid::grid_dimension_2 ( ) const

inlineoverridevirtual

Grid size along the second dimension.

Implements q3::FFTGrid.

Definition at line 114 of file hip.cpp.

References q3::SerialFFTGrid::grid_dimension_2().

◆ grid_dimension_3()

int q3::gpu::HIPGPUFFTGrid::grid_dimension_3 ( ) const

inlineoverridevirtual

Grid size along the third dimension.

Implements q3::FFTGrid.

Definition at line 117 of file hip.cpp.

References q3::SerialFFTGrid::grid_dimension_3().

◆ init()

void q3::gpu::HIPGPUFFTGrid::init	(	int	dimension_1,
		int	dimension_2,
		int	dimension_3,
		bool	use_measured_planning = `false`
	)

inlineoverridevirtual

Initialize with grid dimensions (dimension_1 x dimension_2 x dimension_3).

Implements q3::FFTGrid.

Definition at line 82 of file hip.cpp.

References q3::SerialFFTGrid::init(), and q3::SerialFFTGrid::total_grid_points().

◆ input_buffer()

Complex * q3::gpu::HIPGPUFFTGrid::input_buffer ( ) const

inlineoverridevirtual

Access the FFT grid's internal input buffer for zero-copy transforms.

Callers scatter G-space data directly into this buffer, then call execute_inverse_raw_on_buffers() to transform in-place. This eliminates the memcpy overhead of inverse_raw() / forward().

The buffer has at least local_grid_points() elements of usable space.

Implements q3::FFTGrid.

Definition at line 149 of file hip.cpp.

References q3::SerialFFTGrid::input_buffer().

◆ inverse()

void q3::gpu::HIPGPUFFTGrid::inverse ( Complex * data ) const

inlineoverridevirtual

Inverse FFT: reciprocal-space -> real-space (includes 1/N normalization).

Implements q3::FFTGrid.

Definition at line 133 of file hip.cpp.

References q3::SerialFFTGrid::inverse().

◆ inverse_raw()

void q3::gpu::HIPGPUFFTGrid::inverse_raw ( Complex * data ) const

inlineoverridevirtual

Inverse FFT without 1/N normalization (raw FFTW backward).

Implements q3::FFTGrid.

Definition at line 136 of file hip.cpp.

References q3::SerialFFTGrid::inverse_raw().

◆ inverse_raw_single()

void q3::gpu::HIPGPUFFTGrid::inverse_raw_single ( ComplexSingle * data ) const

inlineoverridevirtual

Inverse FP32 FFT (raw, unnormalized) via hipfftExecC2C with HIPFFT_BACKWARD.

Uses a persistent device buffer allocated in init() to avoid per-call hipMalloc/hipFree overhead.

Implements q3::gpu::GPUFFTGrid.

Definition at line 192 of file hip.cpp.

◆ inverse_raw_single_batch()

void q3::gpu::HIPGPUFFTGrid::inverse_raw_single_batch	(	ComplexSingle *	data,
		int	batch_size
	)		const

inlineoverridevirtual

Batched inverse FP32 FFT (raw, unnormalized) via hipfftPlanMany + hipfftExecC2C.

All batch_size grids are transformed in a single hipFFT kernel launch, eliminating per-band dispatch overhead.

Reimplemented from q3::gpu::GPUFFTGrid.

Definition at line 278 of file hip.cpp.

◆ inverse_transform()

void q3::gpu::HIPGPUFFTGrid::inverse_transform ( Complex * data ) const

inlineoverridevirtual

Inverse FP64 FFT (raw, unnormalized) via hipfftExecZ2Z with HIPFFT_BACKWARD.

Uses a persistent device buffer allocated in init() to avoid per-call hipMalloc/hipFree overhead.

Reimplemented from q3::gpu::GPUFFTGrid.

Definition at line 230 of file hip.cpp.

◆ inverse_transform_batch()

void q3::gpu::HIPGPUFFTGrid::inverse_transform_batch	(	Complex *	data,
		int	batch_size
	)		const

inlineoverridevirtual

Batched inverse FP64 FFT (raw, unnormalized) via hipfftPlanMany + hipfftExecZ2Z.

All batch_size grids are transformed in a single hipFFT kernel launch.

Reimplemented from q3::gpu::GPUFFTGrid.

Definition at line 322 of file hip.cpp.

◆ local_grid_points()

int q3::gpu::HIPGPUFFTGrid::local_grid_points ( ) const

inlineoverridevirtual

Number of real-space grid points stored on the local rank.

Implements q3::FFTGrid.

Definition at line 123 of file hip.cpp.

References q3::SerialFFTGrid::local_grid_points().

◆ local_offset()

int q3::gpu::HIPGPUFFTGrid::local_offset ( ) const

inlineoverridevirtual

Offset into the global grid for the local segment (e.g., start of slab).

Implements q3::FFTGrid.

Definition at line 126 of file hip.cpp.

References q3::SerialFFTGrid::local_offset().

◆ output_buffer()

Complex * q3::gpu::HIPGPUFFTGrid::output_buffer ( ) const

inlineoverridevirtual

Access the FFT grid's internal output buffer for zero-copy transforms.

After execute_inverse_raw_on_buffers(), real-space data is here. After execute_forward_on_buffers(), G-space data is here.

For in-place plans (MPI FFT), this returns the same pointer as input_buffer(). For out-of-place plans (Serial FFT), this is a separate buffer.

Implements q3::FFTGrid.

Definition at line 152 of file hip.cpp.

References q3::SerialFFTGrid::output_buffer().

◆ prepare_forward_from_output()

void q3::gpu::HIPGPUFFTGrid::prepare_forward_from_output ( ) const

inlineoverridevirtual

Prepare the input buffer for a forward transform after modifying the output buffer.

For out-of-place plans (Serial FFT), this copies output_buffer → input_buffer. For in-place plans (MPI FFT), this is a no-op (same buffer).

Call this between the real-space multiply and execute_forward_on_buffers().

Implements q3::FFTGrid.

Definition at line 162 of file hip.cpp.

References q3::SerialFFTGrid::prepare_forward_from_output().

◆ reduce_g_space()

void q3::gpu::HIPGPUFFTGrid::reduce_g_space	(	Complex *	data,
		int	number_of_elements
	)		const

inlineoverridevirtual

Reduce G-space data across the pool to make it redundant on all ranks.

Serial: no-op. MPI: MPI_Allreduce(MPI_IN_PLACE, ..., MPI_SUM, pool_comm).

Implements q3::FFTGrid.

Definition at line 145 of file hip.cpp.

References q3::SerialFFTGrid::reduce_g_space().

◆ total_grid_points()

int q3::gpu::HIPGPUFFTGrid::total_grid_points ( ) const

inlineoverridevirtual

Total number of real-space grid points across all processes.

Implements q3::FFTGrid.

Definition at line 120 of file hip.cpp.

References q3::SerialFFTGrid::total_grid_points().

The documentation for this class was generated from the following file:

src/layer_3_representation/basis/planewave/internal/fft/gpu/hip.cpp

Public Member Functions

Additional Inherited Members

Detailed Description

Constructor & Destructor Documentation

◆ HIPGPUFFTGrid()

◆ ~HIPGPUFFTGrid()

Member Function Documentation

◆ execute_forward_on_buffers()

◆ execute_inverse_raw_on_buffers()

◆ forward()

◆ forward_single()

◆ forward_single_batch()

◆ forward_transform()

◆ forward_transform_batch()

◆ g_vector_to_fft_grid_index()

◆ grid_dimension_1()

◆ grid_dimension_2()

◆ grid_dimension_3()

◆ init()

◆ input_buffer()

◆ inverse()

◆ inverse_raw()

◆ inverse_raw_single()

◆ inverse_raw_single_batch()

◆ inverse_transform()

◆ inverse_transform_batch()

◆ local_grid_points()

◆ local_offset()

◆ output_buffer()

◆ prepare_forward_from_output()

◆ reduce_g_space()

◆ total_grid_points()