#include <metal.hpp>

Inheritance diagram for q3::gpu::MetalGPUFFTGrid:

Public Member Functions
	MetalGPUFFTGrid (std::shared_ptr< execution::gpu::device::GPUDeviceBackend > device_backend=nullptr)

	~MetalGPUFFTGrid () 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

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

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

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

virtual void	forward_transform (Complex *data) const

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

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

virtual void	inverse_transform (Complex *data) const

virtual void	inverse_transform_batch (Complex *data, int batch_size) 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 26 of file metal.hpp.

Constructor & Destructor Documentation

◆ MetalGPUFFTGrid()

q3::gpu::MetalGPUFFTGrid::MetalGPUFFTGrid ( std::shared_ptr< execution::gpu::device::GPUDeviceBackend > device_backend = nullptr )

explicit

◆ ~MetalGPUFFTGrid()

q3::gpu::MetalGPUFFTGrid::~MetalGPUFFTGrid ( )

override

Member Function Documentation

◆ execute_forward_on_buffers()

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

overridevirtual

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.

◆ execute_inverse_raw_on_buffers()

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

overridevirtual

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.

◆ forward()

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

overridevirtual

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

Implements q3::FFTGrid.

◆ forward_single()

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

overridevirtual

Forward single-precision FFT: real-space → reciprocal-space (unnormalized).

Parameters

data	In-out: single-precision complex grid data.

Implements q3::gpu::GPUFFTGrid.

◆ g_vector_to_fft_grid_index()

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

overridevirtual

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

Implements q3::FFTGrid.

◆ grid_dimension_1()

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

overridevirtual

Grid size along the first dimension.

Implements q3::FFTGrid.

◆ grid_dimension_2()

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

overridevirtual

Grid size along the second dimension.

Implements q3::FFTGrid.

◆ grid_dimension_3()

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

overridevirtual

Grid size along the third dimension.

Implements q3::FFTGrid.

◆ init()

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

overridevirtual

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

Implements q3::FFTGrid.

◆ input_buffer()

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

overridevirtual

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.

◆ inverse()

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

overridevirtual

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

Implements q3::FFTGrid.

◆ inverse_raw()

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

overridevirtual

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

Implements q3::FFTGrid.

◆ inverse_raw_single()

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

overridevirtual

Inverse single-precision FFT without normalization (raw backward).

Parameters

data	In-out: single-precision complex grid data.

Implements q3::gpu::GPUFFTGrid.

◆ local_grid_points()

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

overridevirtual

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

Implements q3::FFTGrid.

◆ local_offset()

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

overridevirtual

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

Implements q3::FFTGrid.

◆ output_buffer()

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

overridevirtual

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.

◆ prepare_forward_from_output()

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

overridevirtual

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.

◆ reduce_g_space()

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

overridevirtual

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.

◆ total_grid_points()

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

overridevirtual

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

Implements q3::FFTGrid.

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

src/layer_3_representation/basis/planewave/internal/fft/gpu/metal.hpp

Public Member Functions

Additional Inherited Members

Detailed Description

Constructor & Destructor Documentation

◆ MetalGPUFFTGrid()

◆ ~MetalGPUFFTGrid()

Member Function Documentation

◆ execute_forward_on_buffers()

◆ execute_inverse_raw_on_buffers()

◆ forward()

◆ forward_single()

◆ g_vector_to_fft_grid_index()

◆ grid_dimension_1()

◆ grid_dimension_2()

◆ grid_dimension_3()

◆ init()

◆ input_buffer()

◆ inverse()

◆ inverse_raw()

◆ inverse_raw_single()

◆ local_grid_points()

◆ local_offset()

◆ output_buffer()

◆ prepare_forward_from_output()

◆ reduce_g_space()

◆ total_grid_points()