void save_comp_matrix(comp_matrix *matrix, const char *directory, const char *name) {
size_t err=0;
FILE *fp;
char path[500];
path[0]='\0';
strcat(path, directory);
strcat(path, "/");
strcat(path, name);
strcat(path, ".desp");
fp = fopen(path, "wb+");
check_file_open(fp, path);
err = fwrite(&matrix->siz, sizeof(SA_TYPE), 1, fp);
check_file_write(err, 1, path);
err = fwrite(&matrix->n_desp, sizeof(SA_TYPE), 1, fp);
check_file_write(err, 1, path);
err = fwrite(&matrix->m_desp, sizeof(SA_TYPE), 1, fp);
check_file_write(err, 1, path);
for (SA_TYPE i=0; i<matrix->n_desp; i++) {
err = fwrite(matrix->desp[i], sizeof(SA_TYPE), matrix->m_desp, fp);
check_file_write(err, matrix->m_desp, path);
}
fclose(fp);
#if defined FM_COMP_32 || FM_COMP_64
path[0]='\0';
strcat(path, directory);
strcat(path, "/");
strcat(path, name);
strcat(path, ".count");
fp = fopen(path, "wb+");
check_file_open(fp, path);
err = fwrite(&matrix->n_count, sizeof(SA_TYPE), 1, fp);
check_file_write(err, 1, path);
err = fwrite(&matrix->m_count, sizeof(SA_TYPE), 1, fp);
check_file_write(err, 1, path);
for (SA_TYPE i=0; i<matrix->n_count; i++) {
err = fwrite(matrix->count[i], sizeof(FM_COMP_TYPE), matrix->m_count, fp);
check_file_write(err, matrix->m_count, path);
}
fclose(fp);
#endif
}
void save_ref_vector(ref_vector *vector, const char *directory, const char *name) {
size_t err=0;
FILE *fp;
char path[500];
path[0]='\0';
strcat(path, directory);
strcat(path, "/");
strcat(path, name);
strcat(path, ".vec");
fp = fopen(path, "wb+");
check_file_open(fp, path);
err = fwrite(&vector->n, sizeof(uint64_t), 1, fp);
check_file_write(err, 1, path);
err = fwrite(&vector->dollar, sizeof(uint64_t), 1, fp);
check_file_write(err, 1, path);
err = fwrite(vector->vector, sizeof(uint8_t), vector->n, fp);
check_file_write(err, vector->n, path);
fclose(fp);
}
void save_comp_vector(comp_vector *vector, const char *directory, const char *name) {
size_t err=0;
FILE *fp;
char path[500];
path[0]='\0';
strcat(path, directory);
strcat(path, "/");
strcat(path, name);
strcat(path, ".vec");
fp = fopen(path, "wb+");
check_file_open(fp, path);
err = fwrite(&vector->siz, sizeof(SA_TYPE), 1, fp);
check_file_write(err, 1, path);
err = fwrite(&vector->n, sizeof(SA_TYPE), 1, fp);
check_file_write(err, 1, path);
err = fwrite(&vector->ratio, sizeof(SA_TYPE), 1, fp);
check_file_write(err, 1, path);
err = fwrite(vector->vector, sizeof(SA_TYPE), vector->n, fp);
check_file_write(err, vector->n, path);
fclose(fp);
}
void read_config(char *nucleotides, bool *duplicate_strand, const char *directory) {
if (nucleotides == NULL) {
fprintf(stderr, "%s -> Nucleotides NULL\n", __func__);
exit(EXIT_FAILURE);
}
FILE *fp;
char path[500];
char ds_str[32];
path[0]='\0';
strcat(path, directory);
strcat(path, "/config.txt");
fp = fopen(path, "r");
check_file_open(fp, path);
char *res = fgets(nucleotides, 128, fp);
nucleotides[strlen(nucleotides) - 1] = '\0';
res = fgets(ds_str, 1, fp);
*duplicate_strand = atoi(ds_str);
fclose(fp);
}
void read_comp_vector(comp_vector *vector, const char *directory, const char *name) {
size_t err=0;
FILE *fp;
char path[500];
path[0]='\0';
strcat(path, directory);
strcat(path, "/");
strcat(path, name);
strcat(path, ".vec");
fp = fopen(path, "rb+");
check_file_open(fp, path);
err = fread(&vector->siz, sizeof(SA_TYPE), 1, fp);
check_file_read(err, 1, path);
err = fread(&vector->n, sizeof(SA_TYPE), 1, fp);
check_file_read(err, 1, path);
err = fread(&vector->ratio, sizeof(SA_TYPE), 1, fp);
check_file_read(err, 1, path);
vector->vector = (SA_TYPE *) malloc(vector->n * sizeof(SA_TYPE));
check_malloc(vector->vector, path);
err = fread(vector->vector, sizeof(SA_TYPE), vector->n, fp);
check_file_read(err, vector->n, path);
fclose(fp);
}
void read_vector(vector *vector, const char *directory, const char *name) {
size_t err=0;
FILE *fp;
char path[500]; //TODO: Change to dynamic allocation to avoid buffer overflow
path[0]='\0';
strcat(path, directory);
strcat(path, "/");
strcat(path, name);
strcat(path, ".vec");
fp = fopen(path, "rb+");
check_file_open(fp, path);
err = fread(&vector->n, sizeof(SA_TYPE), 1, fp);
check_file_read(err, 1, path);
vector->vector = (SA_TYPE *) malloc(vector->n * sizeof(SA_TYPE));
check_malloc(vector->vector, path);
err = fread(vector->vector, sizeof(SA_TYPE), vector->n, fp);
check_file_read(err, vector->n, path);
fclose(fp);
}
void read_ref_vector(ref_vector *vector, const char *directory, const char *name) {
size_t err=0;
FILE *fp;
char path[500];
path[0]='\0';
strcat(path, directory);
strcat(path, "/");
strcat(path, name);
strcat(path, ".vec");
fp = fopen(path, "rb+");
check_file_open(fp, path);
err = fread(&vector->n, sizeof(uint64_t), 1, fp);
check_file_read(err, 1, path);
err = fread(&vector->dollar, sizeof(uint64_t), 1, fp);
check_file_read(err, 1, path);
check_file_read(err, 1, path);
vector->vector = (uint8_t *) malloc((vector->n + 1) * sizeof(uint8_t)); //Valgrind errors on dbwt
check_malloc(vector->vector, path);
err = fread(vector->vector, sizeof(uint8_t), vector->n, fp);
check_file_read(err, vector->n, path);
vector->vector[vector->n] = 0; //Valgrind errors on dbwt
fclose(fp);
}
void FileWrite::write(const char* data, std::size_t length)
{
check_file_open();
if( PHYSFS_write(reinterpret_cast<PHYSFS_file*>(mHandle), data, 1, length) != (PHYSFS_sint64)length )
{
BOOST_THROW_EXCEPTION( PhysfsFileException(mFileName) );
}
}
void FileWrite::writeUInt32(uint32_t v)
{
check_file_open();
if( !PHYSFS_writeULE32( reinterpret_cast<PHYSFS_file*>(mHandle), v) )
{
BOOST_THROW_EXCEPTION( PhysfsFileException(mFileName) );
}
}
void File::seek(uint32_t target)
{
check_file_open();
if(!PHYSFS_seek( reinterpret_cast<PHYSFS_file*>(mHandle), target))
{
BOOST_THROW_EXCEPTION( PhysfsFileException(mFileName) );
}
}
float FileRead::readFloat()
{
check_file_open();
float ret;
readRawBytes(reinterpret_cast<char*>(&ret), sizeof(ret));
return ret;
}
char FileRead::readByte()
{
check_file_open();
char ret;
readRawBytes(reinterpret_cast<char*>(&ret), sizeof(ret));
return ret;
}
void load_exome_file(exome *ex, const char *directory) {
FILE *fp;
char path[500];
path[0]='\0';
strcat(path, directory);
strcat(path, "/index");
fp = fopen(path, "r");
check_file_open(fp, path);
char line[MAXLINE];
ex->offset[0]=0;
ex->size = 0;
while (fgets(line, MAXLINE, fp) ) {
if (line[0]=='>') {
int j;
char c = 0;
for(j=0; j<IDMAX-1; j++) {
c = line[j+1];
if (c==' ') break;
ex->chromosome[ex->size*IDMAX+j] = c;
}
ex->chromosome[ex->size*IDMAX+j] = '\0';
sscanf(line + j + 2, "%ju %ju %*s", &ex->start[ex->size], &ex->end[ex->size]);
ex->size++;
if (ex->size >= ex->max_chromosomes) {
ex->max_chromosomes *= 2;
ex->chromosome = (char *)realloc(ex->chromosome,
sizeof(char)*ex->max_chromosomes*IDMAX);
ex->start = (uintmax_t *)realloc(ex->start,
sizeof(uintmax_t)*ex->max_chromosomes);
ex->end = (uintmax_t *)realloc(ex->end,
sizeof(uintmax_t)*ex->max_chromosomes);
ex->offset = (uintmax_t *)realloc(ex->offset,
sizeof(uintmax_t)*ex->max_chromosomes);
}
ex->offset[ex->size] = ex->offset[ex->size-1] + (ex->end[ex->size-1] - ex->start[ex->size-1]+1);
}
}
fclose(fp);
}
uint32_t File::tell() const
{
check_file_open();
PHYSFS_sint64 tp = PHYSFS_tell( reinterpret_cast<PHYSFS_file*> (mHandle) );
if(tp == -1)
BOOST_THROW_EXCEPTION( PhysfsFileException(mFileName) );
return tp;
}
uint32_t File::length() const
{
check_file_open();
PHYSFS_sint64 len = PHYSFS_fileLength( reinterpret_cast<PHYSFS_file*> (mHandle) );
if( len == -1 )
{
BOOST_THROW_EXCEPTION( PhysfsFileException(mFileName) );
}
return len;
}
uint32_t FileRead::readUInt32()
{
check_file_open();
if ( length() - tell() < 4)
{
BOOST_THROW_EXCEPTION( EOFException(mFileName) );
}
PHYSFS_uint32 ret;
if(!PHYSFS_readULE32( reinterpret_cast<PHYSFS_file*>(mHandle), &ret))
{
BOOST_THROW_EXCEPTION( PhysfsFileException(mFileName) );
}
return ret;
}
int creat64(const char *filename, mode_t mode)
{
int e;
funcptr __creat;
DPRINT(("creat64: filename=%s \n", filename));
check_file_open(filename, O_RDWR | O_CREAT);
__creat = load_library_symbol("creat64");
if (__creat == NULL) {
errno = ENOENT;
return -1;
}
DPRINT(("creat64 = %p\n", __creat));
e = __creat(filename, mode);
DPRINT(("creat64: filename=%s e=%d\n", filename, e));
return e;
}
void save_config(char *nucleotides, bool duplicate_strand, const char *directory) {
FILE *fp;
char path[500];
path[0]='\0';
strcat(path, directory);
strcat(path, "/config.txt");
fp = fopen(path, "w");
check_file_open(fp, path);
fputs(nucleotides, fp);
fputc('\n', fp);
fputs((duplicate_strand)?"1":"0", fp);
fputc('\n', fp);
fclose(fp);
}
uint32_t FileRead::readRawBytes( char* target, std::size_t num_of_bytes )
{
check_file_open();
PHYSFS_sint64 num_read = PHYSFS_read(reinterpret_cast<PHYSFS_file*> (mHandle), target, 1, num_of_bytes);
// -1 indicates that reading was not possible
if( num_read == -1)
{
BOOST_THROW_EXCEPTION( PhysfsFileException(mFileName) );
}
if( num_read != (PHYSFS_sint64)num_of_bytes )
{
BOOST_THROW_EXCEPTION ( EOFException(mFileName) );
}
return num_read;
}
/*////////////////////////////////////////////////////////////////
Parse genomic coordinates file, return a map of vectors of intervals
grouped by type of data stream
////////////////////////////////////////////////////////////////*/
unordered_map<string,shared_ptr<vector<TrueGenomicInterval>>>
parseGenomicIntervals(string const & fname) {
ifstream f_in(fname);
check_file_open(f_in, fname);
string line;
unordered_map<string, shared_ptr<vector<TrueGenomicInterval>> > map;
while ( getline(f_in, line) ) {
auto space = line.find(' ');
string suffix = line.substr(0, space);
if (map.find(suffix) == map.end()) {
shared_ptr<vector<TrueGenomicInterval>> intervals(new vector<TrueGenomicInterval>());
map[suffix] = intervals;
}
auto second_space = line.find(' ', space + 1);
unsigned long num_alignments = stoul(line.substr(space + 1, second_space));
map[suffix]->emplace_back( line.substr(second_space + 1), num_alignments );
}
f_in.close();
return map;
}
int __libc_open64(const char *filename, int flags, ...)
{
int e;
funcptr __open;
mode_t mode;
va_list ap;
DPRINT(("__libc_open64: filename=%s \n", filename));
check_file_open(filename, flags);
__open = load_library_symbol("__libc_open64");
if (__open == NULL) {
errno = ENOENT;
return -1;
}
DPRINT(("__libc_open64 = %p\n", __open));
va_start(ap, flags);
mode = va_arg(ap, mode_t);
va_end(ap);
e = __open(filename, flags, mode);
DPRINT(("__libc_open64: filename=%s e=%d\n", filename, e));
return e;
}
void save_exome_file(exome *ex, bool reverse, const char *directory) {
FILE *fp;
char path[500];
path[0]='\0';
strcat(path, directory);
strcat(path, "/index");
fp = fopen(path, "w");
check_file_open(fp, path);
for(uintmax_t i=0; i<ex->size; i++) {
fprintf(fp, ">%s %ju %ju\n", ex->chromosome + i*IDMAX, (uintmax_t) ex->start[i], (uintmax_t) ex->end[i]);
}
if(reverse) {
for(int i=ex->size-1; i>=0; i--) {
fprintf(fp, ">%s %ju %ju\n", ex->chromosome + i*IDMAX, (uintmax_t) ex->start[i], (uintmax_t) ex->end[i]);
}
}
}
void read_comp_matrix(comp_matrix *matrix, const char *directory, const char *name) {
size_t err=0;
FILE *fp;
char path[500];
path[0]='\0';
strcat(path, directory);
strcat(path, "/");
strcat(path, name);
strcat(path, ".desp");
fp = fopen(path, "rb+");
check_file_open(fp, path);
err = fread(&matrix->siz, sizeof(SA_TYPE), 1, fp);
check_file_read(err, 1, path);
err = fread(&matrix->n_desp, sizeof(SA_TYPE), 1, fp);
check_file_read(err, 1, path);
err = fread(&matrix->m_desp, sizeof(SA_TYPE), 1, fp);
check_file_read(err, 1, path);
matrix->desp = (SA_TYPE **) malloc(matrix->n_desp * sizeof(SA_TYPE *));
check_malloc(matrix->desp, path);
for (SA_TYPE i=0; i<matrix->n_desp; i++) {
matrix->desp[i] = (SA_TYPE *) malloc(matrix->m_desp * sizeof(SA_TYPE));
check_malloc(matrix->desp[i], path);
err = fread(matrix->desp[i], sizeof(SA_TYPE), matrix->m_desp, fp);
check_file_read(err, matrix->m_desp, path);
}
fclose(fp);
#if defined FM_COMP_32 || FM_COMP_64
path[0]='\0';
strcat(path, directory);
strcat(path, "/");
strcat(path, name);
strcat(path, ".count");
fp = fopen(path, "rb+");
check_file_open(fp, path);
err = fread(&matrix->n_count, sizeof(SA_TYPE), 1, fp);
check_file_read(err, 1, path);
err = fread(&matrix->m_count, sizeof(SA_TYPE), 1, fp);
check_file_read(err, 1, path);
matrix->count = (FM_COMP_TYPE **) malloc(matrix->n_count * sizeof(FM_COMP_TYPE *));
check_malloc(matrix->count, path);
for (SA_TYPE i=0; i<matrix->n_count; i++){
matrix->count[i] = (FM_COMP_TYPE *) malloc(matrix->m_count * sizeof(FM_COMP_TYPE));
check_malloc(matrix->count[i], path);
err = fread(matrix->count[i], sizeof(FM_COMP_TYPE), matrix->m_count, fp);
check_file_read(err, matrix->m_count, path);
}
fclose(fp);
#endif
}
int main(int argc, char **argv) {
bwt_index backward, forward;
char *Worig;
ref_vector W;
results_list rl_prev, rl_next, rl_prev_i, rl_next_i, rl_final;
uintmax_t read_index=0;
uintmax_t RESULTS, FRAGSIZE;
exome ex;
FILE *queries_file, *output_file;
check_syntax(argc, 7, "exact_search search_file input_dir output_file results_buffer frag_size nucleotides");
timevars();
init_replace_table(argv[6]);
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]);
RESULTS = atoi(argv[4]);
FRAGSIZE = atoi(argv[5]);
if (FRAGSIZE <= 0) {
fprintf(stderr, "Fragsize must be greater than 0\n");
exit(1);
}
tic("Loading FM-Index");
load_bwt_index(NULL, &backward, argv[2], 1, true);
load_bwt_index(NULL, &forward, argv[2], 0, true);
toc();
tic("Preparing search space");
load_exome_file(&ex, argv[2]);
new_results_list(&rl_prev, RESULTS); new_results_list(&rl_prev_i, RESULTS);
new_results_list(&rl_next, RESULTS); new_results_list(&rl_next_i, RESULTS);
new_results_list(&rl_final, RESULTS);
Worig = (char *) malloc( MAXLINE * sizeof(char) );
check_malloc(Worig, "main");
W.vector = (uint8_t *) malloc( MAXLINE * sizeof(uint8_t) );
check_malloc(W.vector, "main");
toc();
intmax_t *k = (intmax_t*)malloc(RESULTS * sizeof(intmax_t));
intmax_t *l = (intmax_t*)malloc(RESULTS * sizeof(intmax_t));
uintmax_t nW_aux;
tic("Sequence mapping");
while(nextFASTAToken(queries_file, Worig, W.vector, &nW_aux)) {
if (!W.vector) {
printf ("Error de lectura de la cadena a buscar\n");
return -1;
}
W.n = nW_aux;
rl_prev.read_index = read_index; rl_prev_i.read_index = read_index;
rl_next.read_index = read_index; rl_next_i.read_index = read_index;
rl_final.read_index = read_index;
BWSearchCPU(W.vector, W.n, &backward, &forward, &rl_prev, &rl_next, &rl_prev_i, &rl_next_i, &rl_final, FRAGSIZE, 1);
write_results(&rl_final, k, l, &ex, &backward, &forward, Worig, nW_aux, 2, output_file);
rl_final.num_results=0;
read_index++;
}
toc();
printf("Memory frees\n");
free(W.vector);
free(Worig);
free(k);
free(l);
free_bwt_index(NULL, &backward);
free_bwt_index(NULL, &forward);
free(rl_prev.list);
free(rl_next.list);
free(rl_prev_i.list);
free(rl_next_i.list);
free(rl_final.list);
fclose(queries_file);
fclose(output_file);
return 0;
}
void encode_reference(ref_vector *X, exome *ex, const char *ref_path, bwt_config_t *bwt_config) {
FILE *ref_file;
ref_file = fopen(ref_path, "r");
check_file_open(ref_file, ref_path);
size_t read, size;
fseek(ref_file, 0, SEEK_END);
read = ftell(ref_file); //Valgrind errors on dbwt
fseek(ref_file, 0, SEEK_SET);
if (bwt_config->duplicate_strand) size = read*2 + 1;
else size = read + 1;
X->vector = (uint8_t *) malloc( size * sizeof(uint8_t) );
check_malloc(X->vector, ref_path);
char *reference = (char *) X->vector;
if (ex !=NULL) ex->size=0;
uintmax_t partial_length=0, total_length=0;
while ( fgets(reference + total_length, MAXLINE, ref_file) ) {
if ( (reference + total_length)[0] == '>') {
if (ex!=NULL) {
if (total_length == 0) {
sscanf(reference + total_length, ">%s ", ex->chromosome + ex->size * IDMAX);
ex->start[ex->size] = 0;
} else {
ex->end[ex->size] = partial_length - 1;
partial_length=0;
if (ex->size == 0)
ex->offset[0] = 0;
else
ex->offset[ex->size] = ex->offset[ex->size-1] + (ex->end[ex->size-1] - ex->start[ex->size-1] + 1);
ex->size++;
if (ex->size >= ex->max_chromosomes) {
ex->max_chromosomes *= 2;
ex->chromosome = (char *)realloc(ex->chromosome,
sizeof(char)*ex->max_chromosomes*IDMAX);
ex->start = (uintmax_t *)realloc(ex->start,
sizeof(uintmax_t)*ex->max_chromosomes);
ex->end = (uintmax_t *)realloc(ex->end,
sizeof(uintmax_t)*ex->max_chromosomes);
ex->offset = (uintmax_t *)realloc(ex->offset,
sizeof(uintmax_t)*ex->max_chromosomes);
}
sscanf(reference + total_length, ">%s ", ex->chromosome + ex->size * IDMAX);
ex->start[ex->size] = 0;
}
}
continue;
}
size_t length = strlen(reference + total_length);
if ((reference + total_length)[length-1]=='\n')
length--;
partial_length += length;
total_length += length;
}
if (ex != NULL) {
ex->end[ex->size] = partial_length - 1;
partial_length=0;
if (ex->size==0)
ex->offset[0] = 0;
else
ex->offset[ex->size] = ex->offset[ex->size-1] + (ex->end[ex->size-1] - ex->start[ex->size-1] + 1);
ex->size++;
}
encode_bases(X->vector, reference, total_length, bwt_config->table);
if (bwt_config->duplicate_strand) {
duplicate_reverse(X->vector, total_length, bwt_config->reverse);
X->n = total_length * 2;
} else {
X->n = total_length;
}
X->dollar = 0;
X->vector[X->n] = 0; //Valgrind errors on dbwt
fclose(ref_file);
}
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;
}
