/**************************************************************************
*            Unix-like crypt(3) Algorithm for Password Encryption
*
*   File    : crypt3.c
*   Purpose : Provides crypt(3) functionality to ANSI C compilers
*             without a need for the crypt library.
*   Author  : Michael Dipperstein
*   Date    : November 3, 1998
*
***************************************************************************
*   The source in this file is heavily borrowed from the crypt3.c file
*   found on several ftp sites on the Internet.  The original source
*   claimed to be BSD, but was not distributed with any BSD license or
*   copyright claims.  I am releasing the source that I have provided into
*   public domain without any restrictions, warranties, or copyright
*   claims of my own.
*
*   The code below has been cleaned and compiles correctly under, gcc,
*   lcc, and Borland's bcc C compilers.  A bug involving the left and
*   right halves of the encrypted data block in the widely published
*   crypt3.c source has been fixed by this version.  All implicit register
*   declarations have been removed, because they generated suboptimal code.
*   All constant data has been explicitly declared as const and all
*   declarations have been given a minimal scope, because I'm paranoid.
*
*   Caution: crypt() returns a pointer to static data.  I left it this way
*            to maintain backward compatibility.  The downside is that
*            successive calls will cause previous results to be lost.
*            This can easily be changed with only minor modifications to
*            the function crypt().
**************************************************************************/

/* Initial permutation */
__device__ static const char IP[] =
{
    58, 50, 42, 34, 26, 18, 10, 2,
    60, 52, 44, 36, 28, 20, 12, 4,
    62, 54, 46, 38, 30, 22, 14, 6,
    64, 56, 48, 40, 32, 24, 16, 8,
    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,
};

/* Final permutation, FP = IP^(-1) */
__device__ static const char FP[] = {
    40, 8, 48, 16, 56, 24, 64, 32,
    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,
};

/**************************************************************************
* Permuted-choice 1 from the key bits to yield C and D.
* Note that bits 8,16... are left out:
* They are intended for a parity check.
**************************************************************************/
__device__ static const char PC1_C[] =
{
    57, 49, 41, 33, 25, 17,  9,
     1, 58, 50, 42, 34, 26, 18,
    10,  2, 59, 51, 43, 35, 27,
    19, 11,  3, 60, 52, 44, 36,
};

__device__ static const char PC1_D[] =
{
    63, 55, 47, 39, 31, 23, 15,
     7, 62, 54, 46, 38, 30, 22,
    14,  6, 61, 53, 45, 37, 29,
    21, 13,  5, 28, 20, 12,  4,
};

/* Sequence of shifts used for the key schedule. */
__device__ static const char shifts[] =
    {1, 1, 2, 2, 2, 2, 2, 2, 1, 2, 2, 2, 2, 2, 2, 1};

/**************************************************************************
* Permuted-choice 2, to pick out the bits from the CD array that generate
* the key schedule.
**************************************************************************/
__device__ static const char PC2_C[] =
{
    14, 17, 11, 24,  1,  5,
     3, 28, 15,  6, 21, 10,
    23, 19, 12,  4, 26,  8,
    16,  7, 27, 20, 13,  2,
};

__device__ static const char PC2_D[] =
{
    41, 52, 31, 37, 47, 55,
    30, 40, 51, 45, 33, 48,
    44, 49, 39, 56, 34, 53,
    46, 42, 50, 36, 29, 32,
};

/* The C and D arrays used to calculate the key schedule. */
__device__ static char C[28];
__device__ static char D[28];

/* The key schedule.  Generated from the key. */
__device__ static char KS[16][48];

/* The E bit-selection table. */
__device__ static char E[48];
__device__ static const char e2[] =
{
    32,  1,  2,  3,  4,  5,
     4,  5,  6,  7,  8,  9,
     8,  9, 10, 11, 12, 13,
    12, 13, 14, 15, 16, 17,
    16, 17, 18, 19, 20, 21,
    20, 21, 22, 23, 24, 25,
    24, 25, 26, 27, 28, 29,
    28, 29, 30, 31, 32,  1,
};

/**************************************************************************
* Function:    setkey
*
* Description: Set up the key schedule from the encryption key.
*
* Inputs:      char *key
*              pointer to 64 character array.  Each character represents a
*              bit in the key.
*
* Returns:     none
**************************************************************************/
__device__ void setkey(char *key)
{
    int i, j, k, temp;

    /**********************************************************************
    * First, generate C and D by permuting the key.  The low order bit of
    * each 8-bit char is not used, so C and D are only 28 bits apiece.
    **********************************************************************/
    for(i = 0; i < 28; i++)
    {
        C[i] = key[PC1_C[i] - 1];
        D[i] = key[PC1_D[i] - 1];
    }

    /**********************************************************************
    * To generate Ki, rotate C and D according to schedule and pick up a
    * permutation using PC2.
    **********************************************************************/
    for(i = 0; i < 16; i++)
    {
        /* rotate */
        for(k = 0; k < shifts[i]; k++)
        {
            temp = C[0];

            for(j = 0; j < 28 - 1; j++)
                C[j] = C[j+1];

            C[27] = temp;
            temp = D[0];
            for(j = 0; j < 28 - 1; j++)
                D[j] = D[j+1];

            D[27] = temp;
        }

        /* get Ki. Note C and D are concatenated */
        for(j = 0; j < 24; j++)
        {
            KS[i][j] = C[PC2_C[j] - 1];
            KS[i][j + 24] = D[PC2_D[j] - 28 -1];
        }
    }

    /* load E with the initial E bit selections */
    for(i=0; i < 48; i++)
        E[i] = e2[i];
}

/**************************************************************************
* The 8 selection functions. For some reason, they give a 0-origin
* index, unlike everything else.
**************************************************************************/

__device__ static const char S[8][64] =
{
    {
        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
    },

    {
        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
    },

    {
        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
    },

    {
         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
    },

    {
         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
    },

    {
        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
    },

    {
         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
    },

    {
        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
    }
};

/**************************************************************************
* P is a permutation on the selected combination of the current L and key.
**************************************************************************/
__device__ static const char P[] =
{
    16,  7, 20, 21,
    29, 12, 28, 17,
     1, 15, 23, 26,
     5, 18, 31, 10,
     2,  8, 24, 14,
    32, 27,  3,  9,
    19, 13, 30,  6,
    22, 11,  4, 25,
};

/* The combination of the key and the input, before selection. */
__device__ static char preS[48];

/**************************************************************************
* Function:    encrypt
*
* Description: Uses DES to encrypt a 64 bit block of data.  Requires
*              setkey to be invoked with the encryption key before it may
*              be used.  The results of the encryption are stored in block.
*
* Inputs:      char *block
*              pointer to 64 character array.  Each character represents a
*              bit in the data block.
*
* Returns:     none
**************************************************************************/
__device__ void encrypt(char *block)
{
    int i, ii, temp, j, k;

    char left[32], right[32]; /* block in two halves */
    char old[32];
    char f[32];

    /* First, permute the bits in the input */
    for(j = 0; j < 32; j++)
        left[j] = block[IP[j] - 1];

    for(;j < 64; j++)
        right[j - 32] = block[IP[j] - 1];

    /* Perform an encryption operation 16 times. */
    for(ii= 0; ii < 16; ii++)
    {
        i = ii;
        /* Save the right array, which will be the new left. */
        for(j = 0; j < 32; j++)
            old[j] = right[j];

        /******************************************************************
        * Expand right to 48 bits using the E selector and
        * exclusive-or with the current key bits.
        ******************************************************************/
        for(j =0 ; j < 48; j++)
            preS[j] = right[E[j] - 1] ^ KS[i][j];

        /******************************************************************
        * The pre-select bits are now considered in 8 groups of 6 bits ea.
        * The 8 selection functions map these 6-bit quantities into 4-bit
        * quantities and the results are permuted to make an f(R, K).
        * The indexing into the selection functions is peculiar;
        * it could be simplified by rewriting the tables.
        ******************************************************************/
        for(j = 0; j < 8; j++)
        {
            temp = 6 * j;
            k = S[j][(preS[temp + 0] << 5) +
                (preS[temp + 1] << 3) +
                (preS[temp + 2] << 2) +
                (preS[temp + 3] << 1) +
                (preS[temp + 4] << 0) +
                (preS[temp + 5] << 4)];

            temp = 4 * j;

            f[temp + 0] = (k >> 3) & 01;
            f[temp + 1] = (k >> 2) & 01;
            f[temp + 2] = (k >> 1) & 01;
            f[temp + 3] = (k >> 0) & 01;
        }

        /******************************************************************
        * The new right is left ^ f(R, K).
        * The f here has to be permuted first, though.
        ******************************************************************/
        for(j = 0; j < 32; j++)
            right[j] = left[j] ^ f[P[j] - 1];

        /* Finally, the new left (the original right) is copied back. */
        for(j = 0; j < 32; j++)
            left[j] = old[j];
    }

    /* The output left and right are reversed. */
    for(j = 0; j < 32; j++)
    {
        temp = left[j];
        left[j] = right[j];
        right[j] = temp;
    }

    /* The final output gets the inverse permutation of the very original. */
    for(j = 0; j < 64; j++)
    {
        i = FP[j];
        if (i < 33)
                block[j] = left[FP[j] - 1];
        else
                block[j] = right[FP[j] - 33];
    }
}

/**************************************************************************
* Function:    crypt
*
* Description: Clone of Unix crypt(3) function.
*
* Inputs:      char *pw
*              pointer to 8 character encryption key (user password)
*              char *salt
*              pointer to 2 character salt used to modify the DES results.
*
* Returns:     Pointer to static array containing the salt concatenated
*              on to the encrypted results.  Same as stored in passwd file.
**************************************************************************/
__global__ void crypt(char* pw, char* salt, char* trip) {
    int i, j, temp;
    char c,
         block[66];             /* 1st store key, then results */
    char iobuf[16];      /* encrypted results */

    for(i = 0; i < 66; i++)
        block[i] = 0;

    /* break pw into 64 bits */
    for(i = 0, c = *pw; c && (i < 64); i++)
    {
        for(j = 0; j < 7; j++, i++)
            block[i] = (c >> (6 - j)) & 01;
        pw++;
        c = *pw;
    }

    /* set key based on pw */
    setkey(block);

    for(i = 0; i < 66; i++)
        block[i] = 0;

    for(i = 0; i < 2; i++)
    {
        /* store salt at beginning of results */
        c = *salt++;
        iobuf[i] = c;

        if(c > 'Z')
            c -= 6;

        if(c > '9')
            c -= 7;

        c -= '.';

        /* use salt to effect the E-bit selection */
        for(j = 0; j < 6; j++)
        {
            if((c >> j) & 01)
            {
                temp = E[6 * i + j];
                E[6 * i +j] = E[6 * i + j + 24];
                E[6 * i + j + 24] = temp;
            }
        }
    }

    /* call DES encryption 25 times using pw as key and initial data = 0 */
    for(i = 0; i < 25; i++)
        encrypt(block);

    /* format encrypted block for standard crypt(3) output */
    for(i=0; i < 11; i++)
    {
        c = 0;
        for(j = 0; j < 6; j++)
        {
            c <<= 1;
            c |= block[6 * i + j];
        }

        c += '.';
        if(c > '9')
            c += 7;

        if(c > 'Z')
            c += 6;

        iobuf[i + 2] = c;
    }

    iobuf[i + 2] = '\0';

    /* prevent premature NULL terminator */
    if(iobuf[1] == '\0')
        iobuf[1] = iobuf[0];

    trip = iobuf;
    return;
}

	.file	"tmpxft_002f1bd5_00000000-6_Crypt.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2012:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2012:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z6setkeyPc
	.type	_Z6setkeyPc, @function
_Z6setkeyPc:
.LFB2008:
	.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
.LFE2008:
	.size	_Z6setkeyPc, .-_Z6setkeyPc
	.globl	_Z7encryptPc
	.type	_Z7encryptPc, @function
_Z7encryptPc:
.LFB2154:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_Z6setkeyPc
	.cfi_endproc
.LFE2154:
	.size	_Z7encryptPc, .-_Z7encryptPc
	.globl	_Z28__device_stub__Z5cryptPcS_S_PcS_S_
	.type	_Z28__device_stub__Z5cryptPcS_S_PcS_S_, @function
_Z28__device_stub__Z5cryptPcS_S_PcS_S_:
.LFB2034:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L6
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z5cryptPcS_S_(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L6:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L8
	call	__stack_chk_fail@PLT
.L8:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2034:
	.size	_Z28__device_stub__Z5cryptPcS_S_PcS_S_, .-_Z28__device_stub__Z5cryptPcS_S_PcS_S_
	.globl	_Z5cryptPcS_S_
	.type	_Z5cryptPcS_S_, @function
_Z5cryptPcS_S_:
.LFB2035:
	.cfi_startproc
	endbr64
	jmp	_Z28__device_stub__Z5cryptPcS_S_PcS_S_
	.cfi_endproc
.LFE2035:
	.size	_Z5cryptPcS_S_, .-_Z5cryptPcS_S_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z5cryptPcS_S_"
.LC1:
	.string	"IP"
.LC2:
	.string	"FP"
.LC3:
	.string	"PC1_C"
.LC4:
	.string	"PC1_D"
.LC5:
	.string	"shifts"
.LC6:
	.string	"PC2_C"
.LC7:
	.string	"PC2_D"
.LC8:
	.string	"C"
.LC9:
	.string	"D"
.LC10:
	.string	"KS"
.LC11:
	.string	"E"
.LC12:
	.string	"e2"
.LC13:
	.string	"S"
.LC14:
	.string	"P"
.LC15:
	.string	"preS"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2037:
	.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	.LC0(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z5cryptPcS_S_(%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
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC1(%rip), %rdx
	movl	$64, %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	.LC2(%rip), %rdx
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	movq	%rdx, %rcx
	movl	$64, %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	.LC3(%rip), %rdx
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	movq	%rdx, %rcx
	movl	$28, %r9d
	leaq	_ZL5PC1_C(%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	.LC4(%rip), %rdx
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	movq	%rdx, %rcx
	movl	$28, %r9d
	leaq	_ZL5PC1_D(%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	.LC5(%rip), %rdx
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	movq	%rdx, %rcx
	movl	$16, %r9d
	leaq	_ZL6shifts(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%r11
	.cfi_def_cfa_offset 24
	popq	%rax
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC6(%rip), %rdx
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	movq	%rdx, %rcx
	movl	$24, %r9d
	leaq	_ZL5PC2_C(%rip), %rsi
	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	.LC7(%rip), %rdx
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	movq	%rdx, %rcx
	movl	$24, %r9d
	leaq	_ZL5PC2_D(%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	.LC8(%rip), %rdx
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	movq	%rdx, %rcx
	movl	$28, %r9d
	leaq	_ZL1C(%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	.LC9(%rip), %rdx
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	movq	%rdx, %rcx
	movl	$28, %r9d
	leaq	_ZL1D(%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	.LC10(%rip), %rdx
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	movq	%rdx, %rcx
	movl	$768, %r9d
	leaq	_ZL2KS(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%r11
	.cfi_def_cfa_offset 24
	popq	%rax
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC11(%rip), %rdx
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	movq	%rdx, %rcx
	movl	$48, %r9d
	leaq	_ZL1E(%rip), %rsi
	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	.LC12(%rip), %rdx
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	movq	%rdx, %rcx
	movl	$48, %r9d
	leaq	_ZL2e2(%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	.LC13(%rip), %rdx
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	movq	%rdx, %rcx
	movl	$512, %r9d
	leaq	_ZL1S(%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	.LC14(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$32, %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	.LC15(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movl	$48, %r9d
	movq	%rdx, %rcx
	leaq	_ZL4preS(%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
.LFE2037:
	.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
	.local	_ZL4preS
	.comm	_ZL4preS,48,32
	.section	.rodata
	.align 32
	.type	_ZL1P, @object
	.size	_ZL1P, 32
_ZL1P:
	.ascii	"\020\007\024\025\035\f\034\021\001\017\027\032\005\022\037\n"
	.ascii	"\002\b\030\016 \033\003\t\023\r\036\006\026\013\004\031"
	.align 32
	.type	_ZL1S, @object
	.size	_ZL1S, 512
_ZL1S:
	.string	"\016\004\r\001\002\017\013\b\003\n\006\f\005\t"
	.string	"\007"
	.string	"\017\007\004\016\002\r\001\n\006\f\013\t\005\003\b\004\001\016\b\r\006\002\013\017\f\t\007\003\n\005"
	.string	"\017\f\b\002\004\t\001\007\005\013\003\016\n"
	.ascii	"\006\r"
	.string	"\017\001\b\016\006\013\003\004\t\007\002\r\f"
	.string	"\005\n\003\r\004\007\017\002\b\016\f"
	.string	"\001\n\006\t\013\005"
	.string	"\016\007\013\n\004\r\001\005\b\f\006\t\003\002\017\r\b\n\001\003\017\004\002\013\006\007\f"
	.ascii	"\005\016\t"
	.string	"\n"
	.string	"\t\016\006\003\017\005\001\r\f\007\013\004\002\b\r\007"
	.string	"\t\003\004\006\n\002\b\005\016\f\013\017\001\r\006\004\t\b\017\003"
	.string	"\013\001\002\f\005\n\016\007\001\n\r"
	.ascii	"\006\t\b\007\004\017\016\003\013\005\002\f"
	.string	"\007\r\016\003"
	.string	"\006\t\n\001\002\b\005\013\f\004\017\r\b\013\005\006\017"
	.string	"\003\004\007\002\f\001\n\016\t\n\006\t"
	.string	"\f\013\007\r\017\001\003\016\005\002\b\004\003\017"
	.ascii	"\006\n\001\r\b\t\004\005\013\f\007\002\016"
	.string	"\002\f\004\001\007\n\013\006\b\005\003\017\r"
	.string	"\016\t\016\013\002\f\004\007\r\001\005"
	.string	"\017\n\003\t\b\006\004\002\001\013\n\r\007\b\017\t\f\005\006\003"
	.string	"\016\013\b\f\007\001\016\002\r\006\017"
	.ascii	"\t\n\004\005\003"
	.string	"\f\001\n\017\t\002\006\b"
	.string	"\r\003\004\016\007\005\013\n\017\004\002\007\f\t\005\006\001\r\016"
	.string	"\013\003\b\t\016\017\005\002\b\f\003\007"
	.string	"\004\n\001\r\013\006\004\003\002\f\t\005\017\n\013\016\001\007\006"
	.ascii	"\b\r"
	.string	"\004\013\002\016\017"
	.string	"\b\r\003\f\t\007\005\n\006\001\r"
	.string	"\013\007\004\t\001\n\016\003\005\f\002\017\b\006\001\004\013\r\f\003\007\016\n\017\006\b"
	.string	"\005\t\002\006\013\r\b\001\004\n\007\t\005"
	.ascii	"\017\016\002\003\f"
	.string	"\r\002\b\004\006\017\013\001\n\t\003\016\005"
	.string	"\f\007\001\017\r\b\n\003\007\004\f\005\006\013"
	.string	"\016\t\002\007\013\004\001\t\f\016\002"
	.string	"\006\n\r\017\003\005\b\002\001\016\007\004\n\b\r\017\f\t"
	.ascii	"\003\005\006\013"
	.align 32
	.type	_ZL2e2, @object
	.size	_ZL2e2, 48
_ZL2e2:
	.ascii	" \001\002\003\004\005\004\005\006\007\b\t\b\t\n\013\f\r\f\r\016"
	.ascii	"\017\020\021\020\021\022\023\024\025\024\025\026\027\030\031"
	.ascii	"\030\031\032\033\034\035\034\035\036\037 \001"
	.local	_ZL1E
	.comm	_ZL1E,48,32
	.local	_ZL2KS
	.comm	_ZL2KS,768,32
	.local	_ZL1D
	.comm	_ZL1D,28,16
	.local	_ZL1C
	.comm	_ZL1C,28,16
	.align 16
	.type	_ZL5PC2_D, @object
	.size	_ZL5PC2_D, 24
_ZL5PC2_D:
	.ascii	")4\037%/7\036(3-!0,1'8\"5.*2$\035 "
	.align 16
	.type	_ZL5PC2_C, @object
	.size	_ZL5PC2_C, 24
_ZL5PC2_C:
	.ascii	"\016\021\013\030\001\005\003\034\017\006\025\n\027\023\f\004"
	.ascii	"\032\b\020\007\033\024\r\002"
	.align 16
	.type	_ZL6shifts, @object
	.size	_ZL6shifts, 16
_ZL6shifts:
	.ascii	"\001\001\002\002\002\002\002\002\001\002\002\002\002\002\002"
	.ascii	"\001"
	.align 16
	.type	_ZL5PC1_D, @object
	.size	_ZL5PC1_D, 28
_ZL5PC1_D:
	.ascii	"?7/'\037\027\017\007>6.&\036\026\016\006=5-%\035\025\r\005\034"
	.ascii	"\024\f\004"
	.align 16
	.type	_ZL5PC1_C, @object
	.size	_ZL5PC1_C, 28
_ZL5PC1_C:
	.ascii	"91)!\031\021\t\001:2*\"\032\022\n\002;3+#\033\023\013\003<4,"
	.ascii	"$"
	.align 32
	.type	_ZL2FP, @object
	.size	_ZL2FP, 64
_ZL2FP:
	.ascii	"(\b0\0208\030@ '\007/\0177\027?\037&\006.\0166\026>\036%\005"
	.ascii	"-\r5\025=\035$\004,\f4\024<\034#\003+\0133\023;\033\"\002*\n"
	.ascii	"2\022:\032!\001)\t1\0219\031"
	.align 32
	.type	_ZL2IP, @object
	.size	_ZL2IP, 64
_ZL2IP:
	.ascii	":2*\"\032\022\n\002<4,$\034\024\f\004>6.&\036\026\016\006@80"
	.ascii	"( \030\020\b91)!\031\021\t\001;3+#\033\023\013\003=5-%\035\025"
	.ascii	"\r\005?7/'\037\027\017\007"
	.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>
/**************************************************************************
*            Unix-like crypt(3) Algorithm for Password Encryption
*
*   File    : crypt3.c
*   Purpose : Provides crypt(3) functionality to ANSI C compilers
*             without a need for the crypt library.
*   Author  : Michael Dipperstein
*   Date    : November 3, 1998
*
***************************************************************************
*   The source in this file is heavily borrowed from the crypt3.c file
*   found on several ftp sites on the Internet.  The original source
*   claimed to be BSD, but was not distributed with any BSD license or
*   copyright claims.  I am releasing the source that I have provided into
*   public domain without any restrictions, warranties, or copyright
*   claims of my own.
*
*   The code below has been cleaned and compiles correctly under, gcc,
*   lcc, and Borland's bcc C compilers.  A bug involving the left and
*   right halves of the encrypted data block in the widely published
*   crypt3.c source has been fixed by this version.  All implicit register
*   declarations have been removed, because they generated suboptimal code.
*   All constant data has been explicitly declared as const and all
*   declarations have been given a minimal scope, because I'm paranoid.
*
*   Caution: crypt() returns a pointer to static data.  I left it this way
*            to maintain backward compatibility.  The downside is that
*            successive calls will cause previous results to be lost.
*            This can easily be changed with only minor modifications to
*            the function crypt().
**************************************************************************/

/* Initial permutation */
__device__ static const char IP[] =
{
    58, 50, 42, 34, 26, 18, 10, 2,
    60, 52, 44, 36, 28, 20, 12, 4,
    62, 54, 46, 38, 30, 22, 14, 6,
    64, 56, 48, 40, 32, 24, 16, 8,
    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,
};

/* Final permutation, FP = IP^(-1) */
__device__ static const char FP[] = {
    40, 8, 48, 16, 56, 24, 64, 32,
    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,
};

/**************************************************************************
* Permuted-choice 1 from the key bits to yield C and D.
* Note that bits 8,16... are left out:
* They are intended for a parity check.
**************************************************************************/
__device__ static const char PC1_C[] =
{
    57, 49, 41, 33, 25, 17,  9,
     1, 58, 50, 42, 34, 26, 18,
    10,  2, 59, 51, 43, 35, 27,
    19, 11,  3, 60, 52, 44, 36,
};

__device__ static const char PC1_D[] =
{
    63, 55, 47, 39, 31, 23, 15,
     7, 62, 54, 46, 38, 30, 22,
    14,  6, 61, 53, 45, 37, 29,
    21, 13,  5, 28, 20, 12,  4,
};

/* Sequence of shifts used for the key schedule. */
__device__ static const char shifts[] =
    {1, 1, 2, 2, 2, 2, 2, 2, 1, 2, 2, 2, 2, 2, 2, 1};

/**************************************************************************
* Permuted-choice 2, to pick out the bits from the CD array that generate
* the key schedule.
**************************************************************************/
__device__ static const char PC2_C[] =
{
    14, 17, 11, 24,  1,  5,
     3, 28, 15,  6, 21, 10,
    23, 19, 12,  4, 26,  8,
    16,  7, 27, 20, 13,  2,
};

__device__ static const char PC2_D[] =
{
    41, 52, 31, 37, 47, 55,
    30, 40, 51, 45, 33, 48,
    44, 49, 39, 56, 34, 53,
    46, 42, 50, 36, 29, 32,
};

/* The C and D arrays used to calculate the key schedule. */
__device__ static char C[28];
__device__ static char D[28];

/* The key schedule.  Generated from the key. */
__device__ static char KS[16][48];

/* The E bit-selection table. */
__device__ static char E[48];
__device__ static const char e2[] =
{
    32,  1,  2,  3,  4,  5,
     4,  5,  6,  7,  8,  9,
     8,  9, 10, 11, 12, 13,
    12, 13, 14, 15, 16, 17,
    16, 17, 18, 19, 20, 21,
    20, 21, 22, 23, 24, 25,
    24, 25, 26, 27, 28, 29,
    28, 29, 30, 31, 32,  1,
};

/**************************************************************************
* Function:    setkey
*
* Description: Set up the key schedule from the encryption key.
*
* Inputs:      char *key
*              pointer to 64 character array.  Each character represents a
*              bit in the key.
*
* Returns:     none
**************************************************************************/
__device__ void setkey(char *key)
{
    int i, j, k, temp;

    /**********************************************************************
    * First, generate C and D by permuting the key.  The low order bit of
    * each 8-bit char is not used, so C and D are only 28 bits apiece.
    **********************************************************************/
    for(i = 0; i < 28; i++)
    {
        C[i] = key[PC1_C[i] - 1];
        D[i] = key[PC1_D[i] - 1];
    }

    /**********************************************************************
    * To generate Ki, rotate C and D according to schedule and pick up a
    * permutation using PC2.
    **********************************************************************/
    for(i = 0; i < 16; i++)
    {
        /* rotate */
        for(k = 0; k < shifts[i]; k++)
        {
            temp = C[0];

            for(j = 0; j < 28 - 1; j++)
                C[j] = C[j+1];

            C[27] = temp;
            temp = D[0];
            for(j = 0; j < 28 - 1; j++)
                D[j] = D[j+1];

            D[27] = temp;
        }

        /* get Ki. Note C and D are concatenated */
        for(j = 0; j < 24; j++)
        {
            KS[i][j] = C[PC2_C[j] - 1];
            KS[i][j + 24] = D[PC2_D[j] - 28 -1];
        }
    }

    /* load E with the initial E bit selections */
    for(i=0; i < 48; i++)
        E[i] = e2[i];
}

/**************************************************************************
* The 8 selection functions. For some reason, they give a 0-origin
* index, unlike everything else.
**************************************************************************/

__device__ static const char S[8][64] =
{
    {
        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
    },

    {
        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
    },

    {
        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
    },

    {
         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
    },

    {
         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
    },

    {
        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
    },

    {
         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
    },

    {
        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
    }
};

/**************************************************************************
* P is a permutation on the selected combination of the current L and key.
**************************************************************************/
__device__ static const char P[] =
{
    16,  7, 20, 21,
    29, 12, 28, 17,
     1, 15, 23, 26,
     5, 18, 31, 10,
     2,  8, 24, 14,
    32, 27,  3,  9,
    19, 13, 30,  6,
    22, 11,  4, 25,
};

/* The combination of the key and the input, before selection. */
__device__ static char preS[48];

/**************************************************************************
* Function:    encrypt
*
* Description: Uses DES to encrypt a 64 bit block of data.  Requires
*              setkey to be invoked with the encryption key before it may
*              be used.  The results of the encryption are stored in block.
*
* Inputs:      char *block
*              pointer to 64 character array.  Each character represents a
*              bit in the data block.
*
* Returns:     none
**************************************************************************/
__device__ void encrypt(char *block)
{
    int i, ii, temp, j, k;

    char left[32], right[32]; /* block in two halves */
    char old[32];
    char f[32];

    /* First, permute the bits in the input */
    for(j = 0; j < 32; j++)
        left[j] = block[IP[j] - 1];

    for(;j < 64; j++)
        right[j - 32] = block[IP[j] - 1];

    /* Perform an encryption operation 16 times. */
    for(ii= 0; ii < 16; ii++)
    {
        i = ii;
        /* Save the right array, which will be the new left. */
        for(j = 0; j < 32; j++)
            old[j] = right[j];

        /******************************************************************
        * Expand right to 48 bits using the E selector and
        * exclusive-or with the current key bits.
        ******************************************************************/
        for(j =0 ; j < 48; j++)
            preS[j] = right[E[j] - 1] ^ KS[i][j];

        /******************************************************************
        * The pre-select bits are now considered in 8 groups of 6 bits ea.
        * The 8 selection functions map these 6-bit quantities into 4-bit
        * quantities and the results are permuted to make an f(R, K).
        * The indexing into the selection functions is peculiar;
        * it could be simplified by rewriting the tables.
        ******************************************************************/
        for(j = 0; j < 8; j++)
        {
            temp = 6 * j;
            k = S[j][(preS[temp + 0] << 5) +
                (preS[temp + 1] << 3) +
                (preS[temp + 2] << 2) +
                (preS[temp + 3] << 1) +
                (preS[temp + 4] << 0) +
                (preS[temp + 5] << 4)];

            temp = 4 * j;

            f[temp + 0] = (k >> 3) & 01;
            f[temp + 1] = (k >> 2) & 01;
            f[temp + 2] = (k >> 1) & 01;
            f[temp + 3] = (k >> 0) & 01;
        }

        /******************************************************************
        * The new right is left ^ f(R, K).
        * The f here has to be permuted first, though.
        ******************************************************************/
        for(j = 0; j < 32; j++)
            right[j] = left[j] ^ f[P[j] - 1];

        /* Finally, the new left (the original right) is copied back. */
        for(j = 0; j < 32; j++)
            left[j] = old[j];
    }

    /* The output left and right are reversed. */
    for(j = 0; j < 32; j++)
    {
        temp = left[j];
        left[j] = right[j];
        right[j] = temp;
    }

    /* The final output gets the inverse permutation of the very original. */
    for(j = 0; j < 64; j++)
    {
        i = FP[j];
        if (i < 33)
                block[j] = left[FP[j] - 1];
        else
                block[j] = right[FP[j] - 33];
    }
}

/**************************************************************************
* Function:    crypt
*
* Description: Clone of Unix crypt(3) function.
*
* Inputs:      char *pw
*              pointer to 8 character encryption key (user password)
*              char *salt
*              pointer to 2 character salt used to modify the DES results.
*
* Returns:     Pointer to static array containing the salt concatenated
*              on to the encrypted results.  Same as stored in passwd file.
**************************************************************************/
__global__ void crypt(char* pw, char* salt, char* trip) {
    int i, j, temp;
    char c,
         block[66];             /* 1st store key, then results */
    char iobuf[16];      /* encrypted results */

    for(i = 0; i < 66; i++)
        block[i] = 0;

    /* break pw into 64 bits */
    for(i = 0, c = *pw; c && (i < 64); i++)
    {
        for(j = 0; j < 7; j++, i++)
            block[i] = (c >> (6 - j)) & 01;
        pw++;
        c = *pw;
    }

    /* set key based on pw */
    setkey(block);

    for(i = 0; i < 66; i++)
        block[i] = 0;

    for(i = 0; i < 2; i++)
    {
        /* store salt at beginning of results */
        c = *salt++;
        iobuf[i] = c;

        if(c > 'Z')
            c -= 6;

        if(c > '9')
            c -= 7;

        c -= '.';

        /* use salt to effect the E-bit selection */
        for(j = 0; j < 6; j++)
        {
            if((c >> j) & 01)
            {
                temp = E[6 * i + j];
                E[6 * i +j] = E[6 * i + j + 24];
                E[6 * i + j + 24] = temp;
            }
        }
    }

    /* call DES encryption 25 times using pw as key and initial data = 0 */
    for(i = 0; i < 25; i++)
        encrypt(block);

    /* format encrypted block for standard crypt(3) output */
    for(i=0; i < 11; i++)
    {
        c = 0;
        for(j = 0; j < 6; j++)
        {
            c <<= 1;
            c |= block[6 * i + j];
        }

        c += '.';
        if(c > '9')
            c += 7;

        if(c > 'Z')
            c += 6;

        iobuf[i + 2] = c;
    }

    iobuf[i + 2] = '\0';

    /* prevent premature NULL terminator */
    if(iobuf[1] == '\0')
        iobuf[1] = iobuf[0];

    trip = iobuf;
    return;
}

	.text
	.file	"Crypt.hip"
	.globl	_Z20__device_stub__cryptPcS_S_  # -- Begin function _Z20__device_stub__cryptPcS_S_
	.type	_Z20__device_stub__cryptPcS_S_,@function
_Z20__device_stub__cryptPcS_S_:         # @_Z20__device_stub__cryptPcS_S_
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	24(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z5cryptPcS_S_, %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_end0:
	.size	_Z20__device_stub__cryptPcS_S_, .Lfunc_end0-_Z20__device_stub__cryptPcS_S_
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB1_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB1_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z5cryptPcS_S_, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end1:
	.size	__hip_module_ctor, .Lfunc_end1-__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	.LBB2_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
.LBB2_2:
	retq
.Lfunc_end2:
	.size	__hip_module_dtor, .Lfunc_end2-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z5cryptPcS_S_,@object          # @_Z5cryptPcS_S_
	.section	.rodata,"a",@progbits
	.globl	_Z5cryptPcS_S_
	.p2align	3, 0x0
_Z5cryptPcS_S_:
	.quad	_Z20__device_stub__cryptPcS_S_
	.size	_Z5cryptPcS_S_, 8

	.type	.L__unnamed_1,@object           # @0
	.section	.rodata.str1.1,"aMS",@progbits,1
.L__unnamed_1:
	.asciz	"_Z5cryptPcS_S_"
	.size	.L__unnamed_1, 15

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

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

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym _Z20__device_stub__cryptPcS_S_
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z5cryptPcS_S_
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z5cryptPcS_S_          ; -- Begin function _Z5cryptPcS_S_
	.globl	_Z5cryptPcS_S_
	.p2align	8
	.type	_Z5cryptPcS_S_,@function
_Z5cryptPcS_S_:                         ; @_Z5cryptPcS_S_
; %bb.0:
	v_mov_b32_e32 v0, 0
	s_mov_b32 s2, 0
.LBB0_1:                                ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_i32 s3, s2, 0x90
	s_add_i32 s2, s2, 1
	scratch_store_b8 off, v0, s3
	s_cmpk_eq_i32 s2, 0x42
	s_cbranch_scc0 .LBB0_1
; %bb.2:                                ; %.preheader90
	s_load_b128 s[0:3], s[0:1], 0x0
	v_mov_b32_e32 v1, 0
	s_mov_b32 s4, 0
	s_waitcnt lgkmcnt(0)
	global_load_u8 v0, v1, s[0:1]
	s_waitcnt vmcnt(0)
	v_cmp_eq_u16_e32 vcc_lo, 0, v0
	s_cbranch_vccnz .LBB0_6
.LBB0_3:                                ; %.preheader89
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_4 Depth 2
	v_bfe_i32 v0, v0, 0, 8
	s_mov_b32 s6, 6
.LBB0_4:                                ;   Parent Loop BB0_3 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_lshrrev_b32_e32 v2, s6, v0
	s_add_i32 s6, s6, -1
	s_mov_b32 s5, s4
	s_add_i32 s7, s4, 0x90
	s_add_i32 s4, s4, 1
	v_and_b32_e32 v2, 1, v2
	s_cmp_lg_u32 s6, -1
	scratch_store_b8 off, v2, s7
	s_cbranch_scc1 .LBB0_4
; %bb.5:                                ;   in Loop: Header=BB0_3 Depth=1
	global_load_u8 v0, v1, s[0:1] offset:1
	s_add_u32 s0, s0, 1
	s_addc_u32 s1, s1, 0
	s_add_i32 s4, s5, 2
	s_cmp_lt_i32 s5, 62
	s_cselect_b32 s5, -1, 0
	s_waitcnt vmcnt(0)
	v_cmp_ne_u16_e32 vcc_lo, 0, v0
	s_and_b32 s5, vcc_lo, s5
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_b32 vcc_lo, exec_lo, s5
	s_cbranch_vccnz .LBB0_3
.LBB0_6:                                ; %._crit_edge
	v_add_nc_u32_e64 v0, 0x90, -1
	s_mov_b64 s[0:1], 0
.LBB0_7:                                ; =>This Inner Loop Header: Depth=1
	s_getpc_b64 s[4:5]
	s_add_u32 s4, s4, _ZL5PC1_C@rel32@lo+4
	s_addc_u32 s5, s5, _ZL5PC1_C@rel32@hi+12
	s_add_u32 s4, s0, s4
	s_addc_u32 s5, s1, s5
	s_getpc_b64 s[6:7]
	s_add_u32 s6, s6, _ZL1C@rel32@lo+4
	s_addc_u32 s7, s7, _ZL1C@rel32@hi+12
	s_add_u32 s6, s0, s6
	s_addc_u32 s7, s1, s7
	s_getpc_b64 s[8:9]
	s_add_u32 s8, s8, _ZL5PC1_D@rel32@lo+4
	s_addc_u32 s9, s9, _ZL5PC1_D@rel32@hi+12
	s_add_u32 s8, s0, s8
	s_addc_u32 s9, s1, s9
	s_clause 0x1
	global_load_i8 v2, v1, s[4:5]
	global_load_i8 v3, v1, s[8:9]
	s_getpc_b64 s[4:5]
	s_add_u32 s4, s4, _ZL1D@rel32@lo+4
	s_addc_u32 s5, s5, _ZL1D@rel32@hi+12
	s_add_u32 s4, s0, s4
	s_addc_u32 s5, s1, s5
	s_add_u32 s0, s0, 1
	s_addc_u32 s1, s1, 0
	s_cmp_lg_u32 s0, 28
	s_waitcnt vmcnt(1)
	v_add_nc_u32_e32 v2, v0, v2
	s_waitcnt vmcnt(0)
	v_add_nc_u32_e32 v3, v0, v3
	s_clause 0x1
	scratch_load_u8 v2, v2, off
	scratch_load_u8 v3, v3, off
	s_waitcnt vmcnt(1)
	global_store_b8 v1, v2, s[6:7]
	s_waitcnt vmcnt(0)
	global_store_b8 v1, v3, s[4:5]
	s_cbranch_scc1 .LBB0_7
; %bb.8:                                ; %.preheader39.i.preheader
	v_mov_b32_e32 v1, 0
	s_mov_b32 s8, 0
	s_getpc_b64 s[0:1]
	s_add_u32 s0, s0, _ZL2KS@rel32@lo+4
	s_addc_u32 s1, s1, _ZL2KS@rel32@hi+12
	s_getpc_b64 s[14:15]
	s_add_u32 s14, s14, _ZL5PC2_D@rel32@lo+4
	s_addc_u32 s15, s15, _ZL5PC2_D@rel32@hi+12
.LBB0_9:                                ; %.preheader39.i
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_10 Depth 2
                                        ;       Child Loop BB0_11 Depth 3
                                        ;       Child Loop BB0_13 Depth 3
                                        ;     Child Loop BB0_16 Depth 2
	v_mov_b32_e32 v2, s8
	s_getpc_b64 s[4:5]
	s_add_u32 s4, s4, _ZL6shifts@rel32@lo+4
	s_addc_u32 s5, s5, _ZL6shifts@rel32@hi+12
	s_mov_b32 s9, 0
	global_load_i8 v2, v2, s[4:5]
	s_waitcnt vmcnt(0)
	v_max_i16 v2, v2, 1
	s_delay_alu instid0(VALU_DEP_1)
	v_and_b32_e32 v2, 0xffff, v2
.LBB0_10:                               ; %.lr.ph.i
                                        ;   Parent Loop BB0_9 Depth=1
                                        ; =>  This Loop Header: Depth=2
                                        ;       Child Loop BB0_11 Depth 3
                                        ;       Child Loop BB0_13 Depth 3
	s_getpc_b64 s[4:5]
	s_add_u32 s4, s4, _ZL1C@rel32@lo+4
	s_addc_u32 s5, s5, _ZL1C@rel32@hi+12
	s_mov_b64 s[6:7], 0
	global_load_u8 v4, v1, s[4:5]
.LBB0_11:                               ;   Parent Loop BB0_9 Depth=1
                                        ;     Parent Loop BB0_10 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	s_add_u32 s10, s6, s4
	s_addc_u32 s11, s7, s5
	s_getpc_b64 s[12:13]
	s_add_u32 s12, s12, _ZL1C@rel32@lo+5
	s_addc_u32 s13, s13, _ZL1C@rel32@hi+13
	s_add_u32 s12, s6, s12
	s_addc_u32 s13, s7, s13
	s_add_u32 s6, s6, 1
	global_load_u8 v3, v1, s[12:13]
	s_addc_u32 s7, s7, 0
	s_cmp_lg_u32 s6, 27
	s_waitcnt vmcnt(0)
	global_store_b8 v1, v3, s[10:11]
	s_cbranch_scc1 .LBB0_11
; %bb.12:                               ;   in Loop: Header=BB0_10 Depth=2
	s_getpc_b64 s[4:5]
	s_add_u32 s4, s4, _ZL1D@rel32@lo+4
	s_addc_u32 s5, s5, _ZL1D@rel32@hi+12
	s_getpc_b64 s[10:11]
	s_add_u32 s10, s10, _ZL1C@rel32@lo+31
	s_addc_u32 s11, s11, _ZL1C@rel32@hi+39
	global_load_u8 v3, v1, s[4:5]
	s_mov_b64 s[6:7], 0
	global_store_b8 v1, v4, s[10:11]
.LBB0_13:                               ;   Parent Loop BB0_9 Depth=1
                                        ;     Parent Loop BB0_10 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	s_add_u32 s10, s6, s4
	s_addc_u32 s11, s7, s5
	s_getpc_b64 s[12:13]
	s_add_u32 s12, s12, _ZL1D@rel32@lo+5
	s_addc_u32 s13, s13, _ZL1D@rel32@hi+13
	s_add_u32 s12, s6, s12
	s_addc_u32 s13, s7, s13
	s_add_u32 s6, s6, 1
	global_load_u8 v4, v1, s[12:13]
	s_addc_u32 s7, s7, 0
	s_cmp_lg_u32 s6, 27
	s_waitcnt vmcnt(0)
	global_store_b8 v1, v4, s[10:11]
	s_cbranch_scc1 .LBB0_13
; %bb.14:                               ;   in Loop: Header=BB0_10 Depth=2
	s_add_i32 s9, s9, 1
	s_getpc_b64 s[4:5]
	s_add_u32 s4, s4, _ZL1D@rel32@lo+31
	s_addc_u32 s5, s5, _ZL1D@rel32@hi+39
	v_cmp_ne_u32_e32 vcc_lo, s9, v2
	global_store_b8 v1, v3, s[4:5]
	s_cbranch_vccnz .LBB0_10
; %bb.15:                               ; %.preheader38.i.preheader
                                        ;   in Loop: Header=BB0_9 Depth=1
	s_mov_b64 s[4:5], 0
.LBB0_16:                               ; %.preheader38.i
                                        ;   Parent Loop BB0_9 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_getpc_b64 s[6:7]
	s_add_u32 s6, s6, _ZL5PC2_C@rel32@lo+4
	s_addc_u32 s7, s7, _ZL5PC2_C@rel32@hi+12
	s_add_u32 s6, s4, s6
	s_addc_u32 s7, s5, s7
	s_getpc_b64 s[10:11]
	s_add_u32 s10, s10, _ZL1C@rel32@lo+3
	s_addc_u32 s11, s11, _ZL1C@rel32@hi+11
	s_add_u32 s12, s0, s4
	s_addc_u32 s13, s1, s5
	global_load_i8 v2, v1, s[6:7]
	s_add_u32 s6, s4, s14
	s_addc_u32 s7, s5, s15
	global_load_i8 v3, v1, s[6:7]
	s_getpc_b64 s[6:7]
	s_add_u32 s6, s6, _ZL1D@rel32@lo-25
	s_addc_u32 s7, s7, _ZL1D@rel32@hi-17
	s_add_u32 s4, s4, 1
	s_addc_u32 s5, s5, 0
	s_cmp_lg_u32 s4, 24
	s_waitcnt vmcnt(1)
	v_bfe_i32 v2, v2, 0, 16
	s_waitcnt vmcnt(0)
	v_bfe_i32 v4, v3, 0, 16
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_ashrrev_i32_e32 v3, 31, v2
	v_add_co_u32 v2, vcc_lo, v2, s10
	v_ashrrev_i32_e32 v5, 31, v4
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e32 v3, vcc_lo, s11, v3, vcc_lo
	v_add_co_u32 v4, vcc_lo, v4, s6
	v_add_co_ci_u32_e32 v5, vcc_lo, s7, v5, vcc_lo
	global_load_u8 v2, v[2:3], off
	global_load_u8 v3, v[4:5], off
	s_waitcnt vmcnt(1)
	global_store_b8 v1, v2, s[12:13]
	s_waitcnt vmcnt(0)
	global_store_b8 v1, v3, s[12:13] offset:24
	s_cbranch_scc1 .LBB0_16
; %bb.17:                               ;   in Loop: Header=BB0_9 Depth=1
	s_add_i32 s8, s8, 1
	s_add_u32 s0, s0, 48
	s_addc_u32 s1, s1, 0
	s_cmp_lg_u32 s8, 16
	s_cbranch_scc1 .LBB0_9
; %bb.18:                               ; %.preheader.i.preheader
	v_mov_b32_e32 v1, 0
	s_mov_b64 s[0:1], 0
.LBB0_19:                               ; %.preheader.i
                                        ; =>This Inner Loop Header: Depth=1
	s_getpc_b64 s[4:5]
	s_add_u32 s4, s4, _ZL2e2@rel32@lo+4
	s_addc_u32 s5, s5, _ZL2e2@rel32@hi+12
	s_add_u32 s4, s0, s4
	s_addc_u32 s5, s1, s5
	global_load_u8 v2, v1, s[4:5]
	s_getpc_b64 s[4:5]
	s_add_u32 s4, s4, _ZL1E@rel32@lo+4
	s_addc_u32 s5, s5, _ZL1E@rel32@hi+12
	s_add_u32 s4, s0, s4
	s_addc_u32 s5, s1, s5
	s_add_u32 s0, s0, 1
	s_addc_u32 s1, s1, 0
	s_cmp_lg_u32 s0, 48
	s_waitcnt vmcnt(0)
	global_store_b8 v1, v2, s[4:5]
	s_cbranch_scc1 .LBB0_19
; %bb.20:                               ; %_Z6setkeyPc.exit.preheader
	v_mov_b32_e32 v1, 0
	s_mov_b32 s0, 0
.LBB0_21:                               ; %_Z6setkeyPc.exit
                                        ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_i32 s1, s0, 0x90
	s_add_i32 s0, s0, 1
	scratch_store_b8 off, v1, s1
	s_cmpk_lg_i32 s0, 0x42
	s_cbranch_scc1 .LBB0_21
; %bb.22:                               ; %.preheader88.preheader
	v_mov_b32_e32 v1, 0
	s_mov_b32 s0, 0
	s_mov_b32 s1, -1
.LBB0_23:                               ; %.preheader88
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_24 Depth 2
	global_load_i8 v2, v1, s[2:3]
	s_xor_b32 s4, s1, -1
	s_waitcnt vmcnt(0)
	v_readfirstlane_b32 s5, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	s_add_i32 s1, s5, 0xfa
	s_cmpk_gt_i32 s5, 0x5a
	s_sext_i32_i8 s1, s1
	s_cselect_b32 s1, s1, s5
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_i32 s5, s1, 0xf9
	s_cmp_gt_i32 s1, 57
	s_cselect_b32 s1, s5, s1
	s_getpc_b64 s[6:7]
	s_add_u32 s6, s6, _ZL1E@rel32@lo+4
	s_addc_u32 s7, s7, _ZL1E@rel32@hi+12
	s_addk_i32 s1, 0xd2
	s_add_u32 s6, s0, s6
	s_sext_i32_i8 s5, s1
	s_addc_u32 s7, 0, s7
	s_mov_b64 s[0:1], 0
.LBB0_24:                               ;   Parent Loop BB0_23 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_lshr_b32 s8, s5, s0
	s_bitcmp0_b32 s8, 0
	s_cbranch_scc1 .LBB0_26
; %bb.25:                               ;   in Loop: Header=BB0_24 Depth=2
	s_add_u32 s8, s6, s0
	s_addc_u32 s9, s7, s1
	s_clause 0x1
	global_load_u8 v2, v1, s[8:9] offset:24
	global_load_u8 v3, v1, s[8:9]
	s_waitcnt vmcnt(1)
	global_store_b8 v1, v2, s[8:9]
	s_waitcnt vmcnt(0)
	global_store_b8 v1, v3, s[8:9] offset:24
.LBB0_26:                               ;   in Loop: Header=BB0_24 Depth=2
	s_add_u32 s0, s0, 1
	s_addc_u32 s1, s1, 0
	s_cmp_lg_u32 s0, 6
	s_cbranch_scc1 .LBB0_24
; %bb.27:                               ;   in Loop: Header=BB0_23 Depth=1
	s_add_u32 s2, s2, 1
	s_addc_u32 s3, s3, 0
	s_mov_b32 s0, 6
	s_and_not1_b32 vcc_lo, exec_lo, s4
	s_mov_b32 s1, 0
	s_cbranch_vccnz .LBB0_23
; %bb.28:                               ; %.preheader87.preheader
	v_add_nc_u32_e64 v1, 48, -1
	v_add_nc_u32_e64 v2, 0x70, -1
	v_or_b32_e64 v3, 0x70, 1
	v_dual_mov_b32 v4, 0 :: v_dual_mov_b32 v5, 16
	s_mov_b32 s2, 0xaaaaaaaa
	s_mov_b32 s10, 0
.LBB0_29:                               ; %.preheader87
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_30 Depth 2
                                        ;     Child Loop BB0_32 Depth 2
                                        ;     Child Loop BB0_34 Depth 2
                                        ;       Child Loop BB0_35 Depth 3
                                        ;       Child Loop BB0_37 Depth 3
                                        ;       Child Loop BB0_39 Depth 3
                                        ;       Child Loop BB0_41 Depth 3
                                        ;       Child Loop BB0_43 Depth 3
                                        ;     Child Loop BB0_46 Depth 2
                                        ;     Child Loop BB0_48 Depth 2
	s_mov_b32 s0, s1
	s_getpc_b64 s[4:5]
	s_add_u32 s4, s4, _ZL2IP@rel32@lo+4
	s_addc_u32 s5, s5, _ZL2IP@rel32@hi+12
.LBB0_30:                               ;   Parent Loop BB0_29 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_i8 v6, v4, s[4:5]
	s_add_i32 s3, s0, 16
	s_add_i32 s0, s0, 1
	s_add_u32 s4, s4, 1
	s_addc_u32 s5, s5, 0
	s_cmp_lg_u32 s0, 32
	s_waitcnt vmcnt(0)
	v_add_nc_u32_e32 v6, v0, v6
	scratch_load_u8 v6, v6, off
	s_waitcnt vmcnt(0)
	scratch_store_b8 off, v6, s3
	s_cbranch_scc1 .LBB0_30
; %bb.31:                               ; %.lr.ph.i80.preheader
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_mov_b32 s0, 0
	s_getpc_b64 s[4:5]
	s_add_u32 s4, s4, _ZL2IP@rel32@lo+36
	s_addc_u32 s5, s5, _ZL2IP@rel32@hi+44
.LBB0_32:                               ; %.lr.ph.i80
                                        ;   Parent Loop BB0_29 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_i8 v6, v4, s[4:5]
	s_add_i32 s3, s0, 48
	s_add_i32 s0, s0, 1
	s_add_u32 s4, s4, 1
	s_addc_u32 s5, s5, 0
	s_cmp_lg_u32 s0, 32
	s_waitcnt vmcnt(0)
	v_add_nc_u32_e32 v6, v0, v6
	scratch_load_u8 v6, v6, off
	s_waitcnt vmcnt(0)
	scratch_store_b8 off, v6, s3
	s_cbranch_scc1 .LBB0_32
; %bb.33:                               ; %.preheader75.i.preheader
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_mov_b32 s0, 0
	s_getpc_b64 s[4:5]
	s_add_u32 s4, s4, _ZL2KS@rel32@lo+4
	s_addc_u32 s5, s5, _ZL2KS@rel32@hi+12
.LBB0_34:                               ; %.preheader75.i
                                        ;   Parent Loop BB0_29 Depth=1
                                        ; =>  This Loop Header: Depth=2
                                        ;       Child Loop BB0_35 Depth 3
                                        ;       Child Loop BB0_37 Depth 3
                                        ;       Child Loop BB0_39 Depth 3
                                        ;       Child Loop BB0_41 Depth 3
                                        ;       Child Loop BB0_43 Depth 3
	s_mov_b32 s3, 0
.LBB0_35:                               ;   Parent Loop BB0_29 Depth=1
                                        ;     Parent Loop BB0_34 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_3) | instid1(SALU_CYCLE_1)
	s_add_i32 s6, s3, 48
	scratch_load_u8 v6, off, s6
	s_add_i32 s6, s3, 0x50
	s_add_i32 s3, s3, 1
	s_cmp_eq_u32 s3, 32
	s_waitcnt vmcnt(0)
	scratch_store_b8 off, v6, s6
	s_cbranch_scc0 .LBB0_35
; %bb.36:                               ; %.preheader74.i
                                        ;   in Loop: Header=BB0_34 Depth=2
	s_mov_b64 s[6:7], 0
.LBB0_37:                               ;   Parent Loop BB0_29 Depth=1
                                        ;     Parent Loop BB0_34 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	s_getpc_b64 s[8:9]
	s_add_u32 s8, s8, _ZL1E@rel32@lo+4
	s_addc_u32 s9, s9, _ZL1E@rel32@hi+12
	s_add_u32 s8, s6, s8
	s_addc_u32 s9, s7, s9
	global_load_i8 v6, v4, s[8:9]
	s_add_u32 s8, s4, s6
	s_addc_u32 s9, s5, s7
	global_load_u8 v7, v4, s[8:9]
	s_getpc_b64 s[8:9]
	s_add_u32 s8, s8, _ZL4preS@rel32@lo+4
	s_addc_u32 s9, s9, _ZL4preS@rel32@hi+12
	s_add_u32 s8, s6, s8
	s_addc_u32 s9, s7, s9
	s_add_u32 s6, s6, 1
	s_addc_u32 s7, s7, 0
	s_cmp_lg_u32 s6, 48
	s_waitcnt vmcnt(1)
	v_add_nc_u32_e32 v6, v1, v6
	scratch_load_u8 v6, v6, off
	s_waitcnt vmcnt(0)
	v_xor_b32_e32 v6, v7, v6
	global_store_b8 v4, v6, s[8:9]
	s_cbranch_scc1 .LBB0_37
; %bb.38:                               ; %.preheader73.i.preheader
                                        ;   in Loop: Header=BB0_34 Depth=2
	v_mov_b32_e32 v6, v3
	s_mov_b64 s[8:9], 0
	s_getpc_b64 s[6:7]
	s_add_u32 s6, s6, _ZL1S@rel32@lo+4
	s_addc_u32 s7, s7, _ZL1S@rel32@hi+12
.LBB0_39:                               ; %.preheader73.i
                                        ;   Parent Loop BB0_29 Depth=1
                                        ;     Parent Loop BB0_34 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	s_getpc_b64 s[12:13]
	s_add_u32 s12, s12, _ZL4preS@rel32@lo+4
	s_addc_u32 s13, s13, _ZL4preS@rel32@hi+12
	s_add_u32 s12, s8, s12
	v_mov_b32_e32 v7, 0
	s_addc_u32 s13, s9, s13
	s_getpc_b64 s[14:15]
	s_add_u32 s14, s14, _ZL4preS@rel32@lo+8
	s_addc_u32 s15, s15, _ZL4preS@rel32@hi+16
	s_add_u32 s14, s8, s14
	s_addc_u32 s15, s9, s15
	s_clause 0x1
	global_load_b32 v8, v4, s[12:13]
	global_load_d16_b16 v7, v7, s[14:15]
	s_add_u32 s8, s8, 6
	s_addc_u32 s9, s9, 0
	s_waitcnt vmcnt(1)
	v_bfe_i32 v9, v8, 0, 8
	v_bfe_i32 v10, v8, 8, 8
	v_bfe_i32 v11, v8, 16, 8
	v_ashrrev_i32_e32 v8, 24, v8
	s_waitcnt vmcnt(0)
	v_bfe_i32 v12, v7, 0, 8
	v_lshlrev_b32_e32 v9, 5, v9
	v_lshlrev_b32_e32 v10, 3, v10
	v_lshlrev_b32_e32 v11, 2, v11
	v_lshlrev_b32_e32 v8, 1, v8
	v_bfe_i32 v7, v7, 8, 8
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v9, v10, v9
	v_lshlrev_b32_e32 v7, 4, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v9, v9, v11
	v_add_nc_u32_e32 v8, v9, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v8, v8, v12
	v_add_nc_u32_e32 v7, v8, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v8, 31, v7
	v_add_co_u32 v7, vcc_lo, s6, v7
	s_add_u32 s6, s6, 64
	v_add_co_ci_u32_e32 v8, vcc_lo, s7, v8, vcc_lo
	s_addc_u32 s7, s7, 0
	s_cmp_lg_u32 s8, 48
	global_load_u8 v7, v[7:8], off
	s_waitcnt vmcnt(0)
	v_lshrrev_b16 v8, 3, v7
	v_lshlrev_b16 v9, 6, v7
	v_lshrrev_b16 v10, 1, v7
	v_and_b32_e32 v7, 1, v7
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_or_b32_e32 v8, v8, v9
	v_and_b32_e32 v9, 1, v10
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshlrev_b16 v7, 8, v7
	v_and_b32_e32 v8, 0x101, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_or_b32_e32 v7, v9, v7
	v_and_b32_e32 v8, 0xffff, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b32_e32 v7, 16, v7
	v_or_b32_e32 v7, v8, v7
	scratch_store_b32 v6, v7, off offset:-1
	v_add_nc_u32_e32 v6, 4, v6
	s_cbranch_scc1 .LBB0_39
; %bb.40:                               ; %.preheader72.i.preheader
                                        ;   in Loop: Header=BB0_34 Depth=2
	s_mov_b32 s3, 0
	s_getpc_b64 s[6:7]
	s_add_u32 s6, s6, _ZL1P@rel32@lo+4
	s_addc_u32 s7, s7, _ZL1P@rel32@hi+12
.LBB0_41:                               ; %.preheader72.i
                                        ;   Parent Loop BB0_29 Depth=1
                                        ;     Parent Loop BB0_34 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	global_load_i8 v6, v4, s[6:7]
	s_add_i32 s8, s3, 16
	s_waitcnt vmcnt(0)
	v_add_nc_u32_e32 v6, v2, v6
	scratch_load_u8 v7, off, s8
	scratch_load_u8 v6, v6, off
	s_add_i32 s8, s3, 48
	s_add_i32 s3, s3, 1
	s_add_u32 s6, s6, 1
	s_addc_u32 s7, s7, 0
	s_cmp_lg_u32 s3, 32
	s_waitcnt vmcnt(0)
	v_xor_b32_e32 v6, v6, v7
	scratch_store_b8 off, v6, s8
	s_cbranch_scc1 .LBB0_41
; %bb.42:                               ; %.preheader71.i.preheader
                                        ;   in Loop: Header=BB0_34 Depth=2
	s_mov_b32 s3, 0
.LBB0_43:                               ; %.preheader71.i
                                        ;   Parent Loop BB0_29 Depth=1
                                        ;     Parent Loop BB0_34 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_3) | instid1(SALU_CYCLE_1)
	s_add_i32 s6, s3, 0x50
	scratch_load_u8 v6, off, s6
	s_add_i32 s6, s3, 16
	s_add_i32 s3, s3, 1
	s_cmp_lg_u32 s3, 32
	s_waitcnt vmcnt(0)
	scratch_store_b8 off, v6, s6
	s_cbranch_scc1 .LBB0_43
; %bb.44:                               ;   in Loop: Header=BB0_34 Depth=2
	s_add_i32 s0, s0, 1
	s_add_u32 s4, s4, 48
	s_addc_u32 s5, s5, 0
	s_cmp_lg_u32 s0, 16
	s_cbranch_scc1 .LBB0_34
; %bb.45:                               ; %.preheader70.i.preheader
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_mov_b32 s0, 0
.LBB0_46:                               ; %.preheader70.i
                                        ;   Parent Loop BB0_29 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_i32 s3, s0, 48
	s_add_i32 s4, s0, 16
	scratch_load_u8 v6, off, s3
	scratch_load_u8 v7, off, s4
	s_add_i32 s0, s0, 1
	s_waitcnt vmcnt(1)
	scratch_store_b8 off, v6, s4
	s_waitcnt vmcnt(0)
	scratch_store_b8 off, v7, s3
	s_cmp_lg_u32 s0, 32
	s_cbranch_scc1 .LBB0_46
; %bb.47:                               ; %.preheader.i81.preheader
                                        ;   in Loop: Header=BB0_29 Depth=1
	v_mov_b32_e32 v6, 0x90
	s_mov_b64 s[4:5], 0
.LBB0_48:                               ; %.preheader.i81
                                        ;   Parent Loop BB0_29 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_getpc_b64 s[6:7]
	s_add_u32 s6, s6, _ZL2FP@rel32@lo+4
	s_addc_u32 s7, s7, _ZL2FP@rel32@hi+12
	s_add_u32 s6, s4, s6
	s_addc_u32 s7, s5, s7
	s_mov_b32 s3, s2
	global_load_i8 v7, v4, s[6:7]
	s_lshr_b64 s[6:7], s[2:3], s4
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s0, s6, 1
	s_cmp_eq_u64 s[0:1], 0
	s_cselect_b32 s0, 0xffffffdf, -1
	s_cselect_b32 s3, -1, 0
	s_add_u32 s4, s4, 1
	v_cndmask_b32_e64 v8, v5, 48, s3
	s_addc_u32 s5, s5, 0
	s_cmp_lg_u32 s4, 64
	s_waitcnt vmcnt(0)
	v_add_nc_u32_e32 v7, s0, v7
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v7, v8, v7
	scratch_load_u8 v7, v7, off
	s_waitcnt vmcnt(0)
	scratch_store_b8 v6, v7, off
	v_add_nc_u32_e32 v6, 1, v6
	s_cbranch_scc1 .LBB0_48
; %bb.49:                               ; %_Z7encryptPc.exit
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_add_i32 s10, s10, 1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_lg_u32 s10, 25
	s_cbranch_scc1 .LBB0_29
; %bb.50:                               ; %.preheader.preheader
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z5cryptPcS_S_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 224
		.amdhsa_kernarg_size 24
		.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 1
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 13
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z5cryptPcS_S_, .Lfunc_end0-_Z5cryptPcS_S_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 2608
; NumSgprs: 18
; NumVgprs: 13
; ScratchSize: 224
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 13
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 1
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	_ZL5PC1_C,@object               ; @_ZL5PC1_C
	.section	.rodata,"a",@progbits
	.p2align	4, 0x0
_ZL5PC1_C:
	.ascii	"91)!\031\021\t\001:2*\"\032\022\n\002;3+#\033\023\013\003<4,$"
	.size	_ZL5PC1_C, 28

	.type	_ZL1C,@object                   ; @_ZL1C
	.local	_ZL1C
	.comm	_ZL1C,28,4
	.type	_ZL5PC1_D,@object               ; @_ZL5PC1_D
	.p2align	4, 0x0
_ZL5PC1_D:
	.ascii	"?7/'\037\027\017\007>6.&\036\026\016\006=5-%\035\025\r\005\034\024\f\004"
	.size	_ZL5PC1_D, 28

	.type	_ZL1D,@object                   ; @_ZL1D
	.local	_ZL1D
	.comm	_ZL1D,28,4
	.type	_ZL6shifts,@object              ; @_ZL6shifts
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0
_ZL6shifts:
	.ascii	"\001\001\002\002\002\002\002\002\001\002\002\002\002\002\002\001"
	.size	_ZL6shifts, 16

	.type	_ZL5PC2_C,@object               ; @_ZL5PC2_C
	.section	.rodata,"a",@progbits
	.p2align	4, 0x0
_ZL5PC2_C:
	.ascii	"\016\021\013\030\001\005\003\034\017\006\025\n\027\023\f\004\032\b\020\007\033\024\r\002"
	.size	_ZL5PC2_C, 24

	.type	_ZL2KS,@object                  ; @_ZL2KS
	.local	_ZL2KS
	.comm	_ZL2KS,768,4
	.type	_ZL5PC2_D,@object               ; @_ZL5PC2_D
	.p2align	4, 0x0
_ZL5PC2_D:
	.ascii	")4\037%/7\036(3-!0,1'8\"5.*2$\035 "
	.size	_ZL5PC2_D, 24

	.type	_ZL2e2,@object                  ; @_ZL2e2
	.p2align	4, 0x0
_ZL2e2:
	.ascii	" \001\002\003\004\005\004\005\006\007\b\t\b\t\n\013\f\r\f\r\016\017\020\021\020\021\022\023\024\025\024\025\026\027\030\031\030\031\032\033\034\035\034\035\036\037 \001"
	.size	_ZL2e2, 48

	.type	_ZL1E,@object                   ; @_ZL1E
	.local	_ZL1E
	.comm	_ZL1E,48,4
	.type	_ZL2IP,@object                  ; @_ZL2IP
	.p2align	4, 0x0
_ZL2IP:
	.ascii	":2*\"\032\022\n\002<4,$\034\024\f\004>6.&\036\026\016\006@80( \030\020\b91)!\031\021\t\001;3+#\033\023\013\003=5-%\035\025\r\005?7/'\037\027\017\007"
	.size	_ZL2IP, 64

	.type	_ZL4preS,@object                ; @_ZL4preS
	.local	_ZL4preS
	.comm	_ZL4preS,48,4
	.type	_ZL1S,@object                   ; @_ZL1S
	.p2align	4, 0x0
_ZL1S:
	.ascii	"\016\004\r\001\002\017\013\b\003\n\006\f\005\t\000\007\000\017\007\004\016\002\r\001\n\006\f\013\t\005\003\b\004\001\016\b\r\006\002\013\017\f\t\007\003\n\005\000\017\f\b\002\004\t\001\007\005\013\003\016\n\000\006\r"
	.ascii	"\017\001\b\016\006\013\003\004\t\007\002\r\f\000\005\n\003\r\004\007\017\002\b\016\f\000\001\n\006\t\013\005\000\016\007\013\n\004\r\001\005\b\f\006\t\003\002\017\r\b\n\001\003\017\004\002\013\006\007\f\000\005\016\t"
	.ascii	"\n\000\t\016\006\003\017\005\001\r\f\007\013\004\002\b\r\007\000\t\003\004\006\n\002\b\005\016\f\013\017\001\r\006\004\t\b\017\003\000\013\001\002\f\005\n\016\007\001\n\r\000\006\t\b\007\004\017\016\003\013\005\002\f"
	.ascii	"\007\r\016\003\000\006\t\n\001\002\b\005\013\f\004\017\r\b\013\005\006\017\000\003\004\007\002\f\001\n\016\t\n\006\t\000\f\013\007\r\017\001\003\016\005\002\b\004\003\017\000\006\n\001\r\b\t\004\005\013\f\007\002\016"
	.ascii	"\002\f\004\001\007\n\013\006\b\005\003\017\r\000\016\t\016\013\002\f\004\007\r\001\005\000\017\n\003\t\b\006\004\002\001\013\n\r\007\b\017\t\f\005\006\003\000\016\013\b\f\007\001\016\002\r\006\017\000\t\n\004\005\003"
	.ascii	"\f\001\n\017\t\002\006\b\000\r\003\004\016\007\005\013\n\017\004\002\007\f\t\005\006\001\r\016\000\013\003\b\t\016\017\005\002\b\f\003\007\000\004\n\001\r\013\006\004\003\002\f\t\005\017\n\013\016\001\007\006\000\b\r"
	.ascii	"\004\013\002\016\017\000\b\r\003\f\t\007\005\n\006\001\r\000\013\007\004\t\001\n\016\003\005\f\002\017\b\006\001\004\013\r\f\003\007\016\n\017\006\b\000\005\t\002\006\013\r\b\001\004\n\007\t\005\000\017\016\002\003\f"
	.ascii	"\r\002\b\004\006\017\013\001\n\t\003\016\005\000\f\007\001\017\r\b\n\003\007\004\f\005\006\013\000\016\t\002\007\013\004\001\t\f\016\002\000\006\n\r\017\003\005\b\002\001\016\007\004\n\b\r\017\f\t\000\003\005\006\013"
	.size	_ZL1S, 512

	.type	_ZL1P,@object                   ; @_ZL1P
	.section	.rodata.cst32,"aM",@progbits,32
	.p2align	4, 0x0
_ZL1P:
	.ascii	"\020\007\024\025\035\f\034\021\001\017\027\032\005\022\037\n\002\b\030\016 \033\003\t\023\r\036\006\026\013\004\031"
	.size	_ZL1P, 32

	.type	_ZL2FP,@object                  ; @_ZL2FP
	.section	.rodata,"a",@progbits
	.p2align	4, 0x0
_ZL2FP:
	.ascii	"(\b0\0208\030@ '\007/\0177\027?\037&\006.\0166\026>\036%\005-\r5\025=\035$\004,\f4\024<\034#\003+\0133\023;\033\"\002*\n2\022:\032!\001)\t1\0219\031"
	.size	_ZL2FP, 64

	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 24
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z5cryptPcS_S_
    .private_segment_fixed_size: 224
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z5cryptPcS_S_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     13
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
