#include <cuda.h>
#include <cuda_fp16.h>

#include <stdio.h>
#include <stdint.h>
#include <zlib.h>
#include <stdlib.h>
#include <time.h>
#include <math.h>
#include <fcntl.h>
#include <limits.h>

typedef struct __ReadSeqList {
	char* sequence;
	unsigned int length;
	struct __ReadSeqList* next;
} ReadSeqList;

typedef struct HashTable {
	unsigned int bits;
	unsigned int count;
	unsigned int read_count;
	unsigned long long int *keys;
    unsigned int *values;
} HashTable;



__device__ const unsigned char seq_nt4_table[256] = { // translate ACGT to 0123
	0, 1, 2, 3,  4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,
	4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,
	4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,
	4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,
	4, 0, 4, 1,  4, 4, 4, 2,  4, 4, 4, 4,  4, 4, 4, 4,
	4, 4, 4, 4,  3, 3, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,
	4, 0, 4, 1,  4, 4, 4, 2,  4, 4, 4, 4,  4, 4, 4, 4,
	4, 4, 4, 4,  3, 3, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,
	4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,
	4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,
	4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,
	4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,
	4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,
	4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,
	4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,
	4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4
};


// funcion para calcular un hash de 64 bits
__device__ unsigned int hash_uint64(unsigned long long int key) {

	key = ~key + (key << 21);
	key = key ^ key >> 24;
	key = (key + (key << 3)) + (key << 8);
	key = key ^ key >> 14;
	key = (key + (key << 2)) + (key << 4);
	key = key ^ key >> 28;
	key = key + (key << 31);
	return (unsigned int)key;
}

HashTable* HashTable_init(unsigned int bits, unsigned int read_count){
    HashTable *ht;
    ht = (HashTable*)calloc(1, sizeof(HashTable));

	ht->read_count = read_count;
	ht->bits = bits;
	ht->count = 0;

    return ht;
}


void HashTable_destory(HashTable *ht) {
	if (!ht) return;
	free(ht);
}


__device__ unsigned int h2b(unsigned int hash, unsigned int bits) {
    return hash * 2654435769U >> (32 - bits);
}

__device__ void hash_insert(HashTable *ht, unsigned long long int kmer) {

	unsigned int iKey, last;
	bool end = false;

    iKey = last = h2b(hash_uint64(kmer), ht->bits);

	while (true)
	{
		unsigned long long int prev = atomicCAS(&(ht->keys[iKey]), 0ULL, kmer);

		if (prev == 0ULL || prev == kmer) {
			atomicAdd(&(ht->values[iKey]), 1U);
			return;
		}

		if(end) return;

		// Collition: Open addressing
		iKey = (iKey + 1U) & ((1U << ht->bits) - 1);

		// loop back
		end = (iKey == last);

	}

}

// insert k-mers in $seq to hash table $ht
__global__ void kernel_count_seq_kmers(HashTable *ht, int k, char **d_reads)
{
	unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x;
	if(tid < ht->read_count) {
        int i, l;
		char *seq = d_reads[tid];

		int len = 100;
        unsigned long long int x[2], mask = (1ULL<<k*2) - 1, shift = (k - 1) * 2;

		for (i = l = 0, x[0] = x[1] = 0; i < len; ++i) {
			int c = seq_nt4_table[(unsigned char)seq[i]];
			if (c < 4) { // not an "N" base
				x[0] = (x[0] << 2 | c) & mask;                  // forward strand
				x[1] = x[1] >> 2 | (unsigned long long int)(3 - c) << shift;  // reverse strand
				if (++l >= k) { // we find a k-mer

					unsigned long long int kmer = x[0] < x[1]? x[0] : x[1];
					hash_insert(ht, kmer); // only add one strand!
				}
			} else l = 0, x[0] = x[1] = 0; // if there is an "N", restart
		}
    }
}

__global__ void kernel_print_hist(const HashTable *ht, unsigned int *cnt_d)
{
	unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x;
	unsigned int pos;

	if(tid < (1U << ht->bits)) {
		if (ht->values[tid] > 0) {
			pos = ht->values[tid] < 256U ? ht->values[tid] : 255;
			atomicAdd(&(cnt_d[pos]), 1U);
		}
	}
}

static int count_file(const char *fn, int k, unsigned int p)
{
	HashTable *ht;
    unsigned int capacity = 1U << p;
    unsigned int cnt[256];
    unsigned int read_count = 0;


    // variables para cuda
	HashTable *ht_d;
	char **reads_d;
	unsigned int *cnt_d;

    FILE * fp;
    char * line = NULL;
    size_t len = 0;
    ssize_t read;

    fp = fopen(fn, "r");
    if (fp == NULL) exit(EXIT_FAILURE);

   
	ReadSeqList *current, *head;
	head = current = NULL;

    while ((read = getline(&line, &len, fp)) != -1) {

        read_count++;

		ReadSeqList *node = (ReadSeqList*)malloc(sizeof(ReadSeqList));
        node->sequence = (char*)malloc(strlen(line));
        strcpy(node->sequence, line);
        node->length = read;
        node->next =NULL;

        if(head == NULL){
            current = head = node;
        } else {
            current = current->next = node;
        }


    }

    fclose(fp);
    if (line) free(line);

    printf("%d\n", read_count);
    unsigned int i;

    char **reads = (char**)malloc(read_count * sizeof(char*)); 

	for(i=0, current = head; current; current=current->next){
        reads[i] = (char*)malloc(current->length);
        sprintf(reads[i], "%s", current->sequence);
        i++;
    }
   

    // inicializar hashtable
	ht = HashTable_init(p, read_count);

    
    unsigned long long int *keys_d;
    unsigned int *values_d;


	// allocate memory in device
	cudaMalloc((void **)&ht_d, sizeof(HashTable));
   	cudaMalloc((void **)&keys_d, capacity * sizeof(unsigned long long int));
	cudaMalloc((void **)&values_d, capacity * sizeof(unsigned int));
	cudaMalloc((void **)&cnt_d, 256 * sizeof(unsigned int));
   	cudaMemset(keys_d, 0ULL, capacity * sizeof(unsigned long long int));
	cudaMemset(values_d, 0, capacity * sizeof(unsigned int));
	cudaMemset(cnt_d, 0, 256 * sizeof(unsigned int));

    
	// copy data to device

    ht->keys = keys_d;
    ht->values = values_d;

	cudaMemcpy(ht_d, ht, sizeof(HashTable), cudaMemcpyHostToDevice);

	char **tmp = (char**)malloc (read_count * sizeof (char*));
    for (int i = 0; i < read_count; i++) {
        cudaMalloc ((void **)&tmp[i], head->length * sizeof (char));
    }

	cudaMalloc((void **)&reads_d, read_count * sizeof(char*));

    cudaMemcpy(reads_d, tmp, read_count * sizeof (char*), cudaMemcpyHostToDevice);
    for (int i = 0; i < read_count; i++) {
        cudaMemcpy(tmp[i], reads[i], head->length * sizeof (char), cudaMemcpyHostToDevice);
    }
    free(tmp);



    printf("total reads: %d\n", read_count);


    // invocar kernels
    unsigned int thr = 1024;


    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);

    cudaEventRecord(start);

	kernel_count_seq_kmers<<<ceil(read_count/thr), thr>>>(ht_d, k, reads_d);
    cudaDeviceSynchronize();

	kernel_print_hist<<<ceil(capacity/thr), thr>>>(ht_d, cnt_d);
    cudaDeviceSynchronize();
    
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    
    float milliseconds = 0;
    cudaEventElapsedTime(&milliseconds, start, stop);
    float seconds = milliseconds / 1000.0f;
    printf("GPU time: %fs\n", seconds);

	cudaMemcpy(ht, ht_d, sizeof(HashTable), cudaMemcpyDeviceToHost);
	cudaMemcpy(ht->keys, keys_d, capacity * sizeof(unsigned long long int), cudaMemcpyDeviceToHost);
	cudaMemcpy(ht->values, values_d, capacity * sizeof(unsigned int), cudaMemcpyDeviceToHost);
	cudaMemcpy(cnt, cnt_d, 256 * sizeof(unsigned int), cudaMemcpyDeviceToHost);


    printf("COUNT: %d\n\n", ht->count);

	for (i = 1; i < 256; ++i)
		printf("%d\t%d\n", i, cnt[i]);


	// limpieza
    cudaFree(reads_d);
	cudaFree(ht_d);
	cudaFree(cnt_d);
	cudaFree(keys_d);
	cudaFree(values_d);

	// limpieza
    i = 0;
	for(current = head; current; current=current->next){
        free(current->sequence);
        free(current);
        free(reads[i]);
        i++;
    }

    free(reads);
	HashTable_destory(ht);
    return 0;
}


int main(int argc, char *argv[])
{
	int k = 31;
    unsigned int p = 27;

    k = (int)strtol(argv[1], NULL, 10);
    p = (unsigned int)strtol(argv[2], NULL, 10);
	count_file(argv[3], k, p);

	return 0;
}

	.file	"tmpxft_003671dc_00000000-6_cuda-slow.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2427:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2427:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z11hash_uint64y
	.type	_Z11hash_uint64y, @function
_Z11hash_uint64y:
.LFB2418:
	.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
.LFE2418:
	.size	_Z11hash_uint64y, .-_Z11hash_uint64y
	.globl	_Z14HashTable_initjj
	.type	_Z14HashTable_initjj, @function
_Z14HashTable_initjj:
.LFB2419:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movl	%esi, %ebp
	movl	$32, %esi
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movl	%edi, %ebx
	movl	$1, %edi
	pushq	%rax
	.cfi_def_cfa_offset 32
	call	calloc@PLT
	xorl	%edx, %edx
	movl	%ebp, 8(%rax)
	movl	%ebx, (%rax)
	movl	%edx, 4(%rax)
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2419:
	.size	_Z14HashTable_initjj, .-_Z14HashTable_initjj
	.globl	_Z17HashTable_destoryP9HashTable
	.type	_Z17HashTable_destoryP9HashTable, @function
_Z17HashTable_destoryP9HashTable:
.LFB2420:
	.cfi_startproc
	endbr64
	testq	%rdi, %rdi
	je	.L6
	jmp	free@PLT
.L6:
	ret
	.cfi_endproc
.LFE2420:
	.size	_Z17HashTable_destoryP9HashTable, .-_Z17HashTable_destoryP9HashTable
	.globl	_Z3h2bjj
	.type	_Z3h2bjj, @function
_Z3h2bjj:
.LFB2421:
	.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
.LFE2421:
	.size	_Z3h2bjj, .-_Z3h2bjj
	.globl	_Z11hash_insertP9HashTabley
	.type	_Z11hash_insertP9HashTabley, @function
_Z11hash_insertP9HashTabley:
.LFB2422:
	.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
.LFE2422:
	.size	_Z11hash_insertP9HashTabley, .-_Z11hash_insertP9HashTabley
	.globl	_Z55__device_stub__Z22kernel_count_seq_kmersP9HashTableiPPcP9HashTableiPPc
	.type	_Z55__device_stub__Z22kernel_count_seq_kmersP9HashTableiPPcP9HashTableiPPc, @function
_Z55__device_stub__Z22kernel_count_seq_kmersP9HashTableiPPcP9HashTableiPPc:
.LFB2449:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movl	%esi, 20(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L13
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z22kernel_count_seq_kmersP9HashTableiPPc(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L13:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L15
	call	__stack_chk_fail@PLT
.L15:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2449:
	.size	_Z55__device_stub__Z22kernel_count_seq_kmersP9HashTableiPPcP9HashTableiPPc, .-_Z55__device_stub__Z22kernel_count_seq_kmersP9HashTableiPPcP9HashTableiPPc
	.globl	_Z22kernel_count_seq_kmersP9HashTableiPPc
	.type	_Z22kernel_count_seq_kmersP9HashTableiPPc, @function
_Z22kernel_count_seq_kmersP9HashTableiPPc:
.LFB2450:
	.cfi_startproc
	endbr64
	jmp	_Z55__device_stub__Z22kernel_count_seq_kmersP9HashTableiPPcP9HashTableiPPc
	.cfi_endproc
.LFE2450:
	.size	_Z22kernel_count_seq_kmersP9HashTableiPPc, .-_Z22kernel_count_seq_kmersP9HashTableiPPc
	.globl	_Z49__device_stub__Z17kernel_print_histPK9HashTablePjPK9HashTablePj
	.type	_Z49__device_stub__Z17kernel_print_histPK9HashTablePjPK9HashTablePj, @function
_Z49__device_stub__Z17kernel_print_histPK9HashTablePjPK9HashTablePj:
.LFB2451:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	movq	%rsi, (%rsp)
	leaq	40(%rsp), %rdi
	leaq	52(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L18
	pushq	32(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z17kernel_print_histPK9HashTablePj(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 144
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L18:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L20
	call	__stack_chk_fail@PLT
.L20:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2451:
	.size	_Z49__device_stub__Z17kernel_print_histPK9HashTablePjPK9HashTablePj, .-_Z49__device_stub__Z17kernel_print_histPK9HashTablePjPK9HashTablePj
	.globl	_Z17kernel_print_histPK9HashTablePj
	.type	_Z17kernel_print_histPK9HashTablePj, @function
_Z17kernel_print_histPK9HashTablePj:
.LFB2452:
	.cfi_startproc
	endbr64
	jmp	_Z49__device_stub__Z17kernel_print_histPK9HashTablePjPK9HashTablePj
	.cfi_endproc
.LFE2452:
	.size	_Z17kernel_print_histPK9HashTablePj, .-_Z17kernel_print_histPK9HashTablePj
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"r"
.LC1:
	.string	"%d\n"
.LC2:
	.string	"%s"
.LC3:
	.string	"total reads: %d\n"
.LC6:
	.string	"GPU time: %fs\n"
.LC7:
	.string	"COUNT: %d\n\n"
.LC8:
	.string	"%d\t%d\n"
	.text
	.type	_ZL10count_filePKcij.isra.0, @function
_ZL10count_filePKcij.isra.0:
.LFB2572:
	.cfi_startproc
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$1192, %rsp
	.cfi_def_cfa_offset 1248
	movl	%edx, 16(%rsp)
	movb	16(%rsp), %cl
	xorl	%edx, %edx
	movl	%esi, 36(%rsp)
	leaq	.LC0(%rip), %rsi
	movq	%fs:40, %rax
	movq	%rax, 1176(%rsp)
	movl	$1, %eax
	movq	%rdx, 80(%rsp)
	movq	%rdx, 88(%rsp)
	sall	%cl, %eax
	movl	%eax, 32(%rsp)
	call	fopen@PLT
	testq	%rax, %rax
	jne	.L40
	movl	$1, %edi
	call	exit@PLT
.L40:
	movq	%rax, %rbx
	leaq	88(%rsp), %rax
	xorl	%ebp, %ebp
	xorl	%r13d, %r13d
	movq	%rax, 8(%rsp)
	xorl	%r12d, %r12d
.L24:
	movq	8(%rsp), %rsi
	leaq	80(%rsp), %rdi
	movq	%rbx, %rdx
	call	getline@PLT
	movq	%rax, %r14
	cmpq	$-1, %rax
	je	.L47
	movl	$24, %edi
	incl	%r12d
	call	malloc@PLT
	movq	80(%rsp), %rsi
	movq	%rax, %r15
	movq	%rsi, %rdi
	movq	%rsi, 40(%rsp)
	call	strlen@PLT
	movq	%rax, %rdi
	movq	%rax, 24(%rsp)
	call	malloc@PLT
	movq	24(%rsp), %rdx
	movq	40(%rsp), %rsi
	movq	%rax, (%r15)
	movq	%rax, %rdi
	call	__strcpy_chk@PLT
	xorl	%eax, %eax
	movl	%r14d, 8(%r15)
	movq	%rax, 16(%r15)
	testq	%rbp, %rbp
	je	.L41
	movq	%r15, 16(%r13)
	jmp	.L25
.L41:
	movq	%r15, %rbp
.L25:
	movq	%r15, %r13
	jmp	.L24
.L47:
	movq	%rbx, %rdi
	call	fclose@PLT
	movq	80(%rsp), %rdi
	testq	%rdi, %rdi
	je	.L27
	call	free@PLT
.L27:
	movl	%r12d, %edx
	leaq	.LC1(%rip), %rsi
	xorl	%eax, %eax
	movq	%rbp, %r13
	movl	$2, %edi
	xorl	%ebx, %ebx
	leaq	.LC2(%rip), %r15
	call	__printf_chk@PLT
	movl	%r12d, %eax
	salq	$3, %rax
	movq	%rax, %rdi
	movq	%rax, 8(%rsp)
	call	malloc@PLT
	movq	%rax, %r14
.L28:
	testq	%r13, %r13
	je	.L48
	movl	8(%r13), %edx
	movq	%rdx, %rdi
	movq	%rdx, 24(%rsp)
	call	malloc@PLT
	movq	0(%r13), %r8
	movq	24(%rsp), %rdx
	movq	%r15, %rcx
	movq	%rax, %rdi
	movl	%ebx, %eax
	movl	$2, %esi
	incl	%ebx
	movq	%rdi, (%r14,%rax,8)
	xorl	%eax, %eax
	call	__sprintf_chk@PLT
	movq	16(%r13), %r13
	jmp	.L28
.L48:
	movl	16(%rsp), %edi
	movl	%r12d, %esi
	call	_Z14HashTable_initjj
	leaq	56(%rsp), %rdi
	movl	$32, %esi
	movq	%rax, %rbx
	call	cudaMalloc@PLT
	movl	32(%rsp), %r13d
	leaq	96(%rsp), %rdi
	leaq	0(,%r13,8), %rax
	movq	%rax, %rsi
	movq	%rax, 16(%rsp)
	call	cudaMalloc@PLT
	leaq	0(,%r13,4), %rax
	leaq	104(%rsp), %rdi
	movq	%rax, %rsi
	movq	%rax, 24(%rsp)
	call	cudaMalloc@PLT
	leaq	72(%rsp), %rdi
	movl	$1024, %esi
	call	cudaMalloc@PLT
	movq	16(%rsp), %rdx
	movq	96(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaMemset@PLT
	movq	24(%rsp), %rdx
	movq	104(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaMemset@PLT
	movq	72(%rsp), %rdi
	movl	$1024, %edx
	xorl	%esi, %esi
	call	cudaMemset@PLT
	movq	96(%rsp), %rax
	movq	56(%rsp), %rdi
	movq	%rbx, %rsi
	movl	$32, %edx
	movl	$1, %ecx
	movq	%rax, 16(%rbx)
	movq	104(%rsp), %rax
	movq	%rax, 24(%rbx)
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rdi
	call	malloc@PLT
	movq	%rax, %r13
	movq	%rax, %rdx
	movq	8(%rsp), %rax
	leaq	(%rax,%r13), %r15
.L30:
	cmpq	%r15, %rdx
	je	.L49
	movl	8(%rbp), %esi
	movq	%rdx, %rdi
	movq	%rdx, 40(%rsp)
	call	cudaMalloc@PLT
	movq	40(%rsp), %rdx
	addq	$8, %rdx
	jmp	.L30
.L49:
	movq	8(%rsp), %rsi
	leaq	64(%rsp), %rdi
	xorl	%r15d, %r15d
	call	cudaMalloc@PLT
	movq	8(%rsp), %rdx
	movq	64(%rsp), %rdi
	movq	%r13, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
.L32:
	cmpl	%r12d, %r15d
	jnb	.L50
	movq	(%r14,%r15,8), %rsi
	movq	0(%r13,%r15,8), %rdi
	movl	$1, %ecx
	incq	%r15
	movl	8(%rbp), %edx
	call	cudaMemcpy@PLT
	jmp	.L32
.L50:
	movq	%r13, %rdi
	call	free@PLT
	movl	%r12d, %edx
	movl	$2, %edi
	xorl	%eax, %eax
	leaq	.LC3(%rip), %rsi
	shrl	$10, %r12d
	call	__printf_chk@PLT
	leaq	112(%rsp), %rdi
	call	cudaEventCreate@PLT
	leaq	120(%rsp), %rdi
	call	cudaEventCreate@PLT
	movq	112(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movl	$4194305, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movabsq	$4294967297, %rax
	salq	$10, %rdx
	movl	$1, %ecx
	movl	%r12d, 128(%rsp)
	movq	%rax, 132(%rsp)
	movq	128(%rsp), %rdi
	movl	136(%rsp), %esi
	movq	%rdx, 140(%rsp)
	movl	$1, 148(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L34
	movq	64(%rsp), %rdx
	movl	36(%rsp), %esi
	movq	56(%rsp), %rdi
	call	_Z55__device_stub__Z22kernel_count_seq_kmersP9HashTableiPPcP9HashTableiPPc
.L34:
	call	cudaDeviceSynchronize@PLT
	movl	32(%rsp), %eax
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$4194305, %edx
	movl	$1, %ecx
	movl	$1, 148(%rsp)
	shrl	$10, %eax
	salq	$10, %rdx
	movl	%eax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 132(%rsp)
	movq	128(%rsp), %rdi
	movl	136(%rsp), %esi
	movq	%rdx, 140(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L35
	movq	72(%rsp), %rsi
	movq	56(%rsp), %rdi
	call	_Z49__device_stub__Z17kernel_print_histPK9HashTablePjPK9HashTablePj
.L35:
	call	cudaDeviceSynchronize@PLT
	movq	120(%rsp), %rdi
	xorl	%esi, %esi
	leaq	152(%rsp), %r13
	movl	$1, %r12d
	leaq	.LC8(%rip), %r15
	call	cudaEventRecord@PLT
	movq	120(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	120(%rsp), %rdx
	movq	112(%rsp), %rsi
	leaq	140(%rsp), %rdi
	movl	$0x00000000, 140(%rsp)
	call	cudaEventElapsedTime@PLT
	leaq	.LC6(%rip), %rsi
	movl	$2, %edi
	movb	$1, %al
	movss	140(%rsp), %xmm0
	divss	.LC5(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	call	__printf_chk@PLT
	movq	56(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdi
	movl	$32, %edx
	call	cudaMemcpy@PLT
	movq	16(%rbx), %rdi
	movq	16(%rsp), %rdx
	movl	$2, %ecx
	movq	96(%rsp), %rsi
	call	cudaMemcpy@PLT
	movq	24(%rbx), %rdi
	movq	24(%rsp), %rdx
	movl	$2, %ecx
	movq	104(%rsp), %rsi
	call	cudaMemcpy@PLT
	movq	72(%rsp), %rsi
	movl	$2, %ecx
	movq	%r13, %rdi
	movl	$1024, %edx
	call	cudaMemcpy@PLT
	movl	4(%rbx), %edx
	movl	$2, %edi
	xorl	%eax, %eax
	leaq	.LC7(%rip), %rsi
	call	__printf_chk@PLT
.L36:
	movl	0(%r13,%r12,4), %ecx
	movl	%r12d, %edx
	movq	%r15, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	incq	%r12
	call	__printf_chk@PLT
	cmpq	$256, %r12
	jne	.L36
	movq	64(%rsp), %rdi
	xorl	%r12d, %r12d
	call	cudaFree@PLT
	movq	56(%rsp), %rdi
	call	cudaFree@PLT
	movq	72(%rsp), %rdi
	call	cudaFree@PLT
	movq	96(%rsp), %rdi
	call	cudaFree@PLT
	movq	104(%rsp), %rdi
	call	cudaFree@PLT
.L37:
	testq	%rbp, %rbp
	je	.L51
	movq	0(%rbp), %rdi
	call	free@PLT
	movq	%rbp, %rdi
	call	free@PLT
	movl	%r12d, %eax
	incl	%r12d
	movq	(%r14,%rax,8), %rdi
	call	free@PLT
	movq	16(%rbp), %rbp
	jmp	.L37
.L51:
	movq	%r14, %rdi
	call	free@PLT
	movq	%rbx, %rdi
	call	_Z17HashTable_destoryP9HashTable
	movq	1176(%rsp), %rax
	subq	%fs:40, %rax
	je	.L39
	call	__stack_chk_fail@PLT
.L39:
	addq	$1192, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.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
	ret
	.cfi_endproc
.LFE2572:
	.size	_ZL10count_filePKcij.isra.0, .-_ZL10count_filePKcij.isra.0
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2424:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rsi, %rbp
	movl	$10, %edx
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	pushq	%rax
	.cfi_def_cfa_offset 32
	movq	8(%rsi), %rdi
	xorl	%esi, %esi
	call	__isoc23_strtol@PLT
	movq	16(%rbp), %rdi
	movl	$10, %edx
	xorl	%esi, %esi
	movq	%rax, %rbx
	call	__isoc23_strtol@PLT
	movq	24(%rbp), %rdi
	movl	%ebx, %esi
	movl	%eax, %edx
	call	_ZL10count_filePKcij.isra.0
	popq	%rdx
	.cfi_def_cfa_offset 24
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2424:
	.size	main, .-main
	.section	.rodata.str1.1
.LC9:
	.string	"_Z17kernel_print_histPK9HashTablePj"
.LC10:
	.string	"_Z22kernel_count_seq_kmersP9HashTableiPPc"
.LC11:
	.string	"seq_nt4_table"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2454:
	.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	.LC9(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z17kernel_print_histPK9HashTablePj(%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	.LC10(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z22kernel_count_seq_kmersP9HashTableiPPc(%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
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC11(%rip), %rdx
	movl	$256, %r9d
	leaq	_ZL13seq_nt4_table(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	call	__cudaRegisterVar@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.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
.LFE2454:
	.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
	.section	.rodata
	.align 32
	.type	_ZL13seq_nt4_table, @object
	.size	_ZL13seq_nt4_table, 256
_ZL13seq_nt4_table:
	.string	""
	.string	"\001\002\003\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004"
	.string	"\004\001\004\004\004\002\004\004\004\004\004\004\004\004\004\004\004\004\003\003\004\004\004\004\004\004\004\004\004\004\004"
	.ascii	"\004\001\004\004\004\002\004\004\004\004\004\004\004\004\004"
	.ascii	"\004\004\004\003\003\004\004\004\004\004\004\004\004\004\004"
	.ascii	"\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004"
	.ascii	"\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004"
	.ascii	"\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004"
	.ascii	"\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004"
	.ascii	"\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004"
	.ascii	"\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004"
	.ascii	"\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004"
	.ascii	"\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004"
	.ascii	"\004\004\004\004\004\004\004\004"
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC5:
	.long	1148846080
	.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 <hip/hip_runtime.h>
#include <hip/hip_fp16.h>

#include <stdio.h>
#include <stdint.h>
#include <zlib.h>
#include <stdlib.h>
#include <time.h>
#include <math.h>
#include <fcntl.h>
#include <limits.h>

typedef struct __ReadSeqList {
	char* sequence;
	unsigned int length;
	struct __ReadSeqList* next;
} ReadSeqList;

typedef struct HashTable {
	unsigned int bits;
	unsigned int count;
	unsigned int read_count;
	unsigned long long int *keys;
    unsigned int *values;
} HashTable;



__device__ const unsigned char seq_nt4_table[256] = { // translate ACGT to 0123
	0, 1, 2, 3,  4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,
	4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,
	4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,
	4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,
	4, 0, 4, 1,  4, 4, 4, 2,  4, 4, 4, 4,  4, 4, 4, 4,
	4, 4, 4, 4,  3, 3, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,
	4, 0, 4, 1,  4, 4, 4, 2,  4, 4, 4, 4,  4, 4, 4, 4,
	4, 4, 4, 4,  3, 3, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,
	4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,
	4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,
	4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,
	4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,
	4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,
	4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,
	4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,
	4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4,  4, 4, 4, 4
};


// funcion para calcular un hash de 64 bits
__device__ unsigned int hash_uint64(unsigned long long int key) {

	key = ~key + (key << 21);
	key = key ^ key >> 24;
	key = (key + (key << 3)) + (key << 8);
	key = key ^ key >> 14;
	key = (key + (key << 2)) + (key << 4);
	key = key ^ key >> 28;
	key = key + (key << 31);
	return (unsigned int)key;
}

HashTable* HashTable_init(unsigned int bits, unsigned int read_count){
    HashTable *ht;
    ht = (HashTable*)calloc(1, sizeof(HashTable));

	ht->read_count = read_count;
	ht->bits = bits;
	ht->count = 0;

    return ht;
}


void HashTable_destory(HashTable *ht) {
	if (!ht) return;
	free(ht);
}


__device__ unsigned int h2b(unsigned int hash, unsigned int bits) {
    return hash * 2654435769U >> (32 - bits);
}

__device__ void hash_insert(HashTable *ht, unsigned long long int kmer) {

	unsigned int iKey, last;
	bool end = false;

    iKey = last = h2b(hash_uint64(kmer), ht->bits);

	while (true)
	{
		unsigned long long int prev = atomicCAS(&(ht->keys[iKey]), 0ULL, kmer);

		if (prev == 0ULL || prev == kmer) {
			atomicAdd(&(ht->values[iKey]), 1U);
			return;
		}

		if(end) return;

		// Collition: Open addressing
		iKey = (iKey + 1U) & ((1U << ht->bits) - 1);

		// loop back
		end = (iKey == last);

	}

}

// insert k-mers in $seq to hash table $ht
__global__ void kernel_count_seq_kmers(HashTable *ht, int k, char **d_reads)
{
	unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x;
	if(tid < ht->read_count) {
        int i, l;
		char *seq = d_reads[tid];

		int len = 100;
        unsigned long long int x[2], mask = (1ULL<<k*2) - 1, shift = (k - 1) * 2;

		for (i = l = 0, x[0] = x[1] = 0; i < len; ++i) {
			int c = seq_nt4_table[(unsigned char)seq[i]];
			if (c < 4) { // not an "N" base
				x[0] = (x[0] << 2 | c) & mask;                  // forward strand
				x[1] = x[1] >> 2 | (unsigned long long int)(3 - c) << shift;  // reverse strand
				if (++l >= k) { // we find a k-mer

					unsigned long long int kmer = x[0] < x[1]? x[0] : x[1];
					hash_insert(ht, kmer); // only add one strand!
				}
			} else l = 0, x[0] = x[1] = 0; // if there is an "N", restart
		}
    }
}

__global__ void kernel_print_hist(const HashTable *ht, unsigned int *cnt_d)
{
	unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x;
	unsigned int pos;

	if(tid < (1U << ht->bits)) {
		if (ht->values[tid] > 0) {
			pos = ht->values[tid] < 256U ? ht->values[tid] : 255;
			atomicAdd(&(cnt_d[pos]), 1U);
		}
	}
}

static int count_file(const char *fn, int k, unsigned int p)
{
	HashTable *ht;
    unsigned int capacity = 1U << p;
    unsigned int cnt[256];
    unsigned int read_count = 0;


    // variables para cuda
	HashTable *ht_d;
	char **reads_d;
	unsigned int *cnt_d;

    FILE * fp;
    char * line = NULL;
    size_t len = 0;
    ssize_t read;

    fp = fopen(fn, "r");
    if (fp == NULL) exit(EXIT_FAILURE);

   
	ReadSeqList *current, *head;
	head = current = NULL;

    while ((read = getline(&line, &len, fp)) != -1) {

        read_count++;

		ReadSeqList *node = (ReadSeqList*)malloc(sizeof(ReadSeqList));
        node->sequence = (char*)malloc(strlen(line));
        strcpy(node->sequence, line);
        node->length = read;
        node->next =NULL;

        if(head == NULL){
            current = head = node;
        } else {
            current = current->next = node;
        }


    }

    fclose(fp);
    if (line) free(line);

    printf("%d\n", read_count);
    unsigned int i;

    char **reads = (char**)malloc(read_count * sizeof(char*)); 

	for(i=0, current = head; current; current=current->next){
        reads[i] = (char*)malloc(current->length);
        sprintf(reads[i], "%s", current->sequence);
        i++;
    }
   

    // inicializar hashtable
	ht = HashTable_init(p, read_count);

    
    unsigned long long int *keys_d;
    unsigned int *values_d;


	// allocate memory in device
	hipMalloc((void **)&ht_d, sizeof(HashTable));
   	hipMalloc((void **)&keys_d, capacity * sizeof(unsigned long long int));
	hipMalloc((void **)&values_d, capacity * sizeof(unsigned int));
	hipMalloc((void **)&cnt_d, 256 * sizeof(unsigned int));
   	hipMemset(keys_d, 0ULL, capacity * sizeof(unsigned long long int));
	hipMemset(values_d, 0, capacity * sizeof(unsigned int));
	hipMemset(cnt_d, 0, 256 * sizeof(unsigned int));

    
	// copy data to device

    ht->keys = keys_d;
    ht->values = values_d;

	hipMemcpy(ht_d, ht, sizeof(HashTable), hipMemcpyHostToDevice);

	char **tmp = (char**)malloc (read_count * sizeof (char*));
    for (int i = 0; i < read_count; i++) {
        hipMalloc ((void **)&tmp[i], head->length * sizeof (char));
    }

	hipMalloc((void **)&reads_d, read_count * sizeof(char*));

    hipMemcpy(reads_d, tmp, read_count * sizeof (char*), hipMemcpyHostToDevice);
    for (int i = 0; i < read_count; i++) {
        hipMemcpy(tmp[i], reads[i], head->length * sizeof (char), hipMemcpyHostToDevice);
    }
    free(tmp);



    printf("total reads: %d\n", read_count);


    // invocar kernels
    unsigned int thr = 1024;


    hipEvent_t start, stop;
    hipEventCreate(&start);
    hipEventCreate(&stop);

    hipEventRecord(start);

	kernel_count_seq_kmers<<<ceil(read_count/thr), thr>>>(ht_d, k, reads_d);
    hipDeviceSynchronize();

	kernel_print_hist<<<ceil(capacity/thr), thr>>>(ht_d, cnt_d);
    hipDeviceSynchronize();
    
    hipEventRecord(stop);
    hipEventSynchronize(stop);
    
    float milliseconds = 0;
    hipEventElapsedTime(&milliseconds, start, stop);
    float seconds = milliseconds / 1000.0f;
    printf("GPU time: %fs\n", seconds);

	hipMemcpy(ht, ht_d, sizeof(HashTable), hipMemcpyDeviceToHost);
	hipMemcpy(ht->keys, keys_d, capacity * sizeof(unsigned long long int), hipMemcpyDeviceToHost);
	hipMemcpy(ht->values, values_d, capacity * sizeof(unsigned int), hipMemcpyDeviceToHost);
	hipMemcpy(cnt, cnt_d, 256 * sizeof(unsigned int), hipMemcpyDeviceToHost);


    printf("COUNT: %d\n\n", ht->count);

	for (i = 1; i < 256; ++i)
		printf("%d\t%d\n", i, cnt[i]);


	// limpieza
    hipFree(reads_d);
	hipFree(ht_d);
	hipFree(cnt_d);
	hipFree(keys_d);
	hipFree(values_d);

	// limpieza
    i = 0;
	for(current = head; current; current=current->next){
        free(current->sequence);
        free(current);
        free(reads[i]);
        i++;
    }

    free(reads);
	HashTable_destory(ht);
    return 0;
}


int main(int argc, char *argv[])
{
	int k = 31;
    unsigned int p = 27;

    k = (int)strtol(argv[1], NULL, 10);
    p = (unsigned int)strtol(argv[2], NULL, 10);
	count_file(argv[3], k, p);

	return 0;
}

	.text
	.file	"cuda-slow.hip"
	.globl	_Z14HashTable_initjj            # -- Begin function _Z14HashTable_initjj
	.type	_Z14HashTable_initjj,@function
_Z14HashTable_initjj:                   # @_Z14HashTable_initjj
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.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 %rbp, -16
	movl	%esi, %ebx
	movl	%edi, %ebp
	movl	$1, %edi
	movl	$32, %esi
	callq	calloc
	movl	%ebx, 8(%rax)
	movl	%ebp, (%rax)
	addq	$8, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z14HashTable_initjj, .Lfunc_end0-_Z14HashTable_initjj
	.cfi_endproc
                                        # -- End function
	.globl	_Z17HashTable_destoryP9HashTable # -- Begin function _Z17HashTable_destoryP9HashTable
	.type	_Z17HashTable_destoryP9HashTable,@function
_Z17HashTable_destoryP9HashTable:       # @_Z17HashTable_destoryP9HashTable
	.cfi_startproc
# %bb.0:
	testq	%rdi, %rdi
	jne	free                            # TAILCALL
# %bb.1:
	retq
.Lfunc_end1:
	.size	_Z17HashTable_destoryP9HashTable, .Lfunc_end1-_Z17HashTable_destoryP9HashTable
	.cfi_endproc
                                        # -- End function
	.globl	_Z37__device_stub__kernel_count_seq_kmersP9HashTableiPPc # -- Begin function _Z37__device_stub__kernel_count_seq_kmersP9HashTableiPPc
	.type	_Z37__device_stub__kernel_count_seq_kmersP9HashTableiPPc,@function
_Z37__device_stub__kernel_count_seq_kmersP9HashTableiPPc: # @_Z37__device_stub__kernel_count_seq_kmersP9HashTableiPPc
	.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	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	12(%rsp), %rcx
	movl	%esi, (%rcx)
	leaq	32(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%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	$_Z22kernel_count_seq_kmersP9HashTableiPPc, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	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_end2:
	.size	_Z37__device_stub__kernel_count_seq_kmersP9HashTableiPPc, .Lfunc_end2-_Z37__device_stub__kernel_count_seq_kmersP9HashTableiPPc
	.cfi_endproc
                                        # -- End function
	.globl	_Z32__device_stub__kernel_print_histPK9HashTablePj # -- Begin function _Z32__device_stub__kernel_print_histPK9HashTablePj
	.type	_Z32__device_stub__kernel_print_histPK9HashTablePj,@function
_Z32__device_stub__kernel_print_histPK9HashTablePj: # @_Z32__device_stub__kernel_print_histPK9HashTablePj
	.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	$80, %rsp
	.cfi_def_cfa_offset 128
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	16(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	64(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%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	$_Z17kernel_print_histPK9HashTablePj, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$96, %rsp
	.cfi_adjust_cfa_offset -96
	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_end3:
	.size	_Z32__device_stub__kernel_print_histPK9HashTablePj, .Lfunc_end3-_Z32__device_stub__kernel_print_histPK9HashTablePj
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function main
.LCPI4_0:
	.long	0x447a0000                      # float 1000
	.text
	.globl	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	$1160, %rsp                     # imm = 0x488
	.cfi_def_cfa_offset 1216
	.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	%rsi, %rbx
	movq	8(%rsi), %rdi
	xorl	%r15d, %r15d
	xorl	%esi, %esi
	movl	$10, %edx
	callq	__isoc23_strtol
	movq	%rax, %r14
	movq	16(%rbx), %rdi
	xorl	%esi, %esi
	movl	$10, %edx
	callq	__isoc23_strtol
	movq	%rax, %r12
	movl	$1, %eax
	movl	%r12d, %ecx
	shll	%cl, %eax
	movq	%rax, 80(%rsp)                  # 8-byte Spill
	movq	24(%rbx), %rdi
	movq	%r15, 32(%rsp)
	xorl	%eax, %eax
	movq	%rax, (%rsp)                    # 8-byte Spill
	movq	%r15, 88(%rsp)
	movl	$.L.str, %esi
	callq	fopen
	testq	%rax, %rax
	je	.LBB4_3
# %bb.1:                                # %.preheader.i
	movq	%rax, %r15
	movq	%r12, 48(%rsp)                  # 8-byte Spill
	movq	%r14, 96(%rsp)                  # 8-byte Spill
	leaq	32(%rsp), %rdi
	leaq	88(%rsp), %rsi
	movq	%rax, %rdx
	callq	getline
	cmpq	$-1, %rax
	je	.LBB4_2
# %bb.4:                                # %.lr.ph.i.preheader
	movq	%rax, %r13
	xorl	%eax, %eax
	movq	%rax, (%rsp)                    # 8-byte Spill
	xorl	%r12d, %r12d
	xorl	%ebp, %ebp
	movq	%r15, 16(%rsp)                  # 8-byte Spill
.LBB4_5:                                # %.lr.ph.i
                                        # =>This Inner Loop Header: Depth=1
	movl	$24, %edi
	callq	malloc
	movq	%rax, %rbx
	movq	32(%rsp), %r15
	movq	%r15, %rdi
	callq	strlen
	movq	%rax, %rdi
	callq	malloc
	movq	%rax, (%rbx)
	movq	%rax, %rdi
	movq	%r15, %rsi
	callq	strcpy
	movl	%r13d, 8(%rbx)
	movq	$0, 16(%rbx)
	movq	%rbx, %r14
	testq	%rbp, %rbp
	je	.LBB4_7
# %bb.6:                                #   in Loop: Header=BB4_5 Depth=1
	movq	%rbx, 16(%r12)
	movq	%rbp, %r14
.LBB4_7:                                #   in Loop: Header=BB4_5 Depth=1
	leaq	32(%rsp), %rdi
	leaq	88(%rsp), %rsi
	movq	16(%rsp), %r15                  # 8-byte Reload
	movq	%r15, %rdx
	callq	getline
	movq	%rax, %r13
	movq	(%rsp), %rax                    # 8-byte Reload
	incl	%eax
	movq	%rax, (%rsp)                    # 8-byte Spill
	movq	%rbx, %r12
	movq	%r14, %rbp
	cmpq	$-1, %r13
	jne	.LBB4_5
	jmp	.LBB4_8
.LBB4_2:
	xorl	%r14d, %r14d
.LBB4_8:                                # %._crit_edge.i
	movq	%r15, %rdi
	callq	fclose
	movq	32(%rsp), %rdi
	testq	%rdi, %rdi
	je	.LBB4_10
# %bb.9:
	callq	free
.LBB4_10:
	xorl	%r13d, %r13d
	movl	$.L.str.1, %edi
	movq	(%rsp), %rbx                    # 8-byte Reload
	movl	%ebx, %esi
	xorl	%eax, %eax
	callq	printf
	movl	%ebx, %ebx
	leaq	(,%rbx,8), %r12
	movq	%r12, %rdi
	callq	malloc
	movq	%rax, %rbp
	testq	%r14, %r14
	je	.LBB4_13
# %bb.11:                               # %.lr.ph123.i.preheader
	movq	%r14, %r15
.LBB4_12:                               # %.lr.ph123.i
                                        # =>This Inner Loop Header: Depth=1
	movl	8(%r15), %edi
	callq	malloc
	movl	%r13d, %ecx
	movq	%rax, (%rbp,%rcx,8)
	movq	(%r15), %rsi
	movq	%rax, %rdi
	callq	strcpy
	incl	%r13d
	movq	16(%r15), %r15
	testq	%r15, %r15
	jne	.LBB4_12
.LBB4_13:                               # %._crit_edge124.i
	movl	$1, %edi
	movl	$32, %esi
	callq	calloc
	movq	%rax, %rcx
	movq	%rax, 16(%rsp)                  # 8-byte Spill
	movq	(%rsp), %rax                    # 8-byte Reload
	movl	%eax, 8(%rcx)
	movq	48(%rsp), %rax                  # 8-byte Reload
	movl	%eax, (%rcx)
	leaq	8(%rsp), %rdi
	movl	$32, %esi
	callq	hipMalloc
	movl	80(%rsp), %r13d                 # 4-byte Reload
	leaq	(,%r13,8), %r15
	leaq	24(%rsp), %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	shlq	$2, %r13
	leaq	64(%rsp), %rdi
	movq	%r13, %rsi
	callq	hipMalloc
	leaq	40(%rsp), %rdi
	movl	$1024, %esi                     # imm = 0x400
	callq	hipMalloc
	leaq	24(%rsp), %rax
	movq	(%rax), %rdi
	xorl	%esi, %esi
	movq	%r15, 104(%rsp)                 # 8-byte Spill
	movq	%r15, %rdx
	callq	hipMemset
	leaq	64(%rsp), %r15
	movq	(%r15), %rdi
	xorl	%esi, %esi
	movq	%r13, 48(%rsp)                  # 8-byte Spill
	movq	%r13, %rdx
	callq	hipMemset
	leaq	40(%rsp), %rax
	movq	(%rax), %rdi
	movl	$1024, %edx                     # imm = 0x400
	xorl	%esi, %esi
	callq	hipMemset
	leaq	24(%rsp), %rax
	movq	(%rax), %rax
	movq	16(%rsp), %rsi                  # 8-byte Reload
	movq	%rax, 16(%rsi)
	movq	(%r15), %rax
	movq	%rax, 24(%rsi)
	leaq	8(%rsp), %rax
	movq	(%rax), %rdi
	movl	$32, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	%r12, %rdi
	callq	malloc
	movq	%rax, %r13
	cmpl	$0, (%rsp)                      # 4-byte Folded Reload
	je	.LBB4_18
# %bb.14:                               # %.lr.ph127.i
	xorl	%r15d, %r15d
.LBB4_15:                               # =>This Inner Loop Header: Depth=1
	leaq	(%r15,%r13), %rdi
	movl	8(%r14), %esi
	callq	hipMalloc
	addq	$8, %r15
	cmpq	%r15, %r12
	jne	.LBB4_15
# %bb.16:                               # %.lr.ph131.i
	leaq	72(%rsp), %r15
	movq	%r15, %rdi
	movq	%r12, %rsi
	callq	hipMalloc
	movq	(%r15), %rdi
	movq	%r13, %rsi
	movq	%r12, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	xorl	%r15d, %r15d
.LBB4_17:                               # =>This Inner Loop Header: Depth=1
	movq	(%r13,%r15,8), %rdi
	movq	(%rbp,%r15,8), %rsi
	movl	8(%r14), %edx
	movl	$1, %ecx
	callq	hipMemcpy
	incq	%r15
	cmpq	%r15, %rbx
	jne	.LBB4_17
	jmp	.LBB4_19
.LBB4_18:                               # %._crit_edge132.critedge.i
	leaq	72(%rsp), %rbx
	movq	%rbx, %rdi
	movq	%r12, %rsi
	callq	hipMalloc
	movq	(%rbx), %rdi
	movq	%r13, %rsi
	movq	%r12, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
.LBB4_19:                               # %._crit_edge132.i
	movabsq	$4294967296, %r12               # imm = 0x100000000
	movq	%r13, %rdi
	callq	free
	movl	$.L.str.3, %edi
	movq	(%rsp), %r15                    # 8-byte Reload
	movl	%r15d, %esi
	xorl	%eax, %eax
	callq	printf
	leaq	112(%rsp), %rbx
	movq	%rbx, %rdi
	callq	hipEventCreate
	leaq	56(%rsp), %rdi
	callq	hipEventCreate
	movq	(%rbx), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	shrl	$10, %r15d
	btsq	$32, %r15
	leaq	1024(%r12), %rdx
	movq	%r15, %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB4_21
# %bb.20:
	movq	8(%rsp), %rdi
	movq	72(%rsp), %rdx
	movq	96(%rsp), %rsi                  # 8-byte Reload
                                        # kill: def $esi killed $esi killed $rsi
	callq	_Z37__device_stub__kernel_count_seq_kmersP9HashTableiPPc
.LBB4_21:
	callq	hipDeviceSynchronize
	movq	80(%rsp), %rdi                  # 8-byte Reload
	shrl	$10, %edi
	orq	%r12, %rdi
	addq	$1024, %r12                     # imm = 0x400
	movl	$1, %esi
	movq	%r12, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	movq	16(%rsp), %r15                  # 8-byte Reload
	jne	.LBB4_23
# %bb.22:
	movq	8(%rsp), %rdi
	movq	40(%rsp), %rsi
	callq	_Z32__device_stub__kernel_print_histPK9HashTablePj
.LBB4_23:
	callq	hipDeviceSynchronize
	movq	56(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	56(%rsp), %rdi
	callq	hipEventSynchronize
	leaq	124(%rsp), %rbx
	movl	$0, (%rbx)
	movq	112(%rsp), %rsi
	movq	56(%rsp), %rdx
	movq	%rbx, %rdi
	callq	hipEventElapsedTime
	movss	(%rbx), %xmm0                   # xmm0 = mem[0],zero,zero,zero
	divss	.LCPI4_0(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	movl	$.L.str.4, %edi
	movb	$1, %al
	callq	printf
	movq	8(%rsp), %rsi
	movl	$32, %edx
	movq	%r15, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	16(%r15), %rdi
	movq	24(%rsp), %rsi
	movq	104(%rsp), %rdx                 # 8-byte Reload
	movl	$2, %ecx
	callq	hipMemcpy
	movq	24(%r15), %rdi
	movq	64(%rsp), %rsi
	movq	48(%rsp), %rdx                  # 8-byte Reload
	movl	$2, %ecx
	callq	hipMemcpy
	movq	40(%rsp), %rsi
	leaq	128(%rsp), %rdi
	movl	$1024, %edx                     # imm = 0x400
	movl	$2, %ecx
	callq	hipMemcpy
	movl	4(%r15), %esi
	movl	$.L.str.5, %edi
	xorl	%eax, %eax
	callq	printf
	movl	$1, %ebx
.LBB4_24:                               # =>This Inner Loop Header: Depth=1
	movl	128(%rsp,%rbx,4), %edx
	movl	$.L.str.6, %edi
	movl	%ebx, %esi
	xorl	%eax, %eax
	callq	printf
	incq	%rbx
	cmpq	$256, %rbx                      # imm = 0x100
	jne	.LBB4_24
# %bb.25:
	movq	72(%rsp), %rdi
	callq	hipFree
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	40(%rsp), %rdi
	callq	hipFree
	movq	24(%rsp), %rdi
	callq	hipFree
	movq	64(%rsp), %rdi
	callq	hipFree
	testq	%r14, %r14
	je	.LBB4_28
# %bb.26:                               # %.lr.ph138.i.preheader
	xorl	%ebx, %ebx
.LBB4_27:                               # %.lr.ph138.i
                                        # =>This Inner Loop Header: Depth=1
	movq	(%r14), %rdi
	callq	free
	movq	%r14, %rdi
	callq	free
	movl	%ebx, %eax
	movq	(%rbp,%rax,8), %rdi
	callq	free
	incl	%ebx
	movq	16(%r14), %r14
	testq	%r14, %r14
	jne	.LBB4_27
.LBB4_28:                               # %_ZL10count_filePKcij.exit
	movq	%rbp, %rdi
	callq	free
	movq	%r15, %rdi
	callq	free
	xorl	%eax, %eax
	addq	$1160, %rsp                     # imm = 0x488
	.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
.LBB4_3:
	.cfi_def_cfa_offset 1216
	movl	$1, %edi
	callq	exit
.Lfunc_end4:
	.size	main, .Lfunc_end4-main
	.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	$_Z22kernel_count_seq_kmersP9HashTableiPPc, %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	$_Z17kernel_print_histPK9HashTablePj, %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	_Z22kernel_count_seq_kmersP9HashTableiPPc,@object # @_Z22kernel_count_seq_kmersP9HashTableiPPc
	.section	.rodata,"a",@progbits
	.globl	_Z22kernel_count_seq_kmersP9HashTableiPPc
	.p2align	3, 0x0
_Z22kernel_count_seq_kmersP9HashTableiPPc:
	.quad	_Z37__device_stub__kernel_count_seq_kmersP9HashTableiPPc
	.size	_Z22kernel_count_seq_kmersP9HashTableiPPc, 8

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

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

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"%d\n"
	.size	.L.str.1, 4

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"total reads: %d\n"
	.size	.L.str.3, 17

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"GPU time: %fs\n"
	.size	.L.str.4, 15

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"COUNT: %d\n\n"
	.size	.L.str.5, 12

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"%d\t%d\n"
	.size	.L.str.6, 7

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z22kernel_count_seq_kmersP9HashTableiPPc"
	.size	.L__unnamed_1, 42

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z17kernel_print_histPK9HashTablePj"
	.size	.L__unnamed_2, 36

	.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

	.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 _Z37__device_stub__kernel_count_seq_kmersP9HashTableiPPc
	.addrsig_sym _Z32__device_stub__kernel_print_histPK9HashTablePj
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z22kernel_count_seq_kmersP9HashTableiPPc
	.addrsig_sym _Z17kernel_print_histPK9HashTablePj
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z22kernel_count_seq_kmersP9HashTableiPPc ; -- Begin function _Z22kernel_count_seq_kmersP9HashTableiPPc
	.globl	_Z22kernel_count_seq_kmersP9HashTableiPPc
	.p2align	8
	.type	_Z22kernel_count_seq_kmersP9HashTableiPPc,@function
_Z22kernel_count_seq_kmersP9HashTableiPPc: ; @_Z22kernel_count_seq_kmersP9HashTableiPPc
; %bb.0:
	s_clause 0x1
	s_load_b64 s[2:3], s[0:1], 0x0
	s_load_b32 s4, s[0:1], 0x24
	s_waitcnt lgkmcnt(0)
	s_load_b32 s5, s[2:3], 0x8
	s_and_b32 s4, s4, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[4:5], null, s15, s4, v[0:1]
	s_mov_b32 s4, exec_lo
	s_waitcnt lgkmcnt(0)
	v_cmpx_gt_u32_e64 s5, v4
	s_cbranch_execz .LBB0_14
; %bb.1:
	s_load_b64 s[4:5], s[0:1], 0x10
	v_dual_mov_b32 v5, 0 :: v_dual_mov_b32 v8, 0
	s_load_b32 s1, s[0:1], 0x8
	v_dual_mov_b32 v9, 0 :: v_dual_mov_b32 v14, 1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshlrev_b64 v[0:1], 3, v[4:5]
	v_mov_b32_e32 v2, v8
	v_mov_b32_e32 v4, 0
	s_mov_b32 s6, 0
	s_mov_b32 s8, 0
	v_mov_b32_e32 v3, v9
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v0, vcc_lo, s4, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s5, v1, vcc_lo
	s_lshl_b32 s0, s1, 1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_lshl_b64 s[4:5], -1, s0
	global_load_b64 v[6:7], v[0:1], off
	v_dual_mov_b32 v0, v8 :: v_dual_mov_b32 v1, v9
	s_not_b64 s[4:5], s[4:5]
	s_add_i32 s7, s0, -2
.LBB0_2:                                ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_5 Depth 2
	s_waitcnt vmcnt(0)
	v_add_co_u32 v10, vcc_lo, v6, s8
	v_add_co_ci_u32_e32 v11, vcc_lo, 0, v7, vcc_lo
	s_getpc_b64 s[10:11]
	s_add_u32 s10, s10, _ZL13seq_nt4_table@rel32@lo+4
	s_addc_u32 s11, s11, _ZL13seq_nt4_table@rel32@hi+12
	s_mov_b32 s9, exec_lo
	v_mov_b32_e32 v13, v9
	global_load_u8 v10, v[10:11], off
	v_dual_mov_b32 v12, v8 :: v_dual_mov_b32 v15, v5
	s_waitcnt vmcnt(0)
	v_and_b32_e32 v10, 0xffff, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v10, s0, v10, s10
	v_add_co_ci_u32_e64 v11, null, 0, s11, s0
	global_load_u8 v16, v[10:11], off
	v_dual_mov_b32 v11, v9 :: v_dual_mov_b32 v10, v8
	s_waitcnt vmcnt(0)
	v_cmpx_gt_u16_e32 4, v16
	s_cbranch_execz .LBB0_13
; %bb.3:                                ;   in Loop: Header=BB0_2 Depth=1
	v_xor_b32_e32 v10, 3, v16
	v_lshlrev_b64 v[2:3], 2, v[2:3]
	v_dual_mov_b32 v11, s6 :: v_dual_and_b32 v12, 0xffff, v16
	v_lshrrev_b64 v[0:1], 2, v[0:1]
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_and_b32_e32 v10, 0xffff, v10
	v_add_nc_u32_e32 v15, 1, v4
	v_or_b32_e32 v2, v2, v12
	s_mov_b32 s10, exec_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_lshlrev_b64 v[12:13], s7, v[10:11]
	v_and_b32_e32 v11, s5, v3
	v_and_b32_e32 v10, s4, v2
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_or_b32_e32 v13, v13, v1
	v_or_b32_e32 v12, v12, v0
	v_cmpx_le_i32_e64 s1, v15
	s_cbranch_execz .LBB0_12
; %bb.4:                                ;   in Loop: Header=BB0_2 Depth=1
	global_load_b32 v20, v5, s[2:3]
	v_cmp_lt_u64_e32 vcc_lo, v[10:11], v[12:13]
	s_mov_b32 s12, 0
	s_mov_b32 s11, 0
	v_dual_cndmask_b32 v1, v13, v11 :: v_dual_cndmask_b32 v0, v12, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_not_b32_e32 v16, v1
	v_lshlrev_b64 v[2:3], 21, v[0:1]
	v_not_b32_e32 v4, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v2, vcc_lo, v2, v4
	v_add_co_ci_u32_e32 v3, vcc_lo, v3, v16, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshrrev_b64 v[16:17], 24, v[2:3]
	v_xor_b32_e32 v2, v16, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v16, v17, v3
	v_mad_u64_u32 v[18:19], null, 0x109, v2, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mov_b32_e32 v2, v19
	v_mad_u64_u32 v[3:4], null, 0x109, v16, v[2:3]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mov_b32_e32 v19, v3
	v_lshrrev_b64 v[2:3], 14, v[18:19]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v2, v2, v18
	v_mad_u64_u32 v[16:17], null, v2, 21, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_mov_b32_e32 v2, v17
	v_xor_b32_e32 v18, v3, v19
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[3:4], null, v18, 21, v[2:3]
	v_mov_b32_e32 v2, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_alignbit_b32 v2, v2, v16, 28
	v_xor_b32_e32 v2, v2, v16
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v2, 0x1e3779b9, v2
	s_waitcnt vmcnt(0)
	v_sub_nc_u32_e32 v3, 32, v20
	v_lshrrev_b32_e32 v16, v3, v2
	s_delay_alu instid0(VALU_DEP_1)
	v_mov_b32_e32 v4, v16
.LBB0_5:                                ;   Parent Loop BB0_2 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_b64 v[17:18], v5, s[2:3] offset:16
	v_lshlrev_b64 v[19:20], 3, v[4:5]
	v_mov_b32_e32 v2, v5
	v_mov_b32_e32 v3, v5
                                        ; implicit-def: $sgpr13
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v17, vcc_lo, v17, v19
	v_add_co_ci_u32_e32 v18, vcc_lo, v18, v20, vcc_lo
	flat_atomic_cmpswap_b64 v[2:3], v[17:18], v[0:3] glc
	s_waitcnt vmcnt(0) lgkmcnt(0)
	v_cmp_ne_u64_e32 vcc_lo, 0, v[2:3]
	v_cmp_ne_u64_e64 s0, v[2:3], v[0:1]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s14, vcc_lo, s0
                                        ; implicit-def: $sgpr0
	s_and_saveexec_b32 s15, s14
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s14, exec_lo, s15
	s_cbranch_execz .LBB0_9
; %bb.6:                                ;   in Loop: Header=BB0_5 Depth=2
	s_xor_b32 s16, s12, -1
	s_mov_b32 s13, -1
	s_mov_b32 s0, 0
	s_and_saveexec_b32 s15, s16
	s_cbranch_execz .LBB0_8
; %bb.7:                                ;   in Loop: Header=BB0_5 Depth=2
	global_load_b32 v2, v5, s[2:3]
	v_add_nc_u32_e32 v3, 1, v4
	s_mov_b32 s0, exec_lo
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_bfe_u32 v4, v3, 0, v2
	v_cmp_eq_u32_e32 vcc_lo, v4, v16
	s_or_not1_b32 s13, vcc_lo, exec_lo
.LBB0_8:                                ; %Flow
                                        ;   in Loop: Header=BB0_5 Depth=2
	s_or_b32 exec_lo, exec_lo, s15
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_b32 s0, s0, exec_lo
	s_and_b32 s13, s13, exec_lo
.LBB0_9:                                ; %Flow40
                                        ;   in Loop: Header=BB0_5 Depth=2
	s_and_not1_saveexec_b32 s14, s14
	s_cbranch_execz .LBB0_11
; %bb.10:                               ;   in Loop: Header=BB0_5 Depth=2
	global_load_b64 v[2:3], v5, s[2:3] offset:24
	v_lshlrev_b64 v[17:18], 2, v[4:5]
	s_and_not1_b32 s13, s13, exec_lo
	s_and_b32 s12, s12, exec_lo
	s_and_not1_b32 s0, s0, exec_lo
	s_or_b32 s13, s13, s12
	s_waitcnt vmcnt(0)
	v_add_co_u32 v2, vcc_lo, v2, v17
	v_add_co_ci_u32_e32 v3, vcc_lo, v3, v18, vcc_lo
	flat_atomic_add_u32 v[2:3], v14
.LBB0_11:                               ;   in Loop: Header=BB0_5 Depth=2
	s_or_b32 exec_lo, exec_lo, s14
	s_xor_b32 s0, s0, -1
	s_mov_b32 s12, s13
	s_and_b32 s0, exec_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s11, s0, s11
	s_and_not1_b32 exec_lo, exec_lo, s11
	s_cbranch_execnz .LBB0_5
.LBB0_12:                               ; %Flow42
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s10
.LBB0_13:                               ; %_Z11hash_insertP9HashTabley.exit
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_4) | instid1(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s9
	v_mov_b32_e32 v0, v12
	v_dual_mov_b32 v2, v10 :: v_dual_mov_b32 v3, v11
	v_dual_mov_b32 v4, v15 :: v_dual_mov_b32 v1, v13
	s_add_i32 s8, s8, 1
	s_cmpk_lg_i32 s8, 0x64
	s_cbranch_scc1 .LBB0_2
.LBB0_14:                               ; %.loopexit
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z22kernel_count_seq_kmersP9HashTableiPPc
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 280
		.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 21
		.amdhsa_next_free_sgpr 17
		.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	_Z22kernel_count_seq_kmersP9HashTableiPPc, .Lfunc_end0-_Z22kernel_count_seq_kmersP9HashTableiPPc
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 952
; NumSgprs: 19
; NumVgprs: 21
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 2
; NumSGPRsForWavesPerEU: 19
; NumVGPRsForWavesPerEU: 21
; 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
	.protected	_Z17kernel_print_histPK9HashTablePj ; -- Begin function _Z17kernel_print_histPK9HashTablePj
	.globl	_Z17kernel_print_histPK9HashTablePj
	.p2align	8
	.type	_Z17kernel_print_histPK9HashTablePj,@function
_Z17kernel_print_histPK9HashTablePj:    ; @_Z17kernel_print_histPK9HashTablePj
; %bb.0:
	s_clause 0x1
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b32 s0, s[0:1], 0x1c
	s_waitcnt lgkmcnt(0)
	s_load_b32 s1, s[4:5], 0x0
	s_and_b32 s0, s0, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s0, v[0:1]
	s_mov_b32 s0, exec_lo
	s_waitcnt lgkmcnt(0)
	v_lshrrev_b32_e32 v0, s1, v1
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v0
	s_cbranch_execz .LBB1_3
; %bb.1:
	s_load_b64 s[0:1], s[4:5], 0x18
	v_mov_b32_e32 v2, 0
	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, s0, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	global_load_b32 v0, v[0:1], off
	s_waitcnt vmcnt(0)
	v_cmp_ne_u32_e32 vcc_lo, 0, v0
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB1_3
; %bb.2:
	v_min_u32_e32 v0, 0xff, v0
	s_delay_alu instid0(VALU_DEP_1)
	v_dual_mov_b32 v1, 1 :: v_dual_lshlrev_b32 v0, 2, v0
	global_atomic_add_u32 v0, v1, s[6:7]
.LBB1_3:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z17kernel_print_histPK9HashTablePj
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 272
		.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 3
		.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_end1:
	.size	_Z17kernel_print_histPK9HashTablePj, .Lfunc_end1-_Z17kernel_print_histPK9HashTablePj
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 184
; NumSgprs: 18
; NumVgprs: 3
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 3
; 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	_ZL13seq_nt4_table,@object      ; @_ZL13seq_nt4_table
	.section	.rodata,"a",@progbits
	.p2align	4, 0x0
_ZL13seq_nt4_table:
	.ascii	"\000\001\002\003\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\000\004\001\004\004\004\002\004\004\004\004\004\004\004\004\004\004\004\004\003\003\004\004\004\004\004\004\004\004\004\004\004\000\004\001\004\004\004\002\004\004\004\004\004\004\004\004\004\004\004\004\003\003\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004\004"
	.size	_ZL13seq_nt4_table, 256

	.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
      - .offset:         8
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         28
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         36
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         38
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         40
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         42
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         44
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         46
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         88
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 280
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z22kernel_count_seq_kmersP9HashTableiPPc
    .private_segment_fixed_size: 0
    .sgpr_count:     19
    .sgpr_spill_count: 0
    .symbol:         _Z22kernel_count_seq_kmersP9HashTableiPPc.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     21
    .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
      - .offset:         16
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         20
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         28
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         30
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         32
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         34
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         36
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         38
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         56
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         80
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 272
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z17kernel_print_histPK9HashTablePj
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z17kernel_print_histPK9HashTablePj.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     3
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
