#include <stdio.h>
#include <stdlib.h>
#include <string.h>
//#include <omp.h>
#include <cuda.h>


const char ARQUIVO_DNA_1[] = {"src/DNA1.txt"};
const char ARQUIVO_DNA_2[] = {"src/DNA2.txt"};

char *DNA, *combinacoesDNA;

int *inicioProcesso, *countCombinacoesDNA;

char combinacoesDNAtxt[24][5] = {"ACGT", "ACTG", "AGCT" , "AGTC", 
                                "ATCG", "ATGC", "CAGT", "CATG", 
                                "CGAT", "CGTA", "CTAG", "CTGA", 
                                "GACT", "GATC", "GCAT", "GCTA", 
                                "GTAC", "GTCA", "TACG", "TAGC", 
                                "TCAG", "TCGA", "TGAC", "TGCA"};
                            
void lerArquivo(const char *nomeArquivo, const int tamanho);
int tamanhoArquivo(const char *nomeArquivo);

__device__ int my_strcmp(const char *str_a, const char *str_b, unsigned len = 5)
{
    int match = 0;
    unsigned i = 0;
    unsigned done = 0;
    while ((i < len) && (match == 0) && !done)
    {
      if ((str_a[i] == 0) || (str_b[i] == 0))
      {
        done = 1;
      } 
      else if (str_a[i] != str_b[i])
      {
        match = i+1;
        if (((int)str_a[i] - (int)str_b[i]) < 0)
        {
        match = 0 - (i + 1);
        }
      }
      i++;
    }
    return match;
}

__global__ void pesquisar(char *dna, char *combinacoes, int *inicio, int *countCombinacoes, int tamanhoDNA, unsigned totalProcessos)
{
    int idProcesso = blockIdx.x * blockDim.x + threadIdx.x;
    if(idProcesso < totalProcessos)
    {
        int i;
    
        char busca[5];
        char combinacao[5];

        int count;
        for(count = inicio[idProcesso]; count < inicio[idProcesso + 1]; count++)
        {
            
            busca[0] = dna[count];
            busca[1] = dna[count + 1];
            busca[2] = dna[count + 2];
            busca[3] = dna[count + 3];
            
            for(i = 0; i < 96; i+=4)
            {
                combinacao[0] = combinacoes[i];
                combinacao[1] = combinacoes[i + 1];
                combinacao[2] = combinacoes[i + 2];
                combinacao[3] = combinacoes[i + 3];
                
                if(my_strcmp(busca, combinacao) == 0)
                {
                    // Add no vetor
                    atomicAdd(&countCombinacoes[i/4],1);
                    break;
                }  
            }
        }
    }
}

__global__ void compararEncontrados(char *dna, char *combinacoes, int *inicio, int *maisEncontrados, int *countmaisEncontrados, int tamanhoDNA, int totalProcessos)
{
    int idProcesso = blockIdx.x * blockDim.x + threadIdx.x;
    if(idProcesso < totalProcessos)
    {

        int i;
    
        char busca[5];
        char combinacao[5];
        
        int count;
        for(count = inicio[idProcesso]; count < inicio[idProcesso + 1]; count++)
        {
            
            busca[0] = dna[count];
            busca[1] = dna[count + 1];
            busca[2] = dna[count + 2];
            busca[3] = dna[count + 3];
            
            for(i = 0; i < 17; i+=4)
            {   
                int tempIndex = maisEncontrados[i/4] * 4;
                combinacao[0] = combinacoes[tempIndex];
                combinacao[1] = combinacoes[tempIndex + 1];
                combinacao[2] = combinacoes[tempIndex + 2];
                combinacao[3] = combinacoes[tempIndex + 3];
                
                if(my_strcmp(busca, combinacao) == 0)
                {
                    // Add no vetor
                    atomicAdd(&countmaisEncontrados[i/4],1);  
                    break;
                }
            }
        }
    }
}


int main(int argc, char *argv[0])
{
    if((argv[1] == NULL) || (argv[2] == NULL))
    {
        printf("\n\nInsira o número de blocos e threads! \n\n");
        return(1);
    }

    int *countMaisEncontrados;

    int numBlocos = atoi(argv[1]);
    int numThreads = atoi(argv[2]);
    int totalProcessos = numBlocos * numThreads;

    const int tamanhoDNA = tamanhoArquivo(ARQUIVO_DNA_1);

    cudaMallocManaged(&DNA, tamanhoDNA * sizeof(char));
    cudaMallocManaged(&combinacoesDNA, 96 * sizeof(char));
    cudaMallocManaged(&inicioProcesso, (totalProcessos + 1) * sizeof(int));
    cudaMallocManaged(&countCombinacoesDNA, 24 * sizeof(int));
    cudaMallocManaged(&countMaisEncontrados, 5 * sizeof(int));
    
    strcpy(combinacoesDNA,"ACGTACTGAGCTAGTCATCGATGCCAGTCATGCGATCGTACTAGCTGAGACTGATCGCATGCTAGTACGTCATACGTAGCTCAGTCGATGACTGCA");
    
    lerArquivo(ARQUIVO_DNA_1, tamanhoDNA);

    // Definindo a posição inicial de cada processo
    int i;
    inicioProcesso[0] = 0;
    for(i = 0; i < totalProcessos; i++)
    {
        inicioProcesso[i + 1] = inicioProcesso[i] + tamanhoDNA / totalProcessos;
    }

    /*
    <<<numDeBlocos, numDeThreadsPorBloco>>>
    gtx 960 16 blocos 64 threads
    gtx 1050 12 blocos 64 threads
    */

    pesquisar <<<numBlocos, numThreads>>> (DNA, combinacoesDNA, inicioProcesso, countCombinacoesDNA, tamanhoDNA, totalProcessos);
    cudaDeviceSynchronize();
    
    int *maisEncontrados;
    cudaMallocManaged(&maisEncontrados, 5 * sizeof(int));
    maisEncontrados[0] = 0;
    // Encontra as 5 combinações mais achadas pegando os seus indices
    for (int count = 0; count < 5; count++)
    {
        if (count == 0)
        {
            for (i = 0; i < 24; i++)
            {
                if (countCombinacoesDNA[maisEncontrados[count]] < countCombinacoesDNA[i])
                {
                    maisEncontrados[count] = i;
                }
            }
        }
        else
        {
            for (i = 0; i < 24; i++)
            {
                if (countCombinacoesDNA[maisEncontrados[count]] < countCombinacoesDNA[i] && countCombinacoesDNA[maisEncontrados[count - 1]] > countCombinacoesDNA[i])
                {
                    maisEncontrados[count] = i;
                }
            }
        }
    }

    lerArquivo(ARQUIVO_DNA_2, tamanhoDNA);

    compararEncontrados<<<numBlocos, numThreads>>>(DNA, combinacoesDNA, inicioProcesso, maisEncontrados, countMaisEncontrados, tamanhoDNA, totalProcessos);
    cudaDeviceSynchronize();

    printf("\n Sequência | Encontrado no DNA1 |\n");
    for(i = 0; i < 24; i++)
    {
        printf("\n %-9s | %-18d |", combinacoesDNAtxt[i], countCombinacoesDNA[i]);
    }

    printf("\n\n\nAs 5 sequências mais encontradas no DNA1 presentes no DNA2\n");
    printf("\n Sequência | DNA1    | DNA2\n");
    for (i = 0; i < 5; i++)
    {
        printf("\n %-9s | %-7d | %d", combinacoesDNAtxt[maisEncontrados[i]], countCombinacoesDNA[maisEncontrados[i]], countMaisEncontrados[i]);
    } 
    printf("\n\n");

    cudaFree(DNA);
    cudaFree(countCombinacoesDNA);
    cudaFree(combinacoesDNA);
    cudaFree(inicioProcesso);
    cudaFree(maisEncontrados);
    cudaFree(countMaisEncontrados);

    return 0;
}

void lerArquivo(const char *nomeArquivo, const int tamanho)
{   
    /*
    omp_set_num_threads(12);
    #pragma omp parallel
    {
        FILE *file;
        file = fopen(nomeArquivo, "r");
    
        int i;
        char busca[2];
        // Leitura do arquivo DNA1
        #pragma omp for private (i) 
        fseek(file, i, SEEK_SET);
        for (i = 0; i < tamanho; i++)
        {
            //fseek(file, i, SEEK_SET);
            fgets(busca, 2, file);
            DNA[i] = busca[0];
            //fscanf(file, "%1c", &DNA[i]);
        }
        fclose(file);
    }
    */
    
    
    FILE *file;
    file = fopen(nomeArquivo,"r");
    
    int i;
    // Leitura do arquivo DNA1
    for(i = 0; i < tamanho; i++)
    {
        fscanf(file, "%1c", &DNA[i]);
    }
    fclose(file);
    
    
}

int tamanhoArquivo(const char *nomeArquivo)
{
    FILE *file;
    file = fopen(nomeArquivo, "r");
    fseek(file, 0, SEEK_END);
    int size;
    size = ftell(file);
    fclose(file);
    return size;
}

	.file	"tmpxft_00342de8_00000000-6_seqgene.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2033:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2033:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z9my_strcmpPKcS0_j
	.type	_Z9my_strcmpPKcS0_j, @function
_Z9my_strcmpPKcS0_j:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2027:
	.size	_Z9my_strcmpPKcS0_j, .-_Z9my_strcmpPKcS0_j
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"r"
.LC1:
	.string	"%1c"
	.text
	.globl	_Z10lerArquivoPKci
	.type	_Z10lerArquivoPKci, @function
_Z10lerArquivoPKci:
.LFB2029:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	leaq	.LC1(%rip), %r13
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	movl	%esi, %r12d
	leaq	.LC0(%rip), %rsi
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	xorl	%ebp, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	pushq	%rdx
	.cfi_def_cfa_offset 48
	call	fopen@PLT
	movq	%rax, %rbx
.L5:
	cmpl	%ebp, %r12d
	jle	.L9
	movq	DNA(%rip), %rdx
	movq	%r13, %rsi
	movq	%rbx, %rdi
	xorl	%eax, %eax
	addq	%rbp, %rdx
	incq	%rbp
	call	__isoc23_fscanf@PLT
	jmp	.L5
.L9:
	popq	%rax
	.cfi_def_cfa_offset 40
	movq	%rbx, %rdi
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	jmp	fclose@PLT
	.cfi_endproc
.LFE2029:
	.size	_Z10lerArquivoPKci, .-_Z10lerArquivoPKci
	.globl	_Z14tamanhoArquivoPKc
	.type	_Z14tamanhoArquivoPKc, @function
_Z14tamanhoArquivoPKc:
.LFB2030:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	leaq	.LC0(%rip), %rsi
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	pushq	%rax
	.cfi_def_cfa_offset 32
	call	fopen@PLT
	movl	$2, %edx
	xorl	%esi, %esi
	movq	%rax, %rbp
	movq	%rax, %rdi
	call	fseek@PLT
	movq	%rbp, %rdi
	call	ftell@PLT
	movq	%rbp, %rdi
	movq	%rax, %rbx
	call	fclose@PLT
	movl	%ebx, %eax
	popq	%rdx
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2030:
	.size	_Z14tamanhoArquivoPKc, .-_Z14tamanhoArquivoPKc
	.globl	_Z37__device_stub__Z9pesquisarPcS_PiS0_ijPcS_PiS0_ij
	.type	_Z37__device_stub__Z9pesquisarPcS_PiS0_ijPcS_PiS0_ij, @function
_Z37__device_stub__Z9pesquisarPcS_PiS0_ijPcS_PiS0_ij:
.LFB2055:
	.cfi_startproc
	endbr64
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	movq	%rdi, 40(%rsp)
	leaq	72(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	84(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	56(%rsp), %rdx
	movq	%rcx, 16(%rsp)
	leaq	64(%rsp), %rcx
	movl	%r8d, 12(%rsp)
	movl	%r9d, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 80(%rsp)
	movq	%rax, 120(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 160(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 72(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L12
	pushq	64(%rsp)
	.cfi_def_cfa_offset 200
	leaq	_Z9pesquisarPcS_PiS0_ij(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 208
	movq	100(%rsp), %rcx
	movl	108(%rsp), %r8d
	movq	88(%rsp), %rsi
	movl	96(%rsp), %edx
	leaq	136(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 200
	popq	%rdx
	.cfi_def_cfa_offset 192
.L12:
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L14
	call	__stack_chk_fail@PLT
.L14:
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2055:
	.size	_Z37__device_stub__Z9pesquisarPcS_PiS0_ijPcS_PiS0_ij, .-_Z37__device_stub__Z9pesquisarPcS_PiS0_ijPcS_PiS0_ij
	.globl	_Z9pesquisarPcS_PiS0_ij
	.type	_Z9pesquisarPcS_PiS0_ij, @function
_Z9pesquisarPcS_PiS0_ij:
.LFB2056:
	.cfi_startproc
	endbr64
	jmp	_Z37__device_stub__Z9pesquisarPcS_PiS0_ijPcS_PiS0_ij
	.cfi_endproc
.LFE2056:
	.size	_Z9pesquisarPcS_PiS0_ij, .-_Z9pesquisarPcS_PiS0_ij
	.globl	_Z51__device_stub__Z19compararEncontradosPcS_PiS0_S0_iiPcS_PiS0_S0_ii
	.type	_Z51__device_stub__Z19compararEncontradosPcS_PiS0_S0_iiPcS_PiS0_S0_ii, @function
_Z51__device_stub__Z19compararEncontradosPcS_PiS0_S0_iiPcS_PiS0_S0_ii:
.LFB2057:
	.cfi_startproc
	endbr64
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	movq	%rdi, 40(%rsp)
	leaq	64(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	76(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	48(%rsp), %rdx
	movq	%rcx, 16(%rsp)
	leaq	56(%rsp), %rcx
	movq	%r8, 8(%rsp)
	movl	%r9d, 4(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 72(%rsp)
	movq	%rax, 112(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	192(%rsp), %rax
	movq	%rax, 160(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L17
	pushq	56(%rsp)
	.cfi_def_cfa_offset 200
	leaq	_Z19compararEncontradosPcS_PiS0_S0_ii(%rip), %rdi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 208
	movq	92(%rsp), %rcx
	movl	100(%rsp), %r8d
	movq	80(%rsp), %rsi
	movl	88(%rsp), %edx
	leaq	128(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 200
	popq	%rdx
	.cfi_def_cfa_offset 192
.L17:
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L19
	call	__stack_chk_fail@PLT
.L19:
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2057:
	.size	_Z51__device_stub__Z19compararEncontradosPcS_PiS0_S0_iiPcS_PiS0_S0_ii, .-_Z51__device_stub__Z19compararEncontradosPcS_PiS0_S0_iiPcS_PiS0_S0_ii
	.globl	_Z19compararEncontradosPcS_PiS0_S0_ii
	.type	_Z19compararEncontradosPcS_PiS0_S0_ii, @function
_Z19compararEncontradosPcS_PiS0_S0_ii:
.LFB2058:
	.cfi_startproc
	endbr64
	jmp	_Z51__device_stub__Z19compararEncontradosPcS_PiS0_S0_iiPcS_PiS0_S0_ii
	.cfi_endproc
.LFE2058:
	.size	_Z19compararEncontradosPcS_PiS0_S0_ii, .-_Z19compararEncontradosPcS_PiS0_S0_ii
	.section	.rodata.str1.1
.LC2:
	.string	"\n\nInsira o n\303\272mero de blocos e threads! \n\n"
.LC3:
	.string	"ACGTACTGAGCTAGTCATCGATGCCAGTCATGCGATCGTACTAGCTGAGACTGATCGCATGCTAGTACGTCATACGTAGCTCAGTCGATGACTGCA"
.LC4:
	.string	"\n Sequ\303\252ncia | Encontrado no DNA1 |\n"
.LC5:
	.string	"\n %-9s | %-18d |"
.LC6:
	.string	"\n\n\nAs 5 sequ\303\252ncias mais encontradas no DNA1 presentes no DNA2\n"
.LC7:
	.string	"\n Sequ\303\252ncia | DNA1    | DNA2\n"
.LC8:
	.string	"\n %-9s | %-7d | %d"
.LC9:
	.string	"\n\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	subq	$48, %rsp
	.cfi_def_cfa_offset 96
	movq	8(%rsi), %rdi
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	testq	%rdi, %rdi
	je	.L23
	movq	16(%rsi), %rbx
	testq	%rbx, %rbx
	jne	.L24
.L23:
	leaq	.LC2(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	$1, %eax
	jmp	.L22
.L24:
	call	atoi@PLT
	leaq	_ZL13ARQUIVO_DNA_1(%rip), %r14
	movq	%rbx, %rdi
	movl	%eax, %r12d
	call	atoi@PLT
	movl	%r12d, %ebp
	movq	%r14, %rdi
	imull	%eax, %ebp
	movl	%eax, %r13d
	call	_Z14tamanhoArquivoPKc
	movl	$1, %edx
	leaq	DNA(%rip), %rdi
	movslq	%eax, %rsi
	movq	%rsi, %rbx
	call	cudaMallocManaged@PLT
	movl	$1, %edx
	movl	$96, %esi
	leaq	combinacoesDNA(%rip), %rdi
	call	cudaMallocManaged@PLT
	leal	1(%rbp), %esi
	movl	$1, %edx
	leaq	inicioProcesso(%rip), %rdi
	movslq	%esi, %rsi
	salq	$2, %rsi
	call	cudaMallocManaged@PLT
	movl	$1, %edx
	movl	$96, %esi
	leaq	countCombinacoesDNA(%rip), %rdi
	call	cudaMallocManaged@PLT
	movq	%rsp, %rdi
	movl	$1, %edx
	movl	$20, %esi
	call	cudaMallocManaged@PLT
	movq	combinacoesDNA(%rip), %rdi
	movl	$97, %ecx
	leaq	.LC3(%rip), %rsi
	rep movsb
	movl	%ebx, %esi
	movq	%r14, %rdi
	call	_Z10lerArquivoPKci
	movq	inicioProcesso(%rip), %rsi
	xorl	%edi, %edi
	xorl	%ecx, %ecx
	movl	%edi, (%rsi)
.L26:
	cmpl	%ecx, %ebp
	jle	.L52
	movl	%ebx, %eax
	cltd
	idivl	%ebp
	addl	(%rsi,%rcx,4), %eax
	movl	%eax, 4(%rsi,%rcx,4)
	incq	%rcx
	jmp	.L26
.L52:
	movabsq	$4294967297, %rax
	movl	%r13d, 28(%rsp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	%rax, 32(%rsp)
	movl	36(%rsp), %ecx
	movq	%rax, 20(%rsp)
	movq	28(%rsp), %rdx
	movl	%r12d, 16(%rsp)
	movl	24(%rsp), %esi
	movq	16(%rsp), %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L28
	movq	countCombinacoesDNA(%rip), %rcx
	movl	%ebp, %r9d
	movl	%ebx, %r8d
	movq	inicioProcesso(%rip), %rdx
	movq	combinacoesDNA(%rip), %rsi
	movq	DNA(%rip), %rdi
	call	_Z37__device_stub__Z9pesquisarPcS_PiS0_ijPcS_PiS0_ij
.L28:
	call	cudaDeviceSynchronize@PLT
	movl	$1, %edx
	movl	$20, %esi
	leaq	8(%rsp), %rdi
	call	cudaMallocManaged@PLT
	movq	8(%rsp), %rax
	xorl	%esi, %esi
	xorl	%edx, %edx
	movq	countCombinacoesDNA(%rip), %rcx
	movl	%esi, (%rax)
.L34:
	xorl	%esi, %esi
	testq	%rdx, %rdx
	jne	.L29
.L31:
	movslq	(%rax), %rsi
	movl	(%rcx,%rdx,4), %edi
	cmpl	%edi, (%rcx,%rsi,4)
	jge	.L30
	movl	%edx, (%rax)
.L30:
	incq	%rdx
	cmpq	$24, %rdx
	jne	.L31
	movl	$1, %edx
	jmp	.L34
.L54:
	movslq	-4(%rax,%rdx,4), %r8
	cmpl	(%rcx,%r8,4), %edi
	jge	.L35
	movl	%esi, (%rax,%rdx,4)
.L35:
	incq	%rsi
	cmpq	$24, %rsi
	je	.L53
.L29:
	movslq	(%rax,%rdx,4), %r8
	movl	(%rcx,%rsi,4), %edi
	cmpl	%edi, (%rcx,%r8,4)
	jge	.L35
	jmp	.L54
.L53:
	incq	%rdx
	cmpq	$5, %rdx
	jne	.L34
	movl	%ebx, %esi
	leaq	_ZL13ARQUIVO_DNA_2(%rip), %rdi
	call	_Z10lerArquivoPKci
	movl	%r13d, 28(%rsp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movabsq	$4294967297, %rax
	movl	%r12d, 16(%rsp)
	movq	%rax, 32(%rsp)
	movl	36(%rsp), %ecx
	movq	%rax, 20(%rsp)
	movq	28(%rsp), %rdx
	movq	16(%rsp), %rdi
	movl	24(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L37
	pushq	%rax
	.cfi_def_cfa_offset 104
	movq	inicioProcesso(%rip), %rdx
	movl	%ebx, %r9d
	movq	combinacoesDNA(%rip), %rsi
	movq	DNA(%rip), %rdi
	pushq	%rbp
	.cfi_def_cfa_offset 112
	movq	16(%rsp), %r8
	movq	24(%rsp), %rcx
	call	_Z51__device_stub__Z19compararEncontradosPcS_PiS0_S0_iiPcS_PiS0_S0_ii
	popq	%rdx
	.cfi_def_cfa_offset 104
	popq	%rcx
	.cfi_def_cfa_offset 96
.L37:
	call	cudaDeviceSynchronize@PLT
	movl	$2, %edi
	xorl	%eax, %eax
	xorl	%ebx, %ebx
	leaq	.LC4(%rip), %rsi
	leaq	combinacoesDNAtxt(%rip), %r12
	call	__printf_chk@PLT
	movq	%r12, %rbp
	leaq	.LC5(%rip), %r13
.L38:
	movq	countCombinacoesDNA(%rip), %rax
	movq	%r12, %rdx
	movq	%r13, %rsi
	movl	$2, %edi
	addq	$5, %r12
	movl	(%rax,%rbx), %ecx
	xorl	%eax, %eax
	addq	$4, %rbx
	call	__printf_chk@PLT
	cmpq	$96, %rbx
	jne	.L38
	leaq	.LC6(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	xorl	%ebx, %ebx
	call	__printf_chk@PLT
	leaq	.LC7(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	leaq	.LC8(%rip), %r12
.L39:
	movq	8(%rsp), %rax
	movq	%r12, %rsi
	movl	$2, %edi
	movslq	(%rax,%rbx), %rdx
	movq	countCombinacoesDNA(%rip), %rax
	movl	(%rax,%rdx,4), %ecx
	movq	(%rsp), %rax
	leaq	(%rdx,%rdx,4), %rdx
	addq	%rbp, %rdx
	movl	(%rax,%rbx), %r8d
	xorl	%eax, %eax
	addq	$4, %rbx
	call	__printf_chk@PLT
	cmpq	$20, %rbx
	jne	.L39
	leaq	.LC9(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	DNA(%rip), %rdi
	call	cudaFree@PLT
	movq	countCombinacoesDNA(%rip), %rdi
	call	cudaFree@PLT
	movq	combinacoesDNA(%rip), %rdi
	call	cudaFree@PLT
	movq	inicioProcesso(%rip), %rdi
	call	cudaFree@PLT
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movq	(%rsp), %rdi
	call	cudaFree@PLT
	xorl	%eax, %eax
.L22:
	movq	40(%rsp), %rdx
	subq	%fs:40, %rdx
	je	.L40
	call	__stack_chk_fail@PLT
.L40:
	addq	$48, %rsp
	.cfi_def_cfa_offset 48
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2028:
	.size	main, .-main
	.section	.rodata.str1.1
.LC10:
	.string	"_Z19compararEncontradosPcS_PiS0_S0_ii"
.LC11:
	.string	"_Z9pesquisarPcS_PiS0_ij"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2060:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC10(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z19compararEncontradosPcS_PiS0_S0_ii(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC11(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z9pesquisarPcS_PiS0_ij(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2060:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.globl	combinacoesDNAtxt
	.data
	.align 32
	.type	combinacoesDNAtxt, @object
	.size	combinacoesDNAtxt, 120
combinacoesDNAtxt:
	.string	"ACGT"
	.string	"ACTG"
	.string	"AGCT"
	.string	"AGTC"
	.string	"ATCG"
	.string	"ATGC"
	.string	"CAGT"
	.string	"CATG"
	.string	"CGAT"
	.string	"CGTA"
	.string	"CTAG"
	.string	"CTGA"
	.string	"GACT"
	.string	"GATC"
	.string	"GCAT"
	.string	"GCTA"
	.string	"GTAC"
	.string	"GTCA"
	.string	"TACG"
	.string	"TAGC"
	.string	"TCAG"
	.string	"TCGA"
	.string	"TGAC"
	.string	"TGCA"
	.globl	countCombinacoesDNA
	.bss
	.align 8
	.type	countCombinacoesDNA, @object
	.size	countCombinacoesDNA, 8
countCombinacoesDNA:
	.zero	8
	.globl	inicioProcesso
	.align 8
	.type	inicioProcesso, @object
	.size	inicioProcesso, 8
inicioProcesso:
	.zero	8
	.globl	combinacoesDNA
	.align 8
	.type	combinacoesDNA, @object
	.size	combinacoesDNA, 8
combinacoesDNA:
	.zero	8
	.globl	DNA
	.align 8
	.type	DNA, @object
	.size	DNA, 8
DNA:
	.zero	8
	.section	.rodata
	.align 8
	.type	_ZL13ARQUIVO_DNA_2, @object
	.size	_ZL13ARQUIVO_DNA_2, 13
_ZL13ARQUIVO_DNA_2:
	.string	"src/DNA2.txt"
	.align 8
	.type	_ZL13ARQUIVO_DNA_1, @object
	.size	_ZL13ARQUIVO_DNA_1, 13
_ZL13ARQUIVO_DNA_1:
	.string	"src/DNA1.txt"
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:

#include <stdio.h>
#include <stdlib.h>
#include <string.h>
//#include <omp.h>
#include <hip/hip_runtime.h>


const char ARQUIVO_DNA_1[] = {"src/DNA1.txt"};
const char ARQUIVO_DNA_2[] = {"src/DNA2.txt"};

char *DNA, *combinacoesDNA;

int *inicioProcesso, *countCombinacoesDNA;

char combinacoesDNAtxt[24][5] = {"ACGT", "ACTG", "AGCT" , "AGTC", 
                                "ATCG", "ATGC", "CAGT", "CATG", 
                                "CGAT", "CGTA", "CTAG", "CTGA", 
                                "GACT", "GATC", "GCAT", "GCTA", 
                                "GTAC", "GTCA", "TACG", "TAGC", 
                                "TCAG", "TCGA", "TGAC", "TGCA"};
                            
void lerArquivo(const char *nomeArquivo, const int tamanho);
int tamanhoArquivo(const char *nomeArquivo);

__device__ int my_strcmp(const char *str_a, const char *str_b, unsigned len = 5)
{
    int match = 0;
    unsigned i = 0;
    unsigned done = 0;
    while ((i < len) && (match == 0) && !done)
    {
      if ((str_a[i] == 0) || (str_b[i] == 0))
      {
        done = 1;
      } 
      else if (str_a[i] != str_b[i])
      {
        match = i+1;
        if (((int)str_a[i] - (int)str_b[i]) < 0)
        {
        match = 0 - (i + 1);
        }
      }
      i++;
    }
    return match;
}

__global__ void pesquisar(char *dna, char *combinacoes, int *inicio, int *countCombinacoes, int tamanhoDNA, unsigned totalProcessos)
{
    int idProcesso = blockIdx.x * blockDim.x + threadIdx.x;
    if(idProcesso < totalProcessos)
    {
        int i;
    
        char busca[5];
        char combinacao[5];

        int count;
        for(count = inicio[idProcesso]; count < inicio[idProcesso + 1]; count++)
        {
            
            busca[0] = dna[count];
            busca[1] = dna[count + 1];
            busca[2] = dna[count + 2];
            busca[3] = dna[count + 3];
            
            for(i = 0; i < 96; i+=4)
            {
                combinacao[0] = combinacoes[i];
                combinacao[1] = combinacoes[i + 1];
                combinacao[2] = combinacoes[i + 2];
                combinacao[3] = combinacoes[i + 3];
                
                if(my_strcmp(busca, combinacao) == 0)
                {
                    // Add no vetor
                    atomicAdd(&countCombinacoes[i/4],1);
                    break;
                }  
            }
        }
    }
}

__global__ void compararEncontrados(char *dna, char *combinacoes, int *inicio, int *maisEncontrados, int *countmaisEncontrados, int tamanhoDNA, int totalProcessos)
{
    int idProcesso = blockIdx.x * blockDim.x + threadIdx.x;
    if(idProcesso < totalProcessos)
    {

        int i;
    
        char busca[5];
        char combinacao[5];
        
        int count;
        for(count = inicio[idProcesso]; count < inicio[idProcesso + 1]; count++)
        {
            
            busca[0] = dna[count];
            busca[1] = dna[count + 1];
            busca[2] = dna[count + 2];
            busca[3] = dna[count + 3];
            
            for(i = 0; i < 17; i+=4)
            {   
                int tempIndex = maisEncontrados[i/4] * 4;
                combinacao[0] = combinacoes[tempIndex];
                combinacao[1] = combinacoes[tempIndex + 1];
                combinacao[2] = combinacoes[tempIndex + 2];
                combinacao[3] = combinacoes[tempIndex + 3];
                
                if(my_strcmp(busca, combinacao) == 0)
                {
                    // Add no vetor
                    atomicAdd(&countmaisEncontrados[i/4],1);  
                    break;
                }
            }
        }
    }
}


int main(int argc, char *argv[0])
{
    if((argv[1] == NULL) || (argv[2] == NULL))
    {
        printf("\n\nInsira o número de blocos e threads! \n\n");
        return(1);
    }

    int *countMaisEncontrados;

    int numBlocos = atoi(argv[1]);
    int numThreads = atoi(argv[2]);
    int totalProcessos = numBlocos * numThreads;

    const int tamanhoDNA = tamanhoArquivo(ARQUIVO_DNA_1);

    hipMallocManaged(&DNA, tamanhoDNA * sizeof(char));
    hipMallocManaged(&combinacoesDNA, 96 * sizeof(char));
    hipMallocManaged(&inicioProcesso, (totalProcessos + 1) * sizeof(int));
    hipMallocManaged(&countCombinacoesDNA, 24 * sizeof(int));
    hipMallocManaged(&countMaisEncontrados, 5 * sizeof(int));
    
    strcpy(combinacoesDNA,"ACGTACTGAGCTAGTCATCGATGCCAGTCATGCGATCGTACTAGCTGAGACTGATCGCATGCTAGTACGTCATACGTAGCTCAGTCGATGACTGCA");
    
    lerArquivo(ARQUIVO_DNA_1, tamanhoDNA);

    // Definindo a posição inicial de cada processo
    int i;
    inicioProcesso[0] = 0;
    for(i = 0; i < totalProcessos; i++)
    {
        inicioProcesso[i + 1] = inicioProcesso[i] + tamanhoDNA / totalProcessos;
    }

    /*
    <<<numDeBlocos, numDeThreadsPorBloco>>>
    gtx 960 16 blocos 64 threads
    gtx 1050 12 blocos 64 threads
    */

    pesquisar <<<numBlocos, numThreads>>> (DNA, combinacoesDNA, inicioProcesso, countCombinacoesDNA, tamanhoDNA, totalProcessos);
    hipDeviceSynchronize();
    
    int *maisEncontrados;
    hipMallocManaged(&maisEncontrados, 5 * sizeof(int));
    maisEncontrados[0] = 0;
    // Encontra as 5 combinações mais achadas pegando os seus indices
    for (int count = 0; count < 5; count++)
    {
        if (count == 0)
        {
            for (i = 0; i < 24; i++)
            {
                if (countCombinacoesDNA[maisEncontrados[count]] < countCombinacoesDNA[i])
                {
                    maisEncontrados[count] = i;
                }
            }
        }
        else
        {
            for (i = 0; i < 24; i++)
            {
                if (countCombinacoesDNA[maisEncontrados[count]] < countCombinacoesDNA[i] && countCombinacoesDNA[maisEncontrados[count - 1]] > countCombinacoesDNA[i])
                {
                    maisEncontrados[count] = i;
                }
            }
        }
    }

    lerArquivo(ARQUIVO_DNA_2, tamanhoDNA);

    compararEncontrados<<<numBlocos, numThreads>>>(DNA, combinacoesDNA, inicioProcesso, maisEncontrados, countMaisEncontrados, tamanhoDNA, totalProcessos);
    hipDeviceSynchronize();

    printf("\n Sequência | Encontrado no DNA1 |\n");
    for(i = 0; i < 24; i++)
    {
        printf("\n %-9s | %-18d |", combinacoesDNAtxt[i], countCombinacoesDNA[i]);
    }

    printf("\n\n\nAs 5 sequências mais encontradas no DNA1 presentes no DNA2\n");
    printf("\n Sequência | DNA1    | DNA2\n");
    for (i = 0; i < 5; i++)
    {
        printf("\n %-9s | %-7d | %d", combinacoesDNAtxt[maisEncontrados[i]], countCombinacoesDNA[maisEncontrados[i]], countMaisEncontrados[i]);
    } 
    printf("\n\n");

    hipFree(DNA);
    hipFree(countCombinacoesDNA);
    hipFree(combinacoesDNA);
    hipFree(inicioProcesso);
    hipFree(maisEncontrados);
    hipFree(countMaisEncontrados);

    return 0;
}

void lerArquivo(const char *nomeArquivo, const int tamanho)
{   
    /*
    omp_set_num_threads(12);
    #pragma omp parallel
    {
        FILE *file;
        file = fopen(nomeArquivo, "r");
    
        int i;
        char busca[2];
        // Leitura do arquivo DNA1
        #pragma omp for private (i) 
        fseek(file, i, SEEK_SET);
        for (i = 0; i < tamanho; i++)
        {
            //fseek(file, i, SEEK_SET);
            fgets(busca, 2, file);
            DNA[i] = busca[0];
            //fscanf(file, "%1c", &DNA[i]);
        }
        fclose(file);
    }
    */
    
    
    FILE *file;
    file = fopen(nomeArquivo,"r");
    
    int i;
    // Leitura do arquivo DNA1
    for(i = 0; i < tamanho; i++)
    {
        fscanf(file, "%1c", &DNA[i]);
    }
    fclose(file);
    
    
}

int tamanhoArquivo(const char *nomeArquivo)
{
    FILE *file;
    file = fopen(nomeArquivo, "r");
    fseek(file, 0, SEEK_END);
    int size;
    size = ftell(file);
    fclose(file);
    return size;
}

	.text
	.file	"seqgene.hip"
	.globl	_Z24__device_stub__pesquisarPcS_PiS0_ij # -- Begin function _Z24__device_stub__pesquisarPcS_PiS0_ij
	.type	_Z24__device_stub__pesquisarPcS_PiS0_ij,@function
_Z24__device_stub__pesquisarPcS_PiS0_ij: # @_Z24__device_stub__pesquisarPcS_PiS0_ij
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$144, %rsp
	.cfi_def_cfa_offset 192
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	56(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	48(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	12(%rsp), %rcx
	movl	%r8d, (%rcx)
	leaq	8(%rsp), %r8
	movl	%r9d, (%r8)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z9pesquisarPcS_PiS0_ij, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$160, %rsp
	.cfi_adjust_cfa_offset -160
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z24__device_stub__pesquisarPcS_PiS0_ij, .Lfunc_end0-_Z24__device_stub__pesquisarPcS_PiS0_ij
	.cfi_endproc
                                        # -- End function
	.globl	_Z34__device_stub__compararEncontradosPcS_PiS0_S0_ii # -- Begin function _Z34__device_stub__compararEncontradosPcS_PiS0_S0_ii
	.type	_Z34__device_stub__compararEncontradosPcS_PiS0_S0_ii,@function
_Z34__device_stub__compararEncontradosPcS_PiS0_S0_ii: # @_Z34__device_stub__compararEncontradosPcS_PiS0_S0_ii
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$160, %rsp
	.cfi_def_cfa_offset 208
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	56(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	48(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	24(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	4(%rsp), %r8
	movl	%r9d, (%r8)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	208(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z19compararEncontradosPcS_PiS0_S0_ii, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$176, %rsp
	.cfi_adjust_cfa_offset -176
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z34__device_stub__compararEncontradosPcS_PiS0_S0_ii, .Lfunc_end1-_Z34__device_stub__compararEncontradosPcS_PiS0_S0_ii
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	8(%rsi), %rdi
	testq	%rdi, %rdi
	je	.LBB2_2
# %bb.1:
	movq	16(%rsi), %rbx
	testq	%rbx, %rbx
	je	.LBB2_2
# %bb.3:
	callq	atoi
	movl	%eax, %r14d
	movq	%rbx, %rdi
	callq	atoi
	movl	%eax, %r15d
	movl	%eax, %ebx
	imull	%r14d, %ebx
	movl	$_ZL13ARQUIVO_DNA_1, %edi
	callq	_Z14tamanhoArquivoPKc
	movl	%eax, %ebp
	movslq	%eax, %r12
	movl	$DNA, %edi
	movq	%r12, %rsi
	movl	$1, %edx
	callq	hipMallocManaged
	movl	$combinacoesDNA, %edi
	movl	$96, %esi
	movl	$1, %edx
	callq	hipMallocManaged
	movslq	%ebx, %r13
	leaq	4(,%r13,4), %rsi
	movl	$inicioProcesso, %edi
	movl	$1, %edx
	callq	hipMallocManaged
	movl	$countCombinacoesDNA, %edi
	movl	$96, %esi
	movl	$1, %edx
	callq	hipMallocManaged
	leaq	16(%rsp), %rdi
	movl	$20, %esi
	movl	$1, %edx
	callq	hipMallocManaged
	movq	combinacoesDNA(%rip), %rdi
	movl	$.L.str.1, %esi
	movl	$97, %edx
	callq	memcpy@PLT
	movl	$_ZL13ARQUIVO_DNA_1, %edi
	movl	%r12d, %esi
	callq	_Z10lerArquivoPKci
	movq	inicioProcesso(%rip), %rcx
	movl	$0, (%rcx)
	testl	%r13d, %r13d
	jle	.LBB2_6
# %bb.4:                                # %.lr.ph
	movl	%ebp, %eax
	cltd
	idivl	%ebx
	movl	%ebx, %edx
	xorl	%esi, %esi
	movl	%eax, %edi
.LBB2_5:                                # =>This Inner Loop Header: Depth=1
	movl	%edi, 4(%rcx,%rsi,4)
	incq	%rsi
	addl	%eax, %edi
	cmpq	%rsi, %rdx
	jne	.LBB2_5
.LBB2_6:                                # %._crit_edge
	movl	%r14d, %r14d
	btsq	$32, %r14
	movl	%r15d, %r15d
	btsq	$32, %r15
	xorl	%r13d, %r13d
	movq	%r14, %rdi
	movl	$1, %esi
	movq	%r15, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_8
# %bb.7:
	movq	DNA(%rip), %rdi
	movq	combinacoesDNA(%rip), %rsi
	movq	inicioProcesso(%rip), %rdx
	movq	countCombinacoesDNA(%rip), %rcx
	movl	%ebp, %r8d
	movl	%ebx, %r9d
	callq	_Z24__device_stub__pesquisarPcS_PiS0_ij
.LBB2_8:
	callq	hipDeviceSynchronize
	leaq	8(%rsp), %r12
	movl	$20, %esi
	movq	%r12, %rdi
	movl	$1, %edx
	callq	hipMallocManaged
	movq	(%r12), %rax
	movl	$0, (%rax)
	movq	countCombinacoesDNA(%rip), %rcx
	xorl	%edx, %edx
.LBB2_9:                                # =>This Loop Header: Depth=1
                                        #     Child Loop BB2_15 Depth 2
                                        #     Child Loop BB2_11 Depth 2
	testq	%rdx, %rdx
	je	.LBB2_10
# %bb.14:                               # %.preheader76
                                        #   in Loop: Header=BB2_9 Depth=1
	movl	(%rax,%rdx,4), %edi
	xorl	%esi, %esi
.LBB2_15:                               #   Parent Loop BB2_9 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movslq	%edi, %r9
	movl	(%rcx,%rsi,4), %r8d
	cmpl	%r8d, (%rcx,%r9,4)
	jge	.LBB2_18
# %bb.16:                               #   in Loop: Header=BB2_15 Depth=2
	movslq	-4(%rax,%rdx,4), %r9
	cmpl	%r8d, (%rcx,%r9,4)
	jle	.LBB2_18
# %bb.17:                               #   in Loop: Header=BB2_15 Depth=2
	movl	%esi, (%rax,%rdx,4)
	movl	%esi, %edi
.LBB2_18:                               #   in Loop: Header=BB2_15 Depth=2
	incq	%rsi
	cmpq	$24, %rsi
	jne	.LBB2_15
	jmp	.LBB2_19
.LBB2_10:                               # %.preheader.preheader
                                        #   in Loop: Header=BB2_9 Depth=1
	xorl	%esi, %esi
	movl	%r13d, %edi
.LBB2_11:                               # %.preheader
                                        #   Parent Loop BB2_9 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movslq	%edi, %r8
	movl	(%rcx,%r8,4), %r8d
	cmpl	(%rcx,%rsi,4), %r8d
	jge	.LBB2_13
# %bb.12:                               #   in Loop: Header=BB2_11 Depth=2
	movl	%esi, (%rax)
	movl	%esi, %r13d
	movl	%esi, %edi
.LBB2_13:                               #   in Loop: Header=BB2_11 Depth=2
	incq	%rsi
	cmpq	$24, %rsi
	jne	.LBB2_11
.LBB2_19:                               # %.loopexit
                                        #   in Loop: Header=BB2_9 Depth=1
	incq	%rdx
	cmpq	$5, %rdx
	jne	.LBB2_9
# %bb.20:
	movl	$_ZL13ARQUIVO_DNA_2, %edi
	movl	%ebp, %esi
	callq	_Z10lerArquivoPKci
	movq	%r14, %rdi
	movl	$1, %esi
	movq	%r15, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_22
# %bb.21:
	movq	DNA(%rip), %rdi
	movq	combinacoesDNA(%rip), %rsi
	movq	inicioProcesso(%rip), %rdx
	movq	8(%rsp), %rcx
	movq	16(%rsp), %r8
	movl	%ebx, (%rsp)
	movl	%ebp, %r9d
	callq	_Z34__device_stub__compararEncontradosPcS_PiS0_S0_ii
.LBB2_22:
	callq	hipDeviceSynchronize
	movl	$.Lstr, %edi
	callq	puts@PLT
	movl	$combinacoesDNAtxt, %ebx
	xorl	%r14d, %r14d
.LBB2_23:                               # =>This Inner Loop Header: Depth=1
	movq	countCombinacoesDNA(%rip), %rax
	movl	(%rax,%r14,4), %edx
	movl	$.L.str.3, %edi
	movq	%rbx, %rsi
	xorl	%eax, %eax
	callq	printf
	incq	%r14
	addq	$5, %rbx
	cmpq	$24, %r14
	jne	.LBB2_23
# %bb.24:
	movl	$.Lstr.1, %edi
	callq	puts@PLT
	movl	$.Lstr.2, %edi
	callq	puts@PLT
	xorl	%ebx, %ebx
.LBB2_25:                               # =>This Inner Loop Header: Depth=1
	movq	8(%rsp), %rax
	movslq	(%rax,%rbx,4), %rax
	leaq	(%rax,%rax,4), %rsi
	addq	$combinacoesDNAtxt, %rsi
	movq	countCombinacoesDNA(%rip), %rcx
	movl	(%rcx,%rax,4), %edx
	movq	16(%rsp), %rax
	movl	(%rax,%rbx,4), %ecx
	movl	$.L.str.6, %edi
	xorl	%eax, %eax
	callq	printf
	incq	%rbx
	cmpq	$5, %rbx
	jne	.LBB2_25
# %bb.26:
	movl	$.Lstr.3, %edi
	callq	puts@PLT
	movq	DNA(%rip), %rdi
	callq	hipFree
	movq	countCombinacoesDNA(%rip), %rdi
	callq	hipFree
	movq	combinacoesDNA(%rip), %rdi
	callq	hipFree
	movq	inicioProcesso(%rip), %rdi
	callq	hipFree
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	xorl	%eax, %eax
	jmp	.LBB2_27
.LBB2_2:
	movl	$.Lstr.4, %edi
	callq	puts@PLT
	movl	$1, %eax
.LBB2_27:
	addq	$24, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end2:
	.size	main, .Lfunc_end2-main
	.cfi_endproc
                                        # -- End function
	.globl	_Z14tamanhoArquivoPKc           # -- Begin function _Z14tamanhoArquivoPKc
	.type	_Z14tamanhoArquivoPKc,@function
_Z14tamanhoArquivoPKc:                  # @_Z14tamanhoArquivoPKc
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	pushq	%rax
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	movl	$.L.str.8, %esi
	callq	fopen
	movq	%rax, %rbx
	movq	%rax, %rdi
	xorl	%esi, %esi
	movl	$2, %edx
	callq	fseek
	movq	%rbx, %rdi
	callq	ftell
	movq	%rax, %r14
	movq	%rbx, %rdi
	callq	fclose
	movl	%r14d, %eax
	addq	$8, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end3:
	.size	_Z14tamanhoArquivoPKc, .Lfunc_end3-_Z14tamanhoArquivoPKc
	.cfi_endproc
                                        # -- End function
	.globl	_Z10lerArquivoPKci              # -- Begin function _Z10lerArquivoPKci
	.type	_Z10lerArquivoPKci,@function
_Z10lerArquivoPKci:                     # @_Z10lerArquivoPKci
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	pushq	%rax
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%esi, %ebp
	movl	$.L.str.8, %esi
	callq	fopen
	movq	%rax, %rbx
	testl	%ebp, %ebp
	jle	.LBB4_3
# %bb.1:                                # %.lr.ph.preheader
	movl	%ebp, %r14d
	xorl	%r15d, %r15d
.LBB4_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movq	DNA(%rip), %rdx
	addq	%r15, %rdx
	movl	$.L.str.9, %esi
	movq	%rbx, %rdi
	xorl	%eax, %eax
	callq	__isoc23_fscanf
	incq	%r15
	cmpq	%r15, %r14
	jne	.LBB4_2
.LBB4_3:                                # %._crit_edge
	movq	%rbx, %rdi
	addq	$8, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	fclose                          # TAILCALL
.Lfunc_end4:
	.size	_Z10lerArquivoPKci, .Lfunc_end4-_Z10lerArquivoPKci
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB5_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB5_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z9pesquisarPcS_PiS0_ij, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z19compararEncontradosPcS_PiS0_S0_ii, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end5:
	.size	__hip_module_ctor, .Lfunc_end5-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB6_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB6_2:
	retq
.Lfunc_end6:
	.size	__hip_module_dtor, .Lfunc_end6-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	DNA,@object                     # @DNA
	.bss
	.globl	DNA
	.p2align	3, 0x0
DNA:
	.quad	0
	.size	DNA, 8

	.type	combinacoesDNA,@object          # @combinacoesDNA
	.globl	combinacoesDNA
	.p2align	3, 0x0
combinacoesDNA:
	.quad	0
	.size	combinacoesDNA, 8

	.type	inicioProcesso,@object          # @inicioProcesso
	.globl	inicioProcesso
	.p2align	3, 0x0
inicioProcesso:
	.quad	0
	.size	inicioProcesso, 8

	.type	countCombinacoesDNA,@object     # @countCombinacoesDNA
	.globl	countCombinacoesDNA
	.p2align	3, 0x0
countCombinacoesDNA:
	.quad	0
	.size	countCombinacoesDNA, 8

	.type	combinacoesDNAtxt,@object       # @combinacoesDNAtxt
	.data
	.globl	combinacoesDNAtxt
	.p2align	4, 0x0
combinacoesDNAtxt:
	.asciz	"ACGT"
	.asciz	"ACTG"
	.asciz	"AGCT"
	.asciz	"AGTC"
	.asciz	"ATCG"
	.asciz	"ATGC"
	.asciz	"CAGT"
	.asciz	"CATG"
	.asciz	"CGAT"
	.asciz	"CGTA"
	.asciz	"CTAG"
	.asciz	"CTGA"
	.asciz	"GACT"
	.asciz	"GATC"
	.asciz	"GCAT"
	.asciz	"GCTA"
	.asciz	"GTAC"
	.asciz	"GTCA"
	.asciz	"TACG"
	.asciz	"TAGC"
	.asciz	"TCAG"
	.asciz	"TCGA"
	.asciz	"TGAC"
	.asciz	"TGCA"
	.size	combinacoesDNAtxt, 120

	.type	_Z9pesquisarPcS_PiS0_ij,@object # @_Z9pesquisarPcS_PiS0_ij
	.section	.rodata,"a",@progbits
	.globl	_Z9pesquisarPcS_PiS0_ij
	.p2align	3, 0x0
_Z9pesquisarPcS_PiS0_ij:
	.quad	_Z24__device_stub__pesquisarPcS_PiS0_ij
	.size	_Z9pesquisarPcS_PiS0_ij, 8

	.type	_Z19compararEncontradosPcS_PiS0_S0_ii,@object # @_Z19compararEncontradosPcS_PiS0_S0_ii
	.globl	_Z19compararEncontradosPcS_PiS0_S0_ii
	.p2align	3, 0x0
_Z19compararEncontradosPcS_PiS0_S0_ii:
	.quad	_Z34__device_stub__compararEncontradosPcS_PiS0_S0_ii
	.size	_Z19compararEncontradosPcS_PiS0_S0_ii, 8

	.type	_ZL13ARQUIVO_DNA_1,@object      # @_ZL13ARQUIVO_DNA_1
_ZL13ARQUIVO_DNA_1:
	.asciz	"src/DNA1.txt"
	.size	_ZL13ARQUIVO_DNA_1, 13

	.type	.L.str.1,@object                # @.str.1
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.1:
	.asciz	"ACGTACTGAGCTAGTCATCGATGCCAGTCATGCGATCGTACTAGCTGAGACTGATCGCATGCTAGTACGTCATACGTAGCTCAGTCGATGACTGCA"
	.size	.L.str.1, 97

	.type	_ZL13ARQUIVO_DNA_2,@object      # @_ZL13ARQUIVO_DNA_2
	.section	.rodata,"a",@progbits
_ZL13ARQUIVO_DNA_2:
	.asciz	"src/DNA2.txt"
	.size	_ZL13ARQUIVO_DNA_2, 13

	.type	.L.str.3,@object                # @.str.3
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.3:
	.asciz	"\n %-9s | %-18d |"
	.size	.L.str.3, 17

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"\n %-9s | %-7d | %d"
	.size	.L.str.6, 19

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"r"
	.size	.L.str.8, 2

	.type	.L.str.9,@object                # @.str.9
.L.str.9:
	.asciz	"%1c"
	.size	.L.str.9, 4

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z9pesquisarPcS_PiS0_ij"
	.size	.L__unnamed_1, 24

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z19compararEncontradosPcS_PiS0_S0_ii"
	.size	.L__unnamed_2, 38

	.type	__hip_fatbin_wrapper,@object    # @__hip_fatbin_wrapper
	.section	.hipFatBinSegment,"a",@progbits
	.p2align	3, 0x0
__hip_fatbin_wrapper:
	.long	1212764230                      # 0x48495046
	.long	1                               # 0x1
	.quad	__hip_fatbin
	.quad	0
	.size	__hip_fatbin_wrapper, 24

	.type	__hip_gpubin_handle,@object     # @__hip_gpubin_handle
	.local	__hip_gpubin_handle
	.comm	__hip_gpubin_handle,8,8
	.section	.init_array,"aw",@init_array
	.p2align	3, 0x0
	.quad	__hip_module_ctor
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.type	.Lstr,@object                   # @str
	.section	.rodata.str1.1,"aMS",@progbits,1
.Lstr:
	.asciz	"\n Sequ\303\252ncia | Encontrado no DNA1 |"
	.size	.Lstr, 36

	.type	.Lstr.1,@object                 # @str.1
.Lstr.1:
	.asciz	"\n\n\nAs 5 sequ\303\252ncias mais encontradas no DNA1 presentes no DNA2"
	.size	.Lstr.1, 63

	.type	.Lstr.2,@object                 # @str.2
.Lstr.2:
	.asciz	"\n Sequ\303\252ncia | DNA1    | DNA2"
	.size	.Lstr.2, 30

	.type	.Lstr.3,@object                 # @str.3
.Lstr.3:
	.asciz	"\n"
	.size	.Lstr.3, 2

	.type	.Lstr.4,@object                 # @str.4
.Lstr.4:
	.asciz	"\n\nInsira o n\303\272mero de blocos e threads! \n"
	.size	.Lstr.4, 42

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym _Z24__device_stub__pesquisarPcS_PiS0_ij
	.addrsig_sym _Z34__device_stub__compararEncontradosPcS_PiS0_S0_ii
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym DNA
	.addrsig_sym combinacoesDNA
	.addrsig_sym inicioProcesso
	.addrsig_sym countCombinacoesDNA
	.addrsig_sym combinacoesDNAtxt
	.addrsig_sym _Z9pesquisarPcS_PiS0_ij
	.addrsig_sym _Z19compararEncontradosPcS_PiS0_S0_ii
	.addrsig_sym _ZL13ARQUIVO_DNA_1
	.addrsig_sym _ZL13ARQUIVO_DNA_2
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z9pesquisarPcS_PiS0_ij ; -- Begin function _Z9pesquisarPcS_PiS0_ij
	.globl	_Z9pesquisarPcS_PiS0_ij
	.p2align	8
	.type	_Z9pesquisarPcS_PiS0_ij,@function
_Z9pesquisarPcS_PiS0_ij:                ; @_Z9pesquisarPcS_PiS0_ij
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x34
	s_load_b32 s3, s[0:1], 0x24
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_mov_b32 s2, exec_lo
	v_cmpx_gt_u32_e64 s3, v1
	s_cbranch_execz .LBB0_16
; %bb.1:
	s_load_b256 s[4:11], s[0:1], 0x0
	v_ashrrev_i32_e32 v2, 31, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 2, v[1:2]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v0, vcc_lo, s8, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s9, v1, vcc_lo
	global_load_b64 v[2:3], v[0:1], off
	s_waitcnt vmcnt(0)
	v_cmp_lt_i32_e32 vcc_lo, v2, v3
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_16
; %bb.2:                                ; %.lr.ph.preheader
	v_and_b32_e32 v4, 0xff, v0
	s_mov_b32 s2, 0
                                        ; implicit-def: $vgpr11
	s_delay_alu instid0(VALU_DEP_1)
	v_dual_mov_b32 v10, 1 :: v_dual_and_b32 v5, 0xffff, v4
.LBB0_3:                                ; %.lr.ph
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_4 Depth 2
                                        ;       Child Loop BB0_5 Depth 3
	v_ashrrev_i32_e32 v4, 31, v2
	v_add_co_u32 v6, vcc_lo, s4, v2
	s_mov_b32 s3, 0
	s_mov_b32 s12, 0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v7, vcc_lo, s5, v4, vcc_lo
                                        ; implicit-def: $sgpr8
                                        ; implicit-def: $sgpr9
	global_load_b32 v4, v[6:7], off
	s_waitcnt vmcnt(0)
	v_lshrrev_b32_e32 v6, 8, v4
	v_lshrrev_b32_e32 v7, 16, v4
	v_lshrrev_b32_e32 v8, 24, v4
	v_and_b32_e32 v4, 0xff, v4
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_lshlrev_b16 v6, 8, v6
	v_and_b32_e32 v7, 0xff, v7
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshlrev_b16 v8, 8, v8
	v_or_b32_e32 v4, v4, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_or_b32_e32 v6, v7, v8
	v_and_b32_e32 v4, 0xffff, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b32_e32 v6, 16, v6
	v_or_b32_e32 v4, v4, v6
.LBB0_4:                                ;   Parent Loop BB0_3 Depth=1
                                        ; =>  This Loop Header: Depth=2
                                        ;       Child Loop BB0_5 Depth 3
	v_mov_b32_e32 v6, s12
	s_mov_b32 s13, 0
	global_load_b32 v6, v6, s[6:7]
	s_waitcnt vmcnt(0)
	v_lshrrev_b32_e32 v7, 8, v6
	v_lshrrev_b32_e32 v8, 16, v6
	v_lshrrev_b32_e32 v9, 24, v6
	v_and_b32_e32 v6, 0xff, v6
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_lshlrev_b16 v7, 8, v7
	v_and_b32_e32 v8, 0xff, v8
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshlrev_b16 v9, 8, v9
	v_or_b32_e32 v6, v6, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_4)
	v_or_b32_e32 v7, v8, v9
	v_and_b32_e32 v8, 0xff, v11
	v_mov_b32_e32 v12, 0
	v_and_b32_e32 v6, 0xffff, v6
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b32_e32 v7, 16, v7
	v_or_b32_e32 v6, v6, v7
	v_and_b32_e32 v7, 0xffff, v8
.LBB0_5:                                ;   Parent Loop BB0_3 Depth=1
                                        ;     Parent Loop BB0_4 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	v_lshlrev_b32_e32 v15, 3, v12
                                        ; implicit-def: $vgpr14
                                        ; implicit-def: $vgpr13
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshrrev_b64 v[8:9], v15, v[4:5]
	v_and_b32_e32 v9, 0xff, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	v_cmp_ne_u16_e32 vcc_lo, 0, v9
	v_add_nc_u32_e32 v9, 1, v12
	s_and_saveexec_b32 s0, vcc_lo
	s_xor_b32 s14, exec_lo, s0
	s_cbranch_execz .LBB0_7
; %bb.6:                                ;   in Loop: Header=BB0_5 Depth=3
	v_lshrrev_b64 v[13:14], v15, v[6:7]
	v_bfe_i32 v8, v8, 0, 8
	v_add_nc_u32_e32 v9, 1, v12
	v_not_b32_e32 v16, v12
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_and_b32_e32 v15, 0xff, v8
	v_and_b32_e32 v14, 0xff, v13
	v_bfe_i32 v13, v13, 0, 8
	v_cmp_eq_u16_e32 vcc_lo, 0, v14
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_cmp_lt_i16_e64 s1, v8, v13
	v_cmp_eq_u16_e64 s0, v15, v14
	v_cndmask_b32_e64 v14, 0, 1, vcc_lo
	v_cndmask_b32_e64 v8, v9, v16, s1
	s_delay_alu instid0(VALU_DEP_3)
	s_or_b32 s0, vcc_lo, s0
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_cndmask_b32_e64 v13, v8, 0, s0
.LBB0_7:                                ; %Flow
                                        ;   in Loop: Header=BB0_5 Depth=3
	s_and_not1_saveexec_b32 s0, s14
; %bb.8:                                ; %._crit_edge
                                        ;   in Loop: Header=BB0_5 Depth=3
	v_dual_mov_b32 v13, 0 :: v_dual_mov_b32 v14, 1
; %bb.9:                                ;   in Loop: Header=BB0_5 Depth=3
	s_or_b32 exec_lo, exec_lo, s0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_or_b32_e32 v8, v14, v13
	v_cmp_lt_u32_e32 vcc_lo, 3, v12
	v_mov_b32_e32 v12, v9
	v_cmp_ne_u32_e64 s0, 0, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s0, vcc_lo, s0
	s_and_b32 s0, exec_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s13, s0, s13
	s_and_not1_b32 exec_lo, exec_lo, s13
	s_cbranch_execnz .LBB0_5
; %bb.10:                               ; %_Z9my_strcmpPKcS0_j.exit
                                        ;   in Loop: Header=BB0_4 Depth=2
	s_or_b32 exec_lo, exec_lo, s13
	s_mov_b32 s1, -1
	s_or_b32 s9, s9, exec_lo
	s_mov_b32 s13, exec_lo
                                        ; implicit-def: $sgpr0
	v_cmpx_ne_u32_e32 0, v13
; %bb.11:                               ;   in Loop: Header=BB0_4 Depth=2
	s_add_i32 s0, s12, 4
	s_cmpk_gt_u32 s12, 0x5b
	s_cselect_b32 s1, -1, 0
	s_and_not1_b32 s9, s9, exec_lo
	s_or_not1_b32 s1, s1, exec_lo
; %bb.12:                               ; %Flow74
                                        ;   in Loop: Header=BB0_4 Depth=2
	s_or_b32 exec_lo, exec_lo, s13
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_b32 s1, exec_lo, s1
	v_mov_b32_e32 v6, s12
	s_or_b32 s3, s1, s3
	s_and_not1_b32 s1, s8, exec_lo
	s_and_b32 s8, s9, exec_lo
	s_mov_b32 s12, s0
	s_or_b32 s8, s1, s8
	s_and_not1_b32 exec_lo, exec_lo, s3
	s_cbranch_execnz .LBB0_4
; %bb.13:                               ; %loop.exit.guard
                                        ;   in Loop: Header=BB0_3 Depth=1
	s_or_b32 exec_lo, exec_lo, s3
	s_and_saveexec_b32 s0, s8
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s0, exec_lo, s0
	s_cbranch_execz .LBB0_15
; %bb.14:                               ;   in Loop: Header=BB0_3 Depth=1
	global_atomic_add_u32 v6, v10, s[10:11]
	global_load_b32 v3, v[0:1], off offset:4
.LBB0_15:                               ; %.loopexit
                                        ;   in Loop: Header=BB0_3 Depth=1
	s_or_b32 exec_lo, exec_lo, s0
	v_add_nc_u32_e32 v2, 1, v2
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_cmp_ge_i32_e32 vcc_lo, v2, v3
	s_or_b32 s2, vcc_lo, s2
	s_and_not1_b32 exec_lo, exec_lo, s2
	s_cbranch_execnz .LBB0_3
.LBB0_16:                               ; %.loopexit37
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z9pesquisarPcS_PiS0_ij
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 296
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 17
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z9pesquisarPcS_PiS0_ij, .Lfunc_end0-_Z9pesquisarPcS_PiS0_ij
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 796
; NumSgprs: 18
; NumVgprs: 17
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 2
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 17
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.protected	_Z19compararEncontradosPcS_PiS0_S0_ii ; -- Begin function _Z19compararEncontradosPcS_PiS0_S0_ii
	.globl	_Z19compararEncontradosPcS_PiS0_S0_ii
	.p2align	8
	.type	_Z19compararEncontradosPcS_PiS0_S0_ii,@function
_Z19compararEncontradosPcS_PiS0_S0_ii:  ; @_Z19compararEncontradosPcS_PiS0_S0_ii
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x3c
	s_load_b32 s3, s[0:1], 0x2c
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_mov_b32 s2, exec_lo
	v_cmpx_gt_i32_e64 s3, v1
	s_cbranch_execz .LBB1_16
; %bb.1:
	s_load_b256 s[4:11], s[0:1], 0x0
	v_ashrrev_i32_e32 v2, 31, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 2, v[1:2]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v0, vcc_lo, s8, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s9, v1, vcc_lo
	global_load_b64 v[2:3], v[0:1], off
	s_waitcnt vmcnt(0)
	v_cmp_lt_i32_e32 vcc_lo, v2, v3
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB1_16
; %bb.2:                                ; %.lr.ph.preheader
	s_load_b64 s[2:3], s[0:1], 0x20
	v_and_b32_e32 v4, 0xff, v0
	s_mov_b32 s9, 0
	s_mov_b32 s12, 0
                                        ; implicit-def: $vgpr11
	s_delay_alu instid0(VALU_DEP_1)
	v_dual_mov_b32 v10, 1 :: v_dual_and_b32 v5, 0xffff, v4
.LBB1_3:                                ; %.lr.ph
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB1_4 Depth 2
                                        ;       Child Loop BB1_5 Depth 3
	v_ashrrev_i32_e32 v4, 31, v2
	v_add_co_u32 v6, vcc_lo, s4, v2
	s_mov_b32 s13, s9
	s_mov_b32 s16, s9
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v7, vcc_lo, s5, v4, vcc_lo
                                        ; implicit-def: $sgpr14
                                        ; implicit-def: $sgpr15
	global_load_b32 v4, v[6:7], off
	s_waitcnt vmcnt(0)
	v_lshrrev_b32_e32 v6, 8, v4
	v_lshrrev_b32_e32 v7, 16, v4
	v_lshrrev_b32_e32 v8, 24, v4
	v_and_b32_e32 v4, 0xff, v4
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_lshlrev_b16 v6, 8, v6
	v_and_b32_e32 v7, 0xff, v7
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshlrev_b16 v8, 8, v8
	v_or_b32_e32 v4, v4, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_or_b32_e32 v6, v7, v8
	v_and_b32_e32 v4, 0xffff, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b32_e32 v6, 16, v6
	v_or_b32_e32 v4, v4, v6
.LBB1_4:                                ;   Parent Loop BB1_3 Depth=1
                                        ; =>  This Loop Header: Depth=2
                                        ;       Child Loop BB1_5 Depth 3
	v_mov_b32_e32 v6, s16
	s_lshr_b32 s8, s16, 2
	s_mov_b32 s17, 0
	global_load_b32 v6, v6, s[10:11]
	s_waitcnt vmcnt(0)
	v_lshlrev_b32_e32 v6, 2, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v7, 31, v6
	v_add_co_u32 v6, vcc_lo, s6, v6
	v_add_co_ci_u32_e32 v7, vcc_lo, s7, v7, vcc_lo
	global_load_b32 v6, v[6:7], off
	s_waitcnt vmcnt(0)
	v_lshrrev_b32_e32 v7, 8, v6
	v_lshrrev_b32_e32 v8, 16, v6
	v_lshrrev_b32_e32 v9, 24, v6
	v_and_b32_e32 v6, 0xff, v6
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_lshlrev_b16 v7, 8, v7
	v_and_b32_e32 v8, 0xff, v8
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshlrev_b16 v9, 8, v9
	v_or_b32_e32 v6, v6, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_4)
	v_or_b32_e32 v7, v8, v9
	v_and_b32_e32 v8, 0xff, v11
	v_mov_b32_e32 v12, 0
	v_and_b32_e32 v6, 0xffff, v6
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b32_e32 v7, 16, v7
	v_or_b32_e32 v6, v6, v7
	v_and_b32_e32 v7, 0xffff, v8
.LBB1_5:                                ;   Parent Loop BB1_3 Depth=1
                                        ;     Parent Loop BB1_4 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	v_lshlrev_b32_e32 v15, 3, v12
                                        ; implicit-def: $vgpr14
                                        ; implicit-def: $vgpr13
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshrrev_b64 v[8:9], v15, v[4:5]
	v_and_b32_e32 v9, 0xff, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	v_cmp_ne_u16_e32 vcc_lo, 0, v9
	v_add_nc_u32_e32 v9, 1, v12
	s_and_saveexec_b32 s0, vcc_lo
	s_xor_b32 s18, exec_lo, s0
	s_cbranch_execz .LBB1_7
; %bb.6:                                ;   in Loop: Header=BB1_5 Depth=3
	v_lshrrev_b64 v[13:14], v15, v[6:7]
	v_bfe_i32 v8, v8, 0, 8
	v_add_nc_u32_e32 v9, 1, v12
	v_not_b32_e32 v16, v12
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_and_b32_e32 v15, 0xff, v8
	v_and_b32_e32 v14, 0xff, v13
	v_bfe_i32 v13, v13, 0, 8
	v_cmp_eq_u16_e32 vcc_lo, 0, v14
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_cmp_lt_i16_e64 s1, v8, v13
	v_cmp_eq_u16_e64 s0, v15, v14
	v_cndmask_b32_e64 v14, 0, 1, vcc_lo
	v_cndmask_b32_e64 v8, v9, v16, s1
	s_delay_alu instid0(VALU_DEP_3)
	s_or_b32 s0, vcc_lo, s0
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_cndmask_b32_e64 v13, v8, 0, s0
.LBB1_7:                                ; %Flow
                                        ;   in Loop: Header=BB1_5 Depth=3
	s_and_not1_saveexec_b32 s0, s18
; %bb.8:                                ; %._crit_edge
                                        ;   in Loop: Header=BB1_5 Depth=3
	v_dual_mov_b32 v13, 0 :: v_dual_mov_b32 v14, 1
; %bb.9:                                ;   in Loop: Header=BB1_5 Depth=3
	s_or_b32 exec_lo, exec_lo, s0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_or_b32_e32 v8, v14, v13
	v_cmp_lt_u32_e32 vcc_lo, 3, v12
	v_mov_b32_e32 v12, v9
	v_cmp_ne_u32_e64 s0, 0, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s0, vcc_lo, s0
	s_and_b32 s0, exec_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s17, s0, s17
	s_and_not1_b32 exec_lo, exec_lo, s17
	s_cbranch_execnz .LBB1_5
; %bb.10:                               ; %_Z9my_strcmpPKcS0_j.exit
                                        ;   in Loop: Header=BB1_4 Depth=2
	s_or_b32 exec_lo, exec_lo, s17
	s_mov_b32 s0, -1
	s_or_b32 s15, s15, exec_lo
	s_mov_b32 s1, exec_lo
	v_cmpx_ne_u32_e32 0, v13
; %bb.11:                               ;   in Loop: Header=BB1_4 Depth=2
	s_add_i32 s17, s16, 4
	s_cmp_gt_u32 s16, 12
	s_mov_b32 s16, s17
	s_cselect_b32 s0, -1, 0
	s_and_not1_b32 s15, s15, exec_lo
	s_or_not1_b32 s0, s0, exec_lo
; %bb.12:                               ; %Flow88
                                        ;   in Loop: Header=BB1_4 Depth=2
	s_or_b32 exec_lo, exec_lo, s1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_4) | instid1(SALU_CYCLE_1)
	s_and_b32 s0, exec_lo, s0
	v_dual_mov_b32 v6, s8 :: v_dual_mov_b32 v7, s9
	s_or_b32 s13, s0, s13
	s_and_not1_b32 s0, s14, exec_lo
	s_and_b32 s1, s15, exec_lo
	s_or_b32 s14, s0, s1
	s_and_not1_b32 exec_lo, exec_lo, s13
	s_cbranch_execnz .LBB1_4
; %bb.13:                               ; %loop.exit.guard
                                        ;   in Loop: Header=BB1_3 Depth=1
	s_or_b32 exec_lo, exec_lo, s13
	s_and_saveexec_b32 s0, s14
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s0, exec_lo, s0
	s_cbranch_execz .LBB1_15
; %bb.14:                               ;   in Loop: Header=BB1_3 Depth=1
	v_lshlrev_b32_e32 v3, 2, v6
	s_waitcnt lgkmcnt(0)
	global_atomic_add_u32 v3, v10, s[2:3]
	global_load_b32 v3, v[0:1], off offset:4
.LBB1_15:                               ; %.loopexit
                                        ;   in Loop: Header=BB1_3 Depth=1
	s_or_b32 exec_lo, exec_lo, s0
	v_add_nc_u32_e32 v2, 1, v2
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_cmp_ge_i32_e32 vcc_lo, v2, v3
	s_or_b32 s12, vcc_lo, s12
	s_and_not1_b32 exec_lo, exec_lo, s12
	s_cbranch_execnz .LBB1_3
.LBB1_16:                               ; %.loopexit42
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z19compararEncontradosPcS_PiS0_S0_ii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 304
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 17
		.amdhsa_next_free_sgpr 19
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end1:
	.size	_Z19compararEncontradosPcS_PiS0_S0_ii, .Lfunc_end1-_Z19compararEncontradosPcS_PiS0_S0_ii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 860
; NumSgprs: 21
; NumVgprs: 17
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 2
; NumSGPRsForWavesPerEU: 21
; NumVGPRsForWavesPerEU: 17
; Occupancy: 16
; WaveLimiterHint : 1
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .offset:         32
        .size:           4
        .value_kind:     by_value
      - .offset:         36
        .size:           4
        .value_kind:     by_value
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         44
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         48
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         52
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         54
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         56
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         58
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         60
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         62
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         104
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 296
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z9pesquisarPcS_PiS0_ij
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z9pesquisarPcS_PiS0_ij.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     17
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         32
        .size:           8
        .value_kind:     global_buffer
      - .offset:         40
        .size:           4
        .value_kind:     by_value
      - .offset:         44
        .size:           4
        .value_kind:     by_value
      - .offset:         48
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         52
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         56
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         60
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         62
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         64
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         66
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         68
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         70
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         104
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         112
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 304
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z19compararEncontradosPcS_PiS0_S0_ii
    .private_segment_fixed_size: 0
    .sgpr_count:     21
    .sgpr_spill_count: 0
    .symbol:         _Z19compararEncontradosPcS_PiS0_S0_ii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     17
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
