// Basado en el codigo entregado durante clases practicas

#include <iostream>
#include <time.h>
#include <cuda_runtime.h>

/*
 *  Lectura Archivo
 */
void Read(float** R, float** G, float** B, int *N, int *S, int** posicionesRompecabezas, const char *filename) {
    printf("Leemos el archivo %s!\n", filename); 
    FILE *fp;
    fp = fopen(filename, "r");
    fscanf(fp, "%d %d\n", N, S);

    int P = (*N) / (*S);
    int imsize = (*N) * (*N);

    float* R1 = new float[imsize];
    float* G1 = new float[imsize];
    float* B1 = new float[imsize];
    int *posicionesTemp = new int[P*P];

    for(int i = 0; i < P*P; i++)
	    fscanf(fp, "%d ", &(posicionesTemp[i]));
	for(int i = 0; i < imsize; i++)
	    fscanf(fp, "%f ", &(R1[i]));
	for(int i = 0; i < imsize; i++)
	    fscanf(fp, "%f ", &(G1[i]));
	for(int i = 0; i < imsize; i++)
	    fscanf(fp, "%f ", &(B1[i]));
    fclose(fp);

    *R = R1; *G = G1; *B = B1; 
    *posicionesRompecabezas = posicionesTemp;
}

/*
 *  Escritura Archivo
 */
void Write(float* R, float* G, float* B, 
	       int M, int N, const char *filename) {
    FILE *fp;
    fp = fopen(filename, "w");
    fprintf(fp, "%d %d\n", M, N);
    for(int i = 0; i < M*N-1; i++)
        fprintf(fp, "%f ", R[i]);
    fprintf(fp, "%f\n", R[M*N-1]);
    for(int i = 0; i < M*N-1; i++)
        fprintf(fp, "%f ", G[i]);
    fprintf(fp, "%f\n", G[M*N-1]);
    for(int i = 0; i < M*N-1; i++)
        fprintf(fp, "%f ", B[i]);
    fprintf(fp, "%f\n", B[M*N-1]);
    fclose(fp);
}


__host__ __device__ int findNewPosition(int P, int position, int* posiciones){

	int newPosition;

	for (int j = 0; j < P*P; j++) {
    	if(position == posiciones[j]) {
        	newPosition = j;
        	break;
    	}
	}

	return newPosition;
}

/*
 *  Procesamiento Imagen CPU
 */
void funcionCPU( float* R, float* G, float* B, float* Rout, float* Gout, float* Bout, 
				 int N, int S, int* posicionesRompecabezas) {

	int P = N/S;
	
	for (int i = 0; i < N*N; i++){
		int xOriginal = i % N;
		int yOriginal = i / N;
		int xRompecabezas = xOriginal / S;
		int yRompecabezas = yOriginal / S;
		int indiceRompecabezas = xRompecabezas + yRompecabezas * P;
		int indiceRompecabezasMezclado =  posicionesRompecabezas[indiceRompecabezas];
		int xRompecabezasMezclado = indiceRompecabezasMezclado % P;
		int yRompecabezasMezclado = indiceRompecabezasMezclado / P;
		int nuevoI = (xOriginal + S * (xRompecabezasMezclado - xRompecabezas)) + N * (yOriginal + S * (yRompecabezasMezclado - yRompecabezas));

		Rout[i] = R[nuevoI];
		Gout[i] = G[nuevoI];
		Bout[i] = B[nuevoI];
	}
}

/*
 *  Procesamiento Imagen GPU
 */
__global__ void kernelGPU( float* R, float* G, float* B, float* Rout, float* Gout, float* Bout, 
						   int N, int S, int* posicionesRompecabezas){
	int i = threadIdx.x + blockDim.x * blockIdx.x;
	int P = N/S;

	if(i < N*N){
		int xOriginal = i % N;
		int yOriginal = i / N;
		int xRompecabezas = xOriginal / S;
		int yRompecabezas = yOriginal / S;
		int indiceRompecabezas = xRompecabezas + yRompecabezas * P;
		int indiceRompecabezasMezclado =  posicionesRompecabezas[indiceRompecabezas];
		int xRompecabezasMezclado = indiceRompecabezasMezclado % P;
		int yRompecabezasMezclado = indiceRompecabezasMezclado / P;
		int nuevoI = (xOriginal + S * (xRompecabezasMezclado - xRompecabezas)) + N * (yOriginal + S * (yRompecabezasMezclado - yRompecabezas));

		Rout[i] = R[nuevoI];
		Gout[i] = G[nuevoI];
		Bout[i] = B[nuevoI];
	} 
}

/*
 *  Codigo Principal
 */
int main(int argc, char **argv){

    /*
     *  Inicializacion
     */
	clock_t t1, t2;
	cudaEvent_t ct1, ct2;
	double ms;
	float dt;
	int N, S;
	int P;
	int *posiciones;
	int *posicionesDev;
    float *Rhost, *Ghost, *Bhost;
    float *Rhostout, *Ghostout, *Bhostout;
    float *Rdev, *Gdev, *Bdev;
    float *Rdevout, *Gdevout, *Bdevout;
    
    //cambiar con los nombres de los archivos correctos
    char names[5][3][30] = {{"img100x100.txt\0", "img100x100CPU.txt\0", "img100x100GPU.txt\0"}, 
    	{"img200x200.txt\0", "img200x200CPU.txt\0", "img200x200GPU.txt\0"},
    	{"img400x400.txt\0", "img400x400CPU.txt\0", "img400x400GPU.txt\0"},
    	{"img800x800.txt\0", "img800x800CPU.txt\0", "img800x800GPU.txt\0"},
    	{"img1600x1600.txt\0", "img1600x1600CPU.txt\0", "img1600x1600GPU.txt\0"}};

    for (int i=0; i<5; i++){
	    Read(&Rhost, &Ghost, &Bhost, &N, &S, &posiciones, names[i][0]); //leemos archivo y reservamos memoria

	    P = N/S;

	    /*
	     *  Parte CPU
	     */
	    Rhostout = new float[N*N];	//reservamos memoria
	    Ghostout = new float[N*N];
	    Bhostout = new float[N*N];

	    t1 = clock();
	    funcionCPU(Rhost, Ghost, Bhost, Rhostout, Ghostout, Bhostout, N, S, posiciones); // Agregar parametros!
	    t2 = clock();
	    ms = 1000.0 * (double)(t2 - t1) / CLOCKS_PER_SEC;
	    std::cout << "Tiempo CPU: " << ms << "[ms]" << std::endl;
	    Write(Rhostout, Ghostout, Bhostout, N, N, names[i][1]);

	    delete[] Rhostout; delete[] Ghostout; delete[] Bhostout;
	    
	    /*
	     *  Parte GPU
	     */

	    int grid_size, block_size = 256;
	    grid_size = (int)ceil((float) N * N / block_size);
	        
	    cudaMalloc((void**)&Rdev, N * N * sizeof(float));
	    cudaMalloc((void**)&Gdev, N * N * sizeof(float));
	    cudaMalloc((void**)&Bdev, N * N * sizeof(float));
	    cudaMalloc((void**)&posicionesDev, P * P * sizeof(float));
	    cudaMemcpy(Rdev, Rhost, N * N * sizeof(float), cudaMemcpyHostToDevice);
	    cudaMemcpy(Gdev, Ghost, N * N * sizeof(float), cudaMemcpyHostToDevice);
	    cudaMemcpy(Bdev, Bhost, N * N * sizeof(float), cudaMemcpyHostToDevice);
	    cudaMemcpy(posicionesDev, posiciones, P * P * sizeof(float), cudaMemcpyHostToDevice);
	        
	    cudaMalloc((void**)&Rdevout, N * N * sizeof(float));
	    cudaMalloc((void**)&Gdevout, N * N * sizeof(float));
	    cudaMalloc((void**)&Bdevout, N * N * sizeof(float));
	    
	    cudaEventCreate(&ct1);
	    cudaEventCreate(&ct2);
	    cudaEventRecord(ct1);
	    kernelGPU<<<grid_size, block_size>>>(Rdev, Gdev, Bdev, Rdevout, Gdevout, Bdevout, N, S, posicionesDev); // Agregar parametros!
	    cudaEventRecord(ct2);
	    cudaEventSynchronize(ct2);
	    cudaEventElapsedTime(&dt, ct1, ct2);
	    std::cout << "Tiempo GPU: " << dt << "[ms]" << std::endl;

	    Rhostout = new float[N*N];
	    Ghostout = new float[N*N];
	    Bhostout = new float[N*N];
	    cudaMemcpy(Rhostout, Rdevout, N * N * sizeof(float), cudaMemcpyDeviceToHost);
	    cudaMemcpy(Ghostout, Gdevout, N * N * sizeof(float), cudaMemcpyDeviceToHost);
	    cudaMemcpy(Bhostout, Bdevout, N * N * sizeof(float), cudaMemcpyDeviceToHost);
	    Write(Rhostout, Ghostout, Bhostout, N, N, names[i][2]);

	    cudaFree(Rdev); cudaFree(Gdev); cudaFree(Bdev);
    	cudaFree(Rdevout); cudaFree(Gdevout); cudaFree(Bdevout);
    	cudaFree(posicionesDev);
    	delete[] Rhost; delete[] Ghost; delete[] Bhost;
    	delete[] Rhostout; delete[] Ghostout; delete[] Bhostout;
    	delete[] posiciones;
	}
	return 0;
}
	.file	"tmpxft_0023f51d_00000000-6_cod.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3642:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3642:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"Leemos el archivo %s!\n"
.LC2:
	.string	"r"
.LC3:
	.string	"%d %d\n"
.LC4:
	.string	"%d "
.LC5:
	.string	"%f "
	.text
	.globl	_Z4ReadPPfS0_S0_PiS1_PS1_PKc
	.type	_Z4ReadPPfS0_S0_PiS1_PS1_PKc, @function
_Z4ReadPPfS0_S0_PiS1_PS1_PKc:
.LFB3635:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	xorl	%eax, %eax
	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
	movq	%r8, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rcx, %rbx
	subq	$56, %rsp
	.cfi_def_cfa_offset 112
	movq	%rdx, 24(%rsp)
	movq	112(%rsp), %rdx
	movq	%rdi, 8(%rsp)
	movl	$2, %edi
	movq	%rsi, 16(%rsp)
	leaq	.LC1(%rip), %rsi
	movq	%r9, 32(%rsp)
	call	__printf_chk@PLT
	movq	112(%rsp), %rdi
	leaq	.LC2(%rip), %rsi
	call	fopen@PLT
	movq	%rbp, %rcx
	leaq	.LC3(%rip), %rsi
	movq	%rbx, %rdx
	movq	%rax, %rdi
	movq	%rax, (%rsp)
	xorl	%eax, %eax
	call	__isoc23_fscanf@PLT
	movl	(%rbx), %ebx
	movl	%ebx, %eax
	imull	%ebx, %ebx
	cltd
	idivl	0(%rbp)
	movslq	%ebx, %rbx
	leaq	0(,%rbx,4), %r12
	movq	%r12, %rdi
	movl	%eax, %ebp
	call	_Znam@PLT
	imull	%ebp, %ebp
	movq	%r12, %rdi
	movq	%rax, %r15
	call	_Znam@PLT
	movq	%r12, %rdi
	movslq	%ebp, %rbp
	movq	%rax, %r14
	call	_Znam@PLT
	salq	$2, %rbp
	movq	%rbp, %rdi
	movq	%rax, %r13
	call	_Znam@PLT
	movq	%rax, %r12
	movq	%rax, %rdx
	addq	%rax, %rbp
.L3:
	cmpq	%rbp, %rdx
	je	.L13
	movq	(%rsp), %rdi
	leaq	.LC4(%rip), %rsi
	xorl	%eax, %eax
	movq	%rdx, 40(%rsp)
	call	__isoc23_fscanf@PLT
	movq	40(%rsp), %rdx
	addq	$4, %rdx
	jmp	.L3
.L13:
	salq	$2, %rbx
	movq	%r15, %rdx
	leaq	(%r15,%rbx), %rbp
.L5:
	cmpq	%rbp, %rdx
	je	.L14
	movq	(%rsp), %rdi
	leaq	.LC5(%rip), %rsi
	xorl	%eax, %eax
	movq	%rdx, 40(%rsp)
	call	__isoc23_fscanf@PLT
	movq	40(%rsp), %rdx
	addq	$4, %rdx
	jmp	.L5
.L14:
	movq	%r14, %rdx
	leaq	(%r14,%rbx), %rbp
.L7:
	cmpq	%rbp, %rdx
	je	.L15
	movq	(%rsp), %rdi
	leaq	.LC5(%rip), %rsi
	xorl	%eax, %eax
	movq	%rdx, 40(%rsp)
	call	__isoc23_fscanf@PLT
	movq	40(%rsp), %rdx
	addq	$4, %rdx
	jmp	.L7
.L15:
	movq	%r13, %rbp
	addq	%r13, %rbx
.L9:
	cmpq	%rbx, %rbp
	je	.L16
	movq	(%rsp), %rdi
	movq	%rbp, %rdx
	leaq	.LC5(%rip), %rsi
	xorl	%eax, %eax
	addq	$4, %rbp
	call	__isoc23_fscanf@PLT
	jmp	.L9
.L16:
	movq	(%rsp), %rdi
	call	fclose@PLT
	movq	8(%rsp), %rax
	movq	%r15, (%rax)
	movq	16(%rsp), %rax
	movq	%r14, (%rax)
	movq	24(%rsp), %rax
	movq	%r13, (%rax)
	movq	32(%rsp), %rax
	movq	%r12, (%rax)
	addq	$56, %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
.LFE3635:
	.size	_Z4ReadPPfS0_S0_PiS1_PS1_PKc, .-_Z4ReadPPfS0_S0_PiS1_PS1_PKc
	.section	.rodata.str1.1
.LC6:
	.string	"w"
.LC7:
	.string	"%f\n"
	.text
	.globl	_Z5WritePfS_S_iiPKc
	.type	_Z5WritePfS_S_iiPKc, @function
_Z5WritePfS_S_iiPKc:
.LFB3636:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movl	%r8d, %r15d
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	xorl	%r14d, %r14d
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movl	%ecx, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rdi, %r12
	movq	%r9, %rdi
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rsi, %rbp
	leaq	.LC6(%rip), %rsi
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rdx, %rbx
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
	call	fopen@PLT
	movl	%r15d, %r8d
	imull	%r13d, %r15d
	movl	%r13d, %ecx
	movq	%rax, 8(%rsp)
	movq	%rax, %rdi
	movl	$2, %esi
	xorl	%eax, %eax
	leaq	.LC3(%rip), %rdx
	call	__fprintf_chk@PLT
	leal	-1(%r15), %r13d
.L18:
	cmpl	%r14d, %r13d
	jle	.L25
	movq	8(%rsp), %rdi
	leaq	.LC5(%rip), %rdx
	movl	$2, %esi
	movb	$1, %al
	cvtss2sd	(%r12,%r14,4), %xmm0
	incq	%r14
	call	__fprintf_chk@PLT
	jmp	.L18
.L25:
	movq	8(%rsp), %rdi
	movslq	%r15d, %r15
	movl	$2, %esi
	movb	$1, %al
	leaq	.LC7(%rip), %rdx
	cvtss2sd	-4(%r12,%r15,4), %xmm0
	leaq	0(,%r15,4), %r14
	xorl	%r12d, %r12d
	call	__fprintf_chk@PLT
	leaq	.LC5(%rip), %r15
.L20:
	cmpl	%r12d, %r13d
	jle	.L26
	movq	8(%rsp), %rdi
	movq	%r15, %rdx
	movl	$2, %esi
	movb	$1, %al
	cvtss2sd	0(%rbp,%r12,4), %xmm0
	incq	%r12
	call	__fprintf_chk@PLT
	jmp	.L20
.L26:
	movq	8(%rsp), %rdi
	movl	$2, %esi
	movb	$1, %al
	xorl	%r15d, %r15d
	leaq	.LC7(%rip), %rdx
	cvtss2sd	-4(%rbp,%r14), %xmm0
	leaq	.LC5(%rip), %rbp
	call	__fprintf_chk@PLT
.L22:
	cmpl	%r15d, %r13d
	jle	.L27
	movq	8(%rsp), %rdi
	movq	%rbp, %rdx
	movl	$2, %esi
	movb	$1, %al
	cvtss2sd	(%rbx,%r15,4), %xmm0
	incq	%r15
	call	__fprintf_chk@PLT
	jmp	.L22
.L27:
	movq	8(%rsp), %rdi
	leaq	.LC7(%rip), %rdx
	movl	$2, %esi
	movb	$1, %al
	cvtss2sd	-4(%rbx,%r14), %xmm0
	call	__fprintf_chk@PLT
	movq	8(%rsp), %rdi
	addq	$24, %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
	jmp	fclose@PLT
	.cfi_endproc
.LFE3636:
	.size	_Z5WritePfS_S_iiPKc, .-_Z5WritePfS_S_iiPKc
	.globl	_Z15findNewPositioniiPi
	.type	_Z15findNewPositioniiPi, @function
_Z15findNewPositioniiPi:
.LFB3637:
	.cfi_startproc
	endbr64
	imull	%edi, %edi
	xorl	%eax, %eax
.L29:
	cmpl	%eax, %edi
	jle	.L28
	leaq	1(%rax), %r8
	cmpl	%esi, (%rdx,%rax,4)
	jne	.L32
	movl	%eax, %ecx
	jmp	.L28
.L32:
	movq	%r8, %rax
	jmp	.L29
.L28:
	movl	%ecx, %eax
	ret
	.cfi_endproc
.LFE3637:
	.size	_Z15findNewPositioniiPi, .-_Z15findNewPositioniiPi
	.globl	_Z10funcionCPUPfS_S_S_S_S_iiPi
	.type	_Z10funcionCPUPfS_S_S_S_S_iiPi, @function
_Z10funcionCPUPfS_S_S_S_S_iiPi:
.LFB3638:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rdi, %r11
	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
	movq	%rcx, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rdx, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movl	56(%rsp), %edi
	movq	%rsi, %rbx
	xorl	%esi, %esi
	movl	64(%rsp), %ecx
	movl	%edi, %eax
	movl	%edi, %r10d
	cltd
	imull	%edi, %r10d
	idivl	%ecx
	movl	%eax, %r13d
.L34:
	movl	%esi, %eax
	cmpl	%esi, %r10d
	jle	.L37
	cltd
	idivl	%edi
	movl	%eax, -8(%rsp)
	movl	%edx, %eax
	movl	%edx, -4(%rsp)
	cltd
	idivl	%ecx
	movl	%eax, %r14d
	movl	-8(%rsp), %eax
	cltd
	idivl	%ecx
	movq	72(%rsp), %rdx
	movl	%eax, %r15d
	movl	%r13d, %eax
	imull	%r15d, %eax
	addl	%r14d, %eax
	cltq
	movl	(%rdx,%rax,4), %eax
	cltd
	idivl	%r13d
	subl	%r15d, %eax
	movl	-8(%rsp), %r15d
	subl	%r14d, %edx
	movl	-4(%rsp), %r14d
	imull	%ecx, %eax
	imull	%ecx, %edx
	addl	%r15d, %eax
	imull	%edi, %eax
	addl	%r14d, %edx
	addl	%edx, %eax
	cltq
	movss	(%r11,%rax,4), %xmm0
	movss	%xmm0, (%r12,%rsi,4)
	movss	(%rbx,%rax,4), %xmm0
	movss	%xmm0, (%r8,%rsi,4)
	movss	0(%rbp,%rax,4), %xmm0
	movss	%xmm0, (%r9,%rsi,4)
	incq	%rsi
	jmp	.L34
.L37:
	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
.LFE3638:
	.size	_Z10funcionCPUPfS_S_S_S_S_iiPi, .-_Z10funcionCPUPfS_S_S_S_S_iiPi
	.globl	_Z42__device_stub__Z9kernelGPUPfS_S_S_S_S_iiPiPfS_S_S_S_S_iiPi
	.type	_Z42__device_stub__Z9kernelGPUPfS_S_S_S_S_iiPiPfS_S_S_S_S_iiPi, @function
_Z42__device_stub__Z9kernelGPUPfS_S_S_S_S_iiPiPfS_S_S_S_S_iiPi:
.LFB3664:
	.cfi_startproc
	endbr64
	subq	$216, %rsp
	.cfi_def_cfa_offset 224
	movq	240(%rsp), %rax
	movq	%rdi, 56(%rsp)
	leaq	80(%rsp), %rdi
	movq	%rsi, 48(%rsp)
	leaq	92(%rsp), %rsi
	movq	%rdx, 40(%rsp)
	leaq	64(%rsp), %rdx
	movq	%rcx, 32(%rsp)
	leaq	72(%rsp), %rcx
	movq	%r8, 24(%rsp)
	movq	%r9, 16(%rsp)
	movq	%rax, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 200(%rsp)
	xorl	%eax, %eax
	leaq	56(%rsp), %rax
	movl	$1, 88(%rsp)
	movq	%rax, 128(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	224(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	232(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 192(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 80(%rsp)
	movq	%rax, 92(%rsp)
	movl	$1, 100(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L38
	pushq	72(%rsp)
	.cfi_def_cfa_offset 232
	leaq	_Z9kernelGPUPfS_S_S_S_S_iiPi(%rip), %rdi
	pushq	72(%rsp)
	.cfi_def_cfa_offset 240
	movq	108(%rsp), %rcx
	movl	116(%rsp), %r8d
	movq	96(%rsp), %rsi
	movl	104(%rsp), %edx
	leaq	144(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 232
	popq	%rdx
	.cfi_def_cfa_offset 224
.L38:
	movq	200(%rsp), %rax
	subq	%fs:40, %rax
	je	.L40
	call	__stack_chk_fail@PLT
.L40:
	addq	$216, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3664:
	.size	_Z42__device_stub__Z9kernelGPUPfS_S_S_S_S_iiPiPfS_S_S_S_S_iiPi, .-_Z42__device_stub__Z9kernelGPUPfS_S_S_S_S_iiPiPfS_S_S_S_S_iiPi
	.globl	_Z9kernelGPUPfS_S_S_S_S_iiPi
	.type	_Z9kernelGPUPfS_S_S_S_S_iiPi, @function
_Z9kernelGPUPfS_S_S_S_S_iiPi:
.LFB3665:
	.cfi_startproc
	endbr64
	jmp	_Z42__device_stub__Z9kernelGPUPfS_S_S_S_S_iiPiPfS_S_S_S_S_iiPi
	.cfi_endproc
.LFE3665:
	.size	_Z9kernelGPUPfS_S_S_S_S_iiPi, .-_Z9kernelGPUPfS_S_S_S_S_iiPi
	.section	.rodata.str1.1
.LC10:
	.string	"Tiempo CPU: "
.LC11:
	.string	"[ms]"
.LC13:
	.string	"Tiempo GPU: "
	.section	.rodata
	.align 32
.LC0:
	.string	"img100x100.txt"
	.string	""
	.zero	14
	.string	"img100x100CPU.txt"
	.string	""
	.zero	11
	.string	"img100x100GPU.txt"
	.string	""
	.zero	11
	.string	"img200x200.txt"
	.string	""
	.zero	14
	.string	"img200x200CPU.txt"
	.string	""
	.zero	11
	.string	"img200x200GPU.txt"
	.string	""
	.zero	11
	.string	"img400x400.txt"
	.string	""
	.zero	14
	.string	"img400x400CPU.txt"
	.string	""
	.zero	11
	.string	"img400x400GPU.txt"
	.string	""
	.zero	11
	.string	"img800x800.txt"
	.string	""
	.zero	14
	.string	"img800x800CPU.txt"
	.string	""
	.zero	11
	.string	"img800x800GPU.txt"
	.string	""
	.zero	11
	.string	"img1600x1600.txt"
	.string	""
	.zero	12
	.string	"img1600x1600CPU.txt"
	.string	""
	.zero	9
	.string	"img1600x1600GPU.txt"
	.string	""
	.zero	9
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB3639:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	leaq	.LC0(%rip), %rsi
	movl	$450, %ecx
	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	$664, %rsp
	.cfi_def_cfa_offset 720
	movq	%fs:40, %rax
	movq	%rax, 648(%rsp)
	xorl	%eax, %eax
	leaq	198(%rsp), %rdi
	leaq	648(%rsp), %rax
	rep movsb
	movq	%rax, 24(%rsp)
	leaq	198(%rsp), %rbx
.L49:
	leaq	56(%rsp), %rcx
	leaq	104(%rsp), %rsi
	leaq	112(%rsp), %rdx
	leaq	96(%rsp), %rdi
	pushq	%r8
	.cfi_def_cfa_offset 728
	pushq	%rbx
	.cfi_def_cfa_offset 736
	leaq	76(%rsp), %r8
	leaq	96(%rsp), %r9
	call	_Z4ReadPPfS0_S0_PiS1_PS1_PKc
	movl	72(%rsp), %edi
	movl	%edi, %eax
	imull	%edi, %edi
	cltd
	idivl	76(%rsp)
	movslq	%edi, %rdi
	salq	$2, %rdi
	movl	%eax, 60(%rsp)
	popq	%r9
	.cfi_def_cfa_offset 728
	popq	%r10
	.cfi_def_cfa_offset 720
	call	_Znam@PLT
	movl	56(%rsp), %edi
	movq	%rax, %r12
	imull	%edi, %edi
	movslq	%edi, %rdi
	salq	$2, %rdi
	call	_Znam@PLT
	movl	56(%rsp), %edi
	movq	%rax, %r13
	imull	%edi, %edi
	movslq	%edi, %rdi
	salq	$2, %rdi
	call	_Znam@PLT
	movq	%rax, 16(%rsp)
	call	clock@PLT
	movq	80(%rsp), %r14
	movq	%r13, %r8
	movq	%r12, %rcx
	movq	%rax, %rbp
	movq	112(%rsp), %rax
	movq	96(%rsp), %r15
	movq	%rax, (%rsp)
	movq	104(%rsp), %rax
	movq	%r15, %rdi
	movq	%rax, 8(%rsp)
	pushq	%r11
	.cfi_def_cfa_offset 728
	pushq	%r14
	.cfi_def_cfa_offset 736
	movl	76(%rsp), %eax
	pushq	%rax
	.cfi_def_cfa_offset 744
	movl	80(%rsp), %eax
	pushq	%rax
	.cfi_def_cfa_offset 752
	movq	48(%rsp), %r9
	movq	32(%rsp), %rdx
	movq	40(%rsp), %rsi
	call	_Z10funcionCPUPfS_S_S_S_S_iiPi
	addq	$32, %rsp
	.cfi_def_cfa_offset 720
	call	clock@PLT
	leaq	.LC10(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	subq	%rbp, %rax
	cvtsi2sdq	%rax, %xmm0
	mulsd	.LC8(%rip), %xmm0
	divsd	.LC9(%rip), %xmm0
	movsd	%xmm0, 32(%rsp)
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsd	32(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	leaq	.LC11(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	call	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@PLT
	movl	56(%rsp), %ecx
	leaq	30(%rbx), %r9
	movq	%r13, %rsi
	movq	16(%rsp), %rdx
	movq	%r12, %rdi
	movl	%ecx, %r8d
	call	_Z5WritePfS_S_iiPKc
	movq	%r12, %rdi
	call	_ZdaPv@PLT
	movq	%r13, %rdi
	call	_ZdaPv@PLT
	movq	16(%rsp), %rdi
	call	_ZdaPv@PLT
	movl	56(%rsp), %ebp
	cvtsi2ssl	%ebp, %xmm0
	imull	%ebp, %ebp
	mulss	%xmm0, %xmm0
	mulss	.LC12(%rip), %xmm0
	call	ceilf@PLT
	movslq	%ebp, %rsi
	leaq	120(%rsp), %rdi
	salq	$2, %rsi
	movss	%xmm0, 16(%rsp)
	call	cudaMalloc@PLT
	movl	56(%rsp), %esi
	leaq	128(%rsp), %rdi
	imull	%esi, %esi
	movslq	%esi, %rsi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	movl	56(%rsp), %esi
	leaq	136(%rsp), %rdi
	imull	%esi, %esi
	movslq	%esi, %rsi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	movl	44(%rsp), %r10d
	leaq	88(%rsp), %rdi
	imull	%r10d, %r10d
	movslq	%r10d, %rbp
	salq	$2, %rbp
	movq	%rbp, %rsi
	call	cudaMalloc@PLT
	movl	56(%rsp), %edx
	movq	120(%rsp), %rdi
	movq	%r15, %rsi
	movl	$1, %ecx
	imull	%edx, %edx
	movslq	%edx, %rdx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movl	56(%rsp), %edx
	movq	8(%rsp), %rsi
	movl	$1, %ecx
	movq	128(%rsp), %rdi
	imull	%edx, %edx
	movslq	%edx, %rdx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movl	56(%rsp), %edx
	movq	(%rsp), %rsi
	movl	$1, %ecx
	movq	136(%rsp), %rdi
	imull	%edx, %edx
	movslq	%edx, %rdx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movq	88(%rsp), %rdi
	movq	%rbp, %rdx
	movq	%r14, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movl	56(%rsp), %esi
	leaq	144(%rsp), %rdi
	imull	%esi, %esi
	movslq	%esi, %rsi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	movl	56(%rsp), %esi
	leaq	152(%rsp), %rdi
	imull	%esi, %esi
	movslq	%esi, %rsi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	movl	56(%rsp), %esi
	leaq	160(%rsp), %rdi
	imull	%esi, %esi
	movslq	%esi, %rsi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	leaq	64(%rsp), %rdi
	call	cudaEventCreate@PLT
	leaq	72(%rsp), %rdi
	call	cudaEventCreate@PLT
	movq	64(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movl	$16777217, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movss	16(%rsp), %xmm0
	salq	$8, %rdx
	movl	$1, %ecx
	cvttss2sil	%xmm0, %eax
	movl	%eax, 172(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 176(%rsp)
	movq	172(%rsp), %rdi
	movl	180(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L44
	pushq	%rax
	.cfi_def_cfa_offset 728
	pushq	96(%rsp)
	.cfi_def_cfa_offset 736
	movl	76(%rsp), %eax
	pushq	%rax
	.cfi_def_cfa_offset 744
	movl	80(%rsp), %eax
	pushq	%rax
	.cfi_def_cfa_offset 752
	movq	192(%rsp), %r9
	movq	184(%rsp), %r8
	movq	176(%rsp), %rcx
	movq	168(%rsp), %rdx
	movq	160(%rsp), %rsi
	movq	152(%rsp), %rdi
	call	_Z42__device_stub__Z9kernelGPUPfS_S_S_S_S_iiPiPfS_S_S_S_S_iiPi
	addq	$32, %rsp
	.cfi_def_cfa_offset 720
.L44:
	movq	72(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	72(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	72(%rsp), %rdx
	movq	64(%rsp), %rsi
	leaq	52(%rsp), %rdi
	call	cudaEventElapsedTime@PLT
	leaq	.LC13(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	cvtss2sd	52(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	leaq	.LC11(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	call	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@PLT
	movl	56(%rsp), %edi
	imull	%edi, %edi
	movslq	%edi, %rdi
	salq	$2, %rdi
	call	_Znam@PLT
	movl	56(%rsp), %edi
	movq	%rax, %r13
	imull	%edi, %edi
	movslq	%edi, %rdi
	salq	$2, %rdi
	call	_Znam@PLT
	movl	56(%rsp), %edi
	movq	%rax, %r12
	imull	%edi, %edi
	movslq	%edi, %rdi
	salq	$2, %rdi
	call	_Znam@PLT
	movl	56(%rsp), %edx
	movl	$2, %ecx
	movq	%r13, %rdi
	movq	144(%rsp), %rsi
	movq	%rax, %rbp
	imull	%edx, %edx
	movslq	%edx, %rdx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movl	56(%rsp), %edx
	movl	$2, %ecx
	movq	%r12, %rdi
	movq	152(%rsp), %rsi
	imull	%edx, %edx
	movslq	%edx, %rdx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movl	56(%rsp), %edx
	movl	$2, %ecx
	movq	%rbp, %rdi
	movq	160(%rsp), %rsi
	imull	%edx, %edx
	movslq	%edx, %rdx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movl	56(%rsp), %ecx
	leaq	60(%rbx), %r9
	movq	%rbp, %rdx
	movq	%r12, %rsi
	movq	%r13, %rdi
	movl	%ecx, %r8d
	call	_Z5WritePfS_S_iiPKc
	movq	120(%rsp), %rdi
	call	cudaFree@PLT
	movq	128(%rsp), %rdi
	call	cudaFree@PLT
	movq	136(%rsp), %rdi
	call	cudaFree@PLT
	movq	144(%rsp), %rdi
	call	cudaFree@PLT
	movq	152(%rsp), %rdi
	call	cudaFree@PLT
	movq	160(%rsp), %rdi
	call	cudaFree@PLT
	movq	88(%rsp), %rdi
	call	cudaFree@PLT
	testq	%r15, %r15
	je	.L45
	movq	%r15, %rdi
	call	_ZdaPv@PLT
.L45:
	cmpq	$0, 8(%rsp)
	je	.L46
	movq	8(%rsp), %rdi
	call	_ZdaPv@PLT
.L46:
	cmpq	$0, (%rsp)
	je	.L47
	movq	(%rsp), %rdi
	call	_ZdaPv@PLT
.L47:
	movq	%r13, %rdi
	call	_ZdaPv@PLT
	movq	%r12, %rdi
	call	_ZdaPv@PLT
	movq	%rbp, %rdi
	call	_ZdaPv@PLT
	testq	%r14, %r14
	je	.L48
	movq	%r14, %rdi
	call	_ZdaPv@PLT
.L48:
	movq	24(%rsp), %rax
	addq	$90, %rbx
	cmpq	%rax, %rbx
	jne	.L49
	movq	648(%rsp), %rax
	subq	%fs:40, %rax
	je	.L50
	call	__stack_chk_fail@PLT
.L50:
	addq	$664, %rsp
	.cfi_def_cfa_offset 56
	xorl	%eax, %eax
	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
.LFE3639:
	.size	main, .-main
	.section	.rodata.str1.1
.LC15:
	.string	"_Z9kernelGPUPfS_S_S_S_S_iiPi"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3667:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 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	.LC15(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z9kernelGPUPfS_S_S_S_S_iiPi(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	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	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE3667:
	.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.cst8,"aM",@progbits,8
	.align 8
.LC8:
	.long	0
	.long	1083129856
	.align 8
.LC9:
	.long	0
	.long	1093567616
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC12:
	.long	998244352
	.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:

// Basado en el codigo entregado durante clases practicas

#include <iostream>
#include <time.h>
#include <hip/hip_runtime.h>

/*
 *  Lectura Archivo
 */
void Read(float** R, float** G, float** B, int *N, int *S, int** posicionesRompecabezas, const char *filename) {
    printf("Leemos el archivo %s!\n", filename); 
    FILE *fp;
    fp = fopen(filename, "r");
    fscanf(fp, "%d %d\n", N, S);

    int P = (*N) / (*S);
    int imsize = (*N) * (*N);

    float* R1 = new float[imsize];
    float* G1 = new float[imsize];
    float* B1 = new float[imsize];
    int *posicionesTemp = new int[P*P];

    for(int i = 0; i < P*P; i++)
	    fscanf(fp, "%d ", &(posicionesTemp[i]));
	for(int i = 0; i < imsize; i++)
	    fscanf(fp, "%f ", &(R1[i]));
	for(int i = 0; i < imsize; i++)
	    fscanf(fp, "%f ", &(G1[i]));
	for(int i = 0; i < imsize; i++)
	    fscanf(fp, "%f ", &(B1[i]));
    fclose(fp);

    *R = R1; *G = G1; *B = B1; 
    *posicionesRompecabezas = posicionesTemp;
}

/*
 *  Escritura Archivo
 */
void Write(float* R, float* G, float* B, 
	       int M, int N, const char *filename) {
    FILE *fp;
    fp = fopen(filename, "w");
    fprintf(fp, "%d %d\n", M, N);
    for(int i = 0; i < M*N-1; i++)
        fprintf(fp, "%f ", R[i]);
    fprintf(fp, "%f\n", R[M*N-1]);
    for(int i = 0; i < M*N-1; i++)
        fprintf(fp, "%f ", G[i]);
    fprintf(fp, "%f\n", G[M*N-1]);
    for(int i = 0; i < M*N-1; i++)
        fprintf(fp, "%f ", B[i]);
    fprintf(fp, "%f\n", B[M*N-1]);
    fclose(fp);
}


__host__ __device__ int findNewPosition(int P, int position, int* posiciones){

	int newPosition;

	for (int j = 0; j < P*P; j++) {
    	if(position == posiciones[j]) {
        	newPosition = j;
        	break;
    	}
	}

	return newPosition;
}

/*
 *  Procesamiento Imagen CPU
 */
void funcionCPU( float* R, float* G, float* B, float* Rout, float* Gout, float* Bout, 
				 int N, int S, int* posicionesRompecabezas) {

	int P = N/S;
	
	for (int i = 0; i < N*N; i++){
		int xOriginal = i % N;
		int yOriginal = i / N;
		int xRompecabezas = xOriginal / S;
		int yRompecabezas = yOriginal / S;
		int indiceRompecabezas = xRompecabezas + yRompecabezas * P;
		int indiceRompecabezasMezclado =  posicionesRompecabezas[indiceRompecabezas];
		int xRompecabezasMezclado = indiceRompecabezasMezclado % P;
		int yRompecabezasMezclado = indiceRompecabezasMezclado / P;
		int nuevoI = (xOriginal + S * (xRompecabezasMezclado - xRompecabezas)) + N * (yOriginal + S * (yRompecabezasMezclado - yRompecabezas));

		Rout[i] = R[nuevoI];
		Gout[i] = G[nuevoI];
		Bout[i] = B[nuevoI];
	}
}

/*
 *  Procesamiento Imagen GPU
 */
__global__ void kernelGPU( float* R, float* G, float* B, float* Rout, float* Gout, float* Bout, 
						   int N, int S, int* posicionesRompecabezas){
	int i = threadIdx.x + blockDim.x * blockIdx.x;
	int P = N/S;

	if(i < N*N){
		int xOriginal = i % N;
		int yOriginal = i / N;
		int xRompecabezas = xOriginal / S;
		int yRompecabezas = yOriginal / S;
		int indiceRompecabezas = xRompecabezas + yRompecabezas * P;
		int indiceRompecabezasMezclado =  posicionesRompecabezas[indiceRompecabezas];
		int xRompecabezasMezclado = indiceRompecabezasMezclado % P;
		int yRompecabezasMezclado = indiceRompecabezasMezclado / P;
		int nuevoI = (xOriginal + S * (xRompecabezasMezclado - xRompecabezas)) + N * (yOriginal + S * (yRompecabezasMezclado - yRompecabezas));

		Rout[i] = R[nuevoI];
		Gout[i] = G[nuevoI];
		Bout[i] = B[nuevoI];
	} 
}

/*
 *  Codigo Principal
 */
int main(int argc, char **argv){

    /*
     *  Inicializacion
     */
	clock_t t1, t2;
	hipEvent_t ct1, ct2;
	double ms;
	float dt;
	int N, S;
	int P;
	int *posiciones;
	int *posicionesDev;
    float *Rhost, *Ghost, *Bhost;
    float *Rhostout, *Ghostout, *Bhostout;
    float *Rdev, *Gdev, *Bdev;
    float *Rdevout, *Gdevout, *Bdevout;
    
    //cambiar con los nombres de los archivos correctos
    char names[5][3][30] = {{"img100x100.txt\0", "img100x100CPU.txt\0", "img100x100GPU.txt\0"}, 
    	{"img200x200.txt\0", "img200x200CPU.txt\0", "img200x200GPU.txt\0"},
    	{"img400x400.txt\0", "img400x400CPU.txt\0", "img400x400GPU.txt\0"},
    	{"img800x800.txt\0", "img800x800CPU.txt\0", "img800x800GPU.txt\0"},
    	{"img1600x1600.txt\0", "img1600x1600CPU.txt\0", "img1600x1600GPU.txt\0"}};

    for (int i=0; i<5; i++){
	    Read(&Rhost, &Ghost, &Bhost, &N, &S, &posiciones, names[i][0]); //leemos archivo y reservamos memoria

	    P = N/S;

	    /*
	     *  Parte CPU
	     */
	    Rhostout = new float[N*N];	//reservamos memoria
	    Ghostout = new float[N*N];
	    Bhostout = new float[N*N];

	    t1 = clock();
	    funcionCPU(Rhost, Ghost, Bhost, Rhostout, Ghostout, Bhostout, N, S, posiciones); // Agregar parametros!
	    t2 = clock();
	    ms = 1000.0 * (double)(t2 - t1) / CLOCKS_PER_SEC;
	    std::cout << "Tiempo CPU: " << ms << "[ms]" << std::endl;
	    Write(Rhostout, Ghostout, Bhostout, N, N, names[i][1]);

	    delete[] Rhostout; delete[] Ghostout; delete[] Bhostout;
	    
	    /*
	     *  Parte GPU
	     */

	    int grid_size, block_size = 256;
	    grid_size = (int)ceil((float) N * N / block_size);
	        
	    hipMalloc((void**)&Rdev, N * N * sizeof(float));
	    hipMalloc((void**)&Gdev, N * N * sizeof(float));
	    hipMalloc((void**)&Bdev, N * N * sizeof(float));
	    hipMalloc((void**)&posicionesDev, P * P * sizeof(float));
	    hipMemcpy(Rdev, Rhost, N * N * sizeof(float), hipMemcpyHostToDevice);
	    hipMemcpy(Gdev, Ghost, N * N * sizeof(float), hipMemcpyHostToDevice);
	    hipMemcpy(Bdev, Bhost, N * N * sizeof(float), hipMemcpyHostToDevice);
	    hipMemcpy(posicionesDev, posiciones, P * P * sizeof(float), hipMemcpyHostToDevice);
	        
	    hipMalloc((void**)&Rdevout, N * N * sizeof(float));
	    hipMalloc((void**)&Gdevout, N * N * sizeof(float));
	    hipMalloc((void**)&Bdevout, N * N * sizeof(float));
	    
	    hipEventCreate(&ct1);
	    hipEventCreate(&ct2);
	    hipEventRecord(ct1);
	    kernelGPU<<<grid_size, block_size>>>(Rdev, Gdev, Bdev, Rdevout, Gdevout, Bdevout, N, S, posicionesDev); // Agregar parametros!
	    hipEventRecord(ct2);
	    hipEventSynchronize(ct2);
	    hipEventElapsedTime(&dt, ct1, ct2);
	    std::cout << "Tiempo GPU: " << dt << "[ms]" << std::endl;

	    Rhostout = new float[N*N];
	    Ghostout = new float[N*N];
	    Bhostout = new float[N*N];
	    hipMemcpy(Rhostout, Rdevout, N * N * sizeof(float), hipMemcpyDeviceToHost);
	    hipMemcpy(Ghostout, Gdevout, N * N * sizeof(float), hipMemcpyDeviceToHost);
	    hipMemcpy(Bhostout, Bdevout, N * N * sizeof(float), hipMemcpyDeviceToHost);
	    Write(Rhostout, Ghostout, Bhostout, N, N, names[i][2]);

	    hipFree(Rdev); hipFree(Gdev); hipFree(Bdev);
    	hipFree(Rdevout); hipFree(Gdevout); hipFree(Bdevout);
    	hipFree(posicionesDev);
    	delete[] Rhost; delete[] Ghost; delete[] Bhost;
    	delete[] Rhostout; delete[] Ghostout; delete[] Bhostout;
    	delete[] posiciones;
	}
	return 0;
}
	.text
	.file	"cod.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.globl	_Z4ReadPPfS0_S0_PiS1_PS1_PKc    # -- Begin function _Z4ReadPPfS0_S0_PiS1_PS1_PKc
	.type	_Z4ReadPPfS0_S0_PiS1_PS1_PKc,@function
_Z4ReadPPfS0_S0_PiS1_PS1_PKc:           # @_Z4ReadPPfS0_S0_PiS1_PS1_PKc
	.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	$72, %rsp
	.cfi_def_cfa_offset 128
	.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	%r9, 64(%rsp)                   # 8-byte Spill
	movq	%r8, %rbx
	movq	%rcx, %r14
	movq	%rdx, 56(%rsp)                  # 8-byte Spill
	movq	%rsi, 48(%rsp)                  # 8-byte Spill
	movq	%rdi, 40(%rsp)                  # 8-byte Spill
	movq	128(%rsp), %r15
	movl	$.L.str, %edi
	movq	%r15, %rsi
	xorl	%eax, %eax
	callq	printf
	movl	$.L.str.1, %esi
	movq	%r15, %rdi
	callq	fopen
	movq	%rax, %r13
	movl	$.L.str.2, %esi
	movq	%rax, %rdi
	movq	%r14, %rdx
	movq	%rbx, %rcx
	xorl	%eax, %eax
	callq	__isoc23_fscanf
	movl	(%r14), %ebp
	movl	%ebp, %eax
	cltd
	idivl	(%rbx)
	movl	%eax, %r14d
	movl	%ebp, %r12d
	imull	%r12d, %r12d
	leaq	(,%r12,4), %r15
	movq	%r15, %rdi
	callq	_Znam
	movq	%rax, 24(%rsp)                  # 8-byte Spill
	movq	%r15, %rdi
	callq	_Znam
	movq	%rax, 16(%rsp)                  # 8-byte Spill
	movq	%r15, %rdi
	callq	_Znam
	movq	%rax, 8(%rsp)                   # 8-byte Spill
	movl	%r14d, %ebx
	imull	%ebx, %ebx
	leaq	(,%rbx,4), %rdi
	callq	_Znam
	movq	%rax, %r15
	testl	%r14d, %r14d
	je	.LBB0_3
# %bb.1:                                # %.lr.ph.preheader
	cmpl	$1, %ebx
	adcl	$0, %ebx
	movq	%r15, %r14
.LBB0_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	$.L.str.3, %esi
	movq	%r13, %rdi
	movq	%r14, %rdx
	xorl	%eax, %eax
	callq	__isoc23_fscanf
	addq	$4, %r14
	decq	%rbx
	jne	.LBB0_2
.LBB0_3:                                # %.preheader52
	movq	%r15, 32(%rsp)                  # 8-byte Spill
	testl	%ebp, %ebp
	je	.LBB0_12
# %bb.4:                                # %.lr.ph55.preheader
	movl	%ebp, %r15d
	cmpl	$1, %r12d
	movl	%r12d, %ebp
	adcl	$0, %ebp
	movq	%rbp, %rbx
	movq	24(%rsp), %r14                  # 8-byte Reload
.LBB0_5:                                # %.lr.ph55
                                        # =>This Inner Loop Header: Depth=1
	movl	$.L.str.4, %esi
	movq	%r13, %rdi
	movq	%r14, %rdx
	xorl	%eax, %eax
	callq	__isoc23_fscanf
	addq	$4, %r14
	decq	%rbx
	jne	.LBB0_5
# %bb.6:                                # %.preheader51
	testl	%r15d, %r15d
	je	.LBB0_12
# %bb.7:
	movq	16(%rsp), %r14                  # 8-byte Reload
.LBB0_8:                                # %.lr.ph57
                                        # =>This Inner Loop Header: Depth=1
	movl	$.L.str.4, %esi
	movq	%r13, %rdi
	movq	%r14, %rdx
	xorl	%eax, %eax
	callq	__isoc23_fscanf
	addq	$4, %r14
	decq	%rbp
	jne	.LBB0_8
# %bb.9:                                # %.preheader
	testl	%r15d, %r15d
	je	.LBB0_12
# %bb.10:                               # %.lr.ph59.preheader
	cmpl	$1, %r12d
	adcl	$0, %r12d
	movq	8(%rsp), %r14                   # 8-byte Reload
.LBB0_11:                               # %.lr.ph59
                                        # =>This Inner Loop Header: Depth=1
	movl	$.L.str.4, %esi
	movq	%r13, %rdi
	movq	%r14, %rdx
	xorl	%eax, %eax
	callq	__isoc23_fscanf
	addq	$4, %r14
	decq	%r12
	jne	.LBB0_11
.LBB0_12:                               # %._crit_edge
	movq	%r13, %rdi
	callq	fclose
	movq	40(%rsp), %rax                  # 8-byte Reload
	movq	24(%rsp), %rcx                  # 8-byte Reload
	movq	%rcx, (%rax)
	movq	48(%rsp), %rax                  # 8-byte Reload
	movq	16(%rsp), %rcx                  # 8-byte Reload
	movq	%rcx, (%rax)
	movq	56(%rsp), %rax                  # 8-byte Reload
	movq	8(%rsp), %rcx                   # 8-byte Reload
	movq	%rcx, (%rax)
	movq	64(%rsp), %rax                  # 8-byte Reload
	movq	32(%rsp), %rcx                  # 8-byte Reload
	movq	%rcx, (%rax)
	addq	$72, %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_end0:
	.size	_Z4ReadPPfS0_S0_PiS1_PS1_PKc, .Lfunc_end0-_Z4ReadPPfS0_S0_PiS1_PS1_PKc
	.cfi_endproc
                                        # -- End function
	.globl	_Z5WritePfS_S_iiPKc             # -- Begin function _Z5WritePfS_S_iiPKc
	.type	_Z5WritePfS_S_iiPKc,@function
_Z5WritePfS_S_iiPKc:                    # @_Z5WritePfS_S_iiPKc
	.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
	movl	%r8d, %r15d
	movl	%ecx, %ebp
	movq	%rdx, %r14
	movq	%rsi, %r12
	movq	%rdi, %r13
	movl	$.L.str.5, %esi
	movq	%r9, %rdi
	callq	fopen
	movq	%rax, %rbx
	movl	$.L.str.2, %esi
	movq	%rax, %rdi
	movl	%ebp, %edx
	movl	%r15d, %ecx
	xorl	%eax, %eax
	callq	fprintf
	imull	%ebp, %r15d
	leal	-1(%r15), %eax
	movl	%eax, 8(%rsp)                   # 4-byte Spill
	movl	%eax, %ebp
	movq	%r15, 16(%rsp)                  # 8-byte Spill
	cmpl	$2, %r15d
	jl	.LBB1_3
# %bb.1:                                # %.lr.ph.preheader
	xorl	%r15d, %r15d
.LBB1_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	xorps	%xmm0, %xmm0
	cvtss2sd	(%r13,%r15,4), %xmm0
	movl	$.L.str.4, %esi
	movq	%rbx, %rdi
	movb	$1, %al
	callq	fprintf
	incq	%r15
	cmpq	%r15, %rbp
	jne	.LBB1_2
.LBB1_3:                                # %._crit_edge
	movslq	8(%rsp), %r15                   # 4-byte Folded Reload
	xorps	%xmm0, %xmm0
	cvtss2sd	(%r13,%r15,4), %xmm0
	movl	$.L.str.6, %esi
	movq	%rbx, %rdi
	movb	$1, %al
	callq	fprintf
	movq	16(%rsp), %r13                  # 8-byte Reload
	cmpl	$1, %r13d
	jle	.LBB1_4
# %bb.9:                                # %.lr.ph43.preheader
	movq	%r15, 8(%rsp)                   # 8-byte Spill
	xorl	%r15d, %r15d
.LBB1_10:                               # %.lr.ph43
                                        # =>This Inner Loop Header: Depth=1
	xorps	%xmm0, %xmm0
	cvtss2sd	(%r12,%r15,4), %xmm0
	movl	$.L.str.4, %esi
	movq	%rbx, %rdi
	movb	$1, %al
	callq	fprintf
	incq	%r15
	cmpq	%r15, %rbp
	jne	.LBB1_10
# %bb.6:                                # %._crit_edge44
	movq	8(%rsp), %rax                   # 8-byte Reload
	xorps	%xmm0, %xmm0
	cvtss2sd	(%r12,%rax,4), %xmm0
	movq	%rax, %r12
	movl	$.L.str.6, %esi
	movq	%rbx, %rdi
	movb	$1, %al
	callq	fprintf
	cmpl	$2, %r13d
	jl	.LBB1_5
# %bb.7:                                # %.lr.ph47.preheader
	xorl	%r15d, %r15d
.LBB1_8:                                # %.lr.ph47
                                        # =>This Inner Loop Header: Depth=1
	xorps	%xmm0, %xmm0
	cvtss2sd	(%r14,%r15,4), %xmm0
	movl	$.L.str.4, %esi
	movq	%rbx, %rdi
	movb	$1, %al
	callq	fprintf
	incq	%r15
	cmpq	%r15, %rbp
	jne	.LBB1_8
	jmp	.LBB1_5
.LBB1_4:                                # %._crit_edge44.thread
	xorps	%xmm0, %xmm0
	cvtss2sd	(%r12,%r15,4), %xmm0
	movl	$.L.str.6, %esi
	movq	%rbx, %rdi
	movb	$1, %al
	callq	fprintf
	movq	%r15, %r12
.LBB1_5:                                # %._crit_edge48
	xorps	%xmm0, %xmm0
	cvtss2sd	(%r14,%r12,4), %xmm0
	movl	$.L.str.6, %esi
	movq	%rbx, %rdi
	movb	$1, %al
	callq	fprintf
	movq	%rbx, %rdi
	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
	jmp	fclose                          # TAILCALL
.Lfunc_end1:
	.size	_Z5WritePfS_S_iiPKc, .Lfunc_end1-_Z5WritePfS_S_iiPKc
	.cfi_endproc
                                        # -- End function
	.globl	_Z15findNewPositioniiPi         # -- Begin function _Z15findNewPositioniiPi
	.type	_Z15findNewPositioniiPi,@function
_Z15findNewPositioniiPi:                # @_Z15findNewPositioniiPi
	.cfi_startproc
# %bb.0:
	testl	%edi, %edi
	je	.LBB2_5
# %bb.1:                                # %.lr.ph.preheader
	movl	%edi, %eax
	imull	%eax, %eax
	cmpl	$1, %eax
	adcl	$0, %eax
	xorl	%ecx, %ecx
.LBB2_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	cmpl	%esi, (%rdx,%rcx,4)
	je	.LBB2_6
# %bb.3:                                #   in Loop: Header=BB2_2 Depth=1
	incq	%rcx
	cmpq	%rcx, %rax
	jne	.LBB2_2
# %bb.4:                                # %._crit_edge
                                        # kill: def $eax killed $eax killed $rax
	retq
.LBB2_5:
	xorl	%eax, %eax
                                        # kill: def $eax killed $eax killed $rax
	retq
.LBB2_6:                                # %._crit_edge.loopexit.split.loop.exit
	movl	%ecx, %eax
                                        # kill: def $eax killed $eax killed $rax
	retq
.Lfunc_end2:
	.size	_Z15findNewPositioniiPi, .Lfunc_end2-_Z15findNewPositioniiPi
	.cfi_endproc
                                        # -- End function
	.globl	_Z10funcionCPUPfS_S_S_S_S_iiPi  # -- Begin function _Z10funcionCPUPfS_S_S_S_S_iiPi
	.type	_Z10funcionCPUPfS_S_S_S_S_iiPi,@function
_Z10funcionCPUPfS_S_S_S_S_iiPi:         # @_Z10funcionCPUPfS_S_S_S_S_iiPi
	.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
	.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	%rdx, %r15
	movl	64(%rsp), %ebx
	movl	56(%rsp), %r11d
	movl	%r11d, %eax
	cltd
	idivl	%ebx
	movl	%eax, -12(%rsp)                 # 4-byte Spill
	testl	%r11d, %r11d
	je	.LBB3_3
# %bb.1:                                # %.lr.ph.preheader
	movl	56(%rsp), %eax
                                        # kill: def $eax killed $eax def $rax
	imull	%eax, %eax
	cmpl	$1, %eax
	adcl	$0, %eax
	movq	%rax, -8(%rsp)                  # 8-byte Spill
	xorl	%r14d, %r14d
.LBB3_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	%r14d, %eax
	cltd
	movl	56(%rsp), %r10d
	idivl	%r10d
	movq	%r9, %r10
	movq	%r8, %r9
	movq	%rcx, %r8
	movq	%rdi, %r11
	movq	%rsi, %rdi
	movq	%r15, %rsi
	movl	%edx, %r15d
	movl	%eax, %ebp
	movl	%edx, %eax
	cltd
	idivl	%ebx
	movl	%eax, %r12d
	movl	%ebp, %eax
	cltd
	idivl	%ebx
	movl	%eax, %r13d
	movl	-12(%rsp), %ecx                 # 4-byte Reload
	imull	%ecx, %eax
	addl	%r12d, %eax
	cltq
	movq	72(%rsp), %rdx
	movl	(%rdx,%rax,4), %eax
	cltd
	idivl	%ecx
	subl	%r12d, %edx
	imull	%ebx, %edx
	addl	%r15d, %edx
	movq	%rsi, %r15
	movq	%rdi, %rsi
	movq	%r11, %rdi
	movq	%r8, %rcx
	movq	%r9, %r8
	movq	%r10, %r9
	subl	%r13d, %eax
	imull	%ebx, %eax
	addl	%ebp, %eax
	imull	56(%rsp), %eax
	addl	%edx, %eax
	cltq
	movss	(%r11,%rax,4), %xmm0            # xmm0 = mem[0],zero,zero,zero
	movss	%xmm0, (%rcx,%r14,4)
	movss	(%rsi,%rax,4), %xmm0            # xmm0 = mem[0],zero,zero,zero
	movss	%xmm0, (%r8,%r14,4)
	movss	(%r15,%rax,4), %xmm0            # xmm0 = mem[0],zero,zero,zero
	movss	%xmm0, (%r10,%r14,4)
	incq	%r14
	cmpq	%r14, -8(%rsp)                  # 8-byte Folded Reload
	jne	.LBB3_2
.LBB3_3:                                # %._crit_edge
	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_end3:
	.size	_Z10funcionCPUPfS_S_S_S_S_iiPi, .Lfunc_end3-_Z10funcionCPUPfS_S_S_S_S_iiPi
	.cfi_endproc
                                        # -- End function
	.globl	_Z24__device_stub__kernelGPUPfS_S_S_S_S_iiPi # -- Begin function _Z24__device_stub__kernelGPUPfS_S_S_S_S_iiPi
	.type	_Z24__device_stub__kernelGPUPfS_S_S_S_S_iiPi,@function
_Z24__device_stub__kernelGPUPfS_S_S_S_S_iiPi: # @_Z24__device_stub__kernelGPUPfS_S_S_S_S_iiPi
	.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	$176, %rsp
	.cfi_def_cfa_offset 224
	.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	16(%rsp), %r8
	movq	%r9, (%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	224(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	232(%rsp), %rax
	movq	%rax, 56(%rbx)
	leaq	240(%rsp), %rax
	movq	%rax, 64(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%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	$_Z9kernelGPUPfS_S_S_S_S_iiPi, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$192, %rsp
	.cfi_adjust_cfa_offset -192
	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_end4:
	.size	_Z24__device_stub__kernelGPUPfS_S_S_S_S_iiPi, .Lfunc_end4-_Z24__device_stub__kernelGPUPfS_S_S_S_S_iiPi
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI5_0:
	.quad	0x408f400000000000              # double 1000
.LCPI5_1:
	.quad	0x412e848000000000              # double 1.0E+6
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0
.LCPI5_2:
	.long	0x3b800000                      # float 0.00390625
	.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	$632, %rsp                      # imm = 0x278
	.cfi_def_cfa_offset 688
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movabsq	$4294967296, %r14               # imm = 0x100000000
	leaq	176(%rsp), %rbx
	movl	$.L__const.main.names, %esi
	movl	$450, %edx                      # imm = 0x1C2
	movq	%rbx, %rdi
	callq	memcpy@PLT
	movl	$5, %eax
	leaq	256(%r14), %rcx
	movq	%rcx, 112(%rsp)                 # 8-byte Spill
.LBB5_1:                                # =>This Inner Loop Header: Depth=1
	movq	%rax, 136(%rsp)                 # 8-byte Spill
	subq	$8, %rsp
	.cfi_adjust_cfa_offset 8
	leaq	168(%rsp), %rdi
	leaq	160(%rsp), %rsi
	leaq	152(%rsp), %rdx
	leaq	8(%rsp), %rcx
	leaq	12(%rsp), %r8
	leaq	176(%rsp), %r9
	pushq	%rbx
	.cfi_adjust_cfa_offset 8
	callq	_Z4ReadPPfS0_S0_PiS1_PS1_PKc
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	(%rsp), %r15d
	movl	%r15d, %eax
	cltd
	idivl	4(%rsp)
	movl	%eax, %r14d
	imull	%r15d, %r15d
	shlq	$2, %r15
	movq	%r15, %rdi
	callq	_Znam
	movq	%rax, %r12
	movq	%r15, %rdi
	callq	_Znam
	movq	%rax, %r13
	movq	%r15, %rdi
	callq	_Znam
	movq	%rax, %rbp
	callq	clock
	movq	%rax, %r15
	movq	160(%rsp), %rdi
	movq	%rdi, 128(%rsp)                 # 8-byte Spill
	movq	152(%rsp), %rsi
	movq	%rsi, 88(%rsp)                  # 8-byte Spill
	movq	144(%rsp), %rdx
	movq	%rdx, 96(%rsp)                  # 8-byte Spill
	movl	(%rsp), %eax
	movl	4(%rsp), %r10d
	movq	168(%rsp), %r11
	movq	%r11, 120(%rsp)                 # 8-byte Spill
	subq	$8, %rsp
	.cfi_adjust_cfa_offset 8
	movq	%r12, %rcx
	movq	%r13, %r8
	movq	%rbp, %r9
	pushq	%r11
	.cfi_adjust_cfa_offset 8
	pushq	%r10
	.cfi_adjust_cfa_offset 8
	pushq	%rax
	.cfi_adjust_cfa_offset 8
	callq	_Z10funcionCPUPfS_S_S_S_S_iiPi
	addq	$32, %rsp
	.cfi_adjust_cfa_offset -32
	callq	clock
	subq	%r15, %rax
	cvtsi2sd	%rax, %xmm0
	mulsd	.LCPI5_0(%rip), %xmm0
	divsd	.LCPI5_1(%rip), %xmm0
	movsd	%xmm0, 8(%rsp)                  # 8-byte Spill
	movl	$_ZSt4cout, %edi
	movl	$.L.str.7, %esi
	movl	$12, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movsd	8(%rsp), %xmm0                  # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r15
	movl	$.L.str.8, %esi
	movl	$4, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	(%r15), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r15, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	(%rsp), %ecx
	movq	%rbx, 8(%rsp)                   # 8-byte Spill
	leaq	30(%rbx), %r9
	movq	%r12, %rdi
	movq	%r13, %rsi
	movq	%rbp, %rdx
	movl	%ecx, %r8d
	callq	_Z5WritePfS_S_iiPKc
	movq	%r12, %rdi
	callq	_ZdaPv
	movq	%r13, %rdi
	callq	_ZdaPv
	movq	%rbp, %rdi
	movq	120(%rsp), %rbp                 # 8-byte Reload
	movq	128(%rsp), %rbx                 # 8-byte Reload
	callq	_ZdaPv
	movl	(%rsp), %r15d
	xorps	%xmm0, %xmm0
	cvtsi2ss	%r15d, %xmm0
	mulss	%xmm0, %xmm0
	mulss	.LCPI5_2(%rip), %xmm0
	callq	ceilf@PLT
	cvttss2si	%xmm0, %r12d
	imull	%r15d, %r15d
	shlq	$2, %r15
	leaq	56(%rsp), %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	movl	(%rsp), %esi
	imull	%esi, %esi
	shlq	$2, %rsi
	leaq	48(%rsp), %rdi
	callq	hipMalloc
	movl	(%rsp), %esi
	imull	%esi, %esi
	shlq	$2, %rsi
	leaq	40(%rsp), %rdi
	callq	hipMalloc
	imull	%r14d, %r14d
	shlq	$2, %r14
	leaq	64(%rsp), %rdi
	movq	%r14, %rsi
	callq	hipMalloc
	movq	56(%rsp), %rdi
	movl	(%rsp), %edx
	imull	%edx, %edx
	shlq	$2, %rdx
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	48(%rsp), %rdi
	movl	(%rsp), %edx
	imull	%edx, %edx
	shlq	$2, %rdx
	movq	88(%rsp), %rsi                  # 8-byte Reload
	movl	$1, %ecx
	callq	hipMemcpy
	movq	40(%rsp), %rdi
	movl	(%rsp), %edx
	imull	%edx, %edx
	shlq	$2, %rdx
	movq	96(%rsp), %rsi                  # 8-byte Reload
	movl	$1, %ecx
	callq	hipMemcpy
	movq	64(%rsp), %rdi
	movq	%rbp, %rsi
	movq	%r14, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movl	(%rsp), %esi
	imull	%esi, %esi
	shlq	$2, %rsi
	leaq	32(%rsp), %rdi
	callq	hipMalloc
	movl	(%rsp), %esi
	imull	%esi, %esi
	shlq	$2, %rsi
	leaq	24(%rsp), %rdi
	callq	hipMalloc
	movl	(%rsp), %esi
	imull	%esi, %esi
	shlq	$2, %rsi
	leaq	16(%rsp), %rdi
	callq	hipMalloc
	leaq	104(%rsp), %rdi
	callq	hipEventCreate
	leaq	72(%rsp), %rdi
	callq	hipEventCreate
	movq	104(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movabsq	$4294967296, %rax               # imm = 0x100000000
	orq	%rax, %r12
	movq	%r12, %rdi
	movl	$1, %esi
	movq	112(%rsp), %rdx                 # 8-byte Reload
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_3
# %bb.2:                                #   in Loop: Header=BB5_1 Depth=1
	movq	56(%rsp), %rdi
	movq	48(%rsp), %rsi
	movq	40(%rsp), %rdx
	movq	32(%rsp), %rcx
	movq	24(%rsp), %r8
	movq	16(%rsp), %r9
	movl	(%rsp), %eax
	movl	4(%rsp), %r10d
	subq	$8, %rsp
	.cfi_adjust_cfa_offset 8
	pushq	72(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	%r10
	.cfi_adjust_cfa_offset 8
	pushq	%rax
	.cfi_adjust_cfa_offset 8
	callq	_Z24__device_stub__kernelGPUPfS_S_S_S_S_iiPi
	addq	$32, %rsp
	.cfi_adjust_cfa_offset -32
.LBB5_3:                                #   in Loop: Header=BB5_1 Depth=1
	movq	72(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	72(%rsp), %rdi
	callq	hipEventSynchronize
	movq	104(%rsp), %rsi
	movq	72(%rsp), %rdx
	leaq	84(%rsp), %rdi
	callq	hipEventElapsedTime
	movl	$_ZSt4cout, %edi
	movl	$.L.str.9, %esi
	movl	$12, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	cvtss2sd	84(%rsp), %xmm0
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r14
	movl	$.L.str.8, %esi
	movl	$4, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	(%r14), %rax
	movq	-24(%rax), %rdi
	addq	%r14, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r14, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	(%rsp), %r15d
	imull	%r15d, %r15d
	shlq	$2, %r15
	movq	%r15, %rdi
	callq	_Znam
	movq	%rax, %r14
	movq	%r15, %rdi
	callq	_Znam
	movq	%rax, %r12
	movq	%r15, %rdi
	callq	_Znam
	movq	%rax, %r13
	movq	32(%rsp), %rsi
	movq	%r14, %rdi
	movq	%r15, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	24(%rsp), %rsi
	movl	(%rsp), %edx
	imull	%edx, %edx
	shlq	$2, %rdx
	movq	%r12, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	16(%rsp), %rsi
	movl	(%rsp), %edx
	imull	%edx, %edx
	shlq	$2, %rdx
	movq	%r13, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movl	(%rsp), %ecx
	movq	8(%rsp), %r15                   # 8-byte Reload
	leaq	60(%r15), %r9
	movq	%r14, %rdi
	movq	%r12, %rsi
	movq	%r13, %rdx
	movl	%ecx, %r8d
	callq	_Z5WritePfS_S_iiPKc
	movq	56(%rsp), %rdi
	callq	hipFree
	movq	48(%rsp), %rdi
	callq	hipFree
	movq	40(%rsp), %rdi
	callq	hipFree
	movq	32(%rsp), %rdi
	callq	hipFree
	movq	24(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	movq	64(%rsp), %rdi
	callq	hipFree
	testq	%rbx, %rbx
	movq	%rbx, %rdi
	movq	%r15, %rbx
	je	.LBB5_5
# %bb.4:                                #   in Loop: Header=BB5_1 Depth=1
	callq	_ZdaPv
.LBB5_5:                                #   in Loop: Header=BB5_1 Depth=1
	movq	88(%rsp), %rdi                  # 8-byte Reload
	testq	%rdi, %rdi
	je	.LBB5_7
# %bb.6:                                #   in Loop: Header=BB5_1 Depth=1
	callq	_ZdaPv
.LBB5_7:                                #   in Loop: Header=BB5_1 Depth=1
	movq	96(%rsp), %rdi                  # 8-byte Reload
	testq	%rdi, %rdi
	je	.LBB5_9
# %bb.8:                                #   in Loop: Header=BB5_1 Depth=1
	callq	_ZdaPv
.LBB5_9:                                #   in Loop: Header=BB5_1 Depth=1
	movq	%r14, %rdi
	callq	_ZdaPv
	movq	%r12, %rdi
	callq	_ZdaPv
	movq	%r13, %rdi
	callq	_ZdaPv
	testq	%rbp, %rbp
	movq	136(%rsp), %r14                 # 8-byte Reload
	je	.LBB5_11
# %bb.10:                               #   in Loop: Header=BB5_1 Depth=1
	movq	%rbp, %rdi
	callq	_ZdaPv
.LBB5_11:                               #   in Loop: Header=BB5_1 Depth=1
	addq	$90, %rbx
	decq	%r14
	movq	%r14, %rax
	jne	.LBB5_1
# %bb.12:
	xorl	%eax, %eax
	addq	$632, %rsp                      # imm = 0x278
	.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_end5:
	.size	main, .Lfunc_end5-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:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB6_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB6_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z9kernelGPUPfS_S_S_S_S_iiPi, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end6:
	.size	__hip_module_ctor, .Lfunc_end6-__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	.LBB7_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
.LBB7_2:
	retq
.Lfunc_end7:
	.size	__hip_module_dtor, .Lfunc_end7-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"Leemos el archivo %s!\n"
	.size	.L.str, 23

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

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

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"%d "
	.size	.L.str.3, 4

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"%f "
	.size	.L.str.4, 4

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"w"
	.size	.L.str.5, 2

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

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

	.type	.L__const.main.names,@object    # @__const.main.names
	.p2align	4, 0x0
.L__const.main.names:
	.asciz	"img100x100.txt\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000"
	.asciz	"img100x100CPU.txt\000\000\000\000\000\000\000\000\000\000\000\000"
	.asciz	"img100x100GPU.txt\000\000\000\000\000\000\000\000\000\000\000\000"
	.asciz	"img200x200.txt\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000"
	.asciz	"img200x200CPU.txt\000\000\000\000\000\000\000\000\000\000\000\000"
	.asciz	"img200x200GPU.txt\000\000\000\000\000\000\000\000\000\000\000\000"
	.asciz	"img400x400.txt\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000"
	.asciz	"img400x400CPU.txt\000\000\000\000\000\000\000\000\000\000\000\000"
	.asciz	"img400x400GPU.txt\000\000\000\000\000\000\000\000\000\000\000\000"
	.asciz	"img800x800.txt\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000"
	.asciz	"img800x800CPU.txt\000\000\000\000\000\000\000\000\000\000\000\000"
	.asciz	"img800x800GPU.txt\000\000\000\000\000\000\000\000\000\000\000\000"
	.asciz	"img1600x1600.txt\000\000\000\000\000\000\000\000\000\000\000\000\000"
	.asciz	"img1600x1600CPU.txt\000\000\000\000\000\000\000\000\000\000"
	.asciz	"img1600x1600GPU.txt\000\000\000\000\000\000\000\000\000\000"
	.size	.L__const.main.names, 450

	.type	.L.str.7,@object                # @.str.7
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.7:
	.asciz	"Tiempo CPU: "
	.size	.L.str.7, 13

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"[ms]"
	.size	.L.str.8, 5

	.type	.L.str.9,@object                # @.str.9
.L.str.9:
	.asciz	"Tiempo GPU: "
	.size	.L.str.9, 13

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z9kernelGPUPfS_S_S_S_S_iiPi"
	.size	.L__unnamed_1, 29

	.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 _Z24__device_stub__kernelGPUPfS_S_S_S_S_iiPi
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z9kernelGPUPfS_S_S_S_S_iiPi
	.addrsig_sym _ZSt4cout
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z9kernelGPUPfS_S_S_S_S_iiPi ; -- Begin function _Z9kernelGPUPfS_S_S_S_S_iiPi
	.globl	_Z9kernelGPUPfS_S_S_S_S_iiPi
	.p2align	8
	.type	_Z9kernelGPUPfS_S_S_S_S_iiPi,@function
_Z9kernelGPUPfS_S_S_S_S_iiPi:           ; @_Z9kernelGPUPfS_S_S_S_S_iiPi
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x4c
	s_load_b64 s[4:5], s[0:1], 0x30
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_mul_i32 s2, s4, s4
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_cmp_gt_i32_e32 vcc_lo, s2, v1
	s_and_saveexec_b32 s2, vcc_lo
	s_cbranch_execz .LBB0_2
; %bb.1:
	s_ashr_i32 s3, s5, 31
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_add_i32 s2, s5, s3
	s_xor_b32 s6, s2, s3
	s_ashr_i32 s2, s4, 31
	v_cvt_f32_u32_e32 v0, s6
	s_add_i32 s7, s4, s2
	s_sub_i32 s9, 0, s6
	s_xor_b32 s7, s7, s2
	s_xor_b32 s10, s2, s3
	v_rcp_iflag_f32_e32 v0, v0
	v_cvt_f32_u32_e32 v2, s7
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v0, 0x4f7ffffe, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_u32_f32_e32 v0, v0
	v_readfirstlane_b32 s8, v0
	v_rcp_iflag_f32_e32 v0, v2
	v_ashrrev_i32_e32 v2, 31, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_mul_i32 s9, s9, s8
	v_add_nc_u32_e32 v4, v1, v2
	s_mul_hi_u32 s9, s8, s9
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_i32 s8, s8, s9
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v0, 0x4f7ffffe, v0
	s_mul_hi_u32 s9, s7, s8
	v_xor_b32_e32 v4, v4, v2
	s_mul_i32 s11, s9, s6
	s_add_i32 s12, s9, 1
	s_sub_i32 s11, s7, s11
	v_cvt_u32_f32_e32 v0, v0
	s_sub_i32 s13, s11, s6
	s_cmp_ge_u32 s11, s6
	s_cselect_b32 s9, s12, s9
	s_cselect_b32 s11, s13, s11
	s_add_i32 s12, s9, 1
	s_cmp_ge_u32 s11, s6
	s_cselect_b32 s9, s12, s9
	s_sub_i32 s11, 0, s7
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v3, s11, v0
	v_mul_hi_u32 v3, v0, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v0, v0, v3
	v_mul_hi_u32 v0, v4, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v3, v0, s7
	v_sub_nc_u32_e32 v3, v4, v3
	v_add_nc_u32_e32 v4, 1, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v5, s7, v3
	v_cmp_le_u32_e32 vcc_lo, s7, v3
	v_dual_cndmask_b32 v0, v0, v4 :: v_dual_cndmask_b32 v3, v3, v5
	v_xor_b32_e32 v5, s2, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_nc_u32_e32 v4, 1, v0
	v_cmp_le_u32_e32 vcc_lo, s7, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v0, v0, v4, vcc_lo
	v_xor_b32_e32 v0, v0, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v0, v0, v5
	v_mul_lo_u32 v3, v0, s4
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v6, v1, v3
	v_ashrrev_i32_e32 v3, 31, v0
	v_ashrrev_i32_e32 v4, 31, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v5, v0, v3
	v_xor_b32_e32 v5, v5, v3
	v_xor_b32_e32 v3, s3, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v8, v5, s8
	v_mul_lo_u32 v10, v8, s6
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v5, v5, v10
	v_add_nc_u32_e32 v10, 1, v8
	v_cmp_le_u32_e32 vcc_lo, s6, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v8, v8, v10 :: v_dual_add_nc_u32 v7, v6, v4
	v_xor_b32_e32 v7, v7, v4
	v_subrev_nc_u32_e32 v10, s6, v5
	v_xor_b32_e32 v4, s3, v4
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_hi_u32 v9, v7, s8
	v_dual_cndmask_b32 v5, v5, v10 :: v_dual_add_nc_u32 v10, 1, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cmp_le_u32_e32 vcc_lo, s6, v5
	v_mul_lo_u32 v11, v9, s6
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v5, v8, v10, vcc_lo
	v_sub_nc_u32_e32 v7, v7, v11
	v_add_nc_u32_e32 v11, 1, v9
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_xor_b32_e32 v5, v5, v3
	v_subrev_nc_u32_e32 v12, s6, v7
	v_cmp_le_u32_e64 s2, s6, v7
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v8, v5, v3
	v_cndmask_b32_e64 v9, v9, v11, s2
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e64 v7, v7, v12, s2
	s_load_b64 s[2:3], s[0:1], 0x38
	v_add_nc_u32_e32 v11, 1, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_cmp_le_u32_e32 vcc_lo, s6, v7
	s_xor_b32 s6, s9, s10
	s_sub_i32 s6, s6, s10
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v7, v9, v11, vcc_lo
	v_xor_b32_e32 v7, v7, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v3, v7, v4
	v_mad_u64_u32 v[4:5], null, v8, s6, v[3:4]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v5, 31, v4
	v_lshlrev_b64 v[4:5], 2, v[4:5]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v4, vcc_lo, s2, v4
	v_add_co_ci_u32_e32 v5, vcc_lo, s3, v5, vcc_lo
	s_ashr_i32 s2, s6, 31
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	s_add_i32 s3, s6, s2
	global_load_b32 v4, v[4:5], off
	s_xor_b32 s3, s3, s2
	v_cvt_f32_u32_e32 v5, s3
	s_sub_i32 s7, 0, s3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_iflag_f32_e32 v5, v5
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v5, 0x4f7ffffe, v5
	v_cvt_u32_f32_e32 v5, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v7, s7, v5
	v_mul_hi_u32 v7, v5, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v5, v5, v7
	s_waitcnt vmcnt(0)
	v_ashrrev_i32_e32 v9, 31, v4
	v_add_nc_u32_e32 v10, v4, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v7, v10, v9
	v_xor_b32_e32 v9, s2, v9
	v_mul_hi_u32 v5, v7, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v10, v5, s3
	v_sub_nc_u32_e32 v7, v7, v10
	v_add_nc_u32_e32 v10, 1, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_subrev_nc_u32_e32 v11, s3, v7
	v_cmp_le_u32_e32 vcc_lo, s3, v7
	v_cndmask_b32_e32 v5, v5, v10, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v7, v7, v11 :: v_dual_add_nc_u32 v10, 1, v5
	v_cmp_le_u32_e32 vcc_lo, s3, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v5, v5, v10, vcc_lo
	v_xor_b32_e32 v5, v5, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v5, v5, v9
	v_mul_lo_u32 v7, v5, s6
	v_sub_nc_u32_e32 v8, v5, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v7, v4, v7
	v_mad_u64_u32 v[4:5], null, v8, s5, v[0:1]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v0, v7, v3
	v_mul_lo_u32 v3, v4, s4
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v0, v0, s5
	s_clause 0x1
	s_load_b256 s[4:11], s[0:1], 0x0
	s_load_b128 s[0:3], s[0:1], 0x20
	v_add3_u32 v3, v0, v6, v3
	v_lshlrev_b64 v[0:1], 2, v[1:2]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v4, 31, v3
	v_lshlrev_b64 v[3:4], 2, v[3:4]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v5, vcc_lo, s4, v3
	v_add_co_ci_u32_e32 v6, vcc_lo, s5, v4, vcc_lo
	global_load_b32 v9, v[5:6], off
	v_add_co_u32 v5, vcc_lo, s10, v0
	v_add_co_ci_u32_e32 v6, vcc_lo, s11, v1, vcc_lo
	v_add_co_u32 v7, vcc_lo, s6, v3
	v_add_co_ci_u32_e32 v8, vcc_lo, s7, v4, vcc_lo
	s_waitcnt vmcnt(0)
	global_store_b32 v[5:6], v9, off
	global_load_b32 v7, v[7:8], off
	v_add_co_u32 v5, vcc_lo, s0, v0
	v_add_co_ci_u32_e32 v6, vcc_lo, s1, v1, vcc_lo
	v_add_co_u32 v2, vcc_lo, s8, v3
	v_add_co_ci_u32_e32 v3, vcc_lo, s9, v4, vcc_lo
	v_add_co_u32 v0, vcc_lo, s2, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s3, v1, vcc_lo
	s_waitcnt vmcnt(0)
	global_store_b32 v[5:6], v7, off
	global_load_b32 v2, v[2:3], off
	s_waitcnt vmcnt(0)
	global_store_b32 v[0:1], v2, off
.LBB0_2:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z9kernelGPUPfS_S_S_S_S_iiPi
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 320
		.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 13
		.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	_Z9kernelGPUPfS_S_S_S_S_iiPi, .Lfunc_end0-_Z9kernelGPUPfS_S_S_S_S_iiPi
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 1088
; NumSgprs: 18
; NumVgprs: 13
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 13
; 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
	.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
      - .address_space:  global
        .offset:         32
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         40
        .size:           8
        .value_kind:     global_buffer
      - .offset:         48
        .size:           4
        .value_kind:     by_value
      - .offset:         52
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         56
        .size:           8
        .value_kind:     global_buffer
      - .offset:         64
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         68
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         72
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         76
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         78
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         80
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         82
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         84
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         86
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         104
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         112
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         120
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         128
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 320
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z9kernelGPUPfS_S_S_S_S_iiPi
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z9kernelGPUPfS_S_S_S_S_iiPi.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     13
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
