forked from bearwithdog/barracuda_beta_hip
-
Notifications
You must be signed in to change notification settings - Fork 0
/
Copy pathbarracuda.cuh
144 lines (96 loc) · 3.14 KB
/
barracuda.cuh
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
#include "hip/hip_runtime.h"
#ifndef BWTALN_CUH
#define BWTALN_CUH
#include "bwt.h"
#include "bwtaln.h"
#include "barracuda.h"
///////////////////////////////////////////////////////////////
// Begin struct (Dag's test)
///////////////////////////////////////////////////////////////
// This struct is for use in CUDA implementation of the functions
// bwa_cal_max_diff() and bwa_approx_mapQ(). In the variable
// "strand_type", bits 1-2 are "type", bit 3 is "strand" in
// corresponding to CPU struct "bwa_seq_t".
//
//typedef struct __align__(16)
typedef struct
{
uint32_t len;
//uint8_t strand_type;
uint8_t strand;
uint8_t type;
uint8_t n_mm;
uint32_t c1;
uint32_t c2;
} bwa_maxdiff_mapQ_t;
///////////////////////////////////////////////////////////////
// End struct (Dag's test)
///////////////////////////////////////////////////////////////
///////////////////////////////////////////////////////////////
// Begin bwa_cal_pac_pos_cuda (Dag's test)
///////////////////////////////////////////////////////////////
void print_cuda_info();
__device__ ubyte_t _bwt_B0(const bwt_t *b, bwtint_t k);
__device__ ubyte_t _bwt_B02(bwtint_t k);
__device__ ubyte_t _rbwt_B02(bwtint_t k);
__device__ ubyte_t _bwt_B03(bwtint_t k, texture<uint4, 1, hipReadModeElementType> *b);
__device__ uint32_t _bwt_bwt(const bwt_t *b, bwtint_t k);
__device__ uint32_t _bwt_bwt2(bwtint_t k);
__device__ uint32_t _rbwt_bwt2(bwtint_t k);
__device__ uint32_t _bwt_bwt3(bwtint_t k, texture<uint4, 1, hipReadModeElementType> *b);
__device__ inline uint32_t* _bwt_occ_intv(const bwt_t *b, bwtint_t k);
__device__ int cuda_bwa_approx_mapQ(const bwa_maxdiff_mapQ_t *p, int mm);
__device__ int cuda_bwa_cal_maxdiff(int l, double err, double thres);
__device__
void update_indices(
int *n_sa_processed,
int *n_sa_remaining,
int *n_sa_in_buf,
int *n_sa_buf_empty);
__device__
void update_indices_in_parallel(
int *n_sa_processed,
int *n_sa_remaining,
int *n_sa_in_buf,
int *n_sa_buf_empty);
__device__
void fetch_read_new_in_parallel(
bwa_maxdiff_mapQ_t *maxdiff_mapQ_buf,
int16_t *sa_origin,
const bwa_maxdiff_mapQ_t *seqs_maxdiff_mapQ_de,
const int offset,
int *n_sa_in_buf,
int *n_sa_buf_empty,
int *n_sa_processed,
int *n_sa_remaining,
int *sa_next_no,
const int n_sa_total);
__device__
void sort_reads(
bwtint_t *sa_buf_arr,
bwa_maxdiff_mapQ_t *maxdiff_mapQ_buf_arr,
int16_t *sa_origin,
bwtint_t *sa_return,
const int *n_sa_in_buf,
int *n_sa_in_buf_prev);
__global__
void cuda_bwa_cal_pac_pos_parallel2(
uint8_t *seqs_mapQ_de,
bwtint_t *seqs_pos_de,
const bwa_maxdiff_mapQ_t *seqs_maxdiff_mapQ_de,
const bwtint_t *seqs_sa_de,
int n_seqs,
int n_block,
int n_seq_per_block,
int block_mod,
int max_mm,
float fnr,
uint4 * bwt_occ_array,
uint4 * rbwt_occ_array,
int * g_log_n_tex,
bwtint_t * bwt_sa_tex,
bwtint_t * rbwt_sa_tex);
///////////////////////////////////////////////////////////////
// End bwa_cal_pac_pos_cuda (Dag's test)
///////////////////////////////////////////////////////////////
#endif // BWTALN_CUH