#include <stdio.h>
#include <stdlib.h>
#include <time.h>

__device__ int IP[64];
__device__ int FP[64];
__device__ int E[48];
__device__ int P[32];
__device__ int SBox[8][64];


// Initial Permutation
int host_IP[64] = {
    57, 49, 41, 33, 25, 17,  9,  1,
    59, 51, 43, 35, 27, 19, 11,  3,
    61, 53, 45, 37, 29, 21, 13,  5,
    63, 55, 47, 39, 31, 23, 15,  7,
    56, 48, 40, 32, 24, 16,  8,  0,
    58, 50, 42, 34, 26, 18, 10,  2,
    60, 52, 44, 36, 28, 20, 12,  4,
    62, 54, 46, 38, 30, 22, 14,  6
};

// Final Permutation
int host_FP[64] = {
    39,  7, 47, 15, 55, 23, 63, 31,
    38,  6, 46, 14, 54, 22, 62, 30,
    37,  5, 45, 13, 53, 21, 61, 29,
    36,  4, 44, 12, 52, 20, 60, 28,
    35,  3, 43, 11, 51, 19, 59, 27,
    34,  2, 42, 10, 50, 18, 58, 26,
    33,  1, 41,  9, 49, 17, 57, 25,
    32,  0, 40,  8, 48, 16, 56, 24
};

// Expansion Function: from 32 bit to 48 bit
int host_E[48] = {
    31,  0,  1,  2,  3,  4,
     3,  4,  5,  6,  7,  8,
     7,  8,  9, 10, 11, 12,
    11, 12, 13, 14, 15, 16,
    15, 16, 17, 18, 19, 20,
    19, 20, 21, 22, 23, 24,
    23, 24, 25, 26, 27, 28,
    27, 28, 29, 30, 31,  0
};

// Permutation
int host_P[32] = {
    15,  6, 19, 20, 28, 11, 27, 16,
     0, 14, 22, 25,  4, 17, 30,  9,
     1,  7, 23, 13, 31, 26,  2,  8,
    18, 12, 29,  5, 21, 10,  3, 24
};

int PC1_LEFT[28] = {
    56, 48, 40, 32, 24, 16,  8,
     0, 57, 49, 41, 33, 25, 17,
     9,  1, 58, 50, 42, 34, 26,
    18, 10,  2, 59, 51, 43, 35,
};

int PC1_RIGHT[28] = {
    62, 54, 46, 38, 30, 22, 14,
     6, 61, 53, 45, 37, 29, 21,
    13,  5, 60, 52, 44, 36, 28,
    20, 12,  4, 27, 19, 11,  3
};

int PC2[48] = {
    13, 16, 10, 23,  0,  4,
     2, 27, 14,  5, 20,  9,
    22, 18, 11,  3, 25,  7,
    15,  6, 26, 19, 12,  1,
    40, 51, 30, 36, 46, 54,
    29, 39, 50, 44, 32, 47,
    43, 48, 38, 55, 33, 52,
    45, 41, 49, 35, 28, 31
};

int Rotations[16] = {1, 1, 2, 2, 2, 2, 2, 2, 1, 2, 2, 2, 2, 2, 2, 1};

// Substitution Boxes
int host_SBox[8][64] = {
    // S1
    {
        14, 4, 13, 1, 2, 15, 11, 8, 3, 10, 6, 12, 5, 9, 0, 7,
        0, 15, 7, 4, 14, 2, 13, 1, 10, 6, 12, 11, 9, 5, 3, 8,
        4, 1, 14, 8, 13, 6, 2, 11, 15, 12, 9, 7, 3, 10, 5, 0,
        15, 12, 8, 2, 4, 9, 1, 7, 5, 11, 3, 14, 10, 0, 6, 13
    },

    // S2
    {
        15, 1, 8, 14, 6, 11, 3, 4, 9, 7, 2, 13, 12, 0, 5, 10,
        3, 13, 4, 7, 15, 2, 8, 14, 12, 0, 1, 10, 6, 9, 11, 5,
        0, 14, 7, 11, 10, 4, 13, 1, 5, 8, 12, 6, 9, 3, 2, 15,
        13, 8, 10, 1, 3, 15, 4, 2, 11, 6, 7, 12, 0, 5, 14, 9
    },

    // S3
    {
        10, 0, 9, 14, 6, 3, 15, 5, 1, 13, 12, 7, 11, 4, 2, 8,
        13, 7, 0, 9, 3, 4, 6, 10, 2, 8, 5, 14, 12, 11, 15, 1,
        13, 6, 4, 9, 8, 15, 3, 0, 11, 1, 2, 12, 5, 10, 14, 7,
        1, 10, 13, 0, 6, 9, 8, 7, 4, 15, 14, 3, 11, 5, 2, 12
    },

    // S4
    {
        7, 13, 14, 3, 0, 6, 9, 10, 1, 2, 8, 5, 11, 12, 4, 15,
        13, 8, 11, 5, 6, 15, 0, 3, 4, 7, 2, 12, 1, 10, 14, 9,
        10, 6, 9, 0, 12, 11, 7, 13, 15, 1, 3, 14, 5, 2, 8, 4,
        3, 15, 0, 6, 10, 1, 13, 8, 9, 4, 5, 11, 12, 7, 2, 14
    },

    // S5
    {
        2, 12, 4, 1, 7, 10, 11, 6, 8, 5, 3, 15, 13, 0, 14, 9,
        14, 11, 2, 12, 4, 7, 13, 1, 5, 0, 15, 10, 3, 9, 8, 6,
        4, 2, 1, 11, 10, 13, 7, 8, 15, 9, 12, 5, 6, 3, 0, 14,
        11, 8, 12, 7, 1, 14, 2, 13, 6, 15, 0, 9, 10, 4, 5, 3
    },

    // S6
    {
        12, 1, 10, 15, 9, 2, 6, 8, 0, 13, 3, 4, 14, 7, 5, 11,
        10, 15, 4, 2, 7, 12, 9, 5, 6, 1, 13, 14, 0, 11, 3, 8,
        9, 14, 15, 5, 2, 8, 12, 3, 7, 0, 4, 10, 1, 13, 11, 6,
        4, 3, 2, 12, 9, 5, 15, 10, 11, 14, 1, 7, 6, 0, 8, 13
    },

    // S7
    {
        4, 11, 2, 14, 15, 0, 8, 13, 3, 12, 9, 7, 5, 10, 6, 1,
        13, 0, 11, 7, 4, 9, 1, 10, 14, 3, 5, 12, 2, 15, 8, 6,
        1, 4, 11, 13, 12, 3, 7, 14, 10, 15, 6, 8, 0, 5, 9, 2,
        6, 11, 13, 8, 1, 4, 10, 7, 9, 5, 0, 15, 14, 2, 3, 12
    },

    // S8
    {
        13, 2, 8, 4, 6, 15, 11, 1, 10, 9, 3, 14, 5, 0, 12, 7,
        1, 15, 13, 8, 10, 3, 7, 4, 12, 5, 6, 11, 0, 14, 9, 2,
        7, 11, 4, 1, 9, 12, 14, 2, 0, 6, 10, 13, 15, 3, 5, 8,
        2, 1, 14, 7, 4, 10, 8, 13, 15, 12, 9, 0, 3, 5, 6, 11
    }
};

__device__
__host__
long long int permutation(long long int data, int data_size, int *table, int table_size) {
    long long int result = 0;
    int i = 0;
    for (; i < table_size; i++) {
        result = (result << 1) + ((data >> (data_size - 1 - table[i])) & 0x1);
    }
    return result;
}

long long int *generate_sub_keys(long long int key, int decrypt) {
    int n_keys = 16;
    long long int *sub_keys = (long long int *) malloc(sizeof(long long int) * n_keys);
    int half_key_length = 28;
    long long int left = permutation(key, 64, PC1_LEFT, half_key_length);
    long long int right = permutation(key, 64, PC1_RIGHT, half_key_length);
    int i = 0;
    for (; i < n_keys; i++) {
        int rotation = Rotations[i];
        left = (((left << rotation) | (left >> (half_key_length - rotation))) & 0xFFFFFFF);
        right = (((right << rotation) | (right >> (half_key_length - rotation))) & 0xFFFFFFF);
        long long int new_key = (left << half_key_length) | right;
        int sub_key_index = (decrypt ? 15 - i : i);
        sub_keys[sub_key_index] = permutation(new_key, half_key_length * 2, PC2, 48);
    }
    return sub_keys;
}

__device__
long long int substitution(long long int data) {
    // data: 48 bit
    long long int result = 0;
    int i = 0;
    for (; i < 8; i++) {
        unsigned int box = data >> (6 * (7 - i)) & 0x3F;
        int outer = ((box & 0x20) >> 4) | (box & 0x1);
        int inner = (box & 0x1E) >> 1;
        result = (result << 4) + SBox[i][(outer << 4) + inner];
    }
    return result;
}

__device__
long int F(unsigned int c, long long int key) {
    long long int lc = c;
    long long int new_c = permutation(lc, 32, E, 48);
    long long int mixed_data = new_c ^ key;
    long long int s_box_result = substitution(mixed_data);
    return permutation(s_box_result, 32, P, 32);
}

__device__
void DES(int index, long long int *MD, long long int *keys) {
    long long int data = permutation(MD[index], 64, IP, 64);
    unsigned int left = data >> 32;
    unsigned int right = (int) data;
    int i = 0;
    for (; i < 16; i++) {
        unsigned int buf = left ^ F(right, keys[i]);
        left = right;
        right = buf;
    }
    data = right;
    data = (data << 32) + left;
    MD[index] = permutation(data, 64, FP, 64);
}

__global__
void kernel_DES(unsigned int quota, unsigned int n_blocks, long long int *MD, long long int *keys) {
    int start_index = (blockIdx.x * blockDim.x + threadIdx.x) * quota;
    int end_index = start_index + quota;
    int i;
    for (i = start_index; i < end_index; i++) {
        if (i >= n_blocks) {
        return;
        }
        DES(i, MD, keys);
    }
}

void runDESCuda(unsigned int n_blocks, long long int *host_MD, long long int *host_sub_keys, int n_cuda_blocks, int n_cuda_threads) {
    unsigned int max_n_threads = 512;
    if (n_cuda_blocks == -1 && n_cuda_threads == -1) {
        n_cuda_threads = max_n_threads;
        n_cuda_blocks = (n_blocks + max_n_threads - 1) / max_n_threads;
    }
    if (n_cuda_threads > max_n_threads) {
        printf("Maximum value of the number of threads is 512. You entered : %d\n", n_cuda_threads);
        return;
    }
    cudaMemcpyToSymbol(IP, host_IP, sizeof(host_IP));
    cudaMemcpyToSymbol(FP, host_FP, sizeof(host_FP));
    cudaMemcpyToSymbol(E, host_E, sizeof(host_E));
    cudaMemcpyToSymbol(P, host_P, sizeof(host_P));
    cudaMemcpyToSymbol(SBox, host_SBox, sizeof(host_SBox));

    long long int *MD, *sub_keys;
    cudaMalloc((void **) &MD, sizeof(long long int) * n_blocks);
    cudaMemcpy(MD, host_MD, sizeof(long long int) * n_blocks, cudaMemcpyHostToDevice);
    cudaMalloc((void **) &sub_keys, sizeof(long long int) * 16);
    cudaMemcpy(sub_keys, host_sub_keys, sizeof(long long int) * 16, cudaMemcpyHostToDevice);
    unsigned int quota = n_blocks / (n_cuda_blocks * n_cuda_threads) + 1;
    printf("%u bytes per threads.\n", quota * 64);
    kernel_DES<<<n_cuda_blocks, n_cuda_threads>>>(quota, n_blocks, MD, sub_keys);
    cudaMemcpy(host_MD, MD, sizeof(long long int) * n_blocks, cudaMemcpyDeviceToHost);

    cudaFree(IP);
    cudaFree(FP);
    cudaFree(E);
    cudaFree(P);
    cudaFree(SBox);
    cudaFree(MD);
    cudaFree(sub_keys);
}

unsigned int n_blocks = 0;
void des_with_file(int decrypt, char *in, char *out, char *key, int n_cuda_blocks, int n_cuda_threads) {
    int buf_size = 8 * n_blocks;
    char *buf = (char *) malloc(sizeof(char) * buf_size);
    FILE *in_fp = fopen(in, "rb");
    if (in_fp == NULL) {
        printf("Can't open the in file :%s\n", in);
        return;
    }
    fread(buf, buf_size, 1, in_fp);
    fclose(in_fp);

    long long int *MD = (long long int *) malloc(sizeof(long long int) * n_blocks);
    int i = 0;
    int j = 0;
    for (i = 0; i < n_blocks; i++) {
        long long int block = 0;
        for (j = 0; j < 8; j++) {
            block = (block << 8) + (buf[(i * 8) + j] & 0xFF);
        }
        MD[i] = block;
    }
    long long int binary_key = 0;
    for (i = 0; i < 8; i++) {
        binary_key = (binary_key << 8) + (key[i] & 0xFF);
    }

    long long int *sub_keys = generate_sub_keys(binary_key, decrypt);
    clock_t begin = clock();
    runDESCuda(n_blocks, MD, sub_keys, n_cuda_blocks, n_cuda_threads);
    clock_t end = clock();
    double time_spent = (double) (end - begin) / CLOCKS_PER_SEC;
    printf("CUDA time: %f\n", time_spent);
    free(sub_keys);

    for (i = 0; i < n_blocks; i++) {
        for (j = 0; j < 8; j++) {
            buf[(i * 8) + (7 - j)] = ((MD[i] >> (j * 8)) & 0xFF);
        }
    }
    FILE *out_fp = fopen(out, "wb");
    if (out_fp == NULL) {
        printf("Can't open the out file :%s\n", out);
        return;
    }
    fwrite(buf, buf_size, 1, out_fp);
    fclose(out_fp);
    free(buf);
    free(MD);
}

void encryption(char *in, char *out, char *key, int n_cuda_blocks, int n_cuda_threads) {
    des_with_file(0, in, out, key, n_cuda_blocks, n_cuda_threads);
}

void decryption(char *in, char *out, char *key, int n_cuda_blocks, int n_cuda_threads) {
    des_with_file(1, in, out, key, n_cuda_blocks, n_cuda_threads);
}

int main(int argc, char** argv) {
    if (argc < 6) {
        printf("usage) ./des.out [e|d] <input_file> <output_file> <n_des_block_size> <n_cuda_blocks> <n_cuda_threads>\n");
        printf("example) ./des.out e in.txt out.txt 1\n");
        return -1;
    }
    int n_cuda_blocks;
    int n_cuda_threads;
    sscanf(argv[5], "%d", &n_blocks);
    sscanf(argv[6], "%d", &n_cuda_blocks);
    sscanf(argv[7], "%d", &n_cuda_threads);
    switch(argv[1][0]) {
        case 'e':
            printf("encryption\n");
            encryption(argv[2], argv[3], argv[4], n_cuda_blocks, n_cuda_threads);
            break;
        case 'd':
            printf("decryption\n");
            decryption(argv[2], argv[3], argv[4], n_cuda_blocks, n_cuda_threads);
            break;
        default:
            printf("mode must be 'e' or 'd'\n");
    }
}

	.file	"tmpxft_00242d17_00000000-6_des.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2039:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2039:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z11permutationxiPii
	.type	_Z11permutationxiPii, @function
_Z11permutationxiPii:
.LFB2027:
	.cfi_startproc
	endbr64
	movq	%rdi, %r9
	movl	%ecx, %r8d
	xorl	%edi, %edi
	xorl	%eax, %eax
	decl	%esi
.L3:
	cmpl	%edi, %r8d
	jle	.L6
	movl	%esi, %ecx
	movq	%r9, %r10
	subl	(%rdx,%rdi,4), %ecx
	addq	%rax, %rax
	sarq	%cl, %r10
	incq	%rdi
	movq	%r10, %rcx
	andl	$1, %ecx
	addq	%rcx, %rax
	jmp	.L3
.L6:
	ret
	.cfi_endproc
.LFE2027:
	.size	_Z11permutationxiPii, .-_Z11permutationxiPii
	.globl	_Z17generate_sub_keysxi
	.type	_Z17generate_sub_keysxi, @function
_Z17generate_sub_keysxi:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	leaq	Rotations(%rip), %r15
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movl	%esi, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rdi, %rbp
	movl	$128, %edi
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	pushq	%rcx
	.cfi_def_cfa_offset 64
	call	malloc@PLT
	movq	%rbp, %rdi
	movl	$28, %ecx
	leaq	PC1_LEFT(%rip), %rdx
	movl	$64, %esi
	movq	%rax, %r12
	call	_Z11permutationxiPii
	leaq	PC1_RIGHT(%rip), %rdx
	movq	%rbp, %rdi
	movl	$28, %ecx
	movl	$64, %esi
	movq	%rax, %rbx
	xorl	%ebp, %ebp
	call	_Z11permutationxiPii
	leaq	PC2(%rip), %rdx
	movq	%rax, %r11
.L10:
	movl	(%r15,%rbp,4), %eax
	movl	$28, %edi
	movq	%rbx, %rsi
	subl	%eax, %edi
	movl	%eax, %ecx
	salq	%cl, %rsi
	movl	%edi, %ecx
	sarq	%cl, %rbx
	movl	%eax, %ecx
	orq	%rbx, %rsi
	movq	%rsi, %rbx
	movq	%r11, %rsi
	salq	%cl, %rsi
	movl	%edi, %ecx
	andl	$268435455, %ebx
	sarq	%cl, %r11
	movq	%rsi, %rax
	movq	%rbx, %rdi
	movl	$48, %ecx
	orq	%r11, %rax
	salq	$28, %rdi
	movl	$56, %esi
	andl	$268435455, %eax
	orq	%rax, %rdi
	movq	%rax, %r11
	movl	$15, %eax
	subl	%ebp, %eax
	testl	%r13d, %r13d
	cmove	%ebp, %eax
	incq	%rbp
	cltq
	leaq	(%r12,%rax,8), %r14
	call	_Z11permutationxiPii
	movq	%rax, (%r14)
	cmpq	$16, %rbp
	jne	.L10
	popq	%rdx
	.cfi_def_cfa_offset 56
	movq	%r12, %rax
	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
.LFE2028:
	.size	_Z17generate_sub_keysxi, .-_Z17generate_sub_keysxi
	.globl	_Z12substitutionx
	.type	_Z12substitutionx, @function
_Z12substitutionx:
.LFB2029:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2029:
	.size	_Z12substitutionx, .-_Z12substitutionx
	.globl	_Z1Fjx
	.type	_Z1Fjx, @function
_Z1Fjx:
.LFB2030:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2030:
	.size	_Z1Fjx, .-_Z1Fjx
	.globl	_Z3DESiPxS_
	.type	_Z3DESiPxS_, @function
_Z3DESiPxS_:
.LFB2031:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2031:
	.size	_Z3DESiPxS_, .-_Z3DESiPxS_
	.globl	_Z34__device_stub__Z10kernel_DESjjPxS_jjPxS_
	.type	_Z34__device_stub__Z10kernel_DESjjPxS_jjPxS_, @function
_Z34__device_stub__Z10kernel_DESjjPxS_jjPxS_:
.LFB2061:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movl	%edi, 28(%rsp)
	leaq	56(%rsp), %rdi
	movl	%esi, 24(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 16(%rsp)
	leaq	40(%rsp), %rdx
	movq	%rcx, 8(%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	8(%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	.L19
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z10kernel_DESjjPxS_(%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
.L19:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L21
	call	__stack_chk_fail@PLT
.L21:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2061:
	.size	_Z34__device_stub__Z10kernel_DESjjPxS_jjPxS_, .-_Z34__device_stub__Z10kernel_DESjjPxS_jjPxS_
	.globl	_Z10kernel_DESjjPxS_
	.type	_Z10kernel_DESjjPxS_, @function
_Z10kernel_DESjjPxS_:
.LFB2062:
	.cfi_startproc
	endbr64
	jmp	_Z34__device_stub__Z10kernel_DESjjPxS_jjPxS_
	.cfi_endproc
.LFE2062:
	.size	_Z10kernel_DESjjPxS_, .-_Z10kernel_DESjjPxS_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"Maximum value of the number of threads is 512. You entered : %d\n"
.LC1:
	.string	"%u bytes per threads.\n"
	.text
	.globl	_Z10runDESCudajPxS_ii
	.type	_Z10runDESCudajPxS_ii, @function
_Z10runDESCudajPxS_ii:
.LFB2032:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rdx, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movl	%edi, %ebx
	subq	$72, %rsp
	.cfi_def_cfa_offset 128
	movq	%rsi, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 56(%rsp)
	movl	%ecx, %eax
	andl	%r8d, %eax
	incl	%eax
	jne	.L25
	leal	511(%rdi), %r9d
	movl	$512, %r13d
	shrl	$9, %r9d
	movl	%r9d, %ebp
	jmp	.L26
.L25:
	movl	%ecx, %ebp
	movl	%r8d, %r13d
	cmpl	$512, %r8d
	jbe	.L26
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L31
	addq	$72, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 56
	movl	%r8d, %edx
	leaq	.LC0(%rip), %rsi
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 48
	movl	$2, %edi
	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	__printf_chk@PLT
.L26:
	.cfi_restore_state
	leaq	_ZL2IP(%rip), %r15
	xorl	%ecx, %ecx
	movl	$1, %r8d
	movl	$256, %edx
	leaq	host_IP(%rip), %rsi
	movq	%r15, %rdi
	leaq	_ZL2FP(%rip), %r14
	call	cudaMemcpyToSymbol@PLT
	xorl	%ecx, %ecx
	movl	$1, %r8d
	movq	%r14, %rdi
	movl	$256, %edx
	leaq	host_FP(%rip), %rsi
	call	cudaMemcpyToSymbol@PLT
	xorl	%ecx, %ecx
	movl	$1, %r8d
	movl	$192, %edx
	leaq	host_E(%rip), %rsi
	leaq	_ZL1E(%rip), %rdi
	call	cudaMemcpyToSymbol@PLT
	xorl	%ecx, %ecx
	movl	$1, %r8d
	movl	$128, %edx
	leaq	host_P(%rip), %rsi
	leaq	_ZL1P(%rip), %rdi
	call	cudaMemcpyToSymbol@PLT
	movl	$1, %r8d
	xorl	%ecx, %ecx
	movl	$2048, %edx
	leaq	host_SBox(%rip), %rsi
	leaq	_ZL4SBox(%rip), %rdi
	call	cudaMemcpyToSymbol@PLT
	movl	%ebx, %eax
	leaq	16(%rsp), %rdi
	salq	$3, %rax
	movq	%rax, %rsi
	movq	%rax, (%rsp)
	call	cudaMalloc@PLT
	movq	(%rsp), %rdx
	movq	8(%rsp), %rsi
	movl	$1, %ecx
	movq	16(%rsp), %rdi
	call	cudaMemcpy@PLT
	leaq	24(%rsp), %rdi
	movl	$128, %esi
	call	cudaMalloc@PLT
	movq	24(%rsp), %rdi
	movq	%r12, %rsi
	movl	$1, %ecx
	movl	$128, %edx
	call	cudaMemcpy@PLT
	movl	%ebp, %ecx
	xorl	%edx, %edx
	movl	%ebx, %eax
	imull	%r13d, %ecx
	leaq	.LC1(%rip), %rsi
	movl	$2, %edi
	divl	%ecx
	leal	1(%rax), %r12d
	xorl	%eax, %eax
	movl	%r12d, %edx
	sall	$6, %edx
	call	__printf_chk@PLT
	movl	%r13d, 44(%rsp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movabsq	$4294967297, %rax
	movl	%ebp, 32(%rsp)
	movq	%rax, 48(%rsp)
	movl	52(%rsp), %ecx
	movq	%rax, 36(%rsp)
	movq	44(%rsp), %rdx
	movq	32(%rsp), %rdi
	movl	40(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L28
	movq	24(%rsp), %rcx
	movq	16(%rsp), %rdx
	movl	%ebx, %esi
	movl	%r12d, %edi
	call	_Z34__device_stub__Z10kernel_DESjjPxS_jjPxS_
.L28:
	movq	(%rsp), %rdx
	movq	16(%rsp), %rsi
	movl	$2, %ecx
	movq	8(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	%r15, %rdi
	call	cudaFree@PLT
	movq	%r14, %rdi
	call	cudaFree@PLT
	leaq	_ZL1E(%rip), %rdi
	call	cudaFree@PLT
	leaq	_ZL1P(%rip), %rdi
	call	cudaFree@PLT
	leaq	_ZL4SBox(%rip), %rdi
	call	cudaFree@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	je	.L29
.L31:
	call	__stack_chk_fail@PLT
.L29:
	addq	$72, %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
.LFE2032:
	.size	_Z10runDESCudajPxS_ii, .-_Z10runDESCudajPxS_ii
	.section	.rodata.str1.1
.LC2:
	.string	"rb"
.LC3:
	.string	"Can't open the in file :%s\n"
.LC5:
	.string	"CUDA time: %f\n"
.LC6:
	.string	"wb"
.LC7:
	.string	"Can't open the out file :%s\n"
	.text
	.globl	_Z13des_with_fileiPcS_S_ii
	.type	_Z13des_with_fileiPcS_S_ii, @function
_Z13des_with_fileiPcS_S_ii:
.LFB2033:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rcx, %r15
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movl	%edi, %r14d
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movl	%r9d, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rsi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$40, %rsp
	.cfi_def_cfa_offset 96
	movl	n_blocks(%rip), %eax
	movq	%rdx, 16(%rsp)
	sall	$3, %eax
	movl	%r8d, 28(%rsp)
	cltq
	movq	%rax, %rdi
	movq	%rax, 8(%rsp)
	call	malloc@PLT
	leaq	.LC2(%rip), %rsi
	movq	%rbp, %rdi
	movq	%rax, %rbx
	call	fopen@PLT
	movq	%rbp, %rdx
	leaq	.LC3(%rip), %rsi
	testq	%rax, %rax
	je	.L46
	movq	8(%rsp), %rdx
	movq	%rax, %r12
	movq	%rax, %r8
	movl	$1, %ecx
	movq	%rbx, %rdi
	movq	%rdx, %rsi
	call	__fread_chk@PLT
	movq	%r12, %rdi
	call	fclose@PLT
	movl	n_blocks(%rip), %r12d
	salq	$3, %r12
	movq	%r12, %rdi
	call	malloc@PLT
	movq	%rax, %rbp
	xorl	%eax, %eax
.L34:
	cmpq	%rax, %r12
	je	.L47
	xorl	%ecx, %ecx
	xorl	%edx, %edx
	leaq	(%rbx,%rax), %rdi
.L35:
	movzbl	(%rdi,%rcx), %esi
	salq	$8, %rdx
	incq	%rcx
	addq	%rsi, %rdx
	cmpq	$8, %rcx
	jne	.L35
	movq	%rdx, 0(%rbp,%rax)
	addq	$8, %rax
	jmp	.L34
.L47:
	xorl	%eax, %eax
	xorl	%edi, %edi
.L36:
	movzbl	(%r15,%rax), %edx
	salq	$8, %rdi
	incq	%rax
	addq	%rdx, %rdi
	cmpq	$8, %rax
	jne	.L36
	movl	%r14d, %esi
	call	_Z17generate_sub_keysxi
	movq	%rax, %r12
	call	clock@PLT
	movl	28(%rsp), %ecx
	movq	%r12, %rdx
	movl	%r13d, %r8d
	movl	n_blocks(%rip), %edi
	movq	%rbp, %rsi
	movq	%rax, %r14
	call	_Z10runDESCudajPxS_ii
	call	clock@PLT
	leaq	.LC5(%rip), %rsi
	movl	$2, %edi
	subq	%r14, %rax
	cvtsi2sdq	%rax, %xmm0
	movb	$1, %al
	divsd	.LC4(%rip), %xmm0
	call	__printf_chk@PLT
	movq	%r12, %rdi
	call	free@PLT
	movl	n_blocks(%rip), %edx
	xorl	%eax, %eax
	salq	$3, %rdx
.L37:
	cmpq	%rax, %rdx
	je	.L48
	movq	0(%rbp,%rax), %r8
	leaq	(%rbx,%rax), %rsi
	xorl	%ecx, %ecx
.L38:
	movq	%r8, %rdi
	decq	%rsi
	sarq	%cl, %rdi
	addl	$8, %ecx
	movb	%dil, 8(%rsi)
	cmpl	$64, %ecx
	jne	.L38
	addq	$8, %rax
	jmp	.L37
.L48:
	movq	16(%rsp), %rdi
	leaq	.LC6(%rip), %rsi
	call	fopen@PLT
	movq	%rax, %r12
	testq	%rax, %rax
	jne	.L40
	movq	16(%rsp), %rdx
	leaq	.LC7(%rip), %rsi
.L46:
	addq	$40, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 56
	movl	$2, %edi
	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
	jmp	__printf_chk@PLT
.L40:
	.cfi_restore_state
	movq	8(%rsp), %rsi
	movq	%rax, %rcx
	movl	$1, %edx
	movq	%rbx, %rdi
	call	fwrite@PLT
	movq	%r12, %rdi
	call	fclose@PLT
	movq	%rbx, %rdi
	call	free@PLT
	addq	$40, %rsp
	.cfi_def_cfa_offset 56
	movq	%rbp, %rdi
	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	free@PLT
	.cfi_endproc
.LFE2033:
	.size	_Z13des_with_fileiPcS_S_ii, .-_Z13des_with_fileiPcS_S_ii
	.globl	_Z10encryptionPcS_S_ii
	.type	_Z10encryptionPcS_S_ii, @function
_Z10encryptionPcS_S_ii:
.LFB2034:
	.cfi_startproc
	endbr64
	movl	%r8d, %r9d
	movl	%ecx, %r8d
	movq	%rdx, %rcx
	movq	%rsi, %rdx
	movq	%rdi, %rsi
	xorl	%edi, %edi
	jmp	_Z13des_with_fileiPcS_S_ii
	.cfi_endproc
.LFE2034:
	.size	_Z10encryptionPcS_S_ii, .-_Z10encryptionPcS_S_ii
	.globl	_Z10decryptionPcS_S_ii
	.type	_Z10decryptionPcS_S_ii, @function
_Z10decryptionPcS_S_ii:
.LFB2035:
	.cfi_startproc
	endbr64
	movl	%r8d, %r9d
	movl	%ecx, %r8d
	movq	%rdx, %rcx
	movq	%rsi, %rdx
	movq	%rdi, %rsi
	movl	$1, %edi
	jmp	_Z13des_with_fileiPcS_S_ii
	.cfi_endproc
.LFE2035:
	.size	_Z10decryptionPcS_S_ii, .-_Z10decryptionPcS_S_ii
	.section	.rodata.str1.1
.LC8:
	.string	"usage) ./des.out [e|d] <input_file> <output_file> <n_des_block_size> <n_cuda_blocks> <n_cuda_threads>\n"
.LC9:
	.string	"example) ./des.out e in.txt out.txt 1\n"
.LC10:
	.string	"%d"
.LC11:
	.string	"encryption\n"
.LC12:
	.string	"decryption\n"
.LC13:
	.string	"mode must be 'e' or 'd'\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2036:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	subq	$24, %rsp
	.cfi_def_cfa_offset 48
	movq	%fs:40, %rax
	movq	%rax, 8(%rsp)
	xorl	%eax, %eax
	cmpl	$5, %edi
	jg	.L52
	leaq	.LC8(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	leaq	.LC9(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	orl	$-1, %eax
	jmp	.L51
.L52:
	movq	40(%rsi), %rdi
	leaq	.LC10(%rip), %rbp
	movq	%rsi, %rbx
	xorl	%eax, %eax
	leaq	n_blocks(%rip), %rdx
	movq	%rbp, %rsi
	call	__isoc23_sscanf@PLT
	movq	48(%rbx), %rdi
	movq	%rsp, %rdx
	movq	%rbp, %rsi
	xorl	%eax, %eax
	call	__isoc23_sscanf@PLT
	movq	56(%rbx), %rdi
	leaq	4(%rsp), %rdx
	xorl	%eax, %eax
	movq	%rbp, %rsi
	call	__isoc23_sscanf@PLT
	movq	8(%rbx), %rax
	movb	(%rax), %al
	cmpb	$100, %al
	je	.L54
	cmpb	$101, %al
	jne	.L55
	leaq	.LC11(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	32(%rbx), %rdx
	movq	24(%rbx), %rsi
	movq	16(%rbx), %rdi
	movl	4(%rsp), %r8d
	movl	(%rsp), %ecx
	call	_Z10encryptionPcS_S_ii
.L56:
	xorl	%eax, %eax
	jmp	.L51
.L54:
	leaq	.LC12(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	32(%rbx), %rdx
	movq	24(%rbx), %rsi
	movq	16(%rbx), %rdi
	movl	4(%rsp), %r8d
	movl	(%rsp), %ecx
	call	_Z10decryptionPcS_S_ii
	jmp	.L56
.L55:
	leaq	.LC13(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	jmp	.L56
.L51:
	movq	8(%rsp), %rdx
	subq	%fs:40, %rdx
	je	.L57
	call	__stack_chk_fail@PLT
.L57:
	addq	$24, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2036:
	.size	main, .-main
	.section	.rodata.str1.1
.LC14:
	.string	"_Z10kernel_DESjjPxS_"
.LC15:
	.string	"IP"
.LC16:
	.string	"FP"
.LC17:
	.string	"E"
.LC18:
	.string	"P"
.LC19:
	.string	"SBox"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2064:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC14(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z10kernel_DESjjPxS_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC15(%rip), %rdx
	movl	$256, %r9d
	leaq	_ZL2IP(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	call	__cudaRegisterVar@PLT
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC16(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$256, %r9d
	leaq	_ZL2FP(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rsi
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC17(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$192, %r9d
	leaq	_ZL1E(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rdi
	.cfi_def_cfa_offset 24
	popq	%r8
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC18(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$128, %r9d
	leaq	_ZL1P(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%r9
	.cfi_def_cfa_offset 24
	popq	%r10
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC19(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movl	$2048, %r9d
	movq	%rdx, %rcx
	leaq	_ZL4SBox(%rip), %rsi
	call	__cudaRegisterVar@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%r11
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2064:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.globl	n_blocks
	.bss
	.align 4
	.type	n_blocks, @object
	.size	n_blocks, 4
n_blocks:
	.zero	4
	.globl	host_SBox
	.data
	.align 32
	.type	host_SBox, @object
	.size	host_SBox, 2048
host_SBox:
	.long	14
	.long	4
	.long	13
	.long	1
	.long	2
	.long	15
	.long	11
	.long	8
	.long	3
	.long	10
	.long	6
	.long	12
	.long	5
	.long	9
	.long	0
	.long	7
	.long	0
	.long	15
	.long	7
	.long	4
	.long	14
	.long	2
	.long	13
	.long	1
	.long	10
	.long	6
	.long	12
	.long	11
	.long	9
	.long	5
	.long	3
	.long	8
	.long	4
	.long	1
	.long	14
	.long	8
	.long	13
	.long	6
	.long	2
	.long	11
	.long	15
	.long	12
	.long	9
	.long	7
	.long	3
	.long	10
	.long	5
	.long	0
	.long	15
	.long	12
	.long	8
	.long	2
	.long	4
	.long	9
	.long	1
	.long	7
	.long	5
	.long	11
	.long	3
	.long	14
	.long	10
	.long	0
	.long	6
	.long	13
	.long	15
	.long	1
	.long	8
	.long	14
	.long	6
	.long	11
	.long	3
	.long	4
	.long	9
	.long	7
	.long	2
	.long	13
	.long	12
	.long	0
	.long	5
	.long	10
	.long	3
	.long	13
	.long	4
	.long	7
	.long	15
	.long	2
	.long	8
	.long	14
	.long	12
	.long	0
	.long	1
	.long	10
	.long	6
	.long	9
	.long	11
	.long	5
	.long	0
	.long	14
	.long	7
	.long	11
	.long	10
	.long	4
	.long	13
	.long	1
	.long	5
	.long	8
	.long	12
	.long	6
	.long	9
	.long	3
	.long	2
	.long	15
	.long	13
	.long	8
	.long	10
	.long	1
	.long	3
	.long	15
	.long	4
	.long	2
	.long	11
	.long	6
	.long	7
	.long	12
	.long	0
	.long	5
	.long	14
	.long	9
	.long	10
	.long	0
	.long	9
	.long	14
	.long	6
	.long	3
	.long	15
	.long	5
	.long	1
	.long	13
	.long	12
	.long	7
	.long	11
	.long	4
	.long	2
	.long	8
	.long	13
	.long	7
	.long	0
	.long	9
	.long	3
	.long	4
	.long	6
	.long	10
	.long	2
	.long	8
	.long	5
	.long	14
	.long	12
	.long	11
	.long	15
	.long	1
	.long	13
	.long	6
	.long	4
	.long	9
	.long	8
	.long	15
	.long	3
	.long	0
	.long	11
	.long	1
	.long	2
	.long	12
	.long	5
	.long	10
	.long	14
	.long	7
	.long	1
	.long	10
	.long	13
	.long	0
	.long	6
	.long	9
	.long	8
	.long	7
	.long	4
	.long	15
	.long	14
	.long	3
	.long	11
	.long	5
	.long	2
	.long	12
	.long	7
	.long	13
	.long	14
	.long	3
	.long	0
	.long	6
	.long	9
	.long	10
	.long	1
	.long	2
	.long	8
	.long	5
	.long	11
	.long	12
	.long	4
	.long	15
	.long	13
	.long	8
	.long	11
	.long	5
	.long	6
	.long	15
	.long	0
	.long	3
	.long	4
	.long	7
	.long	2
	.long	12
	.long	1
	.long	10
	.long	14
	.long	9
	.long	10
	.long	6
	.long	9
	.long	0
	.long	12
	.long	11
	.long	7
	.long	13
	.long	15
	.long	1
	.long	3
	.long	14
	.long	5
	.long	2
	.long	8
	.long	4
	.long	3
	.long	15
	.long	0
	.long	6
	.long	10
	.long	1
	.long	13
	.long	8
	.long	9
	.long	4
	.long	5
	.long	11
	.long	12
	.long	7
	.long	2
	.long	14
	.long	2
	.long	12
	.long	4
	.long	1
	.long	7
	.long	10
	.long	11
	.long	6
	.long	8
	.long	5
	.long	3
	.long	15
	.long	13
	.long	0
	.long	14
	.long	9
	.long	14
	.long	11
	.long	2
	.long	12
	.long	4
	.long	7
	.long	13
	.long	1
	.long	5
	.long	0
	.long	15
	.long	10
	.long	3
	.long	9
	.long	8
	.long	6
	.long	4
	.long	2
	.long	1
	.long	11
	.long	10
	.long	13
	.long	7
	.long	8
	.long	15
	.long	9
	.long	12
	.long	5
	.long	6
	.long	3
	.long	0
	.long	14
	.long	11
	.long	8
	.long	12
	.long	7
	.long	1
	.long	14
	.long	2
	.long	13
	.long	6
	.long	15
	.long	0
	.long	9
	.long	10
	.long	4
	.long	5
	.long	3
	.long	12
	.long	1
	.long	10
	.long	15
	.long	9
	.long	2
	.long	6
	.long	8
	.long	0
	.long	13
	.long	3
	.long	4
	.long	14
	.long	7
	.long	5
	.long	11
	.long	10
	.long	15
	.long	4
	.long	2
	.long	7
	.long	12
	.long	9
	.long	5
	.long	6
	.long	1
	.long	13
	.long	14
	.long	0
	.long	11
	.long	3
	.long	8
	.long	9
	.long	14
	.long	15
	.long	5
	.long	2
	.long	8
	.long	12
	.long	3
	.long	7
	.long	0
	.long	4
	.long	10
	.long	1
	.long	13
	.long	11
	.long	6
	.long	4
	.long	3
	.long	2
	.long	12
	.long	9
	.long	5
	.long	15
	.long	10
	.long	11
	.long	14
	.long	1
	.long	7
	.long	6
	.long	0
	.long	8
	.long	13
	.long	4
	.long	11
	.long	2
	.long	14
	.long	15
	.long	0
	.long	8
	.long	13
	.long	3
	.long	12
	.long	9
	.long	7
	.long	5
	.long	10
	.long	6
	.long	1
	.long	13
	.long	0
	.long	11
	.long	7
	.long	4
	.long	9
	.long	1
	.long	10
	.long	14
	.long	3
	.long	5
	.long	12
	.long	2
	.long	15
	.long	8
	.long	6
	.long	1
	.long	4
	.long	11
	.long	13
	.long	12
	.long	3
	.long	7
	.long	14
	.long	10
	.long	15
	.long	6
	.long	8
	.long	0
	.long	5
	.long	9
	.long	2
	.long	6
	.long	11
	.long	13
	.long	8
	.long	1
	.long	4
	.long	10
	.long	7
	.long	9
	.long	5
	.long	0
	.long	15
	.long	14
	.long	2
	.long	3
	.long	12
	.long	13
	.long	2
	.long	8
	.long	4
	.long	6
	.long	15
	.long	11
	.long	1
	.long	10
	.long	9
	.long	3
	.long	14
	.long	5
	.long	0
	.long	12
	.long	7
	.long	1
	.long	15
	.long	13
	.long	8
	.long	10
	.long	3
	.long	7
	.long	4
	.long	12
	.long	5
	.long	6
	.long	11
	.long	0
	.long	14
	.long	9
	.long	2
	.long	7
	.long	11
	.long	4
	.long	1
	.long	9
	.long	12
	.long	14
	.long	2
	.long	0
	.long	6
	.long	10
	.long	13
	.long	15
	.long	3
	.long	5
	.long	8
	.long	2
	.long	1
	.long	14
	.long	7
	.long	4
	.long	10
	.long	8
	.long	13
	.long	15
	.long	12
	.long	9
	.long	0
	.long	3
	.long	5
	.long	6
	.long	11
	.globl	Rotations
	.align 32
	.type	Rotations, @object
	.size	Rotations, 64
Rotations:
	.long	1
	.long	1
	.long	2
	.long	2
	.long	2
	.long	2
	.long	2
	.long	2
	.long	1
	.long	2
	.long	2
	.long	2
	.long	2
	.long	2
	.long	2
	.long	1
	.globl	PC2
	.align 32
	.type	PC2, @object
	.size	PC2, 192
PC2:
	.long	13
	.long	16
	.long	10
	.long	23
	.long	0
	.long	4
	.long	2
	.long	27
	.long	14
	.long	5
	.long	20
	.long	9
	.long	22
	.long	18
	.long	11
	.long	3
	.long	25
	.long	7
	.long	15
	.long	6
	.long	26
	.long	19
	.long	12
	.long	1
	.long	40
	.long	51
	.long	30
	.long	36
	.long	46
	.long	54
	.long	29
	.long	39
	.long	50
	.long	44
	.long	32
	.long	47
	.long	43
	.long	48
	.long	38
	.long	55
	.long	33
	.long	52
	.long	45
	.long	41
	.long	49
	.long	35
	.long	28
	.long	31
	.globl	PC1_RIGHT
	.align 32
	.type	PC1_RIGHT, @object
	.size	PC1_RIGHT, 112
PC1_RIGHT:
	.long	62
	.long	54
	.long	46
	.long	38
	.long	30
	.long	22
	.long	14
	.long	6
	.long	61
	.long	53
	.long	45
	.long	37
	.long	29
	.long	21
	.long	13
	.long	5
	.long	60
	.long	52
	.long	44
	.long	36
	.long	28
	.long	20
	.long	12
	.long	4
	.long	27
	.long	19
	.long	11
	.long	3
	.globl	PC1_LEFT
	.align 32
	.type	PC1_LEFT, @object
	.size	PC1_LEFT, 112
PC1_LEFT:
	.long	56
	.long	48
	.long	40
	.long	32
	.long	24
	.long	16
	.long	8
	.long	0
	.long	57
	.long	49
	.long	41
	.long	33
	.long	25
	.long	17
	.long	9
	.long	1
	.long	58
	.long	50
	.long	42
	.long	34
	.long	26
	.long	18
	.long	10
	.long	2
	.long	59
	.long	51
	.long	43
	.long	35
	.globl	host_P
	.align 32
	.type	host_P, @object
	.size	host_P, 128
host_P:
	.long	15
	.long	6
	.long	19
	.long	20
	.long	28
	.long	11
	.long	27
	.long	16
	.long	0
	.long	14
	.long	22
	.long	25
	.long	4
	.long	17
	.long	30
	.long	9
	.long	1
	.long	7
	.long	23
	.long	13
	.long	31
	.long	26
	.long	2
	.long	8
	.long	18
	.long	12
	.long	29
	.long	5
	.long	21
	.long	10
	.long	3
	.long	24
	.globl	host_E
	.align 32
	.type	host_E, @object
	.size	host_E, 192
host_E:
	.long	31
	.long	0
	.long	1
	.long	2
	.long	3
	.long	4
	.long	3
	.long	4
	.long	5
	.long	6
	.long	7
	.long	8
	.long	7
	.long	8
	.long	9
	.long	10
	.long	11
	.long	12
	.long	11
	.long	12
	.long	13
	.long	14
	.long	15
	.long	16
	.long	15
	.long	16
	.long	17
	.long	18
	.long	19
	.long	20
	.long	19
	.long	20
	.long	21
	.long	22
	.long	23
	.long	24
	.long	23
	.long	24
	.long	25
	.long	26
	.long	27
	.long	28
	.long	27
	.long	28
	.long	29
	.long	30
	.long	31
	.long	0
	.globl	host_FP
	.align 32
	.type	host_FP, @object
	.size	host_FP, 256
host_FP:
	.long	39
	.long	7
	.long	47
	.long	15
	.long	55
	.long	23
	.long	63
	.long	31
	.long	38
	.long	6
	.long	46
	.long	14
	.long	54
	.long	22
	.long	62
	.long	30
	.long	37
	.long	5
	.long	45
	.long	13
	.long	53
	.long	21
	.long	61
	.long	29
	.long	36
	.long	4
	.long	44
	.long	12
	.long	52
	.long	20
	.long	60
	.long	28
	.long	35
	.long	3
	.long	43
	.long	11
	.long	51
	.long	19
	.long	59
	.long	27
	.long	34
	.long	2
	.long	42
	.long	10
	.long	50
	.long	18
	.long	58
	.long	26
	.long	33
	.long	1
	.long	41
	.long	9
	.long	49
	.long	17
	.long	57
	.long	25
	.long	32
	.long	0
	.long	40
	.long	8
	.long	48
	.long	16
	.long	56
	.long	24
	.globl	host_IP
	.align 32
	.type	host_IP, @object
	.size	host_IP, 256
host_IP:
	.long	57
	.long	49
	.long	41
	.long	33
	.long	25
	.long	17
	.long	9
	.long	1
	.long	59
	.long	51
	.long	43
	.long	35
	.long	27
	.long	19
	.long	11
	.long	3
	.long	61
	.long	53
	.long	45
	.long	37
	.long	29
	.long	21
	.long	13
	.long	5
	.long	63
	.long	55
	.long	47
	.long	39
	.long	31
	.long	23
	.long	15
	.long	7
	.long	56
	.long	48
	.long	40
	.long	32
	.long	24
	.long	16
	.long	8
	.long	0
	.long	58
	.long	50
	.long	42
	.long	34
	.long	26
	.long	18
	.long	10
	.long	2
	.long	60
	.long	52
	.long	44
	.long	36
	.long	28
	.long	20
	.long	12
	.long	4
	.long	62
	.long	54
	.long	46
	.long	38
	.long	30
	.long	22
	.long	14
	.long	6
	.local	_ZL4SBox
	.comm	_ZL4SBox,2048,32
	.local	_ZL1P
	.comm	_ZL1P,128,32
	.local	_ZL1E
	.comm	_ZL1E,192,32
	.local	_ZL2FP
	.comm	_ZL2FP,256,32
	.local	_ZL2IP
	.comm	_ZL2IP,256,32
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC4:
	.long	0
	.long	1093567616
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:


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

__device__ int IP[64];
__device__ int FP[64];
__device__ int E[48];
__device__ int P[32];
__device__ int SBox[8][64];


// Initial Permutation
int host_IP[64] = {
    57, 49, 41, 33, 25, 17,  9,  1,
    59, 51, 43, 35, 27, 19, 11,  3,
    61, 53, 45, 37, 29, 21, 13,  5,
    63, 55, 47, 39, 31, 23, 15,  7,
    56, 48, 40, 32, 24, 16,  8,  0,
    58, 50, 42, 34, 26, 18, 10,  2,
    60, 52, 44, 36, 28, 20, 12,  4,
    62, 54, 46, 38, 30, 22, 14,  6
};

// Final Permutation
int host_FP[64] = {
    39,  7, 47, 15, 55, 23, 63, 31,
    38,  6, 46, 14, 54, 22, 62, 30,
    37,  5, 45, 13, 53, 21, 61, 29,
    36,  4, 44, 12, 52, 20, 60, 28,
    35,  3, 43, 11, 51, 19, 59, 27,
    34,  2, 42, 10, 50, 18, 58, 26,
    33,  1, 41,  9, 49, 17, 57, 25,
    32,  0, 40,  8, 48, 16, 56, 24
};

// Expansion Function: from 32 bit to 48 bit
int host_E[48] = {
    31,  0,  1,  2,  3,  4,
     3,  4,  5,  6,  7,  8,
     7,  8,  9, 10, 11, 12,
    11, 12, 13, 14, 15, 16,
    15, 16, 17, 18, 19, 20,
    19, 20, 21, 22, 23, 24,
    23, 24, 25, 26, 27, 28,
    27, 28, 29, 30, 31,  0
};

// Permutation
int host_P[32] = {
    15,  6, 19, 20, 28, 11, 27, 16,
     0, 14, 22, 25,  4, 17, 30,  9,
     1,  7, 23, 13, 31, 26,  2,  8,
    18, 12, 29,  5, 21, 10,  3, 24
};

int PC1_LEFT[28] = {
    56, 48, 40, 32, 24, 16,  8,
     0, 57, 49, 41, 33, 25, 17,
     9,  1, 58, 50, 42, 34, 26,
    18, 10,  2, 59, 51, 43, 35,
};

int PC1_RIGHT[28] = {
    62, 54, 46, 38, 30, 22, 14,
     6, 61, 53, 45, 37, 29, 21,
    13,  5, 60, 52, 44, 36, 28,
    20, 12,  4, 27, 19, 11,  3
};

int PC2[48] = {
    13, 16, 10, 23,  0,  4,
     2, 27, 14,  5, 20,  9,
    22, 18, 11,  3, 25,  7,
    15,  6, 26, 19, 12,  1,
    40, 51, 30, 36, 46, 54,
    29, 39, 50, 44, 32, 47,
    43, 48, 38, 55, 33, 52,
    45, 41, 49, 35, 28, 31
};

int Rotations[16] = {1, 1, 2, 2, 2, 2, 2, 2, 1, 2, 2, 2, 2, 2, 2, 1};

// Substitution Boxes
int host_SBox[8][64] = {
    // S1
    {
        14, 4, 13, 1, 2, 15, 11, 8, 3, 10, 6, 12, 5, 9, 0, 7,
        0, 15, 7, 4, 14, 2, 13, 1, 10, 6, 12, 11, 9, 5, 3, 8,
        4, 1, 14, 8, 13, 6, 2, 11, 15, 12, 9, 7, 3, 10, 5, 0,
        15, 12, 8, 2, 4, 9, 1, 7, 5, 11, 3, 14, 10, 0, 6, 13
    },

    // S2
    {
        15, 1, 8, 14, 6, 11, 3, 4, 9, 7, 2, 13, 12, 0, 5, 10,
        3, 13, 4, 7, 15, 2, 8, 14, 12, 0, 1, 10, 6, 9, 11, 5,
        0, 14, 7, 11, 10, 4, 13, 1, 5, 8, 12, 6, 9, 3, 2, 15,
        13, 8, 10, 1, 3, 15, 4, 2, 11, 6, 7, 12, 0, 5, 14, 9
    },

    // S3
    {
        10, 0, 9, 14, 6, 3, 15, 5, 1, 13, 12, 7, 11, 4, 2, 8,
        13, 7, 0, 9, 3, 4, 6, 10, 2, 8, 5, 14, 12, 11, 15, 1,
        13, 6, 4, 9, 8, 15, 3, 0, 11, 1, 2, 12, 5, 10, 14, 7,
        1, 10, 13, 0, 6, 9, 8, 7, 4, 15, 14, 3, 11, 5, 2, 12
    },

    // S4
    {
        7, 13, 14, 3, 0, 6, 9, 10, 1, 2, 8, 5, 11, 12, 4, 15,
        13, 8, 11, 5, 6, 15, 0, 3, 4, 7, 2, 12, 1, 10, 14, 9,
        10, 6, 9, 0, 12, 11, 7, 13, 15, 1, 3, 14, 5, 2, 8, 4,
        3, 15, 0, 6, 10, 1, 13, 8, 9, 4, 5, 11, 12, 7, 2, 14
    },

    // S5
    {
        2, 12, 4, 1, 7, 10, 11, 6, 8, 5, 3, 15, 13, 0, 14, 9,
        14, 11, 2, 12, 4, 7, 13, 1, 5, 0, 15, 10, 3, 9, 8, 6,
        4, 2, 1, 11, 10, 13, 7, 8, 15, 9, 12, 5, 6, 3, 0, 14,
        11, 8, 12, 7, 1, 14, 2, 13, 6, 15, 0, 9, 10, 4, 5, 3
    },

    // S6
    {
        12, 1, 10, 15, 9, 2, 6, 8, 0, 13, 3, 4, 14, 7, 5, 11,
        10, 15, 4, 2, 7, 12, 9, 5, 6, 1, 13, 14, 0, 11, 3, 8,
        9, 14, 15, 5, 2, 8, 12, 3, 7, 0, 4, 10, 1, 13, 11, 6,
        4, 3, 2, 12, 9, 5, 15, 10, 11, 14, 1, 7, 6, 0, 8, 13
    },

    // S7
    {
        4, 11, 2, 14, 15, 0, 8, 13, 3, 12, 9, 7, 5, 10, 6, 1,
        13, 0, 11, 7, 4, 9, 1, 10, 14, 3, 5, 12, 2, 15, 8, 6,
        1, 4, 11, 13, 12, 3, 7, 14, 10, 15, 6, 8, 0, 5, 9, 2,
        6, 11, 13, 8, 1, 4, 10, 7, 9, 5, 0, 15, 14, 2, 3, 12
    },

    // S8
    {
        13, 2, 8, 4, 6, 15, 11, 1, 10, 9, 3, 14, 5, 0, 12, 7,
        1, 15, 13, 8, 10, 3, 7, 4, 12, 5, 6, 11, 0, 14, 9, 2,
        7, 11, 4, 1, 9, 12, 14, 2, 0, 6, 10, 13, 15, 3, 5, 8,
        2, 1, 14, 7, 4, 10, 8, 13, 15, 12, 9, 0, 3, 5, 6, 11
    }
};

__device__
__host__
long long int permutation(long long int data, int data_size, int *table, int table_size) {
    long long int result = 0;
    int i = 0;
    for (; i < table_size; i++) {
        result = (result << 1) + ((data >> (data_size - 1 - table[i])) & 0x1);
    }
    return result;
}

long long int *generate_sub_keys(long long int key, int decrypt) {
    int n_keys = 16;
    long long int *sub_keys = (long long int *) malloc(sizeof(long long int) * n_keys);
    int half_key_length = 28;
    long long int left = permutation(key, 64, PC1_LEFT, half_key_length);
    long long int right = permutation(key, 64, PC1_RIGHT, half_key_length);
    int i = 0;
    for (; i < n_keys; i++) {
        int rotation = Rotations[i];
        left = (((left << rotation) | (left >> (half_key_length - rotation))) & 0xFFFFFFF);
        right = (((right << rotation) | (right >> (half_key_length - rotation))) & 0xFFFFFFF);
        long long int new_key = (left << half_key_length) | right;
        int sub_key_index = (decrypt ? 15 - i : i);
        sub_keys[sub_key_index] = permutation(new_key, half_key_length * 2, PC2, 48);
    }
    return sub_keys;
}

__device__
long long int substitution(long long int data) {
    // data: 48 bit
    long long int result = 0;
    int i = 0;
    for (; i < 8; i++) {
        unsigned int box = data >> (6 * (7 - i)) & 0x3F;
        int outer = ((box & 0x20) >> 4) | (box & 0x1);
        int inner = (box & 0x1E) >> 1;
        result = (result << 4) + SBox[i][(outer << 4) + inner];
    }
    return result;
}

__device__
long int F(unsigned int c, long long int key) {
    long long int lc = c;
    long long int new_c = permutation(lc, 32, E, 48);
    long long int mixed_data = new_c ^ key;
    long long int s_box_result = substitution(mixed_data);
    return permutation(s_box_result, 32, P, 32);
}

__device__
void DES(int index, long long int *MD, long long int *keys) {
    long long int data = permutation(MD[index], 64, IP, 64);
    unsigned int left = data >> 32;
    unsigned int right = (int) data;
    int i = 0;
    for (; i < 16; i++) {
        unsigned int buf = left ^ F(right, keys[i]);
        left = right;
        right = buf;
    }
    data = right;
    data = (data << 32) + left;
    MD[index] = permutation(data, 64, FP, 64);
}

__global__
void kernel_DES(unsigned int quota, unsigned int n_blocks, long long int *MD, long long int *keys) {
    int start_index = (blockIdx.x * blockDim.x + threadIdx.x) * quota;
    int end_index = start_index + quota;
    int i;
    for (i = start_index; i < end_index; i++) {
        if (i >= n_blocks) {
        return;
        }
        DES(i, MD, keys);
    }
}

void runDESCuda(unsigned int n_blocks, long long int *host_MD, long long int *host_sub_keys, int n_cuda_blocks, int n_cuda_threads) {
    unsigned int max_n_threads = 512;
    if (n_cuda_blocks == -1 && n_cuda_threads == -1) {
        n_cuda_threads = max_n_threads;
        n_cuda_blocks = (n_blocks + max_n_threads - 1) / max_n_threads;
    }
    if (n_cuda_threads > max_n_threads) {
        printf("Maximum value of the number of threads is 512. You entered : %d\n", n_cuda_threads);
        return;
    }
    hipMemcpyToSymbol(HIP_SYMBOL(IP), host_IP, sizeof(host_IP));
    hipMemcpyToSymbol(HIP_SYMBOL(FP), host_FP, sizeof(host_FP));
    hipMemcpyToSymbol(HIP_SYMBOL(E), host_E, sizeof(host_E));
    hipMemcpyToSymbol(HIP_SYMBOL(P), host_P, sizeof(host_P));
    hipMemcpyToSymbol(HIP_SYMBOL(SBox), host_SBox, sizeof(host_SBox));

    long long int *MD, *sub_keys;
    hipMalloc((void **) &MD, sizeof(long long int) * n_blocks);
    hipMemcpy(MD, host_MD, sizeof(long long int) * n_blocks, hipMemcpyHostToDevice);
    hipMalloc((void **) &sub_keys, sizeof(long long int) * 16);
    hipMemcpy(sub_keys, host_sub_keys, sizeof(long long int) * 16, hipMemcpyHostToDevice);
    unsigned int quota = n_blocks / (n_cuda_blocks * n_cuda_threads) + 1;
    printf("%u bytes per threads.\n", quota * 64);
    kernel_DES<<<n_cuda_blocks, n_cuda_threads>>>(quota, n_blocks, MD, sub_keys);
    hipMemcpy(host_MD, MD, sizeof(long long int) * n_blocks, hipMemcpyDeviceToHost);

    hipFree(IP);
    hipFree(FP);
    hipFree(E);
    hipFree(P);
    hipFree(SBox);
    hipFree(MD);
    hipFree(sub_keys);
}

unsigned int n_blocks = 0;
void des_with_file(int decrypt, char *in, char *out, char *key, int n_cuda_blocks, int n_cuda_threads) {
    int buf_size = 8 * n_blocks;
    char *buf = (char *) malloc(sizeof(char) * buf_size);
    FILE *in_fp = fopen(in, "rb");
    if (in_fp == NULL) {
        printf("Can't open the in file :%s\n", in);
        return;
    }
    fread(buf, buf_size, 1, in_fp);
    fclose(in_fp);

    long long int *MD = (long long int *) malloc(sizeof(long long int) * n_blocks);
    int i = 0;
    int j = 0;
    for (i = 0; i < n_blocks; i++) {
        long long int block = 0;
        for (j = 0; j < 8; j++) {
            block = (block << 8) + (buf[(i * 8) + j] & 0xFF);
        }
        MD[i] = block;
    }
    long long int binary_key = 0;
    for (i = 0; i < 8; i++) {
        binary_key = (binary_key << 8) + (key[i] & 0xFF);
    }

    long long int *sub_keys = generate_sub_keys(binary_key, decrypt);
    clock_t begin = clock();
    runDESCuda(n_blocks, MD, sub_keys, n_cuda_blocks, n_cuda_threads);
    clock_t end = clock();
    double time_spent = (double) (end - begin) / CLOCKS_PER_SEC;
    printf("CUDA time: %f\n", time_spent);
    free(sub_keys);

    for (i = 0; i < n_blocks; i++) {
        for (j = 0; j < 8; j++) {
            buf[(i * 8) + (7 - j)] = ((MD[i] >> (j * 8)) & 0xFF);
        }
    }
    FILE *out_fp = fopen(out, "wb");
    if (out_fp == NULL) {
        printf("Can't open the out file :%s\n", out);
        return;
    }
    fwrite(buf, buf_size, 1, out_fp);
    fclose(out_fp);
    free(buf);
    free(MD);
}

void encryption(char *in, char *out, char *key, int n_cuda_blocks, int n_cuda_threads) {
    des_with_file(0, in, out, key, n_cuda_blocks, n_cuda_threads);
}

void decryption(char *in, char *out, char *key, int n_cuda_blocks, int n_cuda_threads) {
    des_with_file(1, in, out, key, n_cuda_blocks, n_cuda_threads);
}

int main(int argc, char** argv) {
    if (argc < 6) {
        printf("usage) ./des.out [e|d] <input_file> <output_file> <n_des_block_size> <n_cuda_blocks> <n_cuda_threads>\n");
        printf("example) ./des.out e in.txt out.txt 1\n");
        return -1;
    }
    int n_cuda_blocks;
    int n_cuda_threads;
    sscanf(argv[5], "%d", &n_blocks);
    sscanf(argv[6], "%d", &n_cuda_blocks);
    sscanf(argv[7], "%d", &n_cuda_threads);
    switch(argv[1][0]) {
        case 'e':
            printf("encryption\n");
            encryption(argv[2], argv[3], argv[4], n_cuda_blocks, n_cuda_threads);
            break;
        case 'd':
            printf("decryption\n");
            decryption(argv[2], argv[3], argv[4], n_cuda_blocks, n_cuda_threads);
            break;
        default:
            printf("mode must be 'e' or 'd'\n");
    }
}

	.text
	.file	"des.hip"
	.globl	_Z11permutationxiPii            # -- Begin function _Z11permutationxiPii
	.type	_Z11permutationxiPii,@function
_Z11permutationxiPii:                   # @_Z11permutationxiPii
	.cfi_startproc
# %bb.0:
	testl	%ecx, %ecx
	jle	.LBB0_1
# %bb.2:                                # %.lr.ph.preheader
	movl	%ecx, %ecx
	xorl	%r8d, %r8d
	xorl	%eax, %eax
.LBB0_3:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	(%rdx,%r8,4), %r9d
	notl	%r9d
	addl	%esi, %r9d
	xorl	%r10d, %r10d
	btq	%r9, %rdi
	setb	%r10b
	leaq	(%r10,%rax,2), %rax
	incq	%r8
	cmpq	%r8, %rcx
	jne	.LBB0_3
# %bb.4:                                # %._crit_edge
	retq
.LBB0_1:
	xorl	%eax, %eax
	retq
.Lfunc_end0:
	.size	_Z11permutationxiPii, .Lfunc_end0-_Z11permutationxiPii
	.cfi_endproc
                                        # -- End function
	.globl	_Z17generate_sub_keysxi         # -- Begin function _Z17generate_sub_keysxi
	.type	_Z17generate_sub_keysxi,@function
_Z17generate_sub_keysxi:                # @_Z17generate_sub_keysxi
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	pushq	%rax
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%esi, %ebx
	movq	%rdi, %r14
	movl	$128, %edi
	callq	malloc
	movq	$-112, %rcx
	xorl	%edi, %edi
.LBB1_1:                                # %.lr.ph.i
                                        # =>This Inner Loop Header: Depth=1
	movb	$63, %dl
	subb	PC1_LEFT+112(%rcx), %dl
	movzbl	%dl, %edx
	xorl	%esi, %esi
	btq	%rdx, %r14
	setb	%sil
	leaq	(%rsi,%rdi,2), %rdi
	addq	$4, %rcx
	jne	.LBB1_1
# %bb.2:                                # %.lr.ph.i31.preheader
	movq	$-112, %rcx
	xorl	%r8d, %r8d
.LBB1_3:                                # %.lr.ph.i31
                                        # =>This Inner Loop Header: Depth=1
	movb	$63, %dl
	subb	PC1_RIGHT+112(%rcx), %dl
	movzbl	%dl, %edx
	xorl	%esi, %esi
	btq	%rdx, %r14
	setb	%sil
	leaq	(%rsi,%r8,2), %r8
	addq	$4, %rcx
	jne	.LBB1_3
# %bb.4:                                # %_Z11permutationxiPii.exit36.preheader
	xorl	%r9d, %r9d
	movl	$268435455, %r10d               # imm = 0xFFFFFFF
.LBB1_5:                                # =>This Loop Header: Depth=1
                                        #     Child Loop BB1_6 Depth 2
	movb	Rotations(,%r9,4), %dl
	movq	%rdi, %r11
	movl	%edx, %ecx
	shlq	%cl, %r11
	movb	$28, %sil
	subb	%dl, %sil
	movl	%esi, %ecx
	sarq	%cl, %rdi
	orl	%r11d, %edi
	andl	%r10d, %edi
	movq	%r8, %r11
	movl	%edx, %ecx
	shlq	%cl, %r11
	movl	%esi, %ecx
	sarq	%cl, %r8
	orl	%r11d, %r8d
	andl	%r10d, %r8d
	movq	%rdi, %rdx
	shlq	$28, %rdx
	orq	%r8, %rdx
	movq	%r9, %rcx
	xorq	$15, %rcx
	movq	$-192, %r11
	xorl	%esi, %esi
.LBB1_6:                                # %.lr.ph.i37
                                        #   Parent Loop BB1_5 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movb	$55, %bpl
	subb	PC2+192(%r11), %bpl
	movzbl	%bpl, %r14d
	xorl	%r15d, %r15d
	btq	%r14, %rdx
	setb	%r15b
	leaq	(%r15,%rsi,2), %rsi
	addq	$4, %r11
	jne	.LBB1_6
# %bb.7:                                # %_Z11permutationxiPii.exit42
                                        #   in Loop: Header=BB1_5 Depth=1
	testl	%ebx, %ebx
	cmoveq	%r9, %rcx
	movl	%ecx, %ecx
	movq	%rsi, (%rax,%rcx,8)
	incq	%r9
	cmpq	$16, %r9
	jne	.LBB1_5
# %bb.8:
	addq	$8, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z17generate_sub_keysxi, .Lfunc_end1-_Z17generate_sub_keysxi
	.cfi_endproc
                                        # -- End function
	.globl	_Z25__device_stub__kernel_DESjjPxS_ # -- Begin function _Z25__device_stub__kernel_DESjjPxS_
	.type	_Z25__device_stub__kernel_DESjjPxS_,@function
_Z25__device_stub__kernel_DESjjPxS_:    # @_Z25__device_stub__kernel_DESjjPxS_
	.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	12(%rsp), %rax
	movl	%edi, (%rax)
	leaq	8(%rsp), %rdi
	movl	%esi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%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	24(%rsp), %r12
	leaq	16(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z10kernel_DESjjPxS_, %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__kernel_DESjjPxS_, .Lfunc_end2-_Z25__device_stub__kernel_DESjjPxS_
	.cfi_endproc
                                        # -- End function
	.globl	_Z10runDESCudajPxS_ii           # -- Begin function _Z10runDESCudajPxS_ii
	.type	_Z10runDESCudajPxS_ii,@function
_Z10runDESCudajPxS_ii:                  # @_Z10runDESCudajPxS_ii
	.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	$40, %rsp
	.cfi_def_cfa_offset 96
	.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, %eax
	andl	%ecx, %eax
	cmpl	$-1, %eax
	movl	$512, %ebp                      # imm = 0x200
	cmovnel	%r8d, %ebp
	cmpl	$513, %ebp                      # imm = 0x201
	jb	.LBB3_1
# %bb.4:
	movl	$.L.str, %edi
	movl	%ebp, %esi
	xorl	%eax, %eax
	addq	$40, %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	printf                          # TAILCALL
.LBB3_1:
	.cfi_def_cfa_offset 96
	movl	%ecx, %r15d
	movl	%edi, %r14d
	leal	511(%r14), %ecx
	shrl	$9, %ecx
	cmpl	$-1, %eax
	cmovel	%ecx, %r15d
	movl	$IP, %edi
	movq	%rsi, %r12
	movl	$host_IP, %esi
	movq	%rdx, 24(%rsp)                  # 8-byte Spill
	movl	$256, %edx                      # imm = 0x100
	xorl	%ecx, %ecx
	movl	$1, %r8d
	callq	hipMemcpyToSymbol
	movl	$FP, %edi
	movl	$host_FP, %esi
	movl	$256, %edx                      # imm = 0x100
	xorl	%ecx, %ecx
	movl	$1, %r8d
	callq	hipMemcpyToSymbol
	movl	$E, %edi
	movl	$host_E, %esi
	movl	$192, %edx
	xorl	%ecx, %ecx
	movl	$1, %r8d
	callq	hipMemcpyToSymbol
	movl	$P, %edi
	movl	$host_P, %esi
	movl	$128, %edx
	xorl	%ecx, %ecx
	movl	$1, %r8d
	callq	hipMemcpyToSymbol
	movl	$SBox, %edi
	movl	$host_SBox, %esi
	movl	$2048, %edx                     # imm = 0x800
	xorl	%ecx, %ecx
	movl	$1, %r8d
	callq	hipMemcpyToSymbol
	movl	%r14d, %r13d
	shlq	$3, %r13
	leaq	8(%rsp), %rbx
	movq	%rbx, %rdi
	movq	%r13, %rsi
	callq	hipMalloc
	movq	(%rbx), %rdi
	movq	%r12, 32(%rsp)                  # 8-byte Spill
	movq	%r12, %rsi
	movq	%r13, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	16(%rsp), %rbx
	movl	$128, %esi
	movq	%rbx, %rdi
	callq	hipMalloc
	movq	(%rbx), %rdi
	movl	$128, %edx
	movq	24(%rsp), %rsi                  # 8-byte Reload
	movl	$1, %ecx
	callq	hipMemcpy
	movl	%r15d, %ecx
	imull	%ebp, %ecx
	movl	%r14d, %eax
	xorl	%edx, %edx
	divl	%ecx
	movl	%eax, %r12d
	incl	%r12d
	movl	%r12d, %esi
	shll	$6, %esi
	movl	$.L.str.1, %edi
	xorl	%eax, %eax
	callq	printf
	btsq	$32, %r15
	movl	%ebp, %edx
	btsq	$32, %rdx
	movq	%r15, %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB3_3
# %bb.2:
	movq	8(%rsp), %rdx
	movq	16(%rsp), %rcx
	movl	%r12d, %edi
	movl	%r14d, %esi
	callq	_Z25__device_stub__kernel_DESjjPxS_
.LBB3_3:
	movq	8(%rsp), %rsi
	movq	32(%rsp), %rdi                  # 8-byte Reload
	movq	%r13, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movl	$IP, %edi
	callq	hipFree
	movl	$FP, %edi
	callq	hipFree
	movl	$E, %edi
	callq	hipFree
	movl	$P, %edi
	callq	hipFree
	movl	$SBox, %edi
	callq	hipFree
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	addq	$40, %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	_Z10runDESCudajPxS_ii, .Lfunc_end3-_Z10runDESCudajPxS_ii
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function _Z13des_with_fileiPcS_S_ii
.LCPI4_0:
	.quad	0x412e848000000000              # double 1.0E+6
	.text
	.globl	_Z13des_with_fileiPcS_S_ii
	.type	_Z13des_with_fileiPcS_S_ii,@function
_Z13des_with_fileiPcS_S_ii:             # @_Z13des_with_fileiPcS_S_ii
	.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	%r9d, 12(%rsp)                  # 4-byte Spill
	movl	%r8d, 8(%rsp)                   # 4-byte Spill
	movq	%rcx, %rbp
	movq	%rdx, 16(%rsp)                  # 8-byte Spill
	movq	%rsi, %r12
	movl	%edi, %r14d
	movl	n_blocks(%rip), %eax
	shll	$3, %eax
	movslq	%eax, %r15
	movq	%r15, %rdi
	callq	malloc
	movq	%rax, %rbx
	movl	$.L.str.2, %esi
	movq	%r12, %rdi
	callq	fopen
	testq	%rax, %rax
	je	.LBB4_1
# %bb.3:
	movq	%rax, %r13
	movl	$1, %edx
	movq	%rbx, %rdi
	movq	%r15, %rsi
	movq	%rax, %rcx
	callq	fread
	movq	%r13, %rdi
	callq	fclose
	movl	n_blocks(%rip), %r13d
	leaq	(,%r13,8), %rdi
	callq	malloc
	movq	%rax, %r12
	testq	%r13, %r13
	je	.LBB4_8
# %bb.4:                                # %.preheader59.preheader
	xorl	%eax, %eax
	movq	%rbx, %rcx
.LBB4_5:                                # %.preheader59
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB4_6 Depth 2
	xorl	%edx, %edx
	xorl	%esi, %esi
.LBB4_6:                                #   Parent Loop BB4_5 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movq	%rsi, %rdi
	shlq	$8, %rdi
	movzbl	(%rcx,%rdx), %esi
	orq	%rdi, %rsi
	incq	%rdx
	cmpq	$8, %rdx
	jne	.LBB4_6
# %bb.7:                                #   in Loop: Header=BB4_5 Depth=1
	movq	%rsi, (%r12,%rax,8)
	incq	%rax
	addq	$8, %rcx
	cmpq	%r13, %rax
	jne	.LBB4_5
.LBB4_8:                                # %.preheader58.preheader
	xorl	%eax, %eax
	xorl	%edi, %edi
.LBB4_9:                                # %.preheader58
                                        # =>This Inner Loop Header: Depth=1
	movq	%rdi, %rcx
	shlq	$8, %rcx
	movzbl	(%rbp,%rax), %edi
	orq	%rcx, %rdi
	incq	%rax
	cmpq	$8, %rax
	jne	.LBB4_9
# %bb.10:
	movl	%r14d, %esi
	callq	_Z17generate_sub_keysxi
	movq	%rax, %r14
	callq	clock
	movq	%rax, %r13
	movl	n_blocks(%rip), %edi
	movq	%r12, %rsi
	movq	%r14, %rdx
	movl	8(%rsp), %ecx                   # 4-byte Reload
	movl	12(%rsp), %r8d                  # 4-byte Reload
	callq	_Z10runDESCudajPxS_ii
	callq	clock
	subq	%r13, %rax
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI4_0(%rip), %xmm0
	movl	$.L.str.4, %edi
	movb	$1, %al
	callq	printf
	movq	%r14, %rdi
	callq	free
	movl	n_blocks(%rip), %eax
	testq	%rax, %rax
	movq	16(%rsp), %r13                  # 8-byte Reload
	je	.LBB4_15
# %bb.11:                               # %.preheader.preheader
	xorl	%edx, %edx
	xorl	%esi, %esi
.LBB4_12:                               # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB4_13 Depth 2
	movl	%edx, %edx
	movq	(%r12,%rsi,8), %rdi
	movq	%rdx, %r8
	xorl	%ecx, %ecx
.LBB4_13:                               #   Parent Loop BB4_12 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movq	%rdi, %r9
	sarq	%cl, %r9
	movq	%r8, %r10
	xorq	$7, %r10
	movb	%r9b, (%rbx,%r10)
	addq	$8, %rcx
	incq	%r8
	cmpq	$64, %rcx
	jne	.LBB4_13
# %bb.14:                               #   in Loop: Header=BB4_12 Depth=1
	incq	%rsi
	addl	$8, %edx
	cmpq	%rax, %rsi
	jne	.LBB4_12
.LBB4_15:                               # %._crit_edge
	movl	$.L.str.5, %esi
	movq	%r13, %rdi
	callq	fopen
	testq	%rax, %rax
	je	.LBB4_16
# %bb.17:
	movq	%rax, %r14
	movl	$1, %edx
	movq	%rbx, %rdi
	movq	%r15, %rsi
	movq	%rax, %rcx
	callq	fwrite
	movq	%r14, %rdi
	callq	fclose
	movq	%rbx, %rdi
	callq	free
	movq	%r12, %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	free                            # TAILCALL
.LBB4_1:
	.cfi_def_cfa_offset 80
	movl	$.L.str.3, %edi
	movq	%r12, %rsi
	jmp	.LBB4_2
.LBB4_16:
	movl	$.L.str.6, %edi
	movq	%r13, %rsi
.LBB4_2:
	xorl	%eax, %eax
	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	printf                          # TAILCALL
.Lfunc_end4:
	.size	_Z13des_with_fileiPcS_S_ii, .Lfunc_end4-_Z13des_with_fileiPcS_S_ii
	.cfi_endproc
                                        # -- End function
	.globl	_Z10encryptionPcS_S_ii          # -- Begin function _Z10encryptionPcS_S_ii
	.type	_Z10encryptionPcS_S_ii,@function
_Z10encryptionPcS_S_ii:                 # @_Z10encryptionPcS_S_ii
	.cfi_startproc
# %bb.0:
	movl	%r8d, %r9d
	movl	%ecx, %r8d
	movq	%rdx, %rcx
	movq	%rsi, %rdx
	movq	%rdi, %rsi
	xorl	%edi, %edi
	jmp	_Z13des_with_fileiPcS_S_ii      # TAILCALL
.Lfunc_end5:
	.size	_Z10encryptionPcS_S_ii, .Lfunc_end5-_Z10encryptionPcS_S_ii
	.cfi_endproc
                                        # -- End function
	.globl	_Z10decryptionPcS_S_ii          # -- Begin function _Z10decryptionPcS_S_ii
	.type	_Z10decryptionPcS_S_ii,@function
_Z10decryptionPcS_S_ii:                 # @_Z10decryptionPcS_S_ii
	.cfi_startproc
# %bb.0:
	movl	%r8d, %r9d
	movl	%ecx, %r8d
	movq	%rdx, %rcx
	movq	%rsi, %rdx
	movq	%rdi, %rsi
	movl	$1, %edi
	jmp	_Z13des_with_fileiPcS_S_ii      # TAILCALL
.Lfunc_end6:
	.size	_Z10decryptionPcS_S_ii, .Lfunc_end6-_Z10decryptionPcS_S_ii
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$16, %rsp
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -16
	cmpl	$5, %edi
	jg	.LBB7_2
# %bb.1:
	movl	$.Lstr.3, %edi
	callq	puts@PLT
	movl	$.Lstr.4, %edi
	callq	puts@PLT
	movl	$-1, %eax
	jmp	.LBB7_9
.LBB7_2:
	movq	%rsi, %rbx
	movq	40(%rsi), %rdi
	movl	$.L.str.9, %esi
	movl	$n_blocks, %edx
	xorl	%eax, %eax
	callq	__isoc23_sscanf
	movq	48(%rbx), %rdi
	leaq	12(%rsp), %rdx
	movl	$.L.str.9, %esi
	xorl	%eax, %eax
	callq	__isoc23_sscanf
	movq	56(%rbx), %rdi
	leaq	8(%rsp), %rdx
	movl	$.L.str.9, %esi
	xorl	%eax, %eax
	callq	__isoc23_sscanf
	movq	8(%rbx), %rax
	movzbl	(%rax), %eax
	cmpl	$100, %eax
	je	.LBB7_6
# %bb.3:
	cmpl	$101, %eax
	jne	.LBB7_7
# %bb.4:
	movl	$.Lstr.1, %edi
	callq	puts@PLT
	movq	16(%rbx), %rsi
	movq	24(%rbx), %rdx
	movq	32(%rbx), %rcx
	movl	12(%rsp), %r8d
	movl	8(%rsp), %r9d
	xorl	%edi, %edi
	jmp	.LBB7_5
.LBB7_6:
	movl	$.Lstr, %edi
	callq	puts@PLT
	movq	16(%rbx), %rsi
	movq	24(%rbx), %rdx
	movq	32(%rbx), %rcx
	movl	12(%rsp), %r8d
	movl	8(%rsp), %r9d
	movl	$1, %edi
.LBB7_5:
	callq	_Z13des_with_fileiPcS_S_ii
	jmp	.LBB7_8
.LBB7_7:
	movl	$.Lstr.2, %edi
	callq	puts@PLT
.LBB7_8:
	xorl	%eax, %eax
.LBB7_9:
	addq	$16, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end7:
	.size	main, .Lfunc_end7-main
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB8_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB8_2:
	subq	$32, %rsp
	.cfi_adjust_cfa_offset 32
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z10kernel_DESjjPxS_, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	addq	$32, %rsp
	.cfi_adjust_cfa_offset -32
	movl	$IP, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movl	$256, %r9d                      # imm = 0x100
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$FP, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %ecx
	movl	$256, %r9d                      # imm = 0x100
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$E, %esi
	movl	$.L__unnamed_4, %edx
	movl	$.L__unnamed_4, %ecx
	movl	$192, %r9d
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$P, %esi
	movl	$.L__unnamed_5, %edx
	movl	$.L__unnamed_5, %ecx
	movl	$128, %r9d
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$SBox, %esi
	movl	$.L__unnamed_6, %edx
	movl	$.L__unnamed_6, %ecx
	movl	$2048, %r9d                     # imm = 0x800
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$__hip_module_dtor, %edi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end8:
	.size	__hip_module_ctor, .Lfunc_end8-__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	.LBB9_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
.LBB9_2:
	retq
.Lfunc_end9:
	.size	__hip_module_dtor, .Lfunc_end9-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	IP,@object                      # @IP
	.local	IP
	.comm	IP,256,16
	.type	FP,@object                      # @FP
	.local	FP
	.comm	FP,256,16
	.type	E,@object                       # @E
	.local	E
	.comm	E,192,16
	.type	P,@object                       # @P
	.local	P
	.comm	P,128,16
	.type	SBox,@object                    # @SBox
	.local	SBox
	.comm	SBox,2048,16
	.type	host_IP,@object                 # @host_IP
	.data
	.globl	host_IP
	.p2align	4, 0x0
host_IP:
	.long	57                              # 0x39
	.long	49                              # 0x31
	.long	41                              # 0x29
	.long	33                              # 0x21
	.long	25                              # 0x19
	.long	17                              # 0x11
	.long	9                               # 0x9
	.long	1                               # 0x1
	.long	59                              # 0x3b
	.long	51                              # 0x33
	.long	43                              # 0x2b
	.long	35                              # 0x23
	.long	27                              # 0x1b
	.long	19                              # 0x13
	.long	11                              # 0xb
	.long	3                               # 0x3
	.long	61                              # 0x3d
	.long	53                              # 0x35
	.long	45                              # 0x2d
	.long	37                              # 0x25
	.long	29                              # 0x1d
	.long	21                              # 0x15
	.long	13                              # 0xd
	.long	5                               # 0x5
	.long	63                              # 0x3f
	.long	55                              # 0x37
	.long	47                              # 0x2f
	.long	39                              # 0x27
	.long	31                              # 0x1f
	.long	23                              # 0x17
	.long	15                              # 0xf
	.long	7                               # 0x7
	.long	56                              # 0x38
	.long	48                              # 0x30
	.long	40                              # 0x28
	.long	32                              # 0x20
	.long	24                              # 0x18
	.long	16                              # 0x10
	.long	8                               # 0x8
	.long	0                               # 0x0
	.long	58                              # 0x3a
	.long	50                              # 0x32
	.long	42                              # 0x2a
	.long	34                              # 0x22
	.long	26                              # 0x1a
	.long	18                              # 0x12
	.long	10                              # 0xa
	.long	2                               # 0x2
	.long	60                              # 0x3c
	.long	52                              # 0x34
	.long	44                              # 0x2c
	.long	36                              # 0x24
	.long	28                              # 0x1c
	.long	20                              # 0x14
	.long	12                              # 0xc
	.long	4                               # 0x4
	.long	62                              # 0x3e
	.long	54                              # 0x36
	.long	46                              # 0x2e
	.long	38                              # 0x26
	.long	30                              # 0x1e
	.long	22                              # 0x16
	.long	14                              # 0xe
	.long	6                               # 0x6
	.size	host_IP, 256

	.type	host_FP,@object                 # @host_FP
	.globl	host_FP
	.p2align	4, 0x0
host_FP:
	.long	39                              # 0x27
	.long	7                               # 0x7
	.long	47                              # 0x2f
	.long	15                              # 0xf
	.long	55                              # 0x37
	.long	23                              # 0x17
	.long	63                              # 0x3f
	.long	31                              # 0x1f
	.long	38                              # 0x26
	.long	6                               # 0x6
	.long	46                              # 0x2e
	.long	14                              # 0xe
	.long	54                              # 0x36
	.long	22                              # 0x16
	.long	62                              # 0x3e
	.long	30                              # 0x1e
	.long	37                              # 0x25
	.long	5                               # 0x5
	.long	45                              # 0x2d
	.long	13                              # 0xd
	.long	53                              # 0x35
	.long	21                              # 0x15
	.long	61                              # 0x3d
	.long	29                              # 0x1d
	.long	36                              # 0x24
	.long	4                               # 0x4
	.long	44                              # 0x2c
	.long	12                              # 0xc
	.long	52                              # 0x34
	.long	20                              # 0x14
	.long	60                              # 0x3c
	.long	28                              # 0x1c
	.long	35                              # 0x23
	.long	3                               # 0x3
	.long	43                              # 0x2b
	.long	11                              # 0xb
	.long	51                              # 0x33
	.long	19                              # 0x13
	.long	59                              # 0x3b
	.long	27                              # 0x1b
	.long	34                              # 0x22
	.long	2                               # 0x2
	.long	42                              # 0x2a
	.long	10                              # 0xa
	.long	50                              # 0x32
	.long	18                              # 0x12
	.long	58                              # 0x3a
	.long	26                              # 0x1a
	.long	33                              # 0x21
	.long	1                               # 0x1
	.long	41                              # 0x29
	.long	9                               # 0x9
	.long	49                              # 0x31
	.long	17                              # 0x11
	.long	57                              # 0x39
	.long	25                              # 0x19
	.long	32                              # 0x20
	.long	0                               # 0x0
	.long	40                              # 0x28
	.long	8                               # 0x8
	.long	48                              # 0x30
	.long	16                              # 0x10
	.long	56                              # 0x38
	.long	24                              # 0x18
	.size	host_FP, 256

	.type	host_E,@object                  # @host_E
	.globl	host_E
	.p2align	4, 0x0
host_E:
	.long	31                              # 0x1f
	.long	0                               # 0x0
	.long	1                               # 0x1
	.long	2                               # 0x2
	.long	3                               # 0x3
	.long	4                               # 0x4
	.long	3                               # 0x3
	.long	4                               # 0x4
	.long	5                               # 0x5
	.long	6                               # 0x6
	.long	7                               # 0x7
	.long	8                               # 0x8
	.long	7                               # 0x7
	.long	8                               # 0x8
	.long	9                               # 0x9
	.long	10                              # 0xa
	.long	11                              # 0xb
	.long	12                              # 0xc
	.long	11                              # 0xb
	.long	12                              # 0xc
	.long	13                              # 0xd
	.long	14                              # 0xe
	.long	15                              # 0xf
	.long	16                              # 0x10
	.long	15                              # 0xf
	.long	16                              # 0x10
	.long	17                              # 0x11
	.long	18                              # 0x12
	.long	19                              # 0x13
	.long	20                              # 0x14
	.long	19                              # 0x13
	.long	20                              # 0x14
	.long	21                              # 0x15
	.long	22                              # 0x16
	.long	23                              # 0x17
	.long	24                              # 0x18
	.long	23                              # 0x17
	.long	24                              # 0x18
	.long	25                              # 0x19
	.long	26                              # 0x1a
	.long	27                              # 0x1b
	.long	28                              # 0x1c
	.long	27                              # 0x1b
	.long	28                              # 0x1c
	.long	29                              # 0x1d
	.long	30                              # 0x1e
	.long	31                              # 0x1f
	.long	0                               # 0x0
	.size	host_E, 192

	.type	host_P,@object                  # @host_P
	.globl	host_P
	.p2align	4, 0x0
host_P:
	.long	15                              # 0xf
	.long	6                               # 0x6
	.long	19                              # 0x13
	.long	20                              # 0x14
	.long	28                              # 0x1c
	.long	11                              # 0xb
	.long	27                              # 0x1b
	.long	16                              # 0x10
	.long	0                               # 0x0
	.long	14                              # 0xe
	.long	22                              # 0x16
	.long	25                              # 0x19
	.long	4                               # 0x4
	.long	17                              # 0x11
	.long	30                              # 0x1e
	.long	9                               # 0x9
	.long	1                               # 0x1
	.long	7                               # 0x7
	.long	23                              # 0x17
	.long	13                              # 0xd
	.long	31                              # 0x1f
	.long	26                              # 0x1a
	.long	2                               # 0x2
	.long	8                               # 0x8
	.long	18                              # 0x12
	.long	12                              # 0xc
	.long	29                              # 0x1d
	.long	5                               # 0x5
	.long	21                              # 0x15
	.long	10                              # 0xa
	.long	3                               # 0x3
	.long	24                              # 0x18
	.size	host_P, 128

	.type	PC1_LEFT,@object                # @PC1_LEFT
	.globl	PC1_LEFT
	.p2align	4, 0x0
PC1_LEFT:
	.long	56                              # 0x38
	.long	48                              # 0x30
	.long	40                              # 0x28
	.long	32                              # 0x20
	.long	24                              # 0x18
	.long	16                              # 0x10
	.long	8                               # 0x8
	.long	0                               # 0x0
	.long	57                              # 0x39
	.long	49                              # 0x31
	.long	41                              # 0x29
	.long	33                              # 0x21
	.long	25                              # 0x19
	.long	17                              # 0x11
	.long	9                               # 0x9
	.long	1                               # 0x1
	.long	58                              # 0x3a
	.long	50                              # 0x32
	.long	42                              # 0x2a
	.long	34                              # 0x22
	.long	26                              # 0x1a
	.long	18                              # 0x12
	.long	10                              # 0xa
	.long	2                               # 0x2
	.long	59                              # 0x3b
	.long	51                              # 0x33
	.long	43                              # 0x2b
	.long	35                              # 0x23
	.size	PC1_LEFT, 112

	.type	PC1_RIGHT,@object               # @PC1_RIGHT
	.globl	PC1_RIGHT
	.p2align	4, 0x0
PC1_RIGHT:
	.long	62                              # 0x3e
	.long	54                              # 0x36
	.long	46                              # 0x2e
	.long	38                              # 0x26
	.long	30                              # 0x1e
	.long	22                              # 0x16
	.long	14                              # 0xe
	.long	6                               # 0x6
	.long	61                              # 0x3d
	.long	53                              # 0x35
	.long	45                              # 0x2d
	.long	37                              # 0x25
	.long	29                              # 0x1d
	.long	21                              # 0x15
	.long	13                              # 0xd
	.long	5                               # 0x5
	.long	60                              # 0x3c
	.long	52                              # 0x34
	.long	44                              # 0x2c
	.long	36                              # 0x24
	.long	28                              # 0x1c
	.long	20                              # 0x14
	.long	12                              # 0xc
	.long	4                               # 0x4
	.long	27                              # 0x1b
	.long	19                              # 0x13
	.long	11                              # 0xb
	.long	3                               # 0x3
	.size	PC1_RIGHT, 112

	.type	PC2,@object                     # @PC2
	.globl	PC2
	.p2align	4, 0x0
PC2:
	.long	13                              # 0xd
	.long	16                              # 0x10
	.long	10                              # 0xa
	.long	23                              # 0x17
	.long	0                               # 0x0
	.long	4                               # 0x4
	.long	2                               # 0x2
	.long	27                              # 0x1b
	.long	14                              # 0xe
	.long	5                               # 0x5
	.long	20                              # 0x14
	.long	9                               # 0x9
	.long	22                              # 0x16
	.long	18                              # 0x12
	.long	11                              # 0xb
	.long	3                               # 0x3
	.long	25                              # 0x19
	.long	7                               # 0x7
	.long	15                              # 0xf
	.long	6                               # 0x6
	.long	26                              # 0x1a
	.long	19                              # 0x13
	.long	12                              # 0xc
	.long	1                               # 0x1
	.long	40                              # 0x28
	.long	51                              # 0x33
	.long	30                              # 0x1e
	.long	36                              # 0x24
	.long	46                              # 0x2e
	.long	54                              # 0x36
	.long	29                              # 0x1d
	.long	39                              # 0x27
	.long	50                              # 0x32
	.long	44                              # 0x2c
	.long	32                              # 0x20
	.long	47                              # 0x2f
	.long	43                              # 0x2b
	.long	48                              # 0x30
	.long	38                              # 0x26
	.long	55                              # 0x37
	.long	33                              # 0x21
	.long	52                              # 0x34
	.long	45                              # 0x2d
	.long	41                              # 0x29
	.long	49                              # 0x31
	.long	35                              # 0x23
	.long	28                              # 0x1c
	.long	31                              # 0x1f
	.size	PC2, 192

	.type	Rotations,@object               # @Rotations
	.globl	Rotations
	.p2align	4, 0x0
Rotations:
	.long	1                               # 0x1
	.long	1                               # 0x1
	.long	2                               # 0x2
	.long	2                               # 0x2
	.long	2                               # 0x2
	.long	2                               # 0x2
	.long	2                               # 0x2
	.long	2                               # 0x2
	.long	1                               # 0x1
	.long	2                               # 0x2
	.long	2                               # 0x2
	.long	2                               # 0x2
	.long	2                               # 0x2
	.long	2                               # 0x2
	.long	2                               # 0x2
	.long	1                               # 0x1
	.size	Rotations, 64

	.type	host_SBox,@object               # @host_SBox
	.globl	host_SBox
	.p2align	4, 0x0
host_SBox:
	.long	14                              # 0xe
	.long	4                               # 0x4
	.long	13                              # 0xd
	.long	1                               # 0x1
	.long	2                               # 0x2
	.long	15                              # 0xf
	.long	11                              # 0xb
	.long	8                               # 0x8
	.long	3                               # 0x3
	.long	10                              # 0xa
	.long	6                               # 0x6
	.long	12                              # 0xc
	.long	5                               # 0x5
	.long	9                               # 0x9
	.long	0                               # 0x0
	.long	7                               # 0x7
	.long	0                               # 0x0
	.long	15                              # 0xf
	.long	7                               # 0x7
	.long	4                               # 0x4
	.long	14                              # 0xe
	.long	2                               # 0x2
	.long	13                              # 0xd
	.long	1                               # 0x1
	.long	10                              # 0xa
	.long	6                               # 0x6
	.long	12                              # 0xc
	.long	11                              # 0xb
	.long	9                               # 0x9
	.long	5                               # 0x5
	.long	3                               # 0x3
	.long	8                               # 0x8
	.long	4                               # 0x4
	.long	1                               # 0x1
	.long	14                              # 0xe
	.long	8                               # 0x8
	.long	13                              # 0xd
	.long	6                               # 0x6
	.long	2                               # 0x2
	.long	11                              # 0xb
	.long	15                              # 0xf
	.long	12                              # 0xc
	.long	9                               # 0x9
	.long	7                               # 0x7
	.long	3                               # 0x3
	.long	10                              # 0xa
	.long	5                               # 0x5
	.long	0                               # 0x0
	.long	15                              # 0xf
	.long	12                              # 0xc
	.long	8                               # 0x8
	.long	2                               # 0x2
	.long	4                               # 0x4
	.long	9                               # 0x9
	.long	1                               # 0x1
	.long	7                               # 0x7
	.long	5                               # 0x5
	.long	11                              # 0xb
	.long	3                               # 0x3
	.long	14                              # 0xe
	.long	10                              # 0xa
	.long	0                               # 0x0
	.long	6                               # 0x6
	.long	13                              # 0xd
	.long	15                              # 0xf
	.long	1                               # 0x1
	.long	8                               # 0x8
	.long	14                              # 0xe
	.long	6                               # 0x6
	.long	11                              # 0xb
	.long	3                               # 0x3
	.long	4                               # 0x4
	.long	9                               # 0x9
	.long	7                               # 0x7
	.long	2                               # 0x2
	.long	13                              # 0xd
	.long	12                              # 0xc
	.long	0                               # 0x0
	.long	5                               # 0x5
	.long	10                              # 0xa
	.long	3                               # 0x3
	.long	13                              # 0xd
	.long	4                               # 0x4
	.long	7                               # 0x7
	.long	15                              # 0xf
	.long	2                               # 0x2
	.long	8                               # 0x8
	.long	14                              # 0xe
	.long	12                              # 0xc
	.long	0                               # 0x0
	.long	1                               # 0x1
	.long	10                              # 0xa
	.long	6                               # 0x6
	.long	9                               # 0x9
	.long	11                              # 0xb
	.long	5                               # 0x5
	.long	0                               # 0x0
	.long	14                              # 0xe
	.long	7                               # 0x7
	.long	11                              # 0xb
	.long	10                              # 0xa
	.long	4                               # 0x4
	.long	13                              # 0xd
	.long	1                               # 0x1
	.long	5                               # 0x5
	.long	8                               # 0x8
	.long	12                              # 0xc
	.long	6                               # 0x6
	.long	9                               # 0x9
	.long	3                               # 0x3
	.long	2                               # 0x2
	.long	15                              # 0xf
	.long	13                              # 0xd
	.long	8                               # 0x8
	.long	10                              # 0xa
	.long	1                               # 0x1
	.long	3                               # 0x3
	.long	15                              # 0xf
	.long	4                               # 0x4
	.long	2                               # 0x2
	.long	11                              # 0xb
	.long	6                               # 0x6
	.long	7                               # 0x7
	.long	12                              # 0xc
	.long	0                               # 0x0
	.long	5                               # 0x5
	.long	14                              # 0xe
	.long	9                               # 0x9
	.long	10                              # 0xa
	.long	0                               # 0x0
	.long	9                               # 0x9
	.long	14                              # 0xe
	.long	6                               # 0x6
	.long	3                               # 0x3
	.long	15                              # 0xf
	.long	5                               # 0x5
	.long	1                               # 0x1
	.long	13                              # 0xd
	.long	12                              # 0xc
	.long	7                               # 0x7
	.long	11                              # 0xb
	.long	4                               # 0x4
	.long	2                               # 0x2
	.long	8                               # 0x8
	.long	13                              # 0xd
	.long	7                               # 0x7
	.long	0                               # 0x0
	.long	9                               # 0x9
	.long	3                               # 0x3
	.long	4                               # 0x4
	.long	6                               # 0x6
	.long	10                              # 0xa
	.long	2                               # 0x2
	.long	8                               # 0x8
	.long	5                               # 0x5
	.long	14                              # 0xe
	.long	12                              # 0xc
	.long	11                              # 0xb
	.long	15                              # 0xf
	.long	1                               # 0x1
	.long	13                              # 0xd
	.long	6                               # 0x6
	.long	4                               # 0x4
	.long	9                               # 0x9
	.long	8                               # 0x8
	.long	15                              # 0xf
	.long	3                               # 0x3
	.long	0                               # 0x0
	.long	11                              # 0xb
	.long	1                               # 0x1
	.long	2                               # 0x2
	.long	12                              # 0xc
	.long	5                               # 0x5
	.long	10                              # 0xa
	.long	14                              # 0xe
	.long	7                               # 0x7
	.long	1                               # 0x1
	.long	10                              # 0xa
	.long	13                              # 0xd
	.long	0                               # 0x0
	.long	6                               # 0x6
	.long	9                               # 0x9
	.long	8                               # 0x8
	.long	7                               # 0x7
	.long	4                               # 0x4
	.long	15                              # 0xf
	.long	14                              # 0xe
	.long	3                               # 0x3
	.long	11                              # 0xb
	.long	5                               # 0x5
	.long	2                               # 0x2
	.long	12                              # 0xc
	.long	7                               # 0x7
	.long	13                              # 0xd
	.long	14                              # 0xe
	.long	3                               # 0x3
	.long	0                               # 0x0
	.long	6                               # 0x6
	.long	9                               # 0x9
	.long	10                              # 0xa
	.long	1                               # 0x1
	.long	2                               # 0x2
	.long	8                               # 0x8
	.long	5                               # 0x5
	.long	11                              # 0xb
	.long	12                              # 0xc
	.long	4                               # 0x4
	.long	15                              # 0xf
	.long	13                              # 0xd
	.long	8                               # 0x8
	.long	11                              # 0xb
	.long	5                               # 0x5
	.long	6                               # 0x6
	.long	15                              # 0xf
	.long	0                               # 0x0
	.long	3                               # 0x3
	.long	4                               # 0x4
	.long	7                               # 0x7
	.long	2                               # 0x2
	.long	12                              # 0xc
	.long	1                               # 0x1
	.long	10                              # 0xa
	.long	14                              # 0xe
	.long	9                               # 0x9
	.long	10                              # 0xa
	.long	6                               # 0x6
	.long	9                               # 0x9
	.long	0                               # 0x0
	.long	12                              # 0xc
	.long	11                              # 0xb
	.long	7                               # 0x7
	.long	13                              # 0xd
	.long	15                              # 0xf
	.long	1                               # 0x1
	.long	3                               # 0x3
	.long	14                              # 0xe
	.long	5                               # 0x5
	.long	2                               # 0x2
	.long	8                               # 0x8
	.long	4                               # 0x4
	.long	3                               # 0x3
	.long	15                              # 0xf
	.long	0                               # 0x0
	.long	6                               # 0x6
	.long	10                              # 0xa
	.long	1                               # 0x1
	.long	13                              # 0xd
	.long	8                               # 0x8
	.long	9                               # 0x9
	.long	4                               # 0x4
	.long	5                               # 0x5
	.long	11                              # 0xb
	.long	12                              # 0xc
	.long	7                               # 0x7
	.long	2                               # 0x2
	.long	14                              # 0xe
	.long	2                               # 0x2
	.long	12                              # 0xc
	.long	4                               # 0x4
	.long	1                               # 0x1
	.long	7                               # 0x7
	.long	10                              # 0xa
	.long	11                              # 0xb
	.long	6                               # 0x6
	.long	8                               # 0x8
	.long	5                               # 0x5
	.long	3                               # 0x3
	.long	15                              # 0xf
	.long	13                              # 0xd
	.long	0                               # 0x0
	.long	14                              # 0xe
	.long	9                               # 0x9
	.long	14                              # 0xe
	.long	11                              # 0xb
	.long	2                               # 0x2
	.long	12                              # 0xc
	.long	4                               # 0x4
	.long	7                               # 0x7
	.long	13                              # 0xd
	.long	1                               # 0x1
	.long	5                               # 0x5
	.long	0                               # 0x0
	.long	15                              # 0xf
	.long	10                              # 0xa
	.long	3                               # 0x3
	.long	9                               # 0x9
	.long	8                               # 0x8
	.long	6                               # 0x6
	.long	4                               # 0x4
	.long	2                               # 0x2
	.long	1                               # 0x1
	.long	11                              # 0xb
	.long	10                              # 0xa
	.long	13                              # 0xd
	.long	7                               # 0x7
	.long	8                               # 0x8
	.long	15                              # 0xf
	.long	9                               # 0x9
	.long	12                              # 0xc
	.long	5                               # 0x5
	.long	6                               # 0x6
	.long	3                               # 0x3
	.long	0                               # 0x0
	.long	14                              # 0xe
	.long	11                              # 0xb
	.long	8                               # 0x8
	.long	12                              # 0xc
	.long	7                               # 0x7
	.long	1                               # 0x1
	.long	14                              # 0xe
	.long	2                               # 0x2
	.long	13                              # 0xd
	.long	6                               # 0x6
	.long	15                              # 0xf
	.long	0                               # 0x0
	.long	9                               # 0x9
	.long	10                              # 0xa
	.long	4                               # 0x4
	.long	5                               # 0x5
	.long	3                               # 0x3
	.long	12                              # 0xc
	.long	1                               # 0x1
	.long	10                              # 0xa
	.long	15                              # 0xf
	.long	9                               # 0x9
	.long	2                               # 0x2
	.long	6                               # 0x6
	.long	8                               # 0x8
	.long	0                               # 0x0
	.long	13                              # 0xd
	.long	3                               # 0x3
	.long	4                               # 0x4
	.long	14                              # 0xe
	.long	7                               # 0x7
	.long	5                               # 0x5
	.long	11                              # 0xb
	.long	10                              # 0xa
	.long	15                              # 0xf
	.long	4                               # 0x4
	.long	2                               # 0x2
	.long	7                               # 0x7
	.long	12                              # 0xc
	.long	9                               # 0x9
	.long	5                               # 0x5
	.long	6                               # 0x6
	.long	1                               # 0x1
	.long	13                              # 0xd
	.long	14                              # 0xe
	.long	0                               # 0x0
	.long	11                              # 0xb
	.long	3                               # 0x3
	.long	8                               # 0x8
	.long	9                               # 0x9
	.long	14                              # 0xe
	.long	15                              # 0xf
	.long	5                               # 0x5
	.long	2                               # 0x2
	.long	8                               # 0x8
	.long	12                              # 0xc
	.long	3                               # 0x3
	.long	7                               # 0x7
	.long	0                               # 0x0
	.long	4                               # 0x4
	.long	10                              # 0xa
	.long	1                               # 0x1
	.long	13                              # 0xd
	.long	11                              # 0xb
	.long	6                               # 0x6
	.long	4                               # 0x4
	.long	3                               # 0x3
	.long	2                               # 0x2
	.long	12                              # 0xc
	.long	9                               # 0x9
	.long	5                               # 0x5
	.long	15                              # 0xf
	.long	10                              # 0xa
	.long	11                              # 0xb
	.long	14                              # 0xe
	.long	1                               # 0x1
	.long	7                               # 0x7
	.long	6                               # 0x6
	.long	0                               # 0x0
	.long	8                               # 0x8
	.long	13                              # 0xd
	.long	4                               # 0x4
	.long	11                              # 0xb
	.long	2                               # 0x2
	.long	14                              # 0xe
	.long	15                              # 0xf
	.long	0                               # 0x0
	.long	8                               # 0x8
	.long	13                              # 0xd
	.long	3                               # 0x3
	.long	12                              # 0xc
	.long	9                               # 0x9
	.long	7                               # 0x7
	.long	5                               # 0x5
	.long	10                              # 0xa
	.long	6                               # 0x6
	.long	1                               # 0x1
	.long	13                              # 0xd
	.long	0                               # 0x0
	.long	11                              # 0xb
	.long	7                               # 0x7
	.long	4                               # 0x4
	.long	9                               # 0x9
	.long	1                               # 0x1
	.long	10                              # 0xa
	.long	14                              # 0xe
	.long	3                               # 0x3
	.long	5                               # 0x5
	.long	12                              # 0xc
	.long	2                               # 0x2
	.long	15                              # 0xf
	.long	8                               # 0x8
	.long	6                               # 0x6
	.long	1                               # 0x1
	.long	4                               # 0x4
	.long	11                              # 0xb
	.long	13                              # 0xd
	.long	12                              # 0xc
	.long	3                               # 0x3
	.long	7                               # 0x7
	.long	14                              # 0xe
	.long	10                              # 0xa
	.long	15                              # 0xf
	.long	6                               # 0x6
	.long	8                               # 0x8
	.long	0                               # 0x0
	.long	5                               # 0x5
	.long	9                               # 0x9
	.long	2                               # 0x2
	.long	6                               # 0x6
	.long	11                              # 0xb
	.long	13                              # 0xd
	.long	8                               # 0x8
	.long	1                               # 0x1
	.long	4                               # 0x4
	.long	10                              # 0xa
	.long	7                               # 0x7
	.long	9                               # 0x9
	.long	5                               # 0x5
	.long	0                               # 0x0
	.long	15                              # 0xf
	.long	14                              # 0xe
	.long	2                               # 0x2
	.long	3                               # 0x3
	.long	12                              # 0xc
	.long	13                              # 0xd
	.long	2                               # 0x2
	.long	8                               # 0x8
	.long	4                               # 0x4
	.long	6                               # 0x6
	.long	15                              # 0xf
	.long	11                              # 0xb
	.long	1                               # 0x1
	.long	10                              # 0xa
	.long	9                               # 0x9
	.long	3                               # 0x3
	.long	14                              # 0xe
	.long	5                               # 0x5
	.long	0                               # 0x0
	.long	12                              # 0xc
	.long	7                               # 0x7
	.long	1                               # 0x1
	.long	15                              # 0xf
	.long	13                              # 0xd
	.long	8                               # 0x8
	.long	10                              # 0xa
	.long	3                               # 0x3
	.long	7                               # 0x7
	.long	4                               # 0x4
	.long	12                              # 0xc
	.long	5                               # 0x5
	.long	6                               # 0x6
	.long	11                              # 0xb
	.long	0                               # 0x0
	.long	14                              # 0xe
	.long	9                               # 0x9
	.long	2                               # 0x2
	.long	7                               # 0x7
	.long	11                              # 0xb
	.long	4                               # 0x4
	.long	1                               # 0x1
	.long	9                               # 0x9
	.long	12                              # 0xc
	.long	14                              # 0xe
	.long	2                               # 0x2
	.long	0                               # 0x0
	.long	6                               # 0x6
	.long	10                              # 0xa
	.long	13                              # 0xd
	.long	15                              # 0xf
	.long	3                               # 0x3
	.long	5                               # 0x5
	.long	8                               # 0x8
	.long	2                               # 0x2
	.long	1                               # 0x1
	.long	14                              # 0xe
	.long	7                               # 0x7
	.long	4                               # 0x4
	.long	10                              # 0xa
	.long	8                               # 0x8
	.long	13                              # 0xd
	.long	15                              # 0xf
	.long	12                              # 0xc
	.long	9                               # 0x9
	.long	0                               # 0x0
	.long	3                               # 0x3
	.long	5                               # 0x5
	.long	6                               # 0x6
	.long	11                              # 0xb
	.size	host_SBox, 2048

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

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"Maximum value of the number of threads is 512. You entered : %d\n"
	.size	.L.str, 65

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"%u bytes per threads.\n"
	.size	.L.str.1, 23

	.type	n_blocks,@object                # @n_blocks
	.bss
	.globl	n_blocks
	.p2align	2, 0x0
n_blocks:
	.long	0                               # 0x0
	.size	n_blocks, 4

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

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"Can't open the in file :%s\n"
	.size	.L.str.3, 28

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

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"wb"
	.size	.L.str.5, 3

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"Can't open the out file :%s\n"
	.size	.L.str.6, 29

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

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z10kernel_DESjjPxS_"
	.size	.L__unnamed_1, 21

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"IP"
	.size	.L__unnamed_2, 3

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"FP"
	.size	.L__unnamed_3, 3

	.type	.L__unnamed_4,@object           # @3
.L__unnamed_4:
	.asciz	"E"
	.size	.L__unnamed_4, 2

	.type	.L__unnamed_5,@object           # @4
.L__unnamed_5:
	.asciz	"P"
	.size	.L__unnamed_5, 2

	.type	.L__unnamed_6,@object           # @5
.L__unnamed_6:
	.asciz	"SBox"
	.size	.L__unnamed_6, 5

	.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	"decryption"
	.size	.Lstr, 11

	.type	.Lstr.1,@object                 # @str.1
.Lstr.1:
	.asciz	"encryption"
	.size	.Lstr.1, 11

	.type	.Lstr.2,@object                 # @str.2
.Lstr.2:
	.asciz	"mode must be 'e' or 'd'"
	.size	.Lstr.2, 24

	.type	.Lstr.3,@object                 # @str.3
.Lstr.3:
	.asciz	"usage) ./des.out [e|d] <input_file> <output_file> <n_des_block_size> <n_cuda_blocks> <n_cuda_threads>"
	.size	.Lstr.3, 102

	.type	.Lstr.4,@object                 # @str.4
.Lstr.4:
	.asciz	"example) ./des.out e in.txt out.txt 1"
	.size	.Lstr.4, 38

	.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__kernel_DESjjPxS_
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym IP
	.addrsig_sym FP
	.addrsig_sym E
	.addrsig_sym P
	.addrsig_sym SBox
	.addrsig_sym host_IP
	.addrsig_sym host_FP
	.addrsig_sym host_E
	.addrsig_sym host_P
	.addrsig_sym host_SBox
	.addrsig_sym _Z10kernel_DESjjPxS_
	.addrsig_sym n_blocks
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z10kernel_DESjjPxS_    ; -- Begin function _Z10kernel_DESjjPxS_
	.globl	_Z10kernel_DESjjPxS_
	.p2align	8
	.type	_Z10kernel_DESjjPxS_,@function
_Z10kernel_DESjjPxS_:                   ; @_Z10kernel_DESjjPxS_
; %bb.0:
	s_clause 0x1
	s_load_b32 s4, s[0:1], 0x24
	s_load_b64 s[2:3], s[0:1], 0x0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s4, s4, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s4, v[0:1]
	v_mul_lo_u32 v0, v1, s2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v12, s2, v0
	v_cmp_gt_u32_e64 s2, s3, v0
	v_cmp_lt_i32_e32 vcc_lo, v0, v12
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s2, vcc_lo, s2
	s_and_saveexec_b32 s4, s2
	s_cbranch_execz .LBB0_15
; %bb.1:                                ; %.lr.ph.preheader
	s_load_b128 s[4:7], s[0:1], 0x8
	v_mov_b32_e32 v2, 0
	s_mov_b32 s1, 0
	s_mov_b32 s2, 0
.LBB0_2:                                ; %.lr.ph
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_3 Depth 2
                                        ;     Child Loop BB0_5 Depth 2
                                        ;       Child Loop BB0_6 Depth 3
                                        ;       Child Loop BB0_8 Depth 3
                                        ;       Child Loop BB0_10 Depth 3
                                        ;     Child Loop BB0_13 Depth 2
	v_ashrrev_i32_e32 v1, 31, v0
	v_mov_b32_e32 v8, 0
	v_mov_b32_e32 v9, 0
	s_mov_b64 s[8:9], 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[3:4], 3, v[0:1]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v3, vcc_lo, s4, v3
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v4, vcc_lo, s5, v4, vcc_lo
	global_load_b64 v[5:6], v[3:4], off
.LBB0_3:                                ; %.lr.ph.i.i
                                        ;   Parent Loop BB0_2 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_getpc_b64 s[10:11]
	s_add_u32 s10, s10, IP@rel32@lo+4
	s_addc_u32 s11, s11, IP@rel32@hi+12
	s_add_u32 s10, s8, s10
	s_addc_u32 s11, s9, s11
	v_lshlrev_b64 v[8:9], 1, v[8:9]
	s_load_b32 s0, s[10:11], 0x0
	s_waitcnt lgkmcnt(0)
	s_sub_i32 s0, 63, s0
	s_add_u32 s8, s8, 4
	s_waitcnt vmcnt(0)
	v_lshrrev_b64 v[10:11], s0, v[5:6]
	s_addc_u32 s9, s9, 0
	s_cmpk_lg_i32 s8, 0x100
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_and_b32_e32 v1, 1, v10
	v_or_b32_e32 v8, v1, v8
	s_cbranch_scc1 .LBB0_3
; %bb.4:                                ; %_Z11permutationxiPii.exit.i
                                        ;   in Loop: Header=BB0_2 Depth=1
	v_mov_b32_e32 v5, v9
	s_mov_b32 s0, 0
.LBB0_5:                                ;   Parent Loop BB0_2 Depth=1
                                        ; =>  This Loop Header: Depth=2
                                        ;       Child Loop BB0_6 Depth 3
                                        ;       Child Loop BB0_8 Depth 3
                                        ;       Child Loop BB0_10 Depth 3
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	s_lshl_b64 s[8:9], s[0:1], 3
	v_dual_mov_b32 v6, v8 :: v_dual_mov_b32 v7, v9
	s_add_u32 s8, s6, s8
	s_addc_u32 s9, s7, s9
	v_dual_mov_b32 v8, 0 :: v_dual_mov_b32 v1, v6
	global_load_b64 v[10:11], v2, s[8:9]
	v_mov_b32_e32 v9, 0
	s_mov_b64 s[8:9], 0
.LBB0_6:                                ; %.lr.ph.i.i.i
                                        ;   Parent Loop BB0_2 Depth=1
                                        ;     Parent Loop BB0_5 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	s_getpc_b64 s[10:11]
	s_add_u32 s10, s10, E@rel32@lo+4
	s_addc_u32 s11, s11, E@rel32@hi+12
	s_add_u32 s10, s8, s10
	s_addc_u32 s11, s9, s11
	v_lshlrev_b64 v[8:9], 1, v[8:9]
	s_load_b32 s10, s[10:11], 0x0
	s_waitcnt lgkmcnt(0)
	s_sub_i32 s10, 31, s10
	s_add_u32 s8, s8, 4
	v_lshrrev_b64 v[13:14], s10, v[1:2]
	s_addc_u32 s9, s9, 0
	s_cmpk_lg_i32 s8, 0xc0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_and_b32_e32 v13, 1, v13
	v_or_b32_e32 v8, v13, v8
	s_cbranch_scc1 .LBB0_6
; %bb.7:                                ; %_Z11permutationxiPii.exit.i.i
                                        ;   in Loop: Header=BB0_5 Depth=2
	s_waitcnt vmcnt(0)
	v_xor_b32_e32 v11, v9, v11
	s_delay_alu instid0(VALU_DEP_2)
	v_xor_b32_e32 v10, v8, v10
	v_mov_b32_e32 v8, 0
	v_mov_b32_e32 v9, 0
	s_mov_b32 s10, 0
	s_getpc_b64 s[8:9]
	s_add_u32 s8, s8, SBox@rel32@lo+4
	s_addc_u32 s9, s9, SBox@rel32@hi+12
.LBB0_8:                                ;   Parent Loop BB0_2 Depth=1
                                        ;     Parent Loop BB0_5 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	s_xor_b32 s11, s10, 7
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[8:9], 4, v[8:9]
	s_mul_i32 s11, s11, 6
	s_add_i32 s10, s10, 1
	v_ashrrev_i64 v[13:14], s11, v[10:11]
	v_lshlrev_b32_e32 v14, 4, v13
	v_bfe_u32 v15, v13, 1, 4
	v_and_b32_e32 v13, 32, v13
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_and_b32_e32 v14, 16, v14
	v_or3_b32 v13, v14, v13, v15
	s_delay_alu instid0(VALU_DEP_1)
	v_lshlrev_b32_e32 v13, 2, v13
	global_load_b32 v13, v13, s[8:9]
	s_add_u32 s8, s8, 0x100
	s_addc_u32 s9, s9, 0
	s_cmp_lg_u32 s10, 8
	s_waitcnt vmcnt(0)
	v_ashrrev_i32_e32 v14, 31, v13
	v_add_co_u32 v8, vcc_lo, v8, v13
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v9, vcc_lo, v9, v14, vcc_lo
	s_cbranch_scc1 .LBB0_8
; %bb.9:                                ; %.lr.ph.i7.i.i.preheader
                                        ;   in Loop: Header=BB0_5 Depth=2
	v_mov_b32_e32 v10, 0
	v_mov_b32_e32 v11, 0
	s_mov_b64 s[8:9], 0
.LBB0_10:                               ; %.lr.ph.i7.i.i
                                        ;   Parent Loop BB0_2 Depth=1
                                        ;     Parent Loop BB0_5 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	s_getpc_b64 s[10:11]
	s_add_u32 s10, s10, P@rel32@lo+4
	s_addc_u32 s11, s11, P@rel32@hi+12
	s_add_u32 s10, s8, s10
	s_addc_u32 s11, s9, s11
	v_lshlrev_b64 v[10:11], 1, v[10:11]
	s_load_b32 s10, s[10:11], 0x0
	s_waitcnt lgkmcnt(0)
	s_sub_i32 s10, 31, s10
	s_add_u32 s8, s8, 4
	v_lshrrev_b64 v[13:14], s10, v[8:9]
	s_addc_u32 s9, s9, 0
	s_cmpk_lg_i32 s8, 0x80
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_and_b32_e32 v13, 1, v13
	v_or_b32_e32 v10, v13, v10
	s_cbranch_scc1 .LBB0_10
; %bb.11:                               ; %_Z1Fjx.exit.i
                                        ;   in Loop: Header=BB0_5 Depth=2
	s_delay_alu instid0(VALU_DEP_1)
	v_xor_b32_e32 v8, v10, v5
	v_mov_b32_e32 v5, v6
	s_add_i32 s0, s0, 1
	v_mov_b32_e32 v6, v7
	s_cmp_lg_u32 s0, 16
	s_cbranch_scc1 .LBB0_5
; %bb.12:                               ;   in Loop: Header=BB0_2 Depth=1
	v_mov_b32_e32 v5, 0
	v_dual_mov_b32 v7, v1 :: v_dual_mov_b32 v6, 0
	s_mov_b64 s[8:9], 0
.LBB0_13:                               ; %.lr.ph.i19.i
                                        ;   Parent Loop BB0_2 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_getpc_b64 s[10:11]
	s_add_u32 s10, s10, FP@rel32@lo+4
	s_addc_u32 s11, s11, FP@rel32@hi+12
	s_add_u32 s10, s8, s10
	s_addc_u32 s11, s9, s11
	v_lshlrev_b64 v[5:6], 1, v[5:6]
	s_load_b32 s0, s[10:11], 0x0
	s_waitcnt lgkmcnt(0)
	s_sub_i32 s0, 63, s0
	s_add_u32 s8, s8, 4
	v_lshrrev_b64 v[9:10], s0, v[7:8]
	s_addc_u32 s9, s9, 0
	s_cmpk_lg_i32 s8, 0x100
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_and_b32_e32 v1, 1, v9
	v_or_b32_e32 v5, v1, v5
	s_cbranch_scc1 .LBB0_13
; %bb.14:                               ; %_Z3DESiPxS_.exit
                                        ;   in Loop: Header=BB0_2 Depth=1
	v_add_nc_u32_e32 v0, 1, v0
	global_store_b64 v[3:4], v[5:6], off
	v_cmp_ge_i32_e32 vcc_lo, v0, v12
	v_cmp_le_u32_e64 s0, s3, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s0, vcc_lo, s0
	s_and_b32 s0, exec_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s2, s0, s2
	s_and_not1_b32 exec_lo, exec_lo, s2
	s_cbranch_execnz .LBB0_2
.LBB0_15:                               ; %._crit_edge
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z10kernel_DESjjPxS_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 280
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 16
		.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	_Z10kernel_DESjjPxS_, .Lfunc_end0-_Z10kernel_DESjjPxS_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 868
; NumSgprs: 18
; NumVgprs: 16
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 16
; Occupancy: 16
; WaveLimiterHint : 1
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.protected	IP                      ; @IP
	.type	IP,@object
	.section	.bss,"aw",@nobits
	.globl	IP
	.p2align	4, 0x0
IP:
	.zero	256
	.size	IP, 256

	.protected	FP                      ; @FP
	.type	FP,@object
	.globl	FP
	.p2align	4, 0x0
FP:
	.zero	256
	.size	FP, 256

	.protected	E                       ; @E
	.type	E,@object
	.globl	E
	.p2align	4, 0x0
E:
	.zero	192
	.size	E, 192

	.protected	P                       ; @P
	.type	P,@object
	.globl	P
	.p2align	4, 0x0
P:
	.zero	128
	.size	P, 128

	.protected	SBox                    ; @SBox
	.type	SBox,@object
	.globl	SBox
	.p2align	4, 0x0
SBox:
	.zero	2048
	.size	SBox, 2048

	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.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 IP
	.addrsig_sym FP
	.addrsig_sym E
	.addrsig_sym P
	.addrsig_sym SBox
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .offset:         4
        .size:           4
        .value_kind:     by_value
      - .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:     hidden_block_count_x
      - .offset:         28
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         36
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         38
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         40
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         42
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         44
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         46
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         88
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 280
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z10kernel_DESjjPxS_
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z10kernel_DESjjPxS_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     16
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
