#include <cuda_runtime.h>
#include "cublas_v2.h"
#include <curand.h>
#include <curand_kernel.h>
#include <cufft.h>
#include <thrust/device_vector.h>
#include <thrust/device_ptr.h>
#include "fir.hpp"

Include dependency graph for kernels.cuh:

This graph shows which files directly or indirectly include this file:

Go to the source code of this file.

Classes
struct	filter_param

struct	chirp_parameter

struct	tone_parameters
	Descriptor of the mutitone generation. More...

Macros
#define	GPU_KERNELS_INCUDED_h

#define	pi_f 3.14159265358979f

#define	Q_PHASE_alt -1.570796327f

#define	_31_BIT_VALUE 2147483647.5

#define	PFB_DECIM_TPB 64.
	Tune the Thread Per Block used in certain functions. More...

Functions
__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)

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)

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)

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

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

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

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

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

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

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

float2 *	tone_gen (tone_parameters *info, int sampling_rate, float scale=1., bool device=false)

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

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)

Macro Definition Documentation

◆ GPU_KERNELS_INCUDED_h

#define GPU_KERNELS_INCUDED_h

Definition at line 12 of file kernels.cuh.

◆ pi_f

#define pi_f 3.14159265358979f

Definition at line 34 of file kernels.cuh.

Referenced by make_hamming_window(), make_hamming_window< float2 >(), and make_sinc_window().

◆ Q_PHASE_alt

#define Q_PHASE_alt -1.570796327f

Definition at line 35 of file kernels.cuh.

◆ _31_BIT_VALUE

#define _31_BIT_VALUE 2147483647.5

Definition at line 36 of file kernels.cuh.

◆ PFB_DECIM_TPB

#define PFB_DECIM_TPB 64.

Tune the Thread Per Block used in certain functions.

Definition at line 41 of file kernels.cuh.

Referenced by decimate_pfb().

Function Documentation

◆ 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:

◆ 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
	)

Definition at line 395 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_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_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_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:

◆ make_hamming_window()

template<typename T >

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

Definition at line 107 of file kernels.cu.

References pi_f.

◆ make_sinc_window()

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

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:

◆ 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:

◆ print_chirp_params()

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

Definition at line 335 of file kernels.cu.

◆ tone_gen()

float2* tone_gen	(	tone_parameters *	info,
		int	sampling_rate,
		float	scale = `1.`,
		bool	device = `false`
	)

Definition at line 611 of file kernels.cu.

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

◆ 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:

◆ 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.

Classes

Macros

Functions

Macro Definition Documentation

◆ GPU_KERNELS_INCUDED_h

◆ pi_f

◆ Q_PHASE_alt

◆ _31_BIT_VALUE

◆ PFB_DECIM_TPB

Function Documentation

◆ direct_demodulator_fp64()

◆ direct_demodulator_integer()

◆ direct_demodulator_wrapper()

◆ chirp_gen_wrapper()

◆ chirp_demodulator_wrapper()

◆ move_buffer_wrapper()

◆ polyphase_filter_wrapper()

◆ tone_select_wrapper()

◆ make_hamming_window()

◆ make_sinc_window()

◆ make_flat_window()

◆ print_chirp_params()

◆ tone_gen()

◆ decimate_spectra()

◆ cublas_decim()

◆ decimate_pfb()

◆ D_cublas_decim()