#include "kernels.cuh"

Include dependency graph for kernels.cu:

Functions
__global__ void	DIRECT_decimator (uint single_tone_length, size_t total_length, float2 __restrict intput, float2 __restrict output)

__global__ void	direct_demodulator_fp64 (double __restrict tone_frquencies, size_t index_counter, uint single_tone_length, size_t total_length, float2 __restrict intput, float2 *__restrict output)
	Direct demodulation kernel. This kernel takes the raw input from the SDR and separate channels. Note: does not do any filtering. More...

__global__ void	direct_demodulator_integer (int __restrict tone_frequencies, int __restrict tone_phases, int wavetablelen, size_t index_counter, size_t single_tone_length, size_t total_length, float2 __restrict input, float2 __restrict output)
	Integer version of the direct demodulation kernel (numerically more stable). This kernel takes the raw input from the SDR and separate channels. Note: does not do any filtering. More...

void	direct_demodulator_wrapper (int __restrict tone_frequencies, int __restrict tone_phases, int wavetablelen, size_t index_counter, size_t single_tone_length, size_t total_length, float2 __restrict input, float2 __restrict output, cudaStream_t internal_stream)

template<typename T >
T *	make_hamming_window (int length, int side, bool diagnostic, bool host_ret)

template<>
float2 *	make_hamming_window< float2 > (int length, int side, bool diagnostic, bool host_ret)

float2 *	make_flat_window (int length, int side, bool diagnostic)
	Creates a flattop window in the GPU memory. More...

float2 *	make_sinc_window (int length, float fc, bool diagnostic=false, bool host_ret=false)

__global__ void	init_states (curandState *state, int twice_vector_len)

__global__ void	make_rand (curandState state, float2 vector, int len, float scale=1)

void	print_chirp_params (std::string comment, chirp_parameter cp)

__device__ float	modulus (float number, float modulus)

__device__ unsigned int	round_index (unsigned int last_index, unsigned int offset, unsigned int num_f, unsigned int f_len)

__global__ void	chirp_gen (float2 __restrict__ output, unsigned int output_size, chirp_parameter __restrict__ info, unsigned long int last_index, float scale=1)

void	chirp_gen_wrapper (float2 __restrict__ output, unsigned int output_size, chirp_parameter __restrict__ info, unsigned long int last_index, cudaStream_t internal_stream, float scale=1)

__global__ void	chirp_demodulator (float2 __restrict__ input, float2 __restrict__ output, unsigned int output_size, unsigned long int last_index, chirp_parameter *__restrict__ info)

__device__ float	absolute (float2 number)

void	chirp_demodulator_wrapper (float2 __restrict__ input, float2 __restrict__ output, unsigned int output_size, unsigned long int last_index, chirp_parameter *__restrict__ info, cudaStream_t internal_stream)

__global__ void	move_buffer (float2 __restrict__ from, float2 __restrict__ to, int size, int from_offset, int to_offset)

void	move_buffer_wrapper (float2 __restrict__ from, float2 __restrict__ to, int size, int from_offset, int to_offset, cudaStream_t internal_stream)

__global__ void	polyphase_filter (float2 __restrict__ input, float2 __restrict__ output, filter_param *__restrict__ filter_info)

void	polyphase_filter_wrapper (float2 __restrict__ input, float2 __restrict__ output, filter_param *__restrict__ filter_info, cudaStream_t internal_stream)

__global__ void	tone_select (float2 __restrict__ input, float2 __restrict__ output, filter_param *__restrict__ filter_info, int effective_batching)

void	tone_select_wrapper (float2 __restrict__ input, float2 __restrict__ output, filter_param *__restrict__ filter_info, int effective_batching, cudaStream_t internal_stream)

__global__ void	scale_buffer (float2 *__restrict__ input, int input_size, float scale)

float2 *	tone_gen (tone_parameters *info, int sampling_rate, float scale, bool device)

__global__ void	mix_buffers (float2 __restrict__ buffer1, float2 __restrict__ buffer2, float2 *__restrict__ output, int length)

__global__ void	average_spectra (float2 __restrict__ input, float2 __restrict__ output, int decim, int nfft, int input_len)

void	decimate_spectra (float2 __restrict__ input, float2 __restrict__ output, int decim, int nfft, int input_len, int output_len, cudaStream_t stram_f)

__global__ void	accumulate_ffts (float2 __restrict__ input, float2 __restrict__ output, int decim, int nfft, int output_length)

__global__ void	zero_mem (float2 *__restrict__ input, int input_len, float value)

__device__ float	magnitude (float2 sample)

void	cublas_decim (float2 __restrict__ input, float2 __restrict__ output, float2 __restrict__ profile, cuComplex __restrict__ zero, cuComplex __restrict__ one, int ppt, int n_freqs, cublasHandle_t __restrict__ handle)

void	decimate_pfb (float2 __restrict__ input, float2 __restrict__ output, int decim, int nfft, int output_length, cudaStream_t stram_f)

void	D_cublas_decim (double2 __restrict__ input, double2 __restrict__ output, double2 __restrict__ profile, cuDoubleComplex __restrict__ zero, cuDoubleComplex __restrict__ one, int ppt, int n_freqs, cublasHandle_t __restrict__ handle)

__global__ void	double2float (double2 __restrict__ input, float2 __restrict__ output, int length)

__global__ void	float2double (float2 __restrict__ input, double2 __restrict__ output, int length)

Function Documentation

◆ DIRECT_decimator()

__global__ void DIRECT_decimator	(	uint	single_tone_length,
		size_t	total_length,
		float2 *__restrict	intput,
		float2 *__restrict	output
	)

◆ direct_demodulator_fp64()

__global__ void direct_demodulator_fp64	(	double *__restrict	tone_frquencies,
		size_t	index_counter,
		uint	single_tone_length,
		size_t	total_length,
		float2 *__restrict	intput,
		float2 *__restrict	output
	)

Direct demodulation kernel. This kernel takes the raw input from the SDR and separate channels. Note: does not do any filtering.

Definition at line 11 of file kernels.cu.

◆ direct_demodulator_integer()

__global__ void direct_demodulator_integer	(	int *__restrict	tone_frequencies,
		int *__restrict	tone_phases,
		int	wavetablelen,
		size_t	index_counter,
		size_t	single_tone_length,
		size_t	total_length,
		float2 *__restrict	input,
		float2 *__restrict	output
	)

Integer version of the direct demodulation kernel (numerically more stable). This kernel takes the raw input from the SDR and separate channels. Note: does not do any filtering.

Definition at line 45 of file kernels.cu.

◆ direct_demodulator_wrapper()

void direct_demodulator_wrapper	(	int *__restrict	tone_frequencies,
		int *__restrict	tone_phases,
		int	wavetablelen,
		size_t	index_counter,
		size_t	single_tone_length,
		size_t	total_length,
		float2 *__restrict	input,
		float2 *__restrict	output,
		cudaStream_t	internal_stream
	)

Wrapper for the integer direct demodulation. Calls the direct_demodulator_integer() kernel and places it on a given stream.

Definition at line 89 of file kernels.cu.

Referenced by RX_buffer_demodulator::close().

Here is the caller graph for this function:

◆ make_hamming_window()

template<typename T >

float2 * make_hamming_window	(	int	length,
		int	side,
		bool	diagnostic,
		bool	host_ret
	)

Definition at line 107 of file kernels.cu.

References pi_f.

◆ make_hamming_window< float2 >()

template<>

float2* make_hamming_window< float2 >	(	int	length,
		int	side,
		bool	diagnostic,
		bool	host_ret
	)

Definition at line 155 of file kernels.cu.

References pi_f.

◆ make_flat_window()

float2* make_flat_window	(	int	length,
		int	side,
		bool	diagnostic
	)

Creates a flattop window in the GPU memory.

Definition at line 208 of file kernels.cu.

Referenced by RX_buffer_demodulator::RX_buffer_demodulator().

Here is the caller graph for this function:

◆ make_sinc_window()

float2* make_sinc_window	(	int	length,
		float	fc,
		bool	diagnostic = `false`,
		bool	host_ret = `false`
	)

Definition at line 258 of file kernels.cu.

References pi_f.

Referenced by RX_buffer_demodulator::RX_buffer_demodulator().

Here is the caller graph for this function:

◆ init_states()

__global__ void init_states	(	curandState *	state,
		int	twice_vector_len
	)

Definition at line 312 of file kernels.cu.

◆ make_rand()

__global__ void make_rand	(	curandState *	state,
		float2 *	vector,
		int	len,
		float	scale = `1`
	)

Definition at line 323 of file kernels.cu.

◆ print_chirp_params()

void print_chirp_params	(	std::string	comment,
		chirp_parameter	cp
	)

Definition at line 335 of file kernels.cu.

◆ modulus()

__device__ float modulus	(	float	number,
		float	modulus
	)

Definition at line 339 of file kernels.cu.

◆ round_index()

__device__ unsigned int round_index	(	unsigned int	last_index,
		unsigned int	offset,
		unsigned int	num_f,
		unsigned int	f_len
	)

Definition at line 343 of file kernels.cu.

◆ chirp_gen()

__global__ void chirp_gen	(	float2 *__restrict__	output,
		unsigned int	output_size,
		chirp_parameter *__restrict__	info,
		unsigned long int	last_index,
		float	scale = `1`
	)

Definition at line 355 of file kernels.cu.

◆ chirp_gen_wrapper()

void chirp_gen_wrapper	(	float2 *__restrict__	output,
		unsigned int	output_size,
		chirp_parameter *__restrict__	info,
		unsigned long int	last_index,
		cudaStream_t	internal_stream,
		float	scale = `1`
	)

Definition at line 395 of file kernels.cu.

◆ chirp_demodulator()

__global__ void chirp_demodulator	(	float2 *__restrict__	input,
		float2 *__restrict__	output,
		unsigned int	output_size,
		unsigned long int	last_index,
		chirp_parameter *__restrict__	info
	)

Definition at line 411 of file kernels.cu.

◆ absolute()

__device__ float absolute ( float2 number )

Definition at line 450 of file kernels.cu.

◆ chirp_demodulator_wrapper()

void chirp_demodulator_wrapper	(	float2 *__restrict__	input,
		float2 *__restrict__	output,
		unsigned int	output_size,
		unsigned long int	last_index,
		chirp_parameter *__restrict__	info,
		cudaStream_t	internal_stream
	)

Definition at line 452 of file kernels.cu.

Referenced by RX_buffer_demodulator::close().

Here is the caller graph for this function:

◆ move_buffer()

__global__ void move_buffer	(	float2 *__restrict__	from,
		float2 *__restrict__	to,
		int	size,
		int	from_offset,
		int	to_offset
	)

Definition at line 466 of file kernels.cu.

◆ move_buffer_wrapper()

void move_buffer_wrapper	(	float2 *__restrict__	from,
		float2 *__restrict__	to,
		int	size,
		int	from_offset,
		int	to_offset,
		cudaStream_t	internal_stream
	)

Definition at line 482 of file kernels.cu.

Referenced by RX_buffer_demodulator::close().

Here is the caller graph for this function:

◆ polyphase_filter()

__global__ void polyphase_filter	(	float2 *__restrict__	input,
		float2 *__restrict__	output,
		filter_param *__restrict__	filter_info
	)

Definition at line 496 of file kernels.cu.

◆ polyphase_filter_wrapper()

void polyphase_filter_wrapper	(	float2 *__restrict__	input,
		float2 *__restrict__	output,
		filter_param *__restrict__	filter_info,
		cudaStream_t	internal_stream
	)

Definition at line 540 of file kernels.cu.

Referenced by RX_buffer_demodulator::close().

Here is the caller graph for this function:

◆ tone_select()

__global__ void tone_select	(	float2 *__restrict__	input,
		float2 *__restrict__	output,
		filter_param *__restrict__	filter_info,
		int	effective_batching
	)

Definition at line 553 of file kernels.cu.

◆ tone_select_wrapper()

void tone_select_wrapper	(	float2 *__restrict__	input,
		float2 *__restrict__	output,
		filter_param *__restrict__	filter_info,
		int	effective_batching,
		cudaStream_t	internal_stream
	)

Definition at line 578 of file kernels.cu.

Referenced by RX_buffer_demodulator::close().

Here is the caller graph for this function:

◆ scale_buffer()

__global__ void scale_buffer	(	float2 *__restrict__	input,
		int	input_size,
		float	scale
	)

Definition at line 592 of file kernels.cu.

◆ tone_gen()

float2* tone_gen	(	tone_parameters *	info,
		int	sampling_rate,
		float	scale,
		bool	device
	)

Definition at line 611 of file kernels.cu.

References tone_parameters::tone_frquencies, tone_parameters::tones_amplitudes, and tone_parameters::tones_number.

◆ mix_buffers()

__global__ void mix_buffers	(	float2 *__restrict__	buffer1,
		float2 *__restrict__	buffer2,
		float2 *__restrict__	output,
		int	length
	)

Definition at line 709 of file kernels.cu.

◆ average_spectra()

__global__ void average_spectra	(	float2 *__restrict__	input,
		float2 *__restrict__	output,
		int	decim,
		int	nfft,
		int	input_len
	)

Definition at line 726 of file kernels.cu.

◆ decimate_spectra()

void decimate_spectra	(	float2 *__restrict__	input,
		float2 *__restrict__	output,
		int	decim,
		int	nfft,
		int	input_len,
		int	output_len,
		cudaStream_t	stram_f
	)

Definition at line 748 of file kernels.cu.

Referenced by RX_buffer_demodulator::close().

Here is the caller graph for this function:

◆ accumulate_ffts()

__global__ void accumulate_ffts	(	float2 *__restrict__	input,
		float2 *__restrict__	output,
		int	decim,
		int	nfft,
		int	output_length
	)

Definition at line 776 of file kernels.cu.

◆ zero_mem()

__global__ void zero_mem	(	float2 *__restrict__	input,
		int	input_len,
		float	value
	)

Definition at line 815 of file kernels.cu.

◆ magnitude()

__device__ float magnitude ( float2 sample )

Definition at line 827 of file kernels.cu.

◆ cublas_decim()

void cublas_decim	(	float2 *__restrict__	input,
		float2 *__restrict__	output,
		float2 *__restrict__	profile,
		cuComplex *__restrict__	zero,
		cuComplex *__restrict__	one,
		int	ppt,
		int	n_freqs,
		cublasHandle_t *__restrict__	handle
	)

Definition at line 874 of file kernels.cu.

Referenced by RX_buffer_demodulator::close().

Here is the caller graph for this function:

◆ decimate_pfb()

void decimate_pfb	(	float2 *__restrict__	input,
		float2 *__restrict__	output,
		int	decim,
		int	nfft,
		int	output_length,
		cudaStream_t	stram_f
	)

Definition at line 900 of file kernels.cu.

References PFB_DECIM_TPB.

Referenced by RX_buffer_demodulator::close().

Here is the caller graph for this function:

◆ D_cublas_decim()

void D_cublas_decim	(	double2 *__restrict__	input,
		double2 *__restrict__	output,
		double2 *__restrict__	profile,
		cuDoubleComplex *__restrict__	zero,
		cuDoubleComplex *__restrict__	one,
		int	ppt,
		int	n_freqs,
		cublasHandle_t *__restrict__	handle
	)

Definition at line 917 of file kernels.cu.

◆ double2float()

__global__ void double2float	(	double2 *__restrict__	input,
		float2 *__restrict__	output,
		int	length
	)

Definition at line 938 of file kernels.cu.

◆ float2double()

__global__ void float2double	(	float2 *__restrict__	input,
		double2 *__restrict__	output,
		int	length
	)

Definition at line 953 of file kernels.cu.

Functions

Function Documentation

◆ DIRECT_decimator()

◆ direct_demodulator_fp64()

◆ direct_demodulator_integer()

◆ direct_demodulator_wrapper()

◆ make_hamming_window()

◆ make_hamming_window< float2 >()

◆ make_flat_window()

◆ make_sinc_window()

◆ init_states()

◆ make_rand()

◆ print_chirp_params()

◆ modulus()

◆ round_index()

◆ chirp_gen()

◆ chirp_gen_wrapper()

◆ chirp_demodulator()

◆ absolute()

◆ chirp_demodulator_wrapper()

◆ move_buffer()

◆ move_buffer_wrapper()

◆ polyphase_filter()

◆ polyphase_filter_wrapper()

◆ tone_select()

◆ tone_select_wrapper()

◆ scale_buffer()

◆ tone_gen()

◆ mix_buffers()

◆ average_spectra()

◆ decimate_spectra()

◆ accumulate_ffts()

◆ zero_mem()

◆ magnitude()

◆ cublas_decim()

◆ decimate_pfb()

◆ D_cublas_decim()

◆ double2float()

◆ float2double()