File with ReMAS kernels for Nvidia GPUs. More...

Functions
__device__ double	__shfl_downD (double var, unsigned int srcLane, int width=sizeWarp)
	__shfl_downD performs __shfl_down of a double number More...

__inline__ __device__ double	warpReduceSumD (double val)
	warpReduceSumD does double sum reduction within a warp More...

__inline__ __device__ float	warpReduceSumS (float val)
	warpReduceSumD does float sum reduction within a warp More...

__global__ void	kernel_InitDTW (MyType *__restrict__ pV, const int pos, const int size)
	kernel_InitDTW This cuda kernel initializes DTW vector More...

__global__ void	kernel_DTW (const MyType __restrict__ Sequence, MyType __restrict__ pD, const int NSeq, const int Where, const int NST)
	kernel_DTW This cuda kernel performs the Online-DTW process for the current frame More...

__global__ void	kernel_InitSxD (MyType __restrict__ odata, MyType __restrict__ v_SxD, const MyType __restrict__ v_dxState, const int __restrict__ I_SxD, const int blockSize, const bool SizeIsPow2, const int size)
	kernel_InitSxD This cuda kernel sets up the vector SxD. More...

__global__ void	kernel_Sum (MyType __restrict__ odata, const MyType __restrict__ idata, const int blockSize, const bool SizeIsPow2, const int size)
	kernel_Sum This cuda kernel adds the elements of a vector. More...

__global__ void	kernel_Vnorm (MyType *__restrict__ odata)
	kernel_Vnorm This cuda kernel initializes position 0 of a vector More...

__global__ void	kernel_ApplyWindow (MyType __restrict__ X_fft, const short __restrict__ frame, const MyType *__restrict__ v_hanning, const int TTRA, const int NFFT)
	kernel_ApplyWindow scales and set the elements of the audio vector X_fft More...

__global__ void	kernel_UpdateSxD (MyType __restrict__ dest, const MyType ALPHA, const MyType __restrict__ norm, const int size)
	kernel_UpdateSxD This cuda kernel update the elements of SxD vector More...

__global__ void	kernel_CompNorB0 (MyType *__restrict__ norms, const MyType value, const int size)
	kernel_CompNorB0 This cuda kernel computes the norm of a vector when BETA=0 More...

__global__ void	kernel_CompNorB1 (MyType __restrict__ norms, const MyType __restrict__ s_fk, const int NMIDI, const int size)
	kernel_CompNorB1 This cuda kernel computes the norm of a vector when BETA=1 More...

__global__ void	kernel_CompNorBG (MyType __restrict__ norms, MyType __restrict__ ts_fk, const MyType *__restrict__ s_fk, const int NMIDI, const MyType BETA, const int size)
	kernel_CompNorBG This cuda kernel computes the norm of a vector when BETA <> 0 and BETA <> 1 More...

__global__ void	kernel_PowToReal (MyType __restrict__ dest, const MyType __restrict__ src, const MyType ex, const int size)
	kernel_PowToReal This cuda kernel powers the elements of a vector to a real number and stores them in other vector More...

__global__ void	kernel_Modul (MyType __restrict__ dest, const MyType __restrict__ src, const int size)
	kernel_Modul This cuda kernel computes the modulus of elements of a vector and stores them in other vector More...

__global__ void	kernel_Cfreq (MyType __restrict__ dest, const MyType __restrict__ src)
	kernel_Cfreq This cuda kernel computes sqrt(sum of elements of a vector) and stores it in dest[0] More...

__global__ void	kernel_Reduction (MyType *__restrict__ dest, const int size)
	kernel_Reduction This cuda kernel performs a typical sum-reduction of a vector More...

__global__ void	kernel_ReductionPowBeta (MyType *__restrict__ dest, const MyType BETA, const int size)
	kernel_Reduction This cuda kernel performs a typical sum-reduction of a vector More...

__global__ void	__launch_bounds__ (maxThreads, 4) kernel_CompDisB0(MyType *__restrict__ dest
	kernel_CompDisB0 This cuda kernel computes the distortion of a vector when BETA=0 More...

	if (i< size)

__global__ void	kernel_Shift (short *__restrict__ frame, const int TTRAMA, const int TMUEST)
	kernel_Shift shifts the vector elements TMUEST positions on the left More...

__global__ void	kernel_BetaNorm (MyType *__restrict__ vector, const int size)
	kernel__BetaNorm normalized the vector More...

__global__ void	kernel_OneImin (MyType __restrict__ odata, int __restrict__ opos, const MyType *__restrict__ idata, const int blockSize, const bool SizeIsPow2, const int size)
	kernel_OneImin calculates the position of one minimum in a vector More...

__global__ void	kernel_OneIminLast (MyType __restrict__ odata, int __restrict__ opos, const MyType __restrict__ idata, const int __restrict__ ipos, const int blockSize, const bool SizeIsPow2, const int size)
	kernel_OneIminLast used with kernel_OneImin to calculates the position of one minimum in a vector More...

__global__ void	kernel_FirstImin (MyType __restrict__ odata, int __restrict__ opos, const MyType *__restrict__ idata, const int blockSize, const bool SizeIsPow2, const int size)
	kernel_FirstImin calculates the position of the first minimum in a vector More...

__global__ void	kernel_FirstIminLast (MyType __restrict__ odata, int __restrict__ opos, const MyType __restrict__ idata, const int __restrict__ ipos, const int blockSize, const bool SizeIsPow2, const int size)
	kernel_FirstIminLast used with kernel_OneImin to calculates the position of the first minimum in a vector More...

__global__ void	kernel_LastImin (MyType __restrict__ odata, int __restrict__ opos, const MyType *__restrict__ idata, const int blockSize, const bool SizeIsPow2, const int size)
	kernel_LastImin calculates the position of the last minimum in a vector More...

__global__ void	kernel_LastIminLast (MyType __restrict__ odata, int __restrict__ opos, const MyType __restrict__ idata, const int __restrict__ ipos, const int blockSize, const bool SizeIsPow2, const int size)
	kernel_LastIminLast used with kernel_LastImin to calculates the position of the last minimum in a vector More...

Variables
__global__ void const MyType *__restrict__	v_cfreq

__global__ void const MyType __restrict__ const MyType __restrict__	norms

__global__ void const MyType __restrict__ const MyType __restrict__ const MyType *__restrict__	s_fk

__global__ void const MyType __restrict__ const MyType __restrict__ const MyType *__restrict__ const int	NMIDI

__global__ void const MyType __restrict__ const MyType __restrict__ const MyType *__restrict__ const int const int	size

unsigned int	j

unsigned int	stride = i * N_MIDI_PAD

unsigned int	th_row = threadIdx.y

unsigned int	th_col = threadIdx.x

unsigned int	row = i + threadIdx.x

bool	guard = th_row == 0 && row < size && th_col < blockDim.y

MyType	a

MyType	b

MyType	tmp1

__shared__ MyType	sh [sizeWarp]

MyType	tmp2

MyType	tmp3 = (1.0 / (BETA*(BETA-1.0)))

__global__ void const MyType __restrict__ const MyType __restrict__ const MyType __restrict__ const MyType __restrict__	ts_fk

__global__ void const MyType __restrict__ const MyType __restrict__ const MyType __restrict__ const MyType __restrict__ const MyType *__restrict__	tauxi

__global__ void const MyType __restrict__ const MyType __restrict__ const MyType __restrict__ const MyType __restrict__ const MyType *__restrict__ const MyType	BETA

unsigned int	k

MyType	beta1 = BETA-1.0

__shared__ MyType	sh_a [sizeWarp/2]

__shared__ MyType	sh_b [sizeWarp/2]

Detailed Description

File with ReMAS kernels for Nvidia GPUs.

Author: Information Retrieval and Parallel Computing Group, University of Oviedo, Spain; Interdisciplinary Computation and Communication Group, Universitat Politecnica de Valencia, Spain; Signal Processing and Telecommunication Systems Research Group, University of Jaen, Spain.; Contact: remas.nosp@m.pack.nosp@m.@gmai.nosp@m.l.co.nosp@m.m

Date: February 13, 2017

Definition in file kernels.cuh.

Function Documentation

◆ __launch_bounds__()

__global__ void __launch_bounds__	(	maxThreads	,
		4
	)

kernel_CompDisB0 This cuda kernel computes the distortion of a vector when BETA=0

kernel_CompDisBG This cuda kernel computes the distortion of a vector when BETA <> 0 and BETA <> 1

kernel_CompDisB1 This cuda kernel computes the distortion of a vector when BETA=1

Parameters

dest	(out) The output vector
v_cfreq	(in) Vector v_cfreq
norms	(in) Vector norms
s_fk	(in) Vector s_fk
NMIDI	(in) Number of midi notes
size	(in) Size of the vectors

Returns: : Nothing, it is a cuda kernel

Parameters

dest	(out) The output vector
v_cfreq	(in) Vector v_cfreq
norms	(in) Vector norms
s_fk	(in) Vector s_fk
ts_fk	(in) Vector ts_fk
tauxi	(in) Vector tauxi
BETA	(in) BETA value
NMIDI	(in) Number of midi notes
size	(in) Size of the vectors

Returns: : Nothing, it is a cuda kernel

◆ __shfl_downD()

__device__ inline double __shfl_downD	(	double	var,
		unsigned int	srcLane,
		int	width = `sizeWarp`
	)

inline

__shfl_downD performs __shfl_down of a double number

Parameters

var	(inout) The data
srcLane	(in) The lane
width	(in) Width of line

Returns: : The __shfl_down of a double number

Definition at line 43 of file kernels.cuh.

◆ kernel_ApplyWindow()

kernel_ApplyWindow	(	MyType *__restrict__	X_fft,
		const short *__restrict__	frame,
		const MyType *__restrict__	v_hanning,
		const int	TTRA,
		const int	NFFT
	)

kernel_ApplyWindow scales and set the elements of the audio vector X_fft

Parameters

X_fft	(out) The vector to scale and set
frame	(in) The vector with current audio (frame)
v_hanning	(in) The hanning vector
TTRA	(in) Size of the frame
NFFT	(in) Size of the X_fft vector

Returns: : Nothing, it is a cuda kernel

Definition at line 345 of file kernels.cuh.

◆ kernel_BetaNorm()

__global__ void kernel_BetaNorm	(	MyType *__restrict__	vector,
		const int	size
	)

kernel__BetaNorm normalized the vector

Parameters

vector	(out) The vector
size	(in) The vector size

Returns: : Nothing, it is a cuda kernel

Definition at line 901 of file kernels.cuh.

◆ kernel_Cfreq()

__global__ void kernel_Cfreq	(	MyType *__restrict__	dest,
		const MyType *__restrict__	src
	)

kernel_Cfreq This cuda kernel computes sqrt(sum of elements of a vector) and stores it in dest[0]

Parameters

dest	(out) The output vector
src	(in) The input vector

Returns: : Nothing, it is a cuda kernel

Definition at line 527 of file kernels.cuh.

◆ kernel_CompNorB0()

__global__ void kernel_CompNorB0	(	MyType *__restrict__	norms,
		const MyType	value,
		const int	size
	)

kernel_CompNorB0 This cuda kernel computes the norm of a vector when BETA=0

Parameters

norms	(out) The output vector
value	(in) The value for the initialization of the vector
size	(in) Size of the vectors

Returns: : Nothing, it is a cuda kernel

Definition at line 386 of file kernels.cuh.

◆ kernel_CompNorB1()

__global__ void kernel_CompNorB1	(	MyType *__restrict__	norms,
		const MyType *__restrict__	s_fk,
		const int	NMIDI,
		const int	size
	)

kernel_CompNorB1 This cuda kernel computes the norm of a vector when BETA=1

Parameters

norms	(out) The output vector
s_fk	(in) Vector s_fk
NMIDI	(in) Number of midi notes
size	(in) Size of the vectors

Returns: : Nothing, it is a cuda kernel

Definition at line 404 of file kernels.cuh.

◆ kernel_CompNorBG()

__global__ void kernel_CompNorBG	(	MyType *__restrict__	norms,
		MyType *__restrict__	ts_fk,
		const MyType *__restrict__	s_fk,
		const int	NMIDI,
		const MyType	BETA,
		const int	size
	)

kernel_CompNorBG This cuda kernel computes the norm of a vector when BETA <> 0 and BETA <> 1

Parameters

norms	(out) The output vector
ts_fk	(out) Vector ts_fk
s_fk	(in) Vector s_fk
NMIDI	(in) Number of midi notes
BETA	(in) Value of parameter BETA
size	(in) Size of the vectors

Returns: : Nothing, it is a cuda kernel

Definition at line 440 of file kernels.cuh.

◆ kernel_DTW()

__global__ void kernel_DTW	(	const MyType *__restrict__	Sequence,
		MyType *__restrict__	pD,
		const int	NSeq,
		const int	Where,
		const int	NST
	)

kernel_DTW This cuda kernel performs the Online-DTW process for the current frame

Parameters

Sequence	(in) Currente frame
pD	(out) DTW pD vector
NSeq	(in) One referenced position within DTW vectors
Where	(in) Position within DTW vectors
NST	(in) Number of states

Returns: : Nothing, it is a cuda kernel

Definition at line 135 of file kernels.cuh.

◆ kernel_FirstImin()

__global__ void kernel_FirstImin	(	MyType *__restrict__	odata,
		int *__restrict__	opos,
		const MyType *__restrict__	idata,
		const int	blockSize,
		const bool	SizeIsPow2,
		const int	size
	)

kernel_FirstImin calculates the position of the first minimum in a vector

Parameters

odata	(out) Stores the minimums
opos	(out) Stores the positions of the minimums
idata	(in) The vector
blockSize	(in) BlockSize used
SizeIsPow2	(in) True if the vector size is power of 2
size	(in) Size of the vector

Returns: : Nothing, it is a cuda kernel

Definition at line 1099 of file kernels.cuh.

◆ kernel_FirstIminLast()

__global__ void kernel_FirstIminLast	(	MyType *__restrict__	odata,
		int *__restrict__	opos,
		const MyType *__restrict__	idata,
		const int *__restrict__	ipos,
		const int	blockSize,
		const bool	SizeIsPow2,
		const int	size
	)

kernel_FirstIminLast used with kernel_OneImin to calculates the position of the first minimum in a vector

Parameters

odata	(out) Stores the minimums
opos	(out) Stores the positions of the minimums
idata	(in) The vector
ipos	(in) Vector with input positions
blockSize	(in) BlockSize used
SizeIsPow2	(in) True if the vector size is power of 2
size	(in) Size of the vector

Returns: : Nothing, it is a cuda kernel

Definition at line 1192 of file kernels.cuh.

◆ kernel_InitDTW()

__global__ void kernel_InitDTW	(	MyType *__restrict__	pV,
		const int	pos,
		const int	size
	)

kernel_InitDTW This cuda kernel initializes DTW vector

Parameters

pV	(out) DTW pV vector
pos	(in) One special position within the vectors
size	(in) Size of DTW vectors

Returns: : Nothing, it is a cuda kernel

Definition at line 108 of file kernels.cuh.

◆ kernel_InitSxD()

__global__ void kernel_InitSxD	(	MyType *__restrict__	odata,
		MyType *__restrict__	v_SxD,
		const MyType *__restrict__	v_dxState,
		const int *__restrict__	I_SxD,
		const int	blockSize,
		const bool	SizeIsPow2,
		const int	size
	)

kernel_InitSxD This cuda kernel sets up the vector SxD.

Parameters

odata	(inout) Intermedial data vector
v_SxD	(out) v_SxD vector
v_dxState	(out) v_dxState vector
I_SxD	(out) I_SxD vector
blockSize	(in) BlockSize used
SizeIsPow2	(in) True if the vector size is power of 2
size	(in) Size of the vector

Returns: : Nothing, it is a cuda kernel

Definition at line 185 of file kernels.cuh.

◆ kernel_LastImin()

__global__ void kernel_LastImin	(	MyType *__restrict__	odata,
		int *__restrict__	opos,
		const MyType *__restrict__	idata,
		const int	blockSize,
		const bool	SizeIsPow2,
		const int	size
	)

kernel_LastImin calculates the position of the last minimum in a vector

Parameters

odata	(out) Stores the minimums
opos	(out) Stores the positions of the minimums
idata	(in) The vector
blockSize	(in) BlockSize used
SizeIsPow2	(in) True if the vector size is power of 2
size	(in) Size of the vector

Returns: : Nothing, it is a cuda kernel

Definition at line 1283 of file kernels.cuh.

◆ kernel_LastIminLast()

__global__ void kernel_LastIminLast	(	MyType *__restrict__	odata,
		int *__restrict__	opos,
		const MyType *__restrict__	idata,
		const int *__restrict__	ipos,
		const int	blockSize,
		const bool	SizeIsPow2,
		const int	size
	)

kernel_LastIminLast used with kernel_LastImin to calculates the position of the last minimum in a vector

Parameters

odata	(out) Stores the minimums
opos	(out) Stores the positions of the minimums
idata	(in) The vector
ipos	(in) Vector with input positions
blockSize	(in) BlockSize used
SizeIsPow2	(in) True if the vector size is power of 2
size	(in) Size of the vector

Returns: : Nothing, it is a cuda kernel

Definition at line 1376 of file kernels.cuh.

◆ kernel_Modul()

__global__ void kernel_Modul	(	MyType *__restrict__	dest,
		const MyType *__restrict__	src,
		const int	size
	)

kernel_Modul This cuda kernel computes the modulus of elements of a vector and stores them in other vector

Parameters

dest	(out) The output vector
src	(in) The input vector
size	(in) Size of the vectors

Returns: : Nothing, it is a cuda kernel

Definition at line 503 of file kernels.cuh.

◆ kernel_OneImin()

__global__ void kernel_OneImin	(	MyType *__restrict__	odata,
		int *__restrict__	opos,
		const MyType *__restrict__	idata,
		const int	blockSize,
		const bool	SizeIsPow2,
		const int	size
	)

kernel_OneImin calculates the position of one minimum in a vector

Parameters

odata	(out) Stores the minimums
opos	(out) Stores the positions of the minimums
idata	(in) The vector
blockSize	(in) BlockSize used
SizeIsPow2	(in) True if the vector size is power of 2
size	(in) Size of the vector

Returns: : Nothing, it is a cuda kernel

Definition at line 923 of file kernels.cuh.

◆ kernel_OneIminLast()

__global__ void kernel_OneIminLast	(	MyType *__restrict__	odata,
		int *__restrict__	opos,
		const MyType *__restrict__	idata,
		const int *__restrict__	ipos,
		const int	blockSize,
		const bool	SizeIsPow2,
		const int	size
	)

kernel_OneIminLast used with kernel_OneImin to calculates the position of one minimum in a vector

Parameters

odata	(out) Stores the minimums
opos	(out) Stores the positions of the minimums
idata	(in) The vector
ipos	(in) Vector with input positions
blockSize	(in) BlockSize used
SizeIsPow2	(in) True if the vector size is power of 2
size	(in) Size of the vector

Returns: : Nothing, it is a cuda kernel

Definition at line 1012 of file kernels.cuh.

◆ kernel_PowToReal()

__global__ void kernel_PowToReal	(	MyType *__restrict__	dest,
		const MyType *__restrict__	src,
		const MyType	ex,
		const int	size
	)

kernel_PowToReal This cuda kernel powers the elements of a vector to a real number and stores them in other vector

Parameters

dest	(out) The output vector
src	(in) The input vector
ex	(in) Number of midi notes
size	(in) Size of the vectors

Returns: : Nothing, it is a cuda kernel

Definition at line 481 of file kernels.cuh.

◆ kernel_Reduction()

__global__ void kernel_Reduction	(	MyType *__restrict__	dest,
		const int	size
	)

kernel_Reduction This cuda kernel performs a typical sum-reduction of a vector

Parameters

dest	(inout) The vector
size	(in) The vector size

Returns: : Nothing, it is a cuda kernel

Definition at line 560 of file kernels.cuh.

◆ kernel_ReductionPowBeta()

__global__ void kernel_ReductionPowBeta	(	MyType *__restrict__	dest,
		const MyType	BETA,
		const int	size
	)

kernel_Reduction This cuda kernel performs a typical sum-reduction of a vector

Parameters

dest	(inout) The vector
BETA	(in) The value of BETA
size	(in) The vector size

Returns: : Nothing, it is a cuda kernel

Definition at line 586 of file kernels.cuh.

◆ kernel_Shift()

__global__ void kernel_Shift	(	short *__restrict__	frame,
		const int	TTRAMA,
		const int	TMUEST
	)

kernel_Shift shifts the vector elements TMUEST positions on the left

Parameters

frame	(out) The vector
TTRAMA	(in) Number of elements of the TRAMA
TMUEST	(in) Number of elements of the MUESTA

Returns: : Nothing, it is a cuda kernel

Definition at line 880 of file kernels.cuh.

◆ kernel_Sum()

__global__ void kernel_Sum	(	MyType *__restrict__	odata,
		const MyType *__restrict__	idata,
		const int	blockSize,
		const bool	SizeIsPow2,
		const int	size
	)

kernel_Sum This cuda kernel adds the elements of a vector.

Parameters

odata	(inout) Intermedial data vector
idata	(in) vector with values to add
blockSize	(in) BlockSize used
SizeIsPow2	(in) True if the vector size is power of 2
size	(in) Size of the vector

Returns: : Nothing, it is a cuda kernel

Definition at line 260 of file kernels.cuh.

◆ kernel_UpdateSxD()

__global__ void kernel_UpdateSxD	(	MyType *__restrict__	dest,
		const MyType	ALPHA,
		const MyType *__restrict__	norm,
		const int	size
	)

kernel_UpdateSxD This cuda kernel update the elements of SxD vector

Parameters

dest	(inout) The vector SxD
ALPHA	(in) The value of parameter ALPHA
norm	(in) Vector with the norm in position 0
size	(in) Size of the vectors

Returns: : Nothing, it is a cuda kernel

Definition at line 364 of file kernels.cuh.

◆ kernel_Vnorm()

__global__ void kernel_Vnorm ( MyType *__restrict__ odata )

kernel_Vnorm This cuda kernel initializes position 0 of a vector

Parameters

odata (inout) The vector

Returns: : Nothing, it is a cuda kernel

Definition at line 326 of file kernels.cuh.

◆ warpReduceSumD()

__inline__ __device__ double warpReduceSumD ( double val )

warpReduceSumD does double sum reduction within a warp

Parameters

val	(in) The data

Returns: : The sum reduction of double type within a warp

Definition at line 67 of file kernels.cuh.

◆ warpReduceSumS()

__inline__ __device__ double warpReduceSumS ( float val )

warpReduceSumD does float sum reduction within a warp

Parameters

val	(in) The data

Returns: : The sum reduction of float type within a warp

Definition at line 87 of file kernels.cuh.

Variable Documentation

◆ size

__global__ void const MyType *__restrict__ const MyType *__restrict__ const MyType *__restrict__ const MyType *__restrict__ const MyType *__restrict__ const MyType const int const int size

Initial value:

{

unsigned int i = blockIdx.x * blockDim.y + threadIdx.y

Definition at line 625 of file kernels.cuh.

Functions

Variables

Detailed Description

Function Documentation

◆ __launch_bounds__()

◆ __shfl_downD()

◆ kernel_ApplyWindow()

◆ kernel_BetaNorm()

◆ kernel_Cfreq()

◆ kernel_CompNorB0()

◆ kernel_CompNorB1()

◆ kernel_CompNorBG()

◆ kernel_DTW()

◆ kernel_FirstImin()

◆ kernel_FirstIminLast()

◆ kernel_InitDTW()

◆ kernel_InitSxD()

◆ kernel_LastImin()

◆ kernel_LastIminLast()

◆ kernel_Modul()

◆ kernel_OneImin()

◆ kernel_OneIminLast()

◆ kernel_PowToReal()

◆ kernel_Reduction()

◆ kernel_ReductionPowBeta()

◆ kernel_Shift()

◆ kernel_Sum()

◆ kernel_UpdateSxD()

◆ kernel_Vnorm()

◆ warpReduceSumD()

◆ warpReduceSumS()

Variable Documentation

◆ size