bool BWSearch1CPU(uint8_t *W, vector *C, vector *C1, comp_matrix *O, comp_matrix *Oi, result *res, results_list *r_list) {
int16_t start, end, half, n;;
SA_TYPE _k, _l;
_k = res->k;
_l = res->l;
start = res->start;
end = res->end;
n = end - start + 1;
half = n / 2;
result r;
init_result(&r, 0);
bound_result(&r, half, end);
change_result(&r, _k, _l, end);
BWExactSearchBackward(W, C, C1, O, &r);
if (r.k <= r.l) {
r.start = start;
r.pos = half-1;
BWSimpleSearch1Backward(W, C, C1, O, &r, r_list);
if (r.k <= r.l) add_result(&r, r_list); //Match
}
half--;
init_result(&r, 1);
bound_result(&r, start, half);
change_result(&r, _k, _l, start);
BWExactSearchForward(W, C, C1, Oi, &r);
if (r.k <= r.l) {
r.pos = half+1;
r.end = end;
BWSimpleSearch1Forward(W, C, C1, Oi, &r, r_list);
}
return false;
}
void BWSearch1CPU(char *W, vector *C, vector *C1, comp_matrix *O, comp_matrix *Oi, result *res, results_list *r_list) {
unsigned int half, n;
int start, end;
unsigned int _k, _l;
_k = res->k;
_l = res->l;
start = res->start;
end = res->end;
n = end - start + 1;
half = n / 2;
result r;
init_result(&r, 0);
bound_result(&r, half, end);
change_result(&r, _k, _l, end);
BWExactSearchBackward(W, C, C1, O, &r);
if (r.k <= r.l) {
r.start = start;
r.pos = half-1;
BWSimpleSearch1Backward(W, C, C1, O, &r, r_list);
if (r.k <= r.l) add_result(&r, r_list); //Match
}
half--;
init_result(&r, 1);
bound_result(&r, start, half);
change_result(&r, _k, _l, start);
BWExactSearchForward(W, C, C1, Oi, &r);
if (r.k <= r.l) {
r.pos = half+1;
r.end = end;
BWSimpleSearch1Forward(W, C, C1, Oi, &r, r_list);
}
}
bool BWSimpleSearch1Forward(uint8_t *W, vector *C, vector *C1, comp_matrix *O, result *res, results_list *r_list) {
SA_TYPE _k, _l, _k_next, _l_next, _k_aux, _l_aux;
SA_TYPE results, results_next;
int16_t start, end, i;
result r;
start = res->pos;
end = res->end;
_k_next = res->k;
_l_next = res->l;
results_next = _l_next - _k_next;
init_result(&r, 1);
bound_result(&r, start, end);
add_mismatch(&r, MATCH, -1, start);
for(i=start; i<=end; i++) {
_k = _k_next;
_l = _l_next;
//printf("%d:\n", i);
if (_k > _l) {
change_result(res, _k, _l, -1);
return false;
}
BWiteration(_k, _l, _k_next, _l_next, W[i], C, C1, O);
results = results_next;
results_next = _l_next - _k_next;
if (results == results_next) continue;
//Deletion
change_result(&r, _k, _l, i+1);
BWExactSearchForward(W, C, C1, O, &r);
if (r.k<=r.l) {
modify_last_mismatch3(&r, DELETION, -1, i);
add_result(&r, r_list);
}
for (uint8_t b=0;b<nA;b++) {
BWiteration(_k, _l, _k_aux, _l_aux, b, C, C1, O);
//printf("W -> %d, %d - %d\n", b, _k_aux, _l_aux);
if (_k_aux > _l_aux) continue;
//Insertion
change_result(&r, _k_aux, _l_aux, i);
BWExactSearchForward(W, C, C1, O, &r);
if (r.k<=r.l) {
modify_last_mismatch3(&r, INSERTION, b, i);
add_result(&r, r_list);
}
//Mismatch
if (b!=(int)W[i]) {
change_result(&r, _k_aux, _l_aux, i+1);
BWExactSearchForward(W, C, C1, O, &r);
if (r.k<=r.l) {
modify_last_mismatch3(&r, MISMATCH, b, i);
add_result(&r, r_list);
}
}
}
}
//Match at exit in res
change_result(res, _k_next, _l_next, -1);
return false;
}
void BWSimpleSearch1Forward(char *W, vector *C, vector *C1, comp_matrix *O, result *res, results_list *r_list) {
BWiterationVariables();
size_t _k, _l, _k_next, _l_next, _k_aux, _l_aux;
size_t results, results_next;
int start, end, i;
result r;
start = res->pos;
end = res->end;
_k_next = res->k;
_l_next = res->l;
results_next = _l_next - _k_next;
init_result(&r, 1);
bound_result(&r, start, end);
add_mismatch(&r, MATCH, XXX, start);
for(i=start; i<=end; i++) {
_k = _k_next;
_l = _l_next;
//printf("%d:\n", i);
if (_k > _l) {
change_result(res, _k, _l, -1);
return;
}
BWiteration(_k, _l, _k_next, _l_next, W[i], C, C1, O);
results = results_next;
results_next = _l_next - _k_next;
if (results == results_next) continue;
//Deletion
change_result(&r, _k, _l, i+1);
modify_last_mismatch_3(&r, DELETION, XXX, i);
BWExactSearchForward(W, C, C1, O, &r);
if (r.k<=r.l) add_result(&r, r_list);
for (unsigned int b=0;b<nA;b++) {
BWiteration(_k, _l, _k_aux, _l_aux, b, C, C1, O);
//printf("W -> %d, %d - %d\n", b, _k_aux, _l_aux);
if (_k_aux > _l_aux) continue;
//Insertion
change_result(&r, _k_aux, _l_aux, i);
modify_last_mismatch_2(&r, INSERTION, b);
BWExactSearchForward(W, C, C1, O, &r);
if (r.k<=r.l) add_result(&r, r_list);
//Mismatch
if (b!=(unsigned int)W[i]) {
change_result(&r, _k_aux, _l_aux, i+1);
modify_last_mismatch_1(&r, MISMATCH);
BWExactSearchForward(W, C, C1, O, &r);
if (r.k<=r.l) add_result(&r, r_list);
}
}
}
//Match at exit in res
change_result(res, _k_next, _l_next, -1);
}
//TODO: Comprobar que funciona igual y borrar el codigo de esta función
void BWSearch1(char *W, int start, int end, size_t *vec_k, size_t *vec_l, size_t *vec_ki, size_t *vec_li, vector *C, vector *C1, comp_matrix *O, comp_matrix *Oi, results_list *r_list) {
BWiterationVariables();
size_t _k, _l, _ki, _li, _k_aux, _l_aux, _ki_aux, _li_aux;
size_t results, results_last;
unsigned int i, j, half, n;
result r;
n = end - start;
half = n / 2;
init_result(&r, 0);
bound_result(&r, start, end);
if (vec_k[0] <= vec_l[0]) {
change_result(&r, vec_k[0], vec_l[0], -1);
add_result(&r, r_list);
}
add_mismatch(&r, MATCH, XXX, start);
results = vec_l[0] - vec_k[0];
results_last = results;
_k = vec_k[1];
_l = vec_l[1];
results = _l - _k;
//printf("B -> %d: %d -> %d, %u, %u\n", 0, results, results_last, _k, _l);
if (results != results_last) {
//printf("*B -> %d: %d -> %d, %u, %u\n", 0, results, results_last, _k, _l);
//printf("%d: %d -> %d\n", 0, results, results_last);
//Deletion
change_result(&r, _k, _l, -1);
modify_last_mismatch_3(&r, DELETION, XXX, start);
add_result(&r, r_list);
for (size_t b=0;b<nA;b++) {
BWiteration(_k, _l, _k_aux, _l_aux, b, C, C1, O);
//printf("W -> %d, %d - %d\n", b, _k_aux, _l_aux);
if (_k_aux > _l_aux) continue;
//printf("*W -> %d, %d - %d\n", b, _k_aux, _l_aux);
size_t b_w = (size_t) W[start];
//Missmatch
if (b!=b_w) {
change_result(&r, _k_aux, _l_aux, -1);
modify_last_mismatch_2(&r, MISMATCH, b);
add_result(&r, r_list);
}
//Insertion
BWiteration(_k_aux, _l_aux, _k_aux, _l_aux, b_w, C, C1, O);
if (_k_aux <= _l_aux) {
change_result(&r, _k_aux, _l_aux, -1);
modify_last_mismatch_3(&r, 3, INSERTION, b);
add_result(&r, r_list);
}
}
}
for (i=start+2, j=2; j<=half; i++, j++) {
results_last = results;
_k = vec_k[j];
_l = vec_l[j];
results = _l - _k;
//printf("B -> %d: %d -> %d, %u, %u\n", j-1, results, results_last, _k, _l);
if (results == results_last) continue;
//printf("*B -> %d: %d -> %d, %u, %u\n", j-1, results, results_last, _k, _l);
//Deletion
change_result(&r, _k, _l, i-2);
modify_last_mismatch_3(&r, DELETION, XXX, i-1);
BWExactSearchBackward(W, C, C1, O, &r);
if (r.k<=r.l) add_result(&r, r_list);
for (size_t b=0;b<nA;b++) {
BWiteration(_k, _l, _k_aux, _l_aux, b, C, C1, O);
if (_k_aux > _l_aux) continue;
//Insertion
change_result(&r, _k_aux, _l_aux, i-1);
modify_last_mismatch_2(&r, INSERTION, b);
BWExactSearchBackward(W, C, C1, O, &r);
if (r.k<=r.l) add_result(&r, r_list);
//Mismatch
if (b!=(size_t)W[i-1]) {
change_result(&r, _k_aux, _l_aux, i-2);
modify_last_mismatch_1(&r, MISMATCH);
BWExactSearchBackward(W, C, C1, O, &r);
if (r.k<=r.l) add_result(&r, r_list);
}
}
}
//printf("\n");
//TODO: Gestionar bien los errores de la busqueda al revés con Si y restas (precalcular Si con |X| - pos)
half--;
results = vec_li[n] - vec_ki[n];
r.dir=1; //Change direction
results_last = results;
_ki = vec_ki[n-1];
_li = vec_li[n-1];
results = _li - _ki;
//printf("F-> %d: %d -> %d, %u, %u\n", n, results, results_last, _ki, _li);
if (results != results_last) {
//printf("*F -> %d: %d -> %d, %u - %u\n", i+1, results, results_last, _ki, _li);
//Deletion
change_result(&r, _ki, _li, -1);
modify_last_mismatch_3(&r, DELETION, XXX, end);
add_result(&r, r_list);
for (size_t b=0;b<nA;b++) {
BWiteration(_ki, _li, _ki_aux, _li_aux, b, C, C1, Oi);
if (_ki_aux > _li_aux) continue;
size_t b_w = (size_t) W[end];
//Mismatch
if (b!=b_w) {
change_result(&r, _ki_aux, _li_aux, -1);
modify_last_mismatch_2(&r, MISMATCH, b);
add_result(&r, r_list);
}
//Insertion
BWiteration(_ki_aux, _li_aux, _ki_aux, _li_aux, b_w, C, C1, Oi);
//printf("\tI -> %d - %d\n", _ki_aux, _li_aux);
if (_ki_aux <= _li_aux){
change_result(&r, _ki_aux, _li_aux, -1);
modify_last_mismatch_2(&r, INSERTION, b);
add_result(&r, r_list);
}
}
}
for(i=end-2,j=n-2; j>=half; i--, j--) {
results_last = results;
_ki = vec_ki[j];
_li = vec_li[j];
results = _li - _ki;
//printf("F -> %d: %d -> %d, %u - %u\n", i+1, results, results_last, _ki, _li);
if (results == results_last) continue;
//printf("*F -> %d: %d -> %d, %u - %u\n", i+1, results, results_last, _ki, _li);
//TODO: Anadir contador para podar cuando se que ya he encontrado las cadenas que permite la variabilidad en este simbolo.
//Deletion
change_result(&r, _ki, _li, i+2);
modify_last_mismatch_3(&r, DELETION, XXX, i+1);
BWExactSearchForward(W, C, C1, Oi, &r);
if (r.k<=r.l) add_result(&r, r_list);
for (size_t b=0;b<nA;b++) {
BWiteration(_ki, _li, _ki_aux, _li_aux, b, C, C1, Oi);
//printf("W -> %d, %d - %d\n", b, _ki_aux, _li_aux);
if (_ki_aux > _li_aux) continue;
//Insertion
change_result(&r, _ki_aux, _li_aux, i+1);
modify_last_mismatch_2(&r, INSERTION, b);
BWExactSearchForward(W, C, C1, Oi, &r);
if (r.k<=r.l) add_result(&r, r_list);
//Mismatch
if (b!= (size_t) W[i+1]) {
change_result(&r, _ki_aux, _li_aux, i+2);
modify_last_mismatch_1(&r, MISMATCH);
BWExactSearchForward(W, C, C1, Oi, &r);
if (r.k<=r.l) add_result(&r, r_list);
}
}
}
//printf("\n");
}
int main(int argc, char **argv) {
char *h_Worig;
uint8_t *h_We, *d_We;
uint64_t *h_nWe, *d_nWe;
uint32_t *h_k, *h_l, *d_k, *d_l;
uint32_t *h_ki, *h_li, *d_ki, *d_li;
intmax_t *h_k2, *h_l2;
intmax_t *h_ki2, *h_li2;
bwt_index backward, forward;
exome ex;
comp_matrix h_O, d_O, h_Oi, d_Oi;
vector h_C, d_C, h_C1, d_C1;
comp_vector h_S, h_Si;
results_list *r_lists;
uint32_t *k, *l;
cudaSetDevice(0);
cudaError_t error;
if (argc!=8) {
printf("Sintaxis:\n\t%s fichero_bus dir_entrada fichero_sal max_bus_gpu repeticiones max_length nucleotides\n", argv[0]);
return 1;
}
timevars();
init_replace_table(argv[7]);
queries_file = fopen(argv[1], "r");
check_file_open(queries_file, argv[1]);
output_file = fopen(argv[3], "w");
check_file_open(output_file, argv[3]);
MAX_BUS_GPU = atoi(argv[4]);
MAXLINE = atoi(argv[6]);
tic("Cargando FM-Index");
read_vector(&h_C, argv[2], "C");
read_vector(&h_C1, argv[2], "C1");
copy_vector_gpu(&d_C, &h_C);
copy_vector_gpu(&d_C1, &h_C1);
read_comp_matrix_gpu(&h_O, argv[2], "O");
read_comp_matrix_gpu(&h_Oi, argv[2], "Oi");
copy_comp_matrix_gpu(&d_O, &h_O);
copy_comp_matrix_gpu(&d_Oi, &h_Oi);
read_comp_vector(&h_S, argv[2], "S");
read_comp_vector(&h_Si, argv[2], "Si");
backward.C = h_C;
backward.C1 = h_C1;
backward.O = h_O;
backward.S = h_S;
forward.C = h_C;
forward.C1 = h_C1;
forward.O = h_Oi;
forward.S = h_Si;
load_exome_file(&ex, argv[2]);
h_Worig = (char*)malloc(MAX_BUS_GPU * MAXLINE * sizeof(char));
cudaMallocHost((void**) &h_We, MAX_BUS_GPU * MAXLINE * sizeof(uint8_t));
cudaMallocHost((void**) &h_nWe, MAX_BUS_GPU * sizeof(uint64_t));
cudaMalloc((void**) &d_We, MAX_BUS_GPU * MAXLINE * sizeof(uint8_t));
manageCudaError();
cudaMalloc((void**) &d_nWe, MAX_BUS_GPU * sizeof(uint64_t));
manageCudaError();
cudaMallocHost((void**) &h_k, MAX_BUS_GPU * MAXLINE * sizeof(uint32_t));
cudaMallocHost((void**) &h_l, MAX_BUS_GPU * MAXLINE * sizeof(uint32_t));
cudaMallocHost((void**) &h_ki, MAX_BUS_GPU * MAXLINE * sizeof(uint32_t));
cudaMallocHost((void**) &h_li, MAX_BUS_GPU * MAXLINE * sizeof(uint32_t));
cudaMallocHost((void**) &h_k2, MAX_BUS_GPU * MAXLINE * sizeof(intmax_t));
cudaMallocHost((void**) &h_l2, MAX_BUS_GPU * MAXLINE * sizeof(intmax_t));
cudaMallocHost((void**) &h_ki2, MAX_BUS_GPU * MAXLINE * sizeof(intmax_t));
cudaMallocHost((void**) &h_li2, MAX_BUS_GPU * MAXLINE * sizeof(intmax_t));
cudaMalloc((void**) &d_k, MAX_BUS_GPU * MAXLINE * sizeof(uint32_t));
manageCudaError();
cudaMalloc((void**) &d_l, MAX_BUS_GPU * MAXLINE * sizeof(uint32_t));
manageCudaError();
cudaMalloc((void**) &d_ki, MAX_BUS_GPU * MAXLINE * sizeof(uint32_t));
manageCudaError();
cudaMalloc((void**) &d_li, MAX_BUS_GPU * MAXLINE * sizeof(uint32_t));
manageCudaError();
r_lists = (results_list *) malloc(MAX_BUS_GPU * sizeof(results_list));
for (int i=0; i<MAX_BUS_GPU; i++) {
new_results_list(&r_lists[i], RESULTS);
}
k = (uint32_t*)malloc(RESULTS * sizeof(uint32_t));
l = (uint32_t*)malloc(RESULTS * sizeof(uint32_t));
toc();
int TAM_BUS_GPU=0, NUM_BLOQUES_GPU=0;
NUM_REP = atoi(argv[5]);
tic("Leer de disco");
while(nextFASTAToken(queries_file, h_Worig + TAM_BUS_GPU * MAXLINE, h_We + TAM_BUS_GPU * MAXLINE, h_nWe + TAM_BUS_GPU)) {
TAM_BUS_GPU++;
if (TAM_BUS_GPU == MAX_BUS_GPU) break;
}
toc();
NUM_BLOQUES_GPU = (TAM_BUS_GPU / TAM_BLOQUE_GPU);
cudaThreadSynchronize();
tic("CPU -> GPU");
cudaMemcpy(d_We, h_We, TAM_BUS_GPU * MAXLINE * sizeof(uint8_t), cudaMemcpyHostToDevice);
manageCudaError();
cudaMemcpy(d_nWe, h_nWe, TAM_BUS_GPU * sizeof(uint64_t), cudaMemcpyHostToDevice);
manageCudaError();
cudaThreadSynchronize();
toc();
cudaThreadSynchronize();
tic("GPU Kernel");
BWExactSearchBackwardVectorGPUWrapper(NUM_BLOQUES_GPU, TAM_BLOQUE_GPU, d_We, d_nWe, MAXLINE, d_k, d_l, 0, d_O.siz-2, &d_C, &d_C1, &d_O);
BWExactSearchForwardVectorGPUWrapper(NUM_BLOQUES_GPU, TAM_BLOQUE_GPU, d_We, d_nWe, MAXLINE, d_ki, d_li, 0, d_Oi.siz-2, &d_C, &d_C1, &d_Oi);
cudaThreadSynchronize();
toc();
cudaThreadSynchronize();
tic("GPU -> CPU");
cudaMemcpy(h_k, d_k, sizeof(uint32_t) * TAM_BUS_GPU * MAXLINE, cudaMemcpyDeviceToHost);
manageCudaError();
cudaMemcpy(h_l, d_l, sizeof(uint32_t) * TAM_BUS_GPU * MAXLINE, cudaMemcpyDeviceToHost);
manageCudaError();
cudaMemcpy(h_ki, d_ki, sizeof(uint32_t) * TAM_BUS_GPU * MAXLINE, cudaMemcpyDeviceToHost);
manageCudaError();
cudaMemcpy(h_li, d_li, sizeof(uint32_t) * TAM_BUS_GPU * MAXLINE, cudaMemcpyDeviceToHost);
manageCudaError();
cudaThreadSynchronize();
toc();
tic("CPU Vector");
for (int i=0; i<TAM_BUS_GPU; i++) {
BWExactSearchVectorBackward(h_We + MAXLINE*i, 0, h_nWe[i]-1, 0, d_O.siz-2, h_k2 + MAXLINE*i, h_l2 + MAXLINE*i, &backward);
BWExactSearchVectorForward(h_We + MAXLINE*i, 0, h_nWe[i]-1, 0, d_Oi.siz-2, h_ki2 + MAXLINE*i, h_li2 + MAXLINE*i, &forward);
}
toc();
tic("CPU Search 1 error");
for (int i=0; i<TAM_BUS_GPU; i++) {
result res;
init_result(&res, 1);
bound_result(&res, 0, h_nWe[i]-1);
change_result(&res, 0, h_O.siz-2, h_nWe[i]-1);
r_lists[i].num_results = 0;
r_lists[i].read_index = i;
BWSearch1CPU(
h_We + i * MAXLINE,
&backward,
&forward,
&res,
&r_lists[i]);
}
toc();
tic("CPU Search 1 Error Helper");
for (int i=0; i<TAM_BUS_GPU; i++) {
r_lists[i].num_results = 0;
r_lists[i].read_index = i;
BWSearch1GPUHelper(
h_We + i * MAXLINE,
0,
h_nWe[i]-1,
h_k + i*MAXLINE,
h_l + i*MAXLINE,
h_ki + i*MAXLINE,
h_li + i*MAXLINE,
&backward,
&forward,
&r_lists[i]
);
}
toc();
tic("Write results");
for (int i=0; i<TAM_BUS_GPU; i++) {
write_results_gpu(&r_lists[i], k, l, &ex, &backward, &forward, h_Worig + i*MAXLINE, h_nWe[i], 1, output_file);
}
toc();
/*
for (int i=0; i<TAM_BUS_GPU; i++) {
for (int j=0; j<h_nWe[i]; j++) {
if (h_k[i*MAXLINE + j] != h_k2[i*MAXLINE + j]) {
printf("Diferente %d %d\n", i, j);
goto salir;
}
}
}
salir:
*/
/*
for (int i=0; i<h_nWe[0]; i++) {
printf("%u ", h_k[i]);
}
printf("\n");
printf("\n");
for (int i=0; i<h_nWe[0]; i++) {
printf("%u ", h_k2[i]);
}
printf("\n");
*/
cudaFreeHost(h_k);
cudaFree(d_k);
cudaFreeHost(h_l);
cudaFree(d_l);
cudaFreeHost(h_We);
cudaFreeHost(h_nWe);
cudaFree(d_We);
cudaFree(d_nWe);
free(h_C.vector);
cudaFree(d_C.vector);
free(h_C1.vector);
cudaFree(d_C1.vector);
free_comp_matrix_gpu_host(NULL, &h_O);
free_comp_matrix_gpu_device(NULL, &d_O);
fclose(queries_file);
return 0;
}
