#include <stdio.h>
#include <string.h>
#include <stdlib.h>
#include <cuda_runtime.h>

#define ASCII_CHARS 128


int mod(int a, int b) {
    
    return (a%b+b)%b;
}

void printrow(float *msg_v_list) {
    int i;
    for (i = 0; i < 16; i++) {
        printf("PRINTROW: %f\n", msg_v_list[i]);
    }

}

__global__ void mtxEncrypt(float *secretKey, float *msg, float *result, int matrix_dims) {
    int tx = threadIdx.x;
    int ty = threadIdx.y;
    float Pvalue = 0;
    int i;
    for (i = 0; i < 16; i++) {
       // printf("PRINTROW: %f\n", msg[i]);
    }
    for (int k = 0; k < matrix_dims; ++k) {
        float secretElement = secretKey[ty * matrix_dims + k];
        float msgElement = msg[k * matrix_dims + tx];
        
   //     printf("\nMultiplying %f by %f: \n", secretElement, msgElement);
        Pvalue  += secretElement * msgElement;
    }
    
    //printf("PValue: %f    %d     %d\n", Pvalue, (int)Pvalue, ((int)Pvalue%ASCII_CHARS+ASCII_CHARS)%ASCII_CHARS);
    int mod_result = ((int)Pvalue%ASCII_CHARS+ASCII_CHARS)%ASCII_CHARS; // actual modulus function
    result[ty * matrix_dims + tx] = mod_result;
   // printf("result: %f\n", result[ty * matrix_dims + tx]);
}

void secretKey(float **SKey, float **invSKey, int matrix_dims) {
    int k,z;

    float xSKey[4][4] = {{8, 6, 9, 5},
                {6, 9, 5, 10},
                {5, 8, 4, 9},
                {10, 6, 11, 4}};
    float xinvSKey[4][4] = {{-3, 20, -21, 1},
                    {2, -41, 44, 1},
                    {2, -6, 6, -1},
                    {-1, 28, -30, -1}};

    int i, j;
    for (i = 0; i < 4; i++) {
        for (j = 0; j < 4; j++) {
            SKey[0][i*matrix_dims + j] = xSKey[i][j];
            invSKey[0][i*matrix_dims + j] = xinvSKey[i][j];
            printf("%f :  %f \n", i*matrix_dims + j, SKey[0][i*matrix_dims + j]);
        }
    }
}

void pad_msg(char *msg, int matrix_dims) {
    printf("Padding..\n");
    int i;
    unsigned int msg_len = strlen(msg);
    int matrix_size = matrix_dims*matrix_dims;
    int extra_chars = matrix_size - msg_len%matrix_size;

    char *space = (char *)malloc(extra_chars);
    for (i = 0; i < extra_chars; i++) {
        space[i] = ' ';
    }
    strcat(msg, space);
}

void encode_msg(char *msg, float *msg_vector, int matrix_dims) {
    int matrix_size = matrix_dims * matrix_dims;
    int i = 0;
    while (msg[i] != '\0') {
        msg_vector[i] = (int)(msg[i]);
        i++;
    }
}

int main(int argc, char *argv[]) {

    char *msg;
    if (argc > 1) {

        if (fopen(argv[1], "r")) {
            printf("reading file\n\n");
            FILE *msg_file;
            msg_file = fopen(argv[1], "r");
            fseek(msg_file, 0, SEEK_END);
            // ftell() gives current position in the stream
            long msg_file_size = ftell(msg_file);
            // rewind to beginning of file now that we have size
            fseek(msg_file, 0, SEEK_SET);
            // allocate memory for msg var, read file stream into memory
            char *msg_file_text = (char *)malloc(msg_file_size + 1);
            fread(msg_file_text, msg_file_size, 1, msg_file);
            fclose(msg_file);
            // printf("%s", msg_file_text);
            msg = msg_file_text;
        } else {
            msg = (char *)malloc(strlen(argv[1]) + 15);
            strcpy(msg, argv[1]);
        }
    } else {
        const char* jack_msg = "All work and no play makes Jack a dull boy.";
//        const char *jack_msg = "kkkkkkkkkkkkkkkkkkkkkkkkkkkkkkkkkkkkkkkkkkk";
        msg = (char *) malloc(strlen(jack_msg));
        strcpy(msg, jack_msg);
    }
    char *ascii_dict = (char *)malloc(128);

    int matrix_dims = 4;
    int matrix_size = matrix_dims*matrix_dims;


    // malloc() for 4 (int*) pointers to 4 rows, also malloc'd (int)
    int z;
    float **SKey = (float **) malloc(matrix_dims * sizeof(float *));
    for (z = 0; z < matrix_dims; z++) {
        SKey[z] = (float *)malloc(matrix_dims * sizeof(float));
    }
    float **invSKey = (float **)malloc(matrix_dims * sizeof(float *));
    for (z = 0; z < matrix_dims; z++) {
        invSKey[z] = (float *)malloc(matrix_dims * sizeof(float));
    }

    
    secretKey(SKey, invSKey, matrix_dims);

    int i, j;
    printf("%s", msg);
    for (i = 0; i < 4; i++) {
        for (j = 0; j < 4; j++) {
            printf("Skey: %f\n", SKey[0][i*4 + j]);
        }
    }
    // pad message if not divisible by matrix_dims^2

    int msg_size;
    if (strlen(msg) % matrix_size != 0) {
        pad_msg(msg, matrix_dims);
    }
    msg_size = strlen(msg);
    int msg_parts = msg_size / matrix_size;
    printf("%s", msg);

    float *msg_vector = (float *)malloc(strlen(msg)*sizeof(float));
    // message encoded, results stored in msg_vector
    encode_msg(msg, msg_vector, matrix_dims);
    printf("ENCODED!\n");

    float msg_vector_list[msg_parts][matrix_size];
    for (i = 0; i < msg_parts; i++) {
        printf("\n\nCount: %d %d\n", i, msg_size%matrix_size);
        for (j = 0; j < matrix_size; j++) {
            msg_vector_list[i][j] = msg_vector[i*matrix_size + j];
            printf("%f\n", msg_vector_list[i][j], "\n");
        }
    }
    
    
    int nBytes = matrix_size*sizeof(float);
    printf("1\n");
    float **results = (float **)malloc(msg_size * sizeof(float *));
    printf("2\n");
    for (z = 0; z < msg_parts; z++) {
        results[z] = (float *)malloc(matrix_size * sizeof(float));
        memset(results[z], 0, nBytes);
    }
   printf("3\n");

    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, 0);
    cudaSetDevice(0);


    float *secretGpu, *msgGpu, *resultGpu;
    cudaMalloc((void **)&secretGpu, nBytes);


    printf("msg_parts: %d \n", msg_parts);
    printf("\nmsg_vector_list[1]: %d\n", *msg_vector_list[1]);


    cudaMalloc((void **)&resultGpu, nBytes);
    cudaMalloc((void **)&msgGpu, nBytes);


    cudaMemcpy(secretGpu, *SKey, nBytes, cudaMemcpyHostToDevice);
    
    dim3 block(matrix_dims, matrix_dims);
    dim3 grid((matrix_dims+block.x-1)/block.x, (matrix_dims+block.y-1)/block.y);

    for (i = 0; i < msg_parts; i++) {
        cudaMemcpy(msgGpu, msg_vector_list[i], nBytes, cudaMemcpyHostToDevice);
        mtxEncrypt<<<grid, block>>>(secretGpu, msgGpu, resultGpu, matrix_dims);
        cudaMemcpy(results[i], resultGpu, nBytes, cudaMemcpyDeviceToHost);
        cudaDeviceSynchronize();
    }


    for (i = 0; i < msg_parts; i++) {
        for (j = 0; j < matrix_size; j++) {
            printf("%c", (int)results[i][j], char((int)results[i][j]));
        }
    }
    
    float **unEncrypted = (float **)malloc(msg_parts * sizeof(float *));
    for (z = 0; z < msg_parts; z++) {
        unEncrypted[z] = (float *)malloc(matrix_size * sizeof(float));
        memset(unEncrypted[z], 0, nBytes);
    }
     

    cudaMemcpy(secretGpu, *invSKey, nBytes, cudaMemcpyHostToDevice);
    for (i = 0; i < msg_parts; i++) {
        cudaMemcpy(msgGpu, results[i], nBytes, cudaMemcpyHostToDevice);
        mtxEncrypt<<<grid, block>>>(secretGpu, msgGpu, resultGpu, matrix_dims);
        cudaMemcpy(unEncrypted[i], resultGpu, nBytes, cudaMemcpyDeviceToHost);
        cudaDeviceSynchronize();
    }
    for (i = 0; i < msg_parts; i++) {
        for (j = 0; j < matrix_size; j++) {
            printf("%c", (int)unEncrypted[i][j], char((int)unEncrypted[i][j]));
        }
    }
    free(results);
    cudaFree(secretGpu);
    cudaFree(msgGpu);
    cudaFree(resultGpu);
    cudaDeviceReset();

    

}






	.file	"tmpxft_002d1e87_00000000-6_hill.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2035:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2035:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z3modii
	.type	_Z3modii, @function
_Z3modii:
.LFB2027:
	.cfi_startproc
	endbr64
	movl	%edi, %eax
	cltd
	idivl	%esi
	leal	(%rdx,%rsi), %eax
	cltd
	idivl	%esi
	movl	%edx, %eax
	ret
	.cfi_endproc
.LFE2027:
	.size	_Z3modii, .-_Z3modii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC2:
	.string	"PRINTROW: %f\n"
	.text
	.globl	_Z8printrowPf
	.type	_Z8printrowPf, @function
_Z8printrowPf:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	leaq	.LC2(%rip), %r12
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movq	%rdi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	xorl	%ebx, %ebx
.L4:
	cvtss2sd	0(%rbp,%rbx,4), %xmm0
	movq	%r12, %rsi
	movl	$2, %edi
	movb	$1, %al
	call	__printf_chk@PLT
	incq	%rbx
	cmpq	$16, %rbx
	jne	.L4
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2028:
	.size	_Z8printrowPf, .-_Z8printrowPf
	.section	.rodata.str1.1
.LC3:
	.string	"%f :  %f \n"
	.section	.rodata
	.align 32
.LC0:
	.long	1090519040
	.long	1086324736
	.long	1091567616
	.long	1084227584
	.long	1086324736
	.long	1091567616
	.long	1084227584
	.long	1092616192
	.long	1084227584
	.long	1090519040
	.long	1082130432
	.long	1091567616
	.long	1092616192
	.long	1086324736
	.long	1093664768
	.long	1082130432
	.align 32
.LC1:
	.long	-1069547520
	.long	1101004800
	.long	-1045954560
	.long	1065353216
	.long	1073741824
	.long	-1037828096
	.long	1110441984
	.long	1065353216
	.long	1073741824
	.long	-1061158912
	.long	1086324736
	.long	-1082130432
	.long	-1082130432
	.long	1105199104
	.long	-1041235968
	.long	-1082130432
	.text
	.globl	_Z9secretKeyPPfS0_i
	.type	_Z9secretKeyPPfS0_i, @function
_Z9secretKeyPPfS0_i:
.LFB2029:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rdi, %r9
	movq	%rsi, %r10
	movl	$16, %ecx
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	leaq	.LC0(%rip), %rsi
	movslq	%edx, %r15
	xorl	%r8d, %r8d
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movq	%r15, %r14
	xorl	%r13d, %r13d
	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	$200, %rsp
	.cfi_def_cfa_offset 256
	movq	%fs:40, %rax
	movq	%rax, 184(%rsp)
	xorl	%eax, %eax
	leaq	48(%rsp), %rdi
	leaq	0(,%r15,4), %rax
	xorl	%r15d, %r15d
	rep movsl
	leaq	112(%rsp), %rdi
	leaq	.LC1(%rip), %rsi
	movl	$16, %ecx
	rep movsl
	movq	%rax, 8(%rsp)
	leaq	48(%rsp), %r11
.L9:
	movq	%r13, %r12
	movq	%r15, %rbp
	xorl	%ebx, %ebx
.L10:
	movq	(%r9), %rax
	movq	(%r10), %rdx
	movl	$2, %edi
	movq	%r9, 40(%rsp)
	movss	(%r11,%r12), %xmm0
	leaq	.LC3(%rip), %rsi
	movq	%r11, 32(%rsp)
	addq	%rbp, %rax
	movq	%r10, 24(%rsp)
	movss	%xmm0, (%rax)
	movss	112(%rsp,%r12), %xmm0
	addq	$4, %r12
	movl	%r8d, 20(%rsp)
	movss	%xmm0, (%rdx,%rbp)
	leal	(%rbx,%r8), %edx
	cvtss2sd	(%rax), %xmm0
	incl	%ebx
	movb	$1, %al
	addq	$4, %rbp
	call	__printf_chk@PLT
	cmpl	$4, %ebx
	movl	20(%rsp), %r8d
	movq	24(%rsp), %r10
	movq	32(%rsp), %r11
	movq	40(%rsp), %r9
	jne	.L10
	movq	8(%rsp), %rax
	addq	$16, %r13
	addl	%r14d, %r8d
	addq	%rax, %r15
	cmpq	$64, %r13
	jne	.L9
	movq	184(%rsp), %rax
	subq	%fs:40, %rax
	je	.L12
	call	__stack_chk_fail@PLT
.L12:
	addq	$200, %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
.LFE2029:
	.size	_Z9secretKeyPPfS0_i, .-_Z9secretKeyPPfS0_i
	.section	.rodata.str1.1
.LC4:
	.string	"Padding..\n"
	.text
	.globl	_Z7pad_msgPci
	.type	_Z7pad_msgPci, @function
_Z7pad_msgPci:
.LFB2030:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rdi, %rbp
	xorl	%eax, %eax
	movl	$2, %edi
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movl	%esi, %ebx
	leaq	.LC4(%rip), %rsi
	pushq	%rdx
	.cfi_def_cfa_offset 32
	imull	%ebx, %ebx
	call	__printf_chk@PLT
	movq	%rbp, %rdi
	call	strlen@PLT
	xorl	%edx, %edx
	divl	%ebx
	subl	%edx, %ebx
	movslq	%ebx, %rdi
	call	malloc@PLT
	movq	%rax, %rsi
	xorl	%eax, %eax
.L16:
	cmpl	%eax, %ebx
	jle	.L19
	movb	$32, (%rsi,%rax)
	incq	%rax
	jmp	.L16
.L19:
	popq	%rax
	.cfi_def_cfa_offset 24
	movq	%rbp, %rdi
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	strcat@PLT
	.cfi_endproc
.LFE2030:
	.size	_Z7pad_msgPci, .-_Z7pad_msgPci
	.globl	_Z10encode_msgPcPfi
	.type	_Z10encode_msgPcPfi, @function
_Z10encode_msgPcPfi:
.LFB2031:
	.cfi_startproc
	endbr64
	xorl	%edx, %edx
.L21:
	movsbl	(%rdi,%rdx), %eax
	testb	%al, %al
	je	.L23
	cvtsi2ssl	%eax, %xmm0
	movss	%xmm0, (%rsi,%rdx,4)
	incq	%rdx
	jmp	.L21
.L23:
	ret
	.cfi_endproc
.LFE2031:
	.size	_Z10encode_msgPcPfi, .-_Z10encode_msgPcPfi
	.globl	_Z35__device_stub__Z10mtxEncryptPfS_S_iPfS_S_i
	.type	_Z35__device_stub__Z10mtxEncryptPfS_S_iPfS_S_i, @function
_Z35__device_stub__Z10mtxEncryptPfS_S_iPfS_S_i:
.LFB2057:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movl	%ecx, 4(%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L24
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z10mtxEncryptPfS_S_i(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 176
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	120(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L24:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L26
	call	__stack_chk_fail@PLT
.L26:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2057:
	.size	_Z35__device_stub__Z10mtxEncryptPfS_S_iPfS_S_i, .-_Z35__device_stub__Z10mtxEncryptPfS_S_iPfS_S_i
	.globl	_Z10mtxEncryptPfS_S_i
	.type	_Z10mtxEncryptPfS_S_i, @function
_Z10mtxEncryptPfS_S_i:
.LFB2058:
	.cfi_startproc
	endbr64
	jmp	_Z35__device_stub__Z10mtxEncryptPfS_S_iPfS_S_i
	.cfi_endproc
.LFE2058:
	.size	_Z10mtxEncryptPfS_S_i, .-_Z10mtxEncryptPfS_S_i
	.section	.rodata.str1.1
.LC5:
	.string	"r"
.LC6:
	.string	"reading file\n\n"
.LC7:
	.string	"All work and no play makes Jack a dull boy."
.LC8:
	.string	"%s"
.LC9:
	.string	"Skey: %f\n"
.LC10:
	.string	"ENCODED!\n"
.LC11:
	.string	"\n\nCount: %d %d\n"
.LC12:
	.string	"\n"
.LC13:
	.string	"%f\n"
.LC14:
	.string	"1\n"
.LC15:
	.string	"2\n"
.LC16:
	.string	"3\n"
.LC17:
	.string	"msg_parts: %d \n"
.LC18:
	.string	"\nmsg_vector_list[1]: %d\n"
.LC19:
	.string	"%c"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2032:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rsp, %rbp
	.cfi_def_cfa_register 6
	pushq	%r15
	pushq	%r14
	pushq	%r13
	pushq	%r12
	pushq	%rbx
	subq	$1160, %rsp
	.cfi_offset 15, -24
	.cfi_offset 14, -32
	.cfi_offset 13, -40
	.cfi_offset 12, -48
	.cfi_offset 3, -56
	movq	%fs:40, %rax
	movq	%rax, -56(%rbp)
	xorl	%eax, %eax
	decl	%edi
	jle	.L30
	movq	8(%rsi), %rdi
	leaq	.LC5(%rip), %r12
	movq	%rsi, %rbx
	movq	%r12, %rsi
	call	fopen@PLT
	testq	%rax, %rax
	je	.L31
	leaq	.LC6(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	8(%rbx), %rdi
	movq	%r12, %rsi
	call	fopen@PLT
	xorl	%esi, %esi
	movl	$2, %edx
	movq	%rax, %r12
	movq	%rax, %rdi
	call	fseek@PLT
	movq	%r12, %rdi
	call	ftell@PLT
	xorl	%edx, %edx
	xorl	%esi, %esi
	movq	%r12, %rdi
	movq	%rax, %r13
	call	fseek@PLT
	leaq	1(%r13), %r14
	movq	%r14, %rdi
	call	malloc@PLT
	movq	%r12, %r8
	movl	$1, %ecx
	movq	%r13, %rdx
	movq	%r14, %rsi
	movq	%rax, %rdi
	movq	%rax, %rbx
	call	__fread_chk@PLT
	movq	%r12, %rdi
	call	fclose@PLT
	jmp	.L32
.L31:
	movq	8(%rbx), %r13
	movq	%r13, %rdi
	call	strlen@PLT
	leaq	15(%rax), %r12
	movq	%r12, %rdi
	call	malloc@PLT
	movq	%r12, %rdx
	movq	%r13, %rsi
	movq	%rax, %rdi
	movq	%rax, %rbx
	call	__strcpy_chk@PLT
	jmp	.L32
.L30:
	movl	$43, %edi
	call	malloc@PLT
	movl	$43, %ecx
	movl	$44, %edx
	leaq	.LC7(%rip), %rsi
	movq	%rax, %rdi
	movq	%rax, %rbx
	call	__memcpy_chk@PLT
.L32:
	movl	$32, %edi
	xorl	%r13d, %r13d
	call	malloc@PLT
	movq	%rax, %r12
.L33:
	movl	$16, %edi
	call	malloc@PLT
	movq	%rax, (%r12,%r13,8)
	incq	%r13
	cmpq	$4, %r13
	jne	.L33
	movl	$32, %edi
	xorl	%r13d, %r13d
	call	malloc@PLT
	movq	%rax, -1144(%rbp)
.L34:
	movl	$16, %edi
	call	malloc@PLT
	movq	-1144(%rbp), %rcx
	movq	%rax, (%rcx,%r13,8)
	incq	%r13
	cmpq	$4, %r13
	jne	.L34
	movq	%rcx, %rsi
	movl	$4, %edx
	movq	%r12, %rdi
	xorl	%r15d, %r15d
	call	_Z9secretKeyPPfS0_i
	movq	%rbx, %rdx
	movl	$2, %edi
	xorl	%eax, %eax
	leaq	.LC8(%rip), %rsi
	call	__printf_chk@PLT
.L35:
	movq	%r15, %r14
	xorl	%r13d, %r13d
	salq	$4, %r14
.L36:
	movq	(%r12), %rax
	leaq	.LC9(%rip), %rsi
	movl	$2, %edi
	movq	%rax, -1152(%rbp)
	leaq	(%rax,%r13,4), %rax
	incq	%r13
	cvtss2sd	(%rax,%r14), %xmm0
	movb	$1, %al
	call	__printf_chk@PLT
	cmpq	$4, %r13
	jne	.L36
	incq	%r15
	cmpq	$4, %r15
	jne	.L35
	movq	%rbx, %rdi
	call	strlen@PLT
	testb	$15, %al
	je	.L38
	movl	$4, %esi
	movq	%rbx, %rdi
	call	_Z7pad_msgPci
.L38:
	movq	%rbx, %rdi
	call	strlen@PLT
	movl	$16, %ecx
	movl	$2, %edi
	leaq	.LC8(%rip), %rsi
	cltd
	movq	%rax, %r15
	movl	%eax, %r14d
	idivl	%ecx
	movq	%rbx, %rdx
	movl	%eax, %r12d
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	%rbx, %rdi
	call	strlen@PLT
	leaq	0(,%rax,4), %rdi
	call	malloc@PLT
	movl	$4, %edx
	movq	%rbx, %rdi
	movq	%rax, %rsi
	movq	%rax, %r13
	call	_Z10encode_msgPcPfi
	leaq	.LC10(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movslq	%r12d, %rax
	movq	%rsp, %rdx
	movq	%rax, -1168(%rbp)
	salq	$6, %rax
	movq	%rax, %rcx
	andq	$-4096, %rcx
	subq	%rcx, %rdx
.L39:
	cmpq	%rdx, %rsp
	je	.L40
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L39
.L40:
	andl	$4095, %eax
	subq	%rax, %rsp
	testq	%rax, %rax
	je	.L41
	orq	$0, -8(%rsp,%rax)
.L41:
	leaq	3(%rsp), %rax
	movq	%rax, %rbx
	andq	$-4, %rax
	shrq	$2, %rbx
	movq	%rax, -1160(%rbp)
	movl	%r14d, %eax
	xorl	%r14d, %r14d
	movq	%rbx, -1184(%rbp)
	cltd
	movl	$16, %ebx
	idivl	%ebx
	movl	%edx, -1188(%rbp)
.L42:
	movl	%r14d, %edx
	cmpl	%r14d, %r12d
	jle	.L75
	movl	-1188(%rbp), %ecx
	leaq	.LC11(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	-1160(%rbp), %rbx
	movq	%r14, %rax
	salq	$6, %rax
	leaq	0(%r13,%rax), %r9
	addq	%rbx, %rax
	xorl	%ebx, %ebx
	movq	%rax, -1176(%rbp)
.L43:
	movss	(%r9,%rbx), %xmm0
	movq	-1176(%rbp), %rax
	leaq	.LC12(%rip), %rdx
	leaq	.LC13(%rip), %rsi
	movl	$2, %edi
	movq	%r9, -1200(%rbp)
	movss	%xmm0, (%rax,%rbx)
	movb	$1, %al
	addq	$4, %rbx
	cvtss2sd	%xmm0, %xmm0
	call	__printf_chk@PLT
	cmpq	$64, %rbx
	movq	-1200(%rbp), %r9
	jne	.L43
	incq	%r14
	jmp	.L42
.L75:
	leaq	.LC14(%rip), %rsi
	xorl	%eax, %eax
	xorl	%r14d, %r14d
	xorl	%r13d, %r13d
	movl	$2, %edi
	call	__printf_chk@PLT
	movslq	%r15d, %rdi
	salq	$3, %rdi
	call	malloc@PLT
	leaq	.LC15(%rip), %rsi
	movl	$2, %edi
	movq	%rax, %rbx
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L45:
	cmpl	%r14d, %r12d
	jle	.L76
	movl	$64, %edi
	call	malloc@PLT
	movl	$16, %ecx
	movq	%rax, (%rbx,%r14,8)
	movq	%rax, %rdi
	movl	%r13d, %eax
	incq	%r14
	rep stosl
	jmp	.L45
.L76:
	leaq	.LC16(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	xorl	%r14d, %r14d
	call	__printf_chk@PLT
	xorl	%esi, %esi
	leaq	-1088(%rbp), %rdi
	movabsq	$17179869188, %r15
	call	cudaGetDeviceProperties_v2@PLT
	xorl	%edi, %edi
	call	cudaSetDevice@PLT
	leaq	-1136(%rbp), %rdi
	movl	$64, %esi
	call	cudaMalloc@PLT
	movl	%r12d, %edx
	movl	$2, %edi
	xorl	%eax, %eax
	leaq	.LC17(%rip), %rsi
	call	__printf_chk@PLT
	movq	-1184(%rbp), %rax
	movl	$2, %edi
	leaq	.LC18(%rip), %rsi
	cvtss2sd	64(,%rax,4), %xmm0
	movb	$1, %al
	call	__printf_chk@PLT
	leaq	-1120(%rbp), %rdi
	movl	$64, %esi
	call	cudaMalloc@PLT
	leaq	-1128(%rbp), %rdi
	movl	$64, %esi
	call	cudaMalloc@PLT
	movl	$1, %ecx
	movl	$64, %edx
	movq	-1152(%rbp), %rsi
	movq	-1136(%rbp), %rdi
	call	cudaMemcpy@PLT
	movq	-1160(%rbp), %r13
	movabsq	$4294967297, %rax
	movl	$1, -1104(%rbp)
	movq	%rax, -1100(%rbp)
	movl	$1, -1092(%rbp)
.L47:
	cmpl	%r14d, %r12d
	jle	.L77
	movq	-1128(%rbp), %rdi
	movl	$1, %ecx
	movl	$64, %edx
	movq	%r13, %rsi
	call	cudaMemcpy@PLT
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	%r15, %rdx
	movl	-1104(%rbp), %ecx
	movq	-1100(%rbp), %rdi
	movq	%r15, -1112(%rbp)
	movl	-1092(%rbp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L48
	movq	-1120(%rbp), %rdx
	movq	-1128(%rbp), %rsi
	movl	$4, %ecx
	movq	-1136(%rbp), %rdi
	call	_Z35__device_stub__Z10mtxEncryptPfS_S_iPfS_S_i
.L48:
	movq	(%rbx,%r14,8), %rdi
	movq	-1120(%rbp), %rsi
	incq	%r14
	addq	$64, %r13
	movl	$2, %ecx
	movl	$64, %edx
	call	cudaMemcpy@PLT
	call	cudaDeviceSynchronize@PLT
	jmp	.L47
.L77:
	xorl	%r14d, %r14d
	leaq	.LC19(%rip), %r13
.L50:
	cmpl	%r14d, %r12d
	jle	.L78
	xorl	%r15d, %r15d
.L51:
	movq	(%rbx,%r14,8), %rax
	movq	%r13, %rsi
	movl	$2, %edi
	cvttss2sil	(%rax,%r15), %edx
	xorl	%eax, %eax
	addq	$4, %r15
	movsbl	%dl, %ecx
	call	__printf_chk@PLT
	cmpq	$64, %r15
	jne	.L51
	incq	%r14
	jmp	.L50
.L78:
	movq	-1168(%rbp), %rdi
	xorl	%r15d, %r15d
	xorl	%r14d, %r14d
	salq	$3, %rdi
	call	malloc@PLT
	movq	%rax, %r13
.L52:
	cmpl	%r15d, %r12d
	jle	.L79
	movl	$64, %edi
	call	malloc@PLT
	movl	$16, %ecx
	movq	%rax, 0(%r13,%r15,8)
	movq	%rax, %rdi
	movl	%r14d, %eax
	incq	%r15
	rep stosl
	jmp	.L52
.L79:
	movq	-1144(%rbp), %rax
	movq	-1136(%rbp), %rdi
	movl	$1, %ecx
	xorl	%r14d, %r14d
	movabsq	$17179869188, %r15
	movl	$64, %edx
	movq	(%rax), %rsi
	call	cudaMemcpy@PLT
.L54:
	cmpl	%r14d, %r12d
	jle	.L80
	movq	(%rbx,%r14,8), %rsi
	movl	$1, %ecx
	movl	$64, %edx
	movq	-1128(%rbp), %rdi
	call	cudaMemcpy@PLT
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	%r15, %rdx
	movl	-1104(%rbp), %ecx
	movq	-1100(%rbp), %rdi
	movq	%r15, -1112(%rbp)
	movl	-1092(%rbp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L55
	movq	-1120(%rbp), %rdx
	movq	-1128(%rbp), %rsi
	movl	$4, %ecx
	movq	-1136(%rbp), %rdi
	call	_Z35__device_stub__Z10mtxEncryptPfS_S_iPfS_S_i
.L55:
	movq	0(%r13,%r14,8), %rdi
	movq	-1120(%rbp), %rsi
	movl	$2, %ecx
	incq	%r14
	movl	$64, %edx
	call	cudaMemcpy@PLT
	call	cudaDeviceSynchronize@PLT
	jmp	.L54
.L80:
	xorl	%r14d, %r14d
.L57:
	cmpl	%r14d, %r12d
	jle	.L81
	xorl	%r15d, %r15d
.L58:
	movq	0(%r13,%r14,8), %rax
	leaq	.LC19(%rip), %rsi
	movl	$2, %edi
	cvttss2sil	(%rax,%r15), %edx
	xorl	%eax, %eax
	addq	$4, %r15
	movsbl	%dl, %ecx
	call	__printf_chk@PLT
	cmpq	$64, %r15
	jne	.L58
	incq	%r14
	jmp	.L57
.L81:
	movq	%rbx, %rdi
	call	free@PLT
	movq	-1136(%rbp), %rdi
	call	cudaFree@PLT
	movq	-1128(%rbp), %rdi
	call	cudaFree@PLT
	movq	-1120(%rbp), %rdi
	call	cudaFree@PLT
	call	cudaDeviceReset@PLT
	movq	-56(%rbp), %rax
	subq	%fs:40, %rax
	je	.L59
	call	__stack_chk_fail@PLT
.L59:
	leaq	-40(%rbp), %rsp
	xorl	%eax, %eax
	popq	%rbx
	popq	%r12
	popq	%r13
	popq	%r14
	popq	%r15
	popq	%rbp
	.cfi_def_cfa 7, 8
	ret
	.cfi_endproc
.LFE2032:
	.size	main, .-main
	.section	.rodata.str1.1
.LC20:
	.string	"_Z10mtxEncryptPfS_S_i"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2060:
	.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	.LC20(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z10mtxEncryptPfS_S_i(%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
.LFE2060:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:

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

#define ASCII_CHARS 128


int mod(int a, int b) {
    
    return (a%b+b)%b;
}

void printrow(float *msg_v_list) {
    int i;
    for (i = 0; i < 16; i++) {
        printf("PRINTROW: %f\n", msg_v_list[i]);
    }

}

__global__ void mtxEncrypt(float *secretKey, float *msg, float *result, int matrix_dims) {
    int tx = threadIdx.x;
    int ty = threadIdx.y;
    float Pvalue = 0;
    int i;
    for (i = 0; i < 16; i++) {
       // printf("PRINTROW: %f\n", msg[i]);
    }
    for (int k = 0; k < matrix_dims; ++k) {
        float secretElement = secretKey[ty * matrix_dims + k];
        float msgElement = msg[k * matrix_dims + tx];
        
   //     printf("\nMultiplying %f by %f: \n", secretElement, msgElement);
        Pvalue  += secretElement * msgElement;
    }
    
    //printf("PValue: %f    %d     %d\n", Pvalue, (int)Pvalue, ((int)Pvalue%ASCII_CHARS+ASCII_CHARS)%ASCII_CHARS);
    int mod_result = ((int)Pvalue%ASCII_CHARS+ASCII_CHARS)%ASCII_CHARS; // actual modulus function
    result[ty * matrix_dims + tx] = mod_result;
   // printf("result: %f\n", result[ty * matrix_dims + tx]);
}

void secretKey(float **SKey, float **invSKey, int matrix_dims) {
    int k,z;

    float xSKey[4][4] = {{8, 6, 9, 5},
                {6, 9, 5, 10},
                {5, 8, 4, 9},
                {10, 6, 11, 4}};
    float xinvSKey[4][4] = {{-3, 20, -21, 1},
                    {2, -41, 44, 1},
                    {2, -6, 6, -1},
                    {-1, 28, -30, -1}};

    int i, j;
    for (i = 0; i < 4; i++) {
        for (j = 0; j < 4; j++) {
            SKey[0][i*matrix_dims + j] = xSKey[i][j];
            invSKey[0][i*matrix_dims + j] = xinvSKey[i][j];
            printf("%f :  %f \n", i*matrix_dims + j, SKey[0][i*matrix_dims + j]);
        }
    }
}

void pad_msg(char *msg, int matrix_dims) {
    printf("Padding..\n");
    int i;
    unsigned int msg_len = strlen(msg);
    int matrix_size = matrix_dims*matrix_dims;
    int extra_chars = matrix_size - msg_len%matrix_size;

    char *space = (char *)malloc(extra_chars);
    for (i = 0; i < extra_chars; i++) {
        space[i] = ' ';
    }
    strcat(msg, space);
}

void encode_msg(char *msg, float *msg_vector, int matrix_dims) {
    int matrix_size = matrix_dims * matrix_dims;
    int i = 0;
    while (msg[i] != '\0') {
        msg_vector[i] = (int)(msg[i]);
        i++;
    }
}

int main(int argc, char *argv[]) {

    char *msg;
    if (argc > 1) {

        if (fopen(argv[1], "r")) {
            printf("reading file\n\n");
            FILE *msg_file;
            msg_file = fopen(argv[1], "r");
            fseek(msg_file, 0, SEEK_END);
            // ftell() gives current position in the stream
            long msg_file_size = ftell(msg_file);
            // rewind to beginning of file now that we have size
            fseek(msg_file, 0, SEEK_SET);
            // allocate memory for msg var, read file stream into memory
            char *msg_file_text = (char *)malloc(msg_file_size + 1);
            fread(msg_file_text, msg_file_size, 1, msg_file);
            fclose(msg_file);
            // printf("%s", msg_file_text);
            msg = msg_file_text;
        } else {
            msg = (char *)malloc(strlen(argv[1]) + 15);
            strcpy(msg, argv[1]);
        }
    } else {
        const char* jack_msg = "All work and no play makes Jack a dull boy.";
//        const char *jack_msg = "kkkkkkkkkkkkkkkkkkkkkkkkkkkkkkkkkkkkkkkkkkk";
        msg = (char *) malloc(strlen(jack_msg));
        strcpy(msg, jack_msg);
    }
    char *ascii_dict = (char *)malloc(128);

    int matrix_dims = 4;
    int matrix_size = matrix_dims*matrix_dims;


    // malloc() for 4 (int*) pointers to 4 rows, also malloc'd (int)
    int z;
    float **SKey = (float **) malloc(matrix_dims * sizeof(float *));
    for (z = 0; z < matrix_dims; z++) {
        SKey[z] = (float *)malloc(matrix_dims * sizeof(float));
    }
    float **invSKey = (float **)malloc(matrix_dims * sizeof(float *));
    for (z = 0; z < matrix_dims; z++) {
        invSKey[z] = (float *)malloc(matrix_dims * sizeof(float));
    }

    
    secretKey(SKey, invSKey, matrix_dims);

    int i, j;
    printf("%s", msg);
    for (i = 0; i < 4; i++) {
        for (j = 0; j < 4; j++) {
            printf("Skey: %f\n", SKey[0][i*4 + j]);
        }
    }
    // pad message if not divisible by matrix_dims^2

    int msg_size;
    if (strlen(msg) % matrix_size != 0) {
        pad_msg(msg, matrix_dims);
    }
    msg_size = strlen(msg);
    int msg_parts = msg_size / matrix_size;
    printf("%s", msg);

    float *msg_vector = (float *)malloc(strlen(msg)*sizeof(float));
    // message encoded, results stored in msg_vector
    encode_msg(msg, msg_vector, matrix_dims);
    printf("ENCODED!\n");

    float msg_vector_list[msg_parts][matrix_size];
    for (i = 0; i < msg_parts; i++) {
        printf("\n\nCount: %d %d\n", i, msg_size%matrix_size);
        for (j = 0; j < matrix_size; j++) {
            msg_vector_list[i][j] = msg_vector[i*matrix_size + j];
            printf("%f\n", msg_vector_list[i][j], "\n");
        }
    }
    
    
    int nBytes = matrix_size*sizeof(float);
    printf("1\n");
    float **results = (float **)malloc(msg_size * sizeof(float *));
    printf("2\n");
    for (z = 0; z < msg_parts; z++) {
        results[z] = (float *)malloc(matrix_size * sizeof(float));
        memset(results[z], 0, nBytes);
    }
   printf("3\n");

    hipDeviceProp_t deviceProp;
    hipGetDeviceProperties(&deviceProp, 0);
    hipSetDevice(0);


    float *secretGpu, *msgGpu, *resultGpu;
    hipMalloc((void **)&secretGpu, nBytes);


    printf("msg_parts: %d \n", msg_parts);
    printf("\nmsg_vector_list[1]: %d\n", *msg_vector_list[1]);


    hipMalloc((void **)&resultGpu, nBytes);
    hipMalloc((void **)&msgGpu, nBytes);


    hipMemcpy(secretGpu, *SKey, nBytes, hipMemcpyHostToDevice);
    
    dim3 block(matrix_dims, matrix_dims);
    dim3 grid((matrix_dims+block.x-1)/block.x, (matrix_dims+block.y-1)/block.y);

    for (i = 0; i < msg_parts; i++) {
        hipMemcpy(msgGpu, msg_vector_list[i], nBytes, hipMemcpyHostToDevice);
        mtxEncrypt<<<grid, block>>>(secretGpu, msgGpu, resultGpu, matrix_dims);
        hipMemcpy(results[i], resultGpu, nBytes, hipMemcpyDeviceToHost);
        hipDeviceSynchronize();
    }


    for (i = 0; i < msg_parts; i++) {
        for (j = 0; j < matrix_size; j++) {
            printf("%c", (int)results[i][j], char((int)results[i][j]));
        }
    }
    
    float **unEncrypted = (float **)malloc(msg_parts * sizeof(float *));
    for (z = 0; z < msg_parts; z++) {
        unEncrypted[z] = (float *)malloc(matrix_size * sizeof(float));
        memset(unEncrypted[z], 0, nBytes);
    }
     

    hipMemcpy(secretGpu, *invSKey, nBytes, hipMemcpyHostToDevice);
    for (i = 0; i < msg_parts; i++) {
        hipMemcpy(msgGpu, results[i], nBytes, hipMemcpyHostToDevice);
        mtxEncrypt<<<grid, block>>>(secretGpu, msgGpu, resultGpu, matrix_dims);
        hipMemcpy(unEncrypted[i], resultGpu, nBytes, hipMemcpyDeviceToHost);
        hipDeviceSynchronize();
    }
    for (i = 0; i < msg_parts; i++) {
        for (j = 0; j < matrix_size; j++) {
            printf("%c", (int)unEncrypted[i][j], char((int)unEncrypted[i][j]));
        }
    }
    free(results);
    hipFree(secretGpu);
    hipFree(msgGpu);
    hipFree(resultGpu);
    hipDeviceReset();

    

}






	.text
	.file	"hill.hip"
	.globl	_Z3modii                        # -- Begin function _Z3modii
	.type	_Z3modii,@function
_Z3modii:                               # @_Z3modii
	.cfi_startproc
# %bb.0:
                                        # kill: def $esi killed $esi def $rsi
	movl	%edi, %eax
	cltd
	idivl	%esi
                                        # kill: def $edx killed $edx def $rdx
	leal	(%rdx,%rsi), %eax
	cltd
	idivl	%esi
	movl	%edx, %eax
	retq
.Lfunc_end0:
	.size	_Z3modii, .Lfunc_end0-_Z3modii
	.cfi_endproc
                                        # -- End function
	.globl	_Z8printrowPf                   # -- Begin function _Z8printrowPf
	.type	_Z8printrowPf,@function
_Z8printrowPf:                          # @_Z8printrowPf
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	pushq	%rax
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	movq	%rdi, %rbx
	xorl	%r14d, %r14d
.LBB1_1:                                # =>This Inner Loop Header: Depth=1
	xorps	%xmm0, %xmm0
	cvtss2sd	(%rbx,%r14,4), %xmm0
	movl	$.L.str, %edi
	movb	$1, %al
	callq	printf
	incq	%r14
	cmpq	$16, %r14
	jne	.LBB1_1
# %bb.2:
	addq	$8, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z8printrowPf, .Lfunc_end1-_Z8printrowPf
	.cfi_endproc
                                        # -- End function
	.globl	_Z25__device_stub__mtxEncryptPfS_S_i # -- Begin function _Z25__device_stub__mtxEncryptPfS_S_i
	.type	_Z25__device_stub__mtxEncryptPfS_S_i,@function
_Z25__device_stub__mtxEncryptPfS_S_i:   # @_Z25__device_stub__mtxEncryptPfS_S_i
	.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	32(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	24(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	4(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z10mtxEncryptPfS_S_i, %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	_Z25__device_stub__mtxEncryptPfS_S_i, .Lfunc_end2-_Z25__device_stub__mtxEncryptPfS_S_i
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0                          # -- Begin function _Z9secretKeyPPfS0_i
.LCPI3_0:
	.long	0x41000000                      # float 8
	.long	0x40c00000                      # float 6
	.long	0x41100000                      # float 9
	.long	0x40a00000                      # float 5
.LCPI3_1:
	.long	0x40c00000                      # float 6
	.long	0x41100000                      # float 9
	.long	0x40a00000                      # float 5
	.long	0x41200000                      # float 10
.LCPI3_2:
	.long	0x40a00000                      # float 5
	.long	0x41000000                      # float 8
	.long	0x40800000                      # float 4
	.long	0x41100000                      # float 9
.LCPI3_3:
	.long	0x41200000                      # float 10
	.long	0x40c00000                      # float 6
	.long	0x41300000                      # float 11
	.long	0x40800000                      # float 4
.LCPI3_4:
	.long	0xc0400000                      # float -3
	.long	0x41a00000                      # float 20
	.long	0xc1a80000                      # float -21
	.long	0x3f800000                      # float 1
.LCPI3_5:
	.long	0x40000000                      # float 2
	.long	0xc2240000                      # float -41
	.long	0x42300000                      # float 44
	.long	0x3f800000                      # float 1
.LCPI3_6:
	.long	0x40000000                      # float 2
	.long	0xc0c00000                      # float -6
	.long	0x40c00000                      # float 6
	.long	0xbf800000                      # float -1
.LCPI3_7:
	.long	0xbf800000                      # float -1
	.long	0x41e00000                      # float 28
	.long	0xc1f00000                      # float -30
	.long	0xbf800000                      # float -1
	.text
	.globl	_Z9secretKeyPPfS0_i
	.type	_Z9secretKeyPPfS0_i,@function
_Z9secretKeyPPfS0_i:                    # @_Z9secretKeyPPfS0_i
	.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	$184, %rsp
	.cfi_def_cfa_offset 240
	.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, %r15
	movq	%rdi, %r12
	movaps	.LCPI3_0(%rip), %xmm0           # xmm0 = [8.0E+0,6.0E+0,9.0E+0,5.0E+0]
	leaq	112(%rsp), %rax
	movaps	%xmm0, (%rax)
	movaps	.LCPI3_1(%rip), %xmm0           # xmm0 = [6.0E+0,9.0E+0,5.0E+0,1.0E+1]
	movaps	%xmm0, 16(%rax)
	movaps	.LCPI3_2(%rip), %xmm0           # xmm0 = [5.0E+0,8.0E+0,4.0E+0,9.0E+0]
	movaps	%xmm0, 32(%rax)
	movaps	.LCPI3_3(%rip), %xmm0           # xmm0 = [1.0E+1,6.0E+0,1.1E+1,4.0E+0]
	movaps	%xmm0, 48(%rax)
	movaps	.LCPI3_4(%rip), %xmm0           # xmm0 = [-3.0E+0,2.0E+1,-2.1E+1,1.0E+0]
	leaq	48(%rsp), %r13
	movaps	%xmm0, (%r13)
	movaps	.LCPI3_5(%rip), %xmm0           # xmm0 = [2.0E+0,-4.1E+1,4.4E+1,1.0E+0]
	movaps	%xmm0, 16(%r13)
	movaps	.LCPI3_6(%rip), %xmm0           # xmm0 = [2.0E+0,-6.0E+0,6.0E+0,-1.0E+0]
	movaps	%xmm0, 32(%r13)
	movaps	.LCPI3_7(%rip), %xmm0           # xmm0 = [-1.0E+0,2.8E+1,-3.0E+1,-1.0E+0]
	movaps	%xmm0, 48(%r13)
	movl	%edx, 8(%rsp)                   # 4-byte Spill
	movslq	%edx, %rdx
	shlq	$2, %rdx
	movq	%rdx, 16(%rsp)                  # 8-byte Spill
	xorl	%ebx, %ebx
	xorl	%ebp, %ebp
	xorl	%edx, %edx
.LBB3_1:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB3_2 Depth 2
	movq	%rdx, 24(%rsp)                  # 8-byte Spill
	movq	%rbx, 32(%rsp)                  # 8-byte Spill
	movl	%ebp, 12(%rsp)                  # 4-byte Spill
	xorl	%r14d, %r14d
	movq	%rax, 40(%rsp)                  # 8-byte Spill
.LBB3_2:                                #   Parent Loop BB3_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movss	(%rax,%r14,4), %xmm0            # xmm0 = mem[0],zero,zero,zero
	movq	(%r12), %rax
	movss	%xmm0, (%rax,%rbx)
	movss	(%r13,%r14,4), %xmm0            # xmm0 = mem[0],zero,zero,zero
	movq	(%r15), %rcx
	movss	%xmm0, (%rcx,%rbx)
	xorps	%xmm0, %xmm0
	cvtss2sd	(%rax,%rbx), %xmm0
	movl	$.L.str.1, %edi
	movl	%ebp, %esi
	movb	$1, %al
	callq	printf
	movq	40(%rsp), %rax                  # 8-byte Reload
	incq	%r14
	incl	%ebp
	addq	$4, %rbx
	cmpq	$4, %r14
	jne	.LBB3_2
# %bb.3:                                #   in Loop: Header=BB3_1 Depth=1
	movq	24(%rsp), %rdx                  # 8-byte Reload
	incq	%rdx
	movl	12(%rsp), %ebp                  # 4-byte Reload
	addl	8(%rsp), %ebp                   # 4-byte Folded Reload
	movq	32(%rsp), %rbx                  # 8-byte Reload
	addq	16(%rsp), %rbx                  # 8-byte Folded Reload
	addq	$16, %rax
	addq	$16, %r13
	cmpq	$4, %rdx
	jne	.LBB3_1
# %bb.4:
	addq	$184, %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_end3:
	.size	_Z9secretKeyPPfS0_i, .Lfunc_end3-_Z9secretKeyPPfS0_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z7pad_msgPci                   # -- Begin function _Z7pad_msgPci
	.type	_Z7pad_msgPci,@function
_Z7pad_msgPci:                          # @_Z7pad_msgPci
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %rbp, -16
	movl	%esi, %ebp
	movq	%rdi, %rbx
	movl	$.Lstr, %edi
	callq	puts@PLT
	movq	%rbx, %rdi
	callq	strlen
	imull	%ebp, %ebp
                                        # kill: def $eax killed $eax killed $rax
	xorl	%edx, %edx
	divl	%ebp
	subl	%edx, %ebp
	movslq	%ebp, %r14
	movq	%r14, %rdi
	callq	malloc
	testl	%r14d, %r14d
	jle	.LBB4_3
# %bb.1:                                # %.lr.ph.preheader
	xorl	%ecx, %ecx
.LBB4_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movb	$32, (%rax,%rcx)
	incq	%rcx
	cmpq	%r14, %rcx
	jl	.LBB4_2
.LBB4_3:                                # %._crit_edge
	movq	%rbx, %rdi
	movq	%rax, %rsi
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	strcat                          # TAILCALL
.Lfunc_end4:
	.size	_Z7pad_msgPci, .Lfunc_end4-_Z7pad_msgPci
	.cfi_endproc
                                        # -- End function
	.globl	_Z10encode_msgPcPfi             # -- Begin function _Z10encode_msgPcPfi
	.type	_Z10encode_msgPcPfi,@function
_Z10encode_msgPcPfi:                    # @_Z10encode_msgPcPfi
	.cfi_startproc
# %bb.0:
	movb	(%rdi), %al
	testb	%al, %al
	je	.LBB5_3
# %bb.1:                                # %.lr.ph.preheader
	xorl	%ecx, %ecx
.LBB5_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movsbl	%al, %eax
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	movss	%xmm0, (%rsi,%rcx,4)
	movb	1(%rdi,%rcx), %al
	incq	%rcx
	testb	%al, %al
	jne	.LBB5_2
.LBB5_3:                                # %._crit_edge
	retq
.Lfunc_end5:
	.size	_Z10encode_msgPcPfi, .Lfunc_end5-_Z10encode_msgPcPfi
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset %rbp, -16
	movq	%rsp, %rbp
	.cfi_def_cfa_register %rbp
	pushq	%r15
	pushq	%r14
	pushq	%r13
	pushq	%r12
	pushq	%rbx
	subq	$1544, %rsp                     # imm = 0x608
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	cmpl	$2, %edi
	jl	.LBB6_4
# %bb.1:
	movq	%rsi, %rbx
	movq	8(%rsi), %rdi
	movl	$.L.str.3, %esi
	callq	fopen
	testq	%rax, %rax
	je	.LBB6_3
# %bb.2:
	movl	$.Lstr.1, %edi
	callq	puts@PLT
	movq	8(%rbx), %rdi
	movl	$.L.str.3, %esi
	callq	fopen
	movq	%rax, %rbx
	movq	%rax, %rdi
	xorl	%esi, %esi
	movl	$2, %edx
	callq	fseek
	movq	%rbx, %rdi
	callq	ftell
	movq	%rax, %r15
	movq	%rbx, %rdi
	xorl	%esi, %esi
	xorl	%edx, %edx
	callq	fseek
	leaq	1(%r15), %rdi
	callq	malloc
	movq	%rax, %r14
	movl	$1, %edx
	movq	%rax, %rdi
	movq	%r15, %rsi
	movq	%rbx, %rcx
	callq	fread
	movq	%rbx, %rdi
	callq	fclose
	jmp	.LBB6_5
.LBB6_4:
	movl	$43, %edi
	callq	malloc
	movq	%rax, %r14
	movups	.L.str.5(%rip), %xmm0
	movups	%xmm0, (%rax)
	movups	.L.str.5+16(%rip), %xmm0
	movups	%xmm0, 16(%rax)
	movups	.L.str.5+28(%rip), %xmm0
	movups	%xmm0, 28(%rax)
	jmp	.LBB6_5
.LBB6_3:
	movq	8(%rbx), %rbx
	movq	%rbx, %rdi
	callq	strlen
	leaq	15(%rax), %rdi
	callq	malloc
	movq	%rax, %r14
	movq	%rax, %rdi
	movq	%rbx, %rsi
	callq	strcpy
.LBB6_5:
	movl	$32, %edi
	callq	malloc
	movq	%rax, %r15
	xorl	%ebx, %ebx
.LBB6_6:                                # =>This Inner Loop Header: Depth=1
	movl	$16, %edi
	callq	malloc
	movq	%rax, (%r15,%rbx,8)
	incq	%rbx
	cmpq	$4, %rbx
	jne	.LBB6_6
# %bb.7:
	movl	$32, %edi
	callq	malloc
	movq	%rax, %r12
	xorl	%ebx, %ebx
.LBB6_8:                                # =>This Inner Loop Header: Depth=1
	movl	$16, %edi
	callq	malloc
	movq	%rax, (%r12,%rbx,8)
	incq	%rbx
	cmpq	$4, %rbx
	jne	.LBB6_8
# %bb.9:
	movq	%r15, %rdi
	movq	%r12, -104(%rbp)                # 8-byte Spill
	movq	%r12, %rsi
	movl	$4, %edx
	callq	_Z9secretKeyPPfS0_i
	xorl	%ebx, %ebx
	movl	$.L.str.6, %edi
	movq	%r14, %rsi
	xorl	%eax, %eax
	callq	printf
	movq	(%r15), %r15
	movq	%r15, -112(%rbp)                # 8-byte Spill
.LBB6_10:                               # %.preheader197
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB6_11 Depth 2
	xorl	%r12d, %r12d
.LBB6_11:                               #   Parent Loop BB6_10 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	xorps	%xmm0, %xmm0
	cvtss2sd	(%r15,%r12,4), %xmm0
	movl	$.L.str.7, %edi
	movb	$1, %al
	callq	printf
	incq	%r12
	cmpq	$4, %r12
	jne	.LBB6_11
# %bb.12:                               #   in Loop: Header=BB6_10 Depth=1
	incq	%rbx
	addq	$16, %r15
	cmpq	$4, %rbx
	jne	.LBB6_10
# %bb.13:
	movq	%r14, %rdi
	callq	strlen
	testb	$15, %al
	je	.LBB6_15
# %bb.14:
	movq	%r14, %rdi
	movl	$4, %esi
	callq	_Z7pad_msgPci
.LBB6_15:
	movq	%r14, %rdi
	callq	strlen
	leal	15(%rax), %ecx
	testl	%eax, %eax
	movq	%rax, -48(%rbp)                 # 8-byte Spill
	cmovnsl	%eax, %ecx
	sarl	$4, %ecx
	movq	%rcx, -88(%rbp)                 # 8-byte Spill
	movl	$.L.str.6, %edi
	movq	%r14, %rsi
	xorl	%eax, %eax
	callq	printf
	movq	%r14, %rdi
	callq	strlen
	leaq	(,%rax,4), %rdi
	callq	malloc
	movq	%rax, %r15
	movb	(%r14), %al
	testb	%al, %al
	je	.LBB6_18
# %bb.16:                               # %.lr.ph.i.preheader
	xorl	%ecx, %ecx
.LBB6_17:                               # %.lr.ph.i
                                        # =>This Inner Loop Header: Depth=1
	movsbl	%al, %eax
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	movss	%xmm0, (%r15,%rcx,4)
	movb	1(%r14,%rcx), %al
	incq	%rcx
	testb	%al, %al
	jne	.LBB6_17
.LBB6_18:                               # %_Z10encode_msgPcPfi.exit
	movl	$.Lstr.2, %edi
	callq	puts@PLT
	movl	-88(%rbp), %eax                 # 4-byte Reload
	movq	%rsp, %rbx
	movq	%rax, -56(%rbp)                 # 8-byte Spill
	shlq	$6, %rax
	subq	%rax, %rbx
	movq	%rbx, %rsp
	cmpl	$16, -48(%rbp)                  # 4-byte Folded Reload
	jl	.LBB6_23
# %bb.19:                               # %.lr.ph
	movq	-48(%rbp), %rax                 # 8-byte Reload
                                        # kill: def $eax killed $eax killed $rax
	andl	$15, %eax
	movl	%eax, -92(%rbp)                 # 4-byte Spill
	movq	%rbx, %r13
	xorl	%r12d, %r12d
.LBB6_20:                               # =>This Loop Header: Depth=1
                                        #     Child Loop BB6_21 Depth 2
	movl	$.L.str.9, %edi
	movl	%r12d, %esi
	movl	-92(%rbp), %edx                 # 4-byte Reload
	xorl	%eax, %eax
	callq	printf
	xorl	%r14d, %r14d
.LBB6_21:                               #   Parent Loop BB6_20 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movss	(%r15,%r14,4), %xmm0            # xmm0 = mem[0],zero,zero,zero
	movss	%xmm0, (%r13,%r14,4)
	cvtss2sd	%xmm0, %xmm0
	movl	$.L.str.10, %edi
	movl	$.L.str.11, %esi
	movb	$1, %al
	callq	printf
	incq	%r14
	cmpq	$16, %r14
	jne	.LBB6_21
# %bb.22:                               #   in Loop: Header=BB6_20 Depth=1
	incq	%r12
	addq	$64, %r15
	addq	$64, %r13
	cmpq	-56(%rbp), %r12                 # 8-byte Folded Reload
	jne	.LBB6_20
.LBB6_23:                               # %._crit_edge
	movl	$.Lstr.3, %edi
	callq	puts@PLT
	movq	-48(%rbp), %r14                 # 8-byte Reload
	movslq	%r14d, %rdi
	shlq	$3, %rdi
	callq	malloc
	movq	%rax, %r15
	movl	$.Lstr.4, %edi
	callq	puts@PLT
	cmpl	$16, %r14d
	movq	-56(%rbp), %r12                 # 8-byte Reload
	jl	.LBB6_26
# %bb.24:                               # %.lr.ph206.preheader
	xorl	%r14d, %r14d
.LBB6_25:                               # %.lr.ph206
                                        # =>This Inner Loop Header: Depth=1
	movl	$1, %edi
	movl	$64, %esi
	callq	calloc@PLT
	movq	%rax, (%r15,%r14,8)
	incq	%r14
	cmpq	%r14, %r12
	jne	.LBB6_25
.LBB6_26:                               # %._crit_edge207
	movl	$.Lstr.5, %edi
	callq	puts@PLT
	leaq	-1584(%rbp), %rdi
	xorl	%esi, %esi
	callq	hipGetDevicePropertiesR0600
	xorl	%edi, %edi
	callq	hipSetDevice
	leaq	-80(%rbp), %r14
	movl	$64, %esi
	movq	%r14, %rdi
	callq	hipMalloc
	movl	$.L.str.15, %edi
	movq	-88(%rbp), %rsi                 # 8-byte Reload
                                        # kill: def $esi killed $esi killed $rsi
	xorl	%eax, %eax
	callq	printf
	xorps	%xmm0, %xmm0
	cvtss2sd	64(%rbx), %xmm0
	movl	$.L.str.16, %edi
	movb	$1, %al
	callq	printf
	leaq	-64(%rbp), %rdi
	movl	$64, %esi
	callq	hipMalloc
	leaq	-72(%rbp), %rdi
	movl	$64, %esi
	callq	hipMalloc
	movq	(%r14), %rdi
	movl	$64, %edx
	movq	-112(%rbp), %rsi                # 8-byte Reload
	movl	$1, %ecx
	callq	hipMemcpy
	cmpl	$16, -48(%rbp)                  # 4-byte Folded Reload
	jl	.LBB6_50
# %bb.27:                               # %.lr.ph210.preheader
	movabsq	$4294967297, %r12               # imm = 0x100000001
	xorl	%r14d, %r14d
	movq	-56(%rbp), %r13                 # 8-byte Reload
.LBB6_28:                               # %.lr.ph210
                                        # =>This Inner Loop Header: Depth=1
	movq	-72(%rbp), %rdi
	movl	$64, %edx
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	%r12, %rdi
	movl	$1, %esi
	movabsq	$17179869188, %rdx              # imm = 0x400000004
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB6_30
# %bb.29:                               #   in Loop: Header=BB6_28 Depth=1
	movq	-80(%rbp), %rdi
	movq	-72(%rbp), %rsi
	movq	-64(%rbp), %rdx
	movl	$4, %ecx
	callq	_Z25__device_stub__mtxEncryptPfS_S_i
.LBB6_30:                               #   in Loop: Header=BB6_28 Depth=1
	movq	(%r15,%r14,8), %rdi
	movq	-64(%rbp), %rsi
	movl	$64, %edx
	movl	$2, %ecx
	callq	hipMemcpy
	callq	hipDeviceSynchronize
	incq	%r14
	addq	$64, %rbx
	cmpq	%r14, %r13
	jne	.LBB6_28
# %bb.31:                               # %.preheader196
	cmpl	$16, -48(%rbp)                  # 4-byte Folded Reload
	jl	.LBB6_50
# %bb.32:                               # %.preheader195.preheader
	xorl	%ebx, %ebx
.LBB6_33:                               # %.preheader195
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB6_34 Depth 2
	movq	(%r15,%rbx,8), %r14
	xorl	%r13d, %r13d
.LBB6_34:                               #   Parent Loop BB6_33 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	cvttss2si	(%r14,%r13,4), %edi
	callq	putchar@PLT
	incq	%r13
	cmpq	$16, %r13
	jne	.LBB6_34
# %bb.35:                               #   in Loop: Header=BB6_33 Depth=1
	incq	%rbx
	cmpq	-56(%rbp), %rbx                 # 8-byte Folded Reload
	jne	.LBB6_33
# %bb.36:                               # %._crit_edge213
	movq	-88(%rbp), %rdi                 # 8-byte Reload
	shll	$3, %edi
	callq	malloc
	cmpl	$16, -48(%rbp)                  # 4-byte Folded Reload
	jl	.LBB6_50
# %bb.37:                               # %.lr.ph216.preheader
	movq	%rax, %r13
	xorl	%ebx, %ebx
	movq	-56(%rbp), %r14                 # 8-byte Reload
.LBB6_38:                               # %.lr.ph216
                                        # =>This Inner Loop Header: Depth=1
	movl	$1, %edi
	movl	$64, %esi
	callq	calloc@PLT
	movq	%rax, (%r13,%rbx,8)
	incq	%rbx
	cmpq	%rbx, %r14
	jne	.LBB6_38
# %bb.39:                               # %._crit_edge217
	movq	-80(%rbp), %rdi
	movq	-104(%rbp), %rax                # 8-byte Reload
	movq	(%rax), %rsi
	movl	$64, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	cmpl	$16, -48(%rbp)                  # 4-byte Folded Reload
	jl	.LBB6_49
# %bb.40:                               # %.lr.ph220.preheader
	xorl	%r14d, %r14d
.LBB6_41:                               # %.lr.ph220
                                        # =>This Inner Loop Header: Depth=1
	movq	-72(%rbp), %rdi
	movq	(%r15,%r14,8), %rsi
	movl	$64, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	%r12, %rdi
	movl	$1, %esi
	movabsq	$17179869188, %rdx              # imm = 0x400000004
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB6_43
# %bb.42:                               #   in Loop: Header=BB6_41 Depth=1
	movq	-80(%rbp), %rdi
	movq	-72(%rbp), %rsi
	movq	-64(%rbp), %rdx
	movl	$4, %ecx
	callq	_Z25__device_stub__mtxEncryptPfS_S_i
.LBB6_43:                               #   in Loop: Header=BB6_41 Depth=1
	movq	(%r13,%r14,8), %rdi
	movq	-64(%rbp), %rsi
	movl	$64, %edx
	movl	$2, %ecx
	callq	hipMemcpy
	callq	hipDeviceSynchronize
	incq	%r14
	cmpq	%r14, -56(%rbp)                 # 8-byte Folded Reload
	jne	.LBB6_41
# %bb.44:                               # %.preheader194
	cmpl	$16, -48(%rbp)                  # 4-byte Folded Reload
	jl	.LBB6_49
# %bb.45:                               # %.preheader.preheader
	xorl	%ebx, %ebx
.LBB6_46:                               # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB6_47 Depth 2
	movq	(%r13,%rbx,8), %r14
	xorl	%r12d, %r12d
.LBB6_47:                               #   Parent Loop BB6_46 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	cvttss2si	(%r14,%r12,4), %edi
	callq	putchar@PLT
	incq	%r12
	cmpq	$16, %r12
	jne	.LBB6_47
# %bb.48:                               #   in Loop: Header=BB6_46 Depth=1
	incq	%rbx
	cmpq	-56(%rbp), %rbx                 # 8-byte Folded Reload
	jne	.LBB6_46
	jmp	.LBB6_49
.LBB6_50:                               # %._crit_edge217.thread
	movq	-80(%rbp), %rdi
	movq	-104(%rbp), %rax                # 8-byte Reload
	movq	(%rax), %rsi
	movl	$64, %edx
	movl	$1, %ecx
	callq	hipMemcpy
.LBB6_49:                               # %._crit_edge223
	movq	%r15, %rdi
	callq	free
	movq	-80(%rbp), %rdi
	callq	hipFree
	movq	-72(%rbp), %rdi
	callq	hipFree
	movq	-64(%rbp), %rdi
	callq	hipFree
	callq	hipDeviceReset
	xorl	%eax, %eax
	leaq	-40(%rbp), %rsp
	popq	%rbx
	popq	%r12
	popq	%r13
	popq	%r14
	popq	%r15
	popq	%rbp
	.cfi_def_cfa %rsp, 8
	retq
.Lfunc_end6:
	.size	main, .Lfunc_end6-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	.LBB7_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB7_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z10mtxEncryptPfS_S_i, %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_end7:
	.size	__hip_module_ctor, .Lfunc_end7-__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	.LBB8_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
.LBB8_2:
	retq
.Lfunc_end8:
	.size	__hip_module_dtor, .Lfunc_end8-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"PRINTROW: %f\n"
	.size	.L.str, 14

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

	.type	.L.str.1,@object                # @.str.1
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.1:
	.asciz	"%f :  %f \n"
	.size	.L.str.1, 11

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

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"All work and no play makes Jack a dull boy."
	.size	.L.str.5, 44

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"%s"
	.size	.L.str.6, 3

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"Skey: %f\n"
	.size	.L.str.7, 10

	.type	.L.str.9,@object                # @.str.9
.L.str.9:
	.asciz	"\n\nCount: %d %d\n"
	.size	.L.str.9, 16

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

	.type	.L.str.11,@object               # @.str.11
.L.str.11:
	.asciz	"\n"
	.size	.L.str.11, 2

	.type	.L.str.15,@object               # @.str.15
.L.str.15:
	.asciz	"msg_parts: %d \n"
	.size	.L.str.15, 16

	.type	.L.str.16,@object               # @.str.16
.L.str.16:
	.asciz	"\nmsg_vector_list[1]: %d\n"
	.size	.L.str.16, 25

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z10mtxEncryptPfS_S_i"
	.size	.L__unnamed_1, 22

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

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

	.type	.Lstr,@object                   # @str
	.section	.rodata.str1.1,"aMS",@progbits,1
.Lstr:
	.asciz	"Padding.."
	.size	.Lstr, 10

	.type	.Lstr.1,@object                 # @str.1
.Lstr.1:
	.asciz	"reading file\n"
	.size	.Lstr.1, 14

	.type	.Lstr.2,@object                 # @str.2
.Lstr.2:
	.asciz	"ENCODED!"
	.size	.Lstr.2, 9

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

	.type	.Lstr.4,@object                 # @str.4
.Lstr.4:
	.asciz	"2"
	.size	.Lstr.4, 2

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

	.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 _Z25__device_stub__mtxEncryptPfS_S_i
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z10mtxEncryptPfS_S_i
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z10mtxEncryptPfS_S_i   ; -- Begin function _Z10mtxEncryptPfS_S_i
	.globl	_Z10mtxEncryptPfS_S_i
	.p2align	8
	.type	_Z10mtxEncryptPfS_S_i,@function
_Z10mtxEncryptPfS_S_i:                  ; @_Z10mtxEncryptPfS_S_i
; %bb.0:                                ; %.preheader
	s_clause 0x2
	s_load_b32 s2, s[0:1], 0x18
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b64 s[0:1], s[0:1], 0x10
	v_bfe_u32 v1, v0, 10, 10
	v_and_b32_e32 v0, 0x3ff, v0
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_2)
	v_mul_lo_u32 v1, v1, s2
	s_cmp_lt_i32 s2, 1
	s_cbranch_scc1 .LBB0_4
; %bb.1:                                ; %.lr.ph.preheader
	v_mov_b32_e32 v3, 0
	s_mov_b32 s3, s2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mov_b32_e32 v2, v3
	v_lshlrev_b64 v[4:5], 2, v[1:2]
	v_mov_b32_e32 v2, v0
	v_mov_b32_e32 v6, v3
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v4, vcc_lo, s4, v4
	v_add_co_ci_u32_e32 v5, vcc_lo, s5, v5, vcc_lo
.LBB0_2:                                ; %.lr.ph
                                        ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	v_lshlrev_b64 v[7:8], 2, v[2:3]
	v_add_nc_u32_e32 v2, s2, v2
	s_add_i32 s3, s3, -1
	s_cmp_eq_u32 s3, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v7, vcc_lo, s6, v7
	v_add_co_ci_u32_e32 v8, vcc_lo, s7, v8, vcc_lo
	global_load_b32 v9, v[4:5], off
	global_load_b32 v7, v[7:8], off
	v_add_co_u32 v4, vcc_lo, v4, 4
	v_add_co_ci_u32_e32 v5, vcc_lo, 0, v5, vcc_lo
	s_waitcnt vmcnt(0)
	v_fmac_f32_e32 v6, v9, v7
	s_cbranch_scc0 .LBB0_2
; %bb.3:                                ; %._crit_edge.loopexit
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_i32_f32_e32 v2, v6
	v_and_b32_e32 v2, 0x7f, v2
	s_delay_alu instid0(VALU_DEP_1)
	v_cvt_f32_ubyte0_e32 v2, v2
	s_branch .LBB0_5
.LBB0_4:
	v_mov_b32_e32 v2, 0
.LBB0_5:                                ; %Flow
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v0, v1, v0
	v_ashrrev_i32_e32 v1, 31, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 2, v[0:1]
	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_store_b32 v[0:1], v2, off
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z10mtxEncryptPfS_S_i
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 28
		.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 1
		.amdhsa_next_free_vgpr 10
		.amdhsa_next_free_sgpr 8
		.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	_Z10mtxEncryptPfS_S_i, .Lfunc_end0-_Z10mtxEncryptPfS_S_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 288
; NumSgprs: 10
; NumVgprs: 10
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 1
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 10
; NumVGPRsForWavesPerEU: 10
; 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: 1
	.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
      - .offset:         24
        .size:           4
        .value_kind:     by_value
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 28
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z10mtxEncryptPfS_S_i
    .private_segment_fixed_size: 0
    .sgpr_count:     10
    .sgpr_spill_count: 0
    .symbol:         _Z10mtxEncryptPfS_S_i.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     10
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
