A collection of most GPU parallel primitives that is implemented as CUDA kernel (most map and expand operations). More...

#include "codec.cu"
#include <limits.h>

Macros
#define	MIN3(a, b, c) ((a) < (b) ? ((a) < (c) ? (a) : (c)) : ((b) < (c) ? (b) : (c)))

Functions
template<typename T >
T	transfer_last_element (T *deviceArr, int n)
	transfer last element of the GPU array to main memory. More...

__global__ void	cal_combination_len (Int3 input, int distance, int output, int n)
	precalculate the number of positions required in the output array of generate combination operation. More...

__global__ void	cal_pair_len (int inputRange, int outputRange, int n)
	precalculate the output range of generate pair operation. More...

__global__ void	cal_pair_len_diag (int inputRange, int outputRange, int n)
	precalculate the output range of pair generation in diagonal position for overlap mode. More...

__global__ void	cal_pair_len_nondiag (Int2 pairs, int seqOffset, int *outputRange, int n)
	precalculate the output range of pair generation in diagonal position for overlap mode. More...

__global__ void	cal_pair_len_lowerbound (int indexes, int inputOffsets, int *outputLengths, int lowerbound, int n)
	precalculate the number of positions required in the output array of generate pair operation with lower bound constratint. More...

__global__ void	generate_pairs (int indexes, Int2 outputs, int inputOffsets, int outputOffsets, int *lesserIndex, int lowerbound, int carry, int n)
	combinatorially generate pairs of indexes within the same group. More...

__global__ void	generate_smaller_index (int indexes, int outputs, int inputOffsets, int outputOffsets, int carry, int n)
	combinatorially generate pairs of indexes within the same group but record only the partial output. More...

__device__ char	levenshtein_static (char *allStr, unsigned int start1, unsigned int start2, int len1, int len2)
	calculate Levenshtein distance of 2 strings in GPU where len1,len2<=18. More...

__device__ char	levenshtein (char *allStr, unsigned int start1, unsigned int start2, int len1, int len2)
	calculate Levenshtein distance of 2 strings in GPU where len1>18 or len2>18. More...

__device__ char	hamming (char *allStr, unsigned int start1, unsigned int start2, int len1, int len2)
	calculate Hamming distance of 2 strings in GPU. More...

__global__ void	cal_distance (char allStr, unsigned int offsets, Int2 index, SeqInfo seqInfo, int seqOffset, int distance, char measure, char distanceOutput, char *flagOutput, int n, int seqLen)
	calculate distances of strings from given pairs and flag ones exceeding the threshold. More...

__global__ void	cal_distance (char allStr, unsigned int offsets, Int2 index, int distance, char measure, char distanceOutput, char *flagOutput, int n, int seqLen)

__global__ void	pair2rep (Int2 pairs, Int2 indexOut, size_t freqOut, SeqInfo seqInfo, int inputOffsets, int outputOffsets, int n)
	turning pairs and frequencies from sequence format to repertoire format. More...

__global__ void	gen_assignment (int matrix, int output, int nBit, int nRow, int nColumn)
	expand operation part of solving bin packing for 2D buffer. More...

__global__ void	gen_bounds (size_t matrix, int keyOut, int *valueOut, int nBit, int valueMax, int nRow, int nColumn)
	expand operation part of solving bin packing for lower bound. More...

__global__ void	flag_lowerbound (int valueInput, int valueOffsets, char *output, int lowerbound, int n)
	flag data to be removed after the lower bound has been processed. More...

__global__ void	make_row_index (int *output, int n, int nRepeat)
	utility to generate keys for matrix processing. More...

__global__ void	toSizeT (int input, size_t output, int n)
	utility to cast types. More...

__global__ void	init_overlap_output (SeqInfo info, Int2 indexOut, size_t freqOut, int inputOffsets, int *outputOffsets, int n)
	generate initial output of overlap mode. More...

__global__ void	toInt3 (char inputs, unsigned int offsets, Int3 *output, int n)
	generate Int3 sequence representation from allStr More...

__global__ void	toInt3 (char inputs, unsigned int offsets, SeqInfo seqInfo, Int3 output, int n)
	generate Int3 sequence representation from allStr, filtered by seqInfo More...

Variables
const size_t	MAX = INT_MAX

Detailed Description

A collection of most GPU parallel primitives that is implemented as CUDA kernel (most map and expand operations).

Follows Facade design pattern.

Macro Definition Documentation

◆ MIN3

#define MIN3	(	a,
		b,
		c
	)	((a) < (b) ? ((a) < (c) ? (a) : (c)) : ((b) < (c) ? (b) : (c)))

Function Documentation

◆ cal_combination_len()

__global__ void cal_combination_len	(	Int3 *	input,
		int	distance,
		int *	output,
		int	n
	)

precalculate the number of positions required in the output array of generate combination operation.

Parameters

input	sequences to generate combination
distance	Levenshtein threshold
output	position output for each sequence
n	array length of input and output

◆ cal_distance() [1/2]

__global__ void cal_distance	(	char *	allStr,
		unsigned int *	offsets,
		Int2 *	index,
		SeqInfo *	seqInfo,
		int *	seqOffset,
		int	distance,
		char	measure,
		char *	distanceOutput,
		char *	flagOutput,
		int	n,
		int	seqLen
	)

calculate distances of strings from given pairs and flag ones exceeding the threshold.

Parameters

allStr	container of all sequences
offsets	start/end position of each sequence
index	pairs of sequence to calculate
seqInfo	info to lookup originalIndex
seqOffset	offset to lookup seqInfo
distance	Levenshtein/Hamming distance threshold
measure	enum representing Levenshtein/Hamming
distanceOutput	output distance, maybe null
flagOutput	array output flag
n	array length of index
seqLen	array length of seq

◆ cal_distance() [2/2]

__global__ void cal_distance	(	char *	allStr,
		unsigned int *	offsets,
		Int2 *	index,
		int	distance,
		char	measure,
		char *	distanceOutput,
		char *	flagOutput,
		int	n,
		int	seqLen
	)

◆ cal_pair_len()

__global__ void cal_pair_len	(	int *	inputRange,
		int *	outputRange,
		int	n
	)

precalculate the output range of generate pair operation.

Parameters

inputRange	range of each input group
outputRange	range of each output group
n	array length of inputRange and outputRange

◆ cal_pair_len_diag()

__global__ void cal_pair_len_diag	(	int *	inputRange,
		int *	outputRange,
		int	n
	)

precalculate the output range of pair generation in diagonal position for overlap mode.

Parameters

inputRange	range of each input group
outputRange	range of each output group
n	array length of inputRange and outputRange

◆ cal_pair_len_lowerbound()

__global__ void cal_pair_len_lowerbound	(	int *	indexes,
		int *	inputOffsets,
		int *	outputLengths,
		int	lowerbound,
		int	n
	)

precalculate the number of positions required in the output array of generate pair operation with lower bound constratint.

Parameters

indexes	value of seqIndexes to generate pair
inputOffsets	group offsets
outputLengths	output position requirement
lowerbound	the processing limit for the indexes
n	array length

◆ cal_pair_len_nondiag()

__global__ void cal_pair_len_nondiag	(	Int2 *	pairs,
		int *	seqOffset,
		int *	outputRange,
		int	n
	)

precalculate the output range of pair generation in diagonal position for overlap mode.

Parameters

pairs	pairs of sequences
seqOffset	range of each input group
outputRange	range of each output group
n	array length of inputRange and outputRange

◆ flag_lowerbound()

__global__ void flag_lowerbound	(	int *	valueInput,
		int *	valueOffsets,
		char *	output,
		int	lowerbound,
		int	n
	)

flag data to be removed after the lower bound has been processed.

This includes both useless group and processed rows.

Parameters

valueInput	seqIndex input
valueOffsets	group offset
output	flag output
lowerbound	the lowerbound used
n	array length of valueOffsets

◆ gen_assignment()

__global__ void gen_assignment	(	int *	matrix,
		int *	output,
		int	nBit,
		int	nRow,
		int	nColumn
	)

expand operation part of solving bin packing for 2D buffer.

Parameters

matrix	statistics of all chunks where each row record the histogram count of each chunk and nRow=nChunk
output	assignment of each chunk to the bins
nBit	bin capacity expressed in log2 form
nRow	number of rows of the matrix
nColumn	number of columns of the matrix

◆ gen_bounds()

__global__ void gen_bounds	(	size_t *	matrix,
		int *	keyOut,
		int *	valueOut,
		int	nBit,
		int	valueMax,
		int	nRow,
		int	nColumn
	)

expand operation part of solving bin packing for lower bound.

Parameters

matrix	statistics of all chunks where each row record the histogram count of each chunk and nRow=nChunk
keyOut	the regrouping of each bin
valueOut	the upper bound of each grouped bin
nBit	bin capacity expressed in log2 form
valueMax	last sequence index
nRow	number of rows of the matrix
nColumn	number of columns of the matrix

◆ generate_pairs()

__global__ void generate_pairs	(	int *	indexes,
		Int2 *	outputs,
		int *	inputOffsets,
		int *	outputOffsets,
		int *	lesserIndex,
		int	lowerbound,
		int	carry,
		int	n
	)

combinatorially generate pairs of indexes within the same group.

Parameters

indexes	value of seqIndexes to generate pair
outputs	pairs output
inputOffsets	precalculated group offsets
outputOffsets	precalculated output position requirement
lesserIndex	by partial output for histogram
lowerbound	the processing limit for the indexes
carry	latest offset from previous chunk in the stream
n	array length of inputOffsets and outputOffsets

◆ generate_smaller_index()

__global__ void generate_smaller_index	(	int *	indexes,
		int *	outputs,
		int *	inputOffsets,
		int *	outputOffsets,
		int	carry,
		int	n
	)

combinatorially generate pairs of indexes within the same group but record only the partial output.

Parameters

indexes	value of seqIndexes to generate pair
outputs	smaller index output
inputOffsets	precalculated group offsets
outputOffsets	precalculated output position requirement
carry	latest offset from previous chunk in the stream
n	array length of inputOffsets and outputOffsets

◆ hamming()

__device__ char hamming	(	char *	allStr,
		unsigned int	start1,
		unsigned int	start2,
		int	len1,
		int	len2
	)

calculate Hamming distance of 2 strings in GPU.

Parameters

allStr	database of all sequences
start1	start index of string one
start2	start index of string two
len1	length of string one
len2	length of string two

◆ init_overlap_output()

__global__ void init_overlap_output	(	SeqInfo *	info,
		Int2 *	indexOut,
		size_t *	freqOut,
		int *	inputOffsets,
		int *	outputOffsets,
		int	n
	)

generate initial output of overlap mode.

Parameters

info	information of each sequence
indexOut	repertoire pair output
freqOut	frequency output for the pair
inputOffsets	index range of input to operate on
outputOffsets	index range of output to operate on
n	length of inputOffset and outputOffset

◆ levenshtein()

__device__ char levenshtein	(	char *	allStr,
		unsigned int	start1,
		unsigned int	start2,
		int	len1,
		int	len2
	)

calculate Levenshtein distance of 2 strings in GPU where len1>18 or len2>18.

Parameters

allStr	database of all sequences
start1	start index of string one
start2	start index of string two
len1	length of string one
len2	length of string two

◆ levenshtein_static()

__device__ char levenshtein_static	(	char *	allStr,
		unsigned int	start1,
		unsigned int	start2,
		int	len1,
		int	len2
	)

calculate Levenshtein distance of 2 strings in GPU where len1,len2<=18.

Parameters

allStr	database of all sequences
start1	start index of string one
start2	start index of string two
len1	length of string one
len2	length of string two

◆ make_row_index()

__global__ void make_row_index	(	int *	output,
		int	n,
		int	nRepeat
	)

utility to generate keys for matrix processing.

Parameters

output	key output with range 0 to n-1 each repeating nRepeat time
n	number of rows
nRepeat	number of columns

◆ pair2rep()

__global__ void pair2rep	(	Int2 *	pairs,
		Int2 *	indexOut,
		size_t *	freqOut,
		SeqInfo *	seqInfo,
		int *	inputOffsets,
		int *	outputOffsets,
		int	n
	)

turning pairs and frequencies from sequence format to repertoire format.

Parameters

pairs	pair result from nearest neighbor search
indexOut	repertiore pair output
freqOut	frequency output
seqInfo	information of each CDR3 sequence
inputOffsets	seqInfo Offset
outputOffsets	output range of indexOut and freqOut
n	number of pairs

◆ toInt3() [1/2]

__global__ void toInt3	(	char *	inputs,
		unsigned int *	offsets,
		Int3 *	output,
		int	n
	)

generate Int3 sequence representation from allStr

Parameters

inputs	container of all sequences
offsets	start/end position of each sequence
output	array of Int3 result with the length of n
n	length of offsets

◆ toInt3() [2/2]

__global__ void toInt3	(	char *	inputs,
		unsigned int *	offsets,
		SeqInfo *	seqInfo,
		Int3 *	output,
		int	n
	)

generate Int3 sequence representation from allStr, filtered by seqInfo

Parameters

inputs	container of all sequences
offsets	start/end position of each sequence
seqInfo	selection of strings to generate output, denoted by originalIndex
output	array of Int3 result with the length of n
n	length of seqInfo

◆ toSizeT()

__global__ void toSizeT	(	int *	input,
		size_t *	output,
		int	n
	)

utility to cast types.

Parameters

input	input array
output	output array
n	number of rows

◆ transfer_last_element()

template<typename T >

T transfer_last_element	(	T *	deviceArr,
		int	n
	)

transfer last element of the GPU array to main memory.

Parameters

deviceArr	the GPU array
n	array length

Variable Documentation

◆ MAX

const size_t MAX = INT_MAX

Macros

Functions

Variables

Detailed Description

Macro Definition Documentation

◆ MIN3

Function Documentation

◆ cal_combination_len()

◆ cal_distance() [1/2]

◆ cal_distance() [2/2]

◆ cal_pair_len()

◆ cal_pair_len_diag()

◆ cal_pair_len_lowerbound()

◆ cal_pair_len_nondiag()

◆ flag_lowerbound()

◆ gen_assignment()

◆ gen_bounds()

◆ generate_pairs()

◆ generate_smaller_index()

◆ hamming()

◆ init_overlap_output()

◆ levenshtein()

◆ levenshtein_static()

◆ make_row_index()

◆ pair2rep()

◆ toInt3() [1/2]

◆ toInt3() [2/2]

◆ toSizeT()

◆ transfer_last_element()

Variable Documentation

◆ MAX