#include<bits/stdc++.h>

using namespace std;

typedef unsigned long long ull;
#define MAX 1000000 //10 e 6

ull LnRnBlocks[17*2]; // from l0r0 to l16r16

ull CnDnBlocks[17*2]; //from c0d0 to c16d16

ull keysBlocks[16];  //from key[1] = k0 to key[16] = k15

ull allCipherDES[MAX];

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

int PC1[56] = {
   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,
   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
};

int PC2[48] = {
   14, 17, 11, 24,  1,  5,
    3, 28, 15,  6, 21, 10,
   23, 19, 12,  4, 26,  8,
   16,  7, 27, 20, 13,  2,
   41, 52, 31, 37, 47, 55,
   30, 40, 51, 45, 33, 48,
   44, 49, 39, 56, 34, 53,
   46, 42, 50, 36, 29, 32
};

int IniPer[64] = {
    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
};

int reverseIniPer[64] ={
    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
};

int Expansion[48] ={
    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
};

int Pbox[32] ={
    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
};

int Sbox[512] = { //8*4*16
    
    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
 };

ull iniKey[8] = {
    0x13,0x34,0x57,0x79,0x9B,0xBC,0xDF,0xF1};

ull message[8] = {    
    0x01,0x23,0x45,0x67,0x89,0xAB,0xCD,0xEF};



__host__ __device__ ull generateKeyPlus(ull *d_iniKey, int *d_PC1){
    ull keyPlus=0L;
    for(int i=56-1;i>=0;i--){
        if( d_iniKey[ d_PC1[i]/8 ] & (1 << ( ( 64-d_PC1[i]) % 8 ) ) ){
            keyPlus|=( 1LL<< (55-i)*1L );
        }
    }
    return keyPlus;
}

__host__ __device__ ull* splitKeyPlus(ull keyPlus){
    ull c0=0L, d0=0L;

    for(int i=0;i<28;i++){
        if(keyPlus & (1LL<<i*1L ) ) d0|=(1LL<<i*1L);
        if(keyPlus & (1LL<< i*1L +28L ) ) c0|=(1LL<<i*1L);
    }
    ull new_array[2]= {c0,d0};
    return new_array;
}

__host__ __device__ ull* splitIniPer(ull codeIniPer){
    ull l0=0L, r0=0L;

    for(int i=0;i<32;i++){
        if(codeIniPer & (1LL<<i*1L ) ) r0|=(1LL<<i*1L);
        if(codeIniPer & (1LL<< i*1L +32L ) ) l0|=(1LL<<i*1L);
    }
    ull new_array[2] ={ l0,r0};
    return new_array;
}

__host__ __device__ void generateCnDnBlocks( ull *seedKey,ull *d_CnDnBlocks, ull *d_Rotations){
    d_CnDnBlocks[0] = seedKey[0];
    d_CnDnBlocks[1] = seedKey[1];
    ull cn ,dn, newCn, newDn;
    ull getOnCn, getOnDn;

    for(int i=1;i<=16;i++){
        
        getOnCn = 0L;
        getOnDn = 0L;

        cn = d_CnDnBlocks[(i*2+0)-2];
        dn = d_CnDnBlocks[(i*2+1)-2];

        for(ull j=0;j< d_Rotations[ i-1 ];j++){

            if(  cn & (1 << (27-j) )  ) getOnCn|= 1LL << (d_Rotations[ i-1 ]==1 ? j: 1-j) ;
            if(  dn & (1 << (27-j) )  ) getOnDn|= 1LL << (d_Rotations[ i-1 ]==1 ? j: 1-j);

        }

        newCn = cn << d_Rotations[ i-1 ];
        newDn = dn << d_Rotations[ i-1 ];

        for(ull j=0; j< d_Rotations[ i-1 ] ;j++){
            newCn &= ~(1<< (28+j) );
            newDn &= ~(1<< (28+j) );
        }
        newCn |= ( getOnCn );
        newDn |= ( getOnDn );
        d_CnDnBlocks[ (i*2)+0 ] = newCn;
        d_CnDnBlocks[ (i*2)+1 ] = newDn;
    }

}

__host__ __device__ ull joinCnDn(ull cn, ull dn){ return (cn<<28) | dn; }

__host__ __device__ void generateKeysBlocks(ull *d_CnDnBlocks, int *d_PC2,ull *d_keysBlocks){
    ull cnDn, keyn;

    for(int i=1;i<=16;i++){
        cnDn = joinCnDn( d_CnDnBlocks[i*2+0], d_CnDnBlocks[i*2+1] );
        keyn = 0L;

        for(int j=48-1;j>=0;j--){
            if( cnDn & ( 1LL << (56-d_PC2[j])*1L ) ) {
                 keyn|= ( 1LL<< (47-j)*1L );
            }
        }
        d_keysBlocks[i-1] = keyn;
    }

}

__host__ __device__ ull generateIniPer(int *d_IniPer, ull *d_message){
    ull keyPlus=0L;

    for(int i=64-1;i>=0;i--){
        if( d_message[ (d_IniPer[i]/8) >=8 ? 7: (d_IniPer[i]/8) ]  & (1LL << ( ( 64-d_IniPer[i]) % 8 ) ) ){
            keyPlus|=( 1LL<< (63-i)*1L );
        }
    }
    
    return keyPlus;
}

__host__ __device__ ull expandRn(ull Rn,int *d_Expansion){
    //from a Rn 32 bit to a Kn 48 bit
    ull exRn=0L;
    for(int j=48-1;j>=0;j--){
            if( Rn & ( 1LL << (32-d_Expansion[j])*1L ) ) {
                 exRn|= ( 1LL<< (47-j)*1L );
            }
    }
    return exRn;
}

__host__ __device__ ull xorOperation(ull En, ull Kn){
    return (Kn ^ En);
}

__host__ __device__ ull getSboxNumber(int Bn, int k, int *d_Sbox){

    int row=0,col=0;
    if( Bn & 1<<0 ) row |= ( 1<<0);
    if( Bn & 1<<6-1) row |=( 1<<1 );

    for(int i=1;i<=4;i++){
        if( Bn & 1<<i ) col |=(1<<(i-1));
    }

    return d_Sbox[ (k*4*16)+(row*16)+col ];
}


__host__ __device__ ull generateSboxCombination(ull Bn,int *d_Sbox){

    int Bbox[8];
    ull sbBox[8];
    int number=0, k=7;
    ull snBnChain=0L;
    ull step=28L;

    for(int i=0;i<=48;i++ ){

        if( i%6==0 && i>=6){
            

            Bbox[ k-- ] = number;
            number = 0;
        }
        if( Bn & (1LL<<i*1L) ){
            number |= (1LL<< ( (i%6)*1L ) );
        }
    }

    for(int i=0;i<8;i++){
        sbBox[i] = getSboxNumber( Bbox[i], i,d_Sbox);

       
    }

    
    for(int i=0;i<8;i++){
        snBnChain |= (sbBox[i]<< step);
        step-=4;
    }

    return snBnChain;

}

__host__ __device__ ull generateFalgorithm(ull snBn, int *d_Pbox){

    ull fn=0L;
    for(int j=32-1;j>=0;j--){
            if( snBn & ( 1LL << (32-d_Pbox[j])*1L ) ) {
                 fn|= ( 1LL<< (31-j)*1L );
            }
    }
    return fn;
}

__host__ __device__ void generateLnRnBlocks(ull *L0R0,ull *d_LnRnBlocks, ull *d_keysBlocks,int *d_Expansion, int *d_Sbox,int *d_Pbox){

    d_LnRnBlocks[0] = L0R0[0];
    d_LnRnBlocks[1] = L0R0[1];
    ull fn;

    for(int time=1; time<=16;time++){

        ull Ln = d_LnRnBlocks[ (time*2+0)-2 ];
        ull Rn = d_LnRnBlocks[ (time*2+1)-2 ];
       
        ull snBn = 
            generateSboxCombination( xorOperation( expandRn( Rn, d_Expansion ),d_keysBlocks[ time-1 ] ),d_Sbox );
        
        fn = generateFalgorithm(snBn,d_Pbox);

        d_LnRnBlocks[ (time*2+0) ] = Rn;
        d_LnRnBlocks[ (time*2+1) ] = (Ln ^ fn);

    }

}

__host__ __device__ ull reverseLnRn( ull *LnRn, ull *LnRn1){
    ull Ln = *LnRn;
    ull Rn = *LnRn1;

    return ( Rn<<32L) | Ln;
}

__host__ __device__ ull generateCipherMessage( ull RnLn, int *d_reverseIniPer ){

    ull cipher=0L;
    for(int j=64-1;j>=0;j--){
            if( RnLn & ( 1LL << (64-d_reverseIniPer[j])*1L ) ) {
                 cipher|= ( 1LL<< (63-j)*1L );
            }
    }
    return cipher;
}


__global__ void cipherDES(
    ull *d_LnRnBlocks,
    ull *d_CnDnBlocks,
    ull *d_keysBlocks,
    ull *d_allCipherDES,
    ull *d_Rotations,
    int *d_PC1,
    int *d_PC2,
    int *d_IniPer,
    int *d_reverseIniPer,
    int *d_Expansion,
    int *d_Pbox,
    int *d_Sbox,
    ull *d_iniKey,
    ull *d_message,
    ull *d_result
){
    ull *keyHalves = splitKeyPlus( generateKeyPlus(d_iniKey,d_PC1) );
    generateCnDnBlocks( keyHalves,d_CnDnBlocks,d_Rotations );
    generateKeysBlocks(d_CnDnBlocks,d_PC2,d_keysBlocks);
    ull *iniPerHalves = splitIniPer(generateIniPer(d_IniPer,d_message) ); //got L0 and R0
    

    generateLnRnBlocks( iniPerHalves,d_LnRnBlocks, d_keysBlocks, d_Expansion, d_Sbox, d_Pbox);

    ull revLnRn = reverseLnRn( &d_LnRnBlocks[16*2+0],&d_LnRnBlocks[16*2+1] );
   
    ull cipherMessage = generateCipherMessage( revLnRn,d_reverseIniPer );
    
    int iindex = threadIdx.x + blockIdx.x * blockDim.x;
    d_result[iindex] = cipherMessage;
    // printf("Index %i - Hex Cipher: %llX\n", iindex, cipherMessage);
   
}

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

    int THREADS_BLOCK;
    
    sscanf (argv[1], "%i", &THREADS_BLOCK);
    

    //host and devices copies
    ull *d_LnRnBlocks;//17 size
    ull *d_CnDnBlocks;//17 size
    ull *d_keysBlocks;//16 size
    ull *d_allCipherDES;//10^6 size
    ull *d_Rotations;//16 size
    int *d_PC1;//56 size
    int *d_PC2;//48 size
    int *d_IniPer;//64 size
    int *d_reverseIniPer;//64 size
    int *d_Expansion;//48 size
    int *d_Pbox;//32 size
    int *d_Sbox;//8*8*16 size [8][4][16]
    ull *d_iniKey;//8 size
    ull *d_message;//8 size
    ull *result; //MAX SIZE
    ull *d_result;//MAX SIZE

    //size of host and device copies
    int sd_LnRnBlocks = 17 * 2 * sizeof(ull);
    int sd_CnDnBlocks = 17 * 2 * sizeof(ull);
    int sd_keysBlocks = 16 * sizeof(ull);
    int sd_allCipherDES = 1000000 * sizeof(ull);
    int sd_Rotations = 16 * sizeof(ull);
    int sd_PC1 = 56 * sizeof(int);
    int sd_PC2 = 48 * sizeof(int);
    int sd_IniPer = 64 * sizeof(int);
    int sd_reverseIniPer = 64 * sizeof(int);
    int sd_Expansion = 48 * sizeof(int);
    int sd_Pbox = 32 * sizeof(int);
    int sd_Sbox = 512 * sizeof(int);
    int sd_iniKey = 8 * sizeof(ull);
    int sd_message = 8 * sizeof(ull);
    int sd_result = MAX * sizeof(ull);

    //alloc space for device copies
    cudaMalloc( (void **)&d_LnRnBlocks, sd_LnRnBlocks );
    cudaMalloc( (void **)&d_CnDnBlocks, sd_CnDnBlocks );
    cudaMalloc( (void **)&d_keysBlocks, sd_keysBlocks );
    cudaMalloc( (void **)&d_allCipherDES, sd_allCipherDES );
    cudaMalloc( (void **)&d_Rotations, sd_Rotations );
    cudaMalloc( (void **)&d_PC1, sd_PC1 );
    cudaMalloc( (void **)&d_PC2, sd_PC2 );
    cudaMalloc( (void **)&d_IniPer, sd_IniPer );
    cudaMalloc( (void **)&d_reverseIniPer, sd_reverseIniPer );
    cudaMalloc( (void **)&d_Expansion, sd_Expansion );
    cudaMalloc( (void **)&d_Pbox, sd_Pbox );
    cudaMalloc( (void **)&d_Sbox, sd_Sbox );
    cudaMalloc( (void **)&d_iniKey, sd_iniKey );
    cudaMalloc( (void **)&d_message, sd_message );
    cudaMalloc( (void **)&d_result, sd_result );

    //alloc space for host copies
    result = (ull *) malloc(sd_result);

    //copy inputs to device
    cudaMemcpy(d_LnRnBlocks,LnRnBlocks,sd_LnRnBlocks,cudaMemcpyHostToDevice);
    cudaMemcpy(d_CnDnBlocks,CnDnBlocks,sd_CnDnBlocks,cudaMemcpyHostToDevice);
    cudaMemcpy(d_keysBlocks,keysBlocks,sd_keysBlocks,cudaMemcpyHostToDevice);
    cudaMemcpy(d_allCipherDES,allCipherDES,sd_allCipherDES,cudaMemcpyHostToDevice);
    cudaMemcpy(d_Rotations,Rotations,sd_Rotations,cudaMemcpyHostToDevice);
    cudaMemcpy(d_PC1,PC1,sd_PC1,cudaMemcpyHostToDevice);
    cudaMemcpy(d_PC2,PC2,sd_PC2,cudaMemcpyHostToDevice);
    cudaMemcpy(d_IniPer,IniPer,sd_IniPer,cudaMemcpyHostToDevice);
    cudaMemcpy(d_reverseIniPer,reverseIniPer,sd_reverseIniPer,cudaMemcpyHostToDevice);
    cudaMemcpy(d_Expansion,Expansion,sd_Expansion,cudaMemcpyHostToDevice);
    cudaMemcpy(d_Pbox,Pbox,sd_Pbox,cudaMemcpyHostToDevice);
    cudaMemcpy(d_Sbox,Sbox,sd_Sbox,cudaMemcpyHostToDevice);
    cudaMemcpy(d_iniKey,iniKey,sd_iniKey,cudaMemcpyHostToDevice);
    cudaMemcpy(d_message,message,sd_message,cudaMemcpyHostToDevice);
   


    int global_work_size = ceil(MAX/THREADS_BLOCK);
    int local_work_size = THREADS_BLOCK;
    if(THREADS_BLOCK > 1000 ){
         global_work_size = MAX - (MAX%THREADS_BLOCK);
         global_work_size = ceil( global_work_size/THREADS_BLOCK);
         local_work_size = THREADS_BLOCK;
    }
    //launch kernel     ----------- HERES THE MAGIC ---------
    cipherDES<<<global_work_size,local_work_size>>>(
        d_LnRnBlocks,
        d_CnDnBlocks,
        d_keysBlocks,
        d_allCipherDES,
        d_Rotations,
        d_PC1,
        d_PC2,
        d_IniPer,
        d_reverseIniPer,
        d_Expansion,
        d_Pbox,
        d_Sbox,
        d_iniKey,
        d_message,
        d_result
    );


    //copy result back to host
    cudaMemcpy(result,d_result,sd_result,cudaMemcpyDeviceToHost);

    //print result

    // for(int i=0;i<MAX;i++){
    //     printf("cipher: %i\t%llX\n",i,result[i]);
    //     fflush(stdout);
    // }


    //free host space
    free(result);

    //free cuda space
    cudaFree(d_LnRnBlocks);
    cudaFree(d_CnDnBlocks);
    cudaFree(d_keysBlocks);
    cudaFree(d_allCipherDES);
    cudaFree(d_Rotations);
    cudaFree(d_PC1);
    cudaFree(d_PC2);
    cudaFree(d_IniPer);
    cudaFree(d_reverseIniPer);
    cudaFree(d_Expansion);
    cudaFree(d_Pbox);
    cudaFree(d_Sbox);
    cudaFree(d_iniKey);
    cudaFree(d_message);

    return 0;
}


// L16 01000011010000100011001000110100    1128411700
// R16 00001010010011001101100110010101    172808597

// revL16R16 0000101001001100110110011001010101000011010000100011001000110100  742207273711055412

// cipher 1000010111101000000100110101010000001111000010101011010000000101     9648983453391827973
	.file	"tmpxft_002cb171_00000000-6_cctest.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB10843:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE10843:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z15generateKeyPlusPyPi
	.type	_Z15generateKeyPlusPyPi, @function
_Z15generateKeyPlusPyPi:
.LFB10825:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rdi, %r8
	movq	%rsi, %r9
	xorl	%edi, %edi
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movl	$55, %esi
	movl	$55, %ebx
	movl	$8, %r10d
	movl	$1, %r11d
	movl	$1, %ebp
.L4:
	movl	(%r9,%rsi,4), %eax
	movl	(%r9,%rsi,4), %ecx
	cltd
	negl	%ecx
	idivl	%r10d
	andl	$7, %ecx
	movslq	%eax, %rdx
	movl	%r11d, %eax
	sall	%cl, %eax
	cltq
	andq	(%r8,%rdx,8), %rax
	je	.L3
	movl	%ebx, %ecx
	movq	%rbp, %rax
	subl	%esi, %ecx
	salq	%cl, %rax
	orq	%rax, %rdi
.L3:
	subq	$1, %rsi
	jnb	.L4
	movq	%rdi, %rax
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE10825:
	.size	_Z15generateKeyPlusPyPi, .-_Z15generateKeyPlusPyPi
	.globl	_Z12splitKeyPlusy
	.type	_Z12splitKeyPlusy, @function
_Z12splitKeyPlusy:
.LFB10826:
	.cfi_startproc
	endbr64
	xorl	%eax, %eax
	ret
	.cfi_endproc
.LFE10826:
	.size	_Z12splitKeyPlusy, .-_Z12splitKeyPlusy
	.globl	_Z11splitIniPery
	.type	_Z11splitIniPery, @function
_Z11splitIniPery:
.LFB12559:
	.cfi_startproc
	endbr64
	xorl	%eax, %eax
	ret
	.cfi_endproc
.LFE12559:
	.size	_Z11splitIniPery, .-_Z11splitIniPery
	.globl	_Z18generateCnDnBlocksPyS_S_
	.type	_Z18generateCnDnBlocksPyS_S_, @function
_Z18generateCnDnBlocksPyS_S_:
.LFB10828:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	xorl	%r8d, %r8d
	movl	$1, %r9d
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movl	$27, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rdx, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	(%rdi), %rax
	movq	%rax, (%rsi)
	movq	8(%rdi), %rax
	movq	%rax, 8(%rsi)
.L24:
	movq	(%rsi), %rdx
	movq	8(%rsi), %rax
	xorl	%edi, %edi
	xorl	%r11d, %r11d
	movq	0(%rbp,%r8,8), %r10
	xorl	%ebx, %ebx
.L14:
	cmpq	%rdi, %r10
	je	.L33
	movl	%r13d, %ecx
	movl	%r9d, %r15d
	movl	%edi, %r14d
	subl	%edi, %ecx
	sall	%cl, %r15d
	movslq	%r15d, %r15
	testq	%rdx, %r15
	je	.L15
	movl	%edi, %ecx
	movq	%r10, %r12
	salq	%cl, %r12
	movq	%r12, %rcx
	cmpq	$1, %r10
	je	.L17
	movl	%r9d, %ecx
	movl	$1, %r12d
	subl	%edi, %ecx
	salq	%cl, %r12
	movq	%r12, %rcx
.L17:
	orq	%rcx, %rbx
.L15:
	testq	%rax, %r15
	je	.L18
	movl	%r14d, %ecx
	movq	%r10, %r15
	salq	%cl, %r15
	movq	%r15, %rcx
	cmpq	$1, %r10
	je	.L20
	movl	%r9d, %ecx
	movl	$1, %r12d
	subl	%r14d, %ecx
	salq	%cl, %r12
	movq	%r12, %rcx
.L20:
	orq	%rcx, %r11
.L18:
	incq	%rdi
	jmp	.L14
.L33:
	movl	%edi, %ecx
	xorl	%r10d, %r10d
	salq	%cl, %rdx
	salq	%cl, %rax
.L22:
	cmpq	%r10, %rdi
	je	.L34
	leal	28(%r10), %ecx
	movl	%r9d, %r14d
	incq	%r10
	sall	%cl, %r14d
	movl	%r14d, %ecx
	notl	%ecx
	movslq	%ecx, %rcx
	andq	%rcx, %rdx
	andq	%rcx, %rax
	jmp	.L22
.L34:
	orq	%rbx, %rdx
	orq	%r11, %rax
	incq	%r8
	addq	$16, %rsi
	movq	%rdx, (%rsi)
	movq	%rax, 8(%rsi)
	cmpq	$16, %r8
	jne	.L24
	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
.LFE10828:
	.size	_Z18generateCnDnBlocksPyS_S_, .-_Z18generateCnDnBlocksPyS_S_
	.globl	_Z8joinCnDnyy
	.type	_Z8joinCnDnyy, @function
_Z8joinCnDnyy:
.LFB10829:
	.cfi_startproc
	endbr64
	salq	$28, %rdi
	movq	%rdi, %rax
	orq	%rsi, %rax
	ret
	.cfi_endproc
.LFE10829:
	.size	_Z8joinCnDnyy, .-_Z8joinCnDnyy
	.globl	_Z18generateKeysBlocksPyPiS_
	.type	_Z18generateKeysBlocksPyPiS_, @function
_Z18generateKeysBlocksPyPiS_:
.LFB10830:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	leaq	16(%rdi), %rax
	movq	%rsi, %r10
	movq	%rdx, %r8
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	xorl	%edi, %edi
	movl	$1, %r11d
	movl	$47, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movl	$56, %ebx
.L39:
	movq	(%rax), %rsi
	movl	$47, %edx
	xorl	%r9d, %r9d
	salq	$28, %rsi
	orq	8(%rax), %rsi
.L38:
	movl	%ebx, %ecx
	movq	%r11, %r14
	subl	(%r10,%rdx,4), %ecx
	salq	%cl, %r14
	testq	%rsi, %r14
	je	.L37
	movl	%ebp, %ecx
	movq	%r11, %r14
	subl	%edx, %ecx
	salq	%cl, %r14
	orq	%r14, %r9
.L37:
	subq	$1, %rdx
	jnb	.L38
	movq	%r9, (%r8,%rdi,8)
	incq	%rdi
	addq	$16, %rax
	cmpq	$16, %rdi
	jne	.L39
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE10830:
	.size	_Z18generateKeysBlocksPyPiS_, .-_Z18generateKeysBlocksPyPiS_
	.globl	_Z14generateIniPerPiPy
	.type	_Z14generateIniPerPiPy, @function
_Z14generateIniPerPiPy:
.LFB10831:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movl	$1, %ebx
	movq	%rdi, %r8
	movq	%rsi, %r9
	xorl	%edi, %edi
	movl	$63, %esi
	movl	$8, %r10d
	salq	$63, %rbx
	movl	$1, %r11d
.L49:
	movl	(%r8,%rsi,4), %ecx
	movl	$56, %eax
	cmpl	$63, %ecx
	jg	.L47
	movl	%ecx, %eax
	cltd
	idivl	%r10d
	cltq
	salq	$3, %rax
.L47:
	negl	%ecx
	movq	%r11, %rdx
	andl	$7, %ecx
	salq	%cl, %rdx
	andq	(%r9,%rax), %rdx
	je	.L48
	movq	%rbx, %rax
	movl	%esi, %ecx
	shrq	%cl, %rax
	orq	%rax, %rdi
.L48:
	subq	$1, %rsi
	jnb	.L49
	movq	%rdi, %rax
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE10831:
	.size	_Z14generateIniPerPiPy, .-_Z14generateIniPerPiPy
	.globl	_Z8expandRnyPi
	.type	_Z8expandRnyPi, @function
_Z8expandRnyPi:
.LFB10832:
	.cfi_startproc
	endbr64
	movl	$47, %eax
	xorl	%edx, %edx
	movl	$32, %r9d
	movl	$1, %r8d
	movl	$47, %r10d
.L58:
	movl	%r9d, %ecx
	movq	%r8, %r11
	subl	(%rsi,%rax,4), %ecx
	salq	%cl, %r11
	testq	%rdi, %r11
	je	.L57
	movl	%r10d, %ecx
	movq	%r8, %r11
	subl	%eax, %ecx
	salq	%cl, %r11
	orq	%r11, %rdx
.L57:
	subq	$1, %rax
	jnb	.L58
	movq	%rdx, %rax
	ret
	.cfi_endproc
.LFE10832:
	.size	_Z8expandRnyPi, .-_Z8expandRnyPi
	.globl	_Z12xorOperationyy
	.type	_Z12xorOperationyy, @function
_Z12xorOperationyy:
.LFB10833:
	.cfi_startproc
	endbr64
	movq	%rsi, %rax
	xorq	%rdi, %rax
	ret
	.cfi_endproc
.LFE10833:
	.size	_Z12xorOperationyy, .-_Z12xorOperationyy
	.globl	_Z13getSboxNumberiiPi
	.type	_Z13getSboxNumberiiPi, @function
_Z13getSboxNumberiiPi:
.LFB10834:
	.cfi_startproc
	endbr64
	movl	%edi, %r9d
	andl	$1, %r9d
	testl	$32, %edi
	je	.L65
	orl	$2, %r9d
.L65:
	movl	$1, %eax
	xorl	%r8d, %r8d
	movl	$1, %r10d
.L67:
	btl	%eax, %edi
	jnc	.L66
	leal	-1(%rax), %ecx
	movl	%r10d, %r11d
	sall	%cl, %r11d
	orl	%r11d, %r8d
.L66:
	incl	%eax
	cmpl	$5, %eax
	jne	.L67
	leal	(%r9,%rsi,4), %eax
	sall	$4, %eax
	addl	%r8d, %eax
	cltq
	movslq	(%rdx,%rax,4), %rax
	ret
	.cfi_endproc
.LFE10834:
	.size	_Z13getSboxNumberiiPi, .-_Z13getSboxNumberiiPi
	.globl	_Z23generateSboxCombinationyPi
	.type	_Z23generateSboxCombinationyPi, @function
_Z23generateSboxCombinationyPi:
.LFB10835:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rdi, %r9
	movq	%rsi, %rbp
	movl	$7, %r8d
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	xorl	%esi, %esi
	xorl	%edi, %edi
	movl	$6, %r10d
	movl	$1, %r11d
	subq	$120, %rsp
	.cfi_def_cfa_offset 144
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
.L78:
	movl	%esi, %eax
	cltd
	idivl	%r10d
	testl	%edx, %edx
	jne	.L76
	cmpl	$5, %esi
	jle	.L76
	movslq	%r8d, %rax
	decl	%r8d
	movl	%edi, 8(%rsp,%rax,4)
	xorl	%edi, %edi
.L76:
	movq	%r11, %rax
	movl	%esi, %ecx
	salq	%cl, %rax
	testq	%r9, %rax
	je	.L77
	movq	%r11, %rax
	movl	%edx, %ecx
	salq	%cl, %rax
	orl	%eax, %edi
.L77:
	incl	%esi
	cmpl	$49, %esi
	jne	.L78
	xorl	%ebx, %ebx
.L79:
	movl	8(%rsp,%rbx,4), %edi
	movq	%rbp, %rdx
	movl	%ebx, %esi
	call	_Z13getSboxNumberiiPi
	movq	%rax, %rdx
	leaq	40(%rsp), %rax
	movq	%rdx, (%rax,%rbx,8)
	incq	%rbx
	cmpq	$8, %rbx
	jne	.L79
	movl	$28, %ecx
	xorl	%edx, %edx
.L80:
	movq	(%rax), %rsi
	addq	$8, %rax
	salq	%cl, %rsi
	subq	$4, %rcx
	orq	%rsi, %rdx
	cmpq	$-4, %rcx
	jne	.L80
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L81
	call	__stack_chk_fail@PLT
.L81:
	addq	$120, %rsp
	.cfi_def_cfa_offset 24
	movq	%rdx, %rax
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE10835:
	.size	_Z23generateSboxCombinationyPi, .-_Z23generateSboxCombinationyPi
	.globl	_Z18generateFalgorithmyPi
	.type	_Z18generateFalgorithmyPi, @function
_Z18generateFalgorithmyPi:
.LFB10836:
	.cfi_startproc
	endbr64
	movl	$31, %eax
	xorl	%edx, %edx
	movl	$32, %r9d
	movl	$1, %r8d
	movl	$31, %r10d
.L97:
	movl	%r9d, %ecx
	movq	%r8, %r11
	subl	(%rsi,%rax,4), %ecx
	salq	%cl, %r11
	testq	%rdi, %r11
	je	.L96
	movl	%r10d, %ecx
	movq	%r8, %r11
	subl	%eax, %ecx
	salq	%cl, %r11
	orq	%r11, %rdx
.L96:
	subq	$1, %rax
	jnb	.L97
	movq	%rdx, %rax
	ret
	.cfi_endproc
.LFE10836:
	.size	_Z18generateFalgorithmyPi, .-_Z18generateFalgorithmyPi
	.globl	_Z18generateLnRnBlocksPyS_S_PiS0_S0_
	.type	_Z18generateLnRnBlocksPyS_S_PiS0_S0_, @function
_Z18generateLnRnBlocksPyS_S_PiS0_S0_:
.LFB10837:
	.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
	movq	%rcx, %r14
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	xorl	%r12d, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rsi, %rbx
	subq	$40, %rsp
	.cfi_def_cfa_offset 96
	movq	(%rdi), %rcx
	movq	%r8, 8(%rsp)
	movq	%rcx, (%rsi)
	movq	8(%rdi), %rax
	movq	%r9, 16(%rsp)
	movq	%rax, 8(%rsi)
.L103:
	movq	8(%rbx), %r13
	movq	(%rdx,%r12,8), %r15
	movq	%r14, %rsi
	movq	%rdx, 24(%rsp)
	movq	(%rbx), %rbp
	incq	%r12
	addq	$16, %rbx
	movq	%r13, %rdi
	call	_Z8expandRnyPi
	movq	8(%rsp), %rsi
	movq	%r15, %rdi
	xorq	%rax, %rdi
	call	_Z23generateSboxCombinationyPi
	movq	16(%rsp), %rsi
	movq	%rax, %rdi
	call	_Z18generateFalgorithmyPi
	movq	%r13, (%rbx)
	movq	24(%rsp), %rdx
	xorq	%rax, %rbp
	movq	%rbp, 8(%rbx)
	cmpq	$16, %r12
	jne	.L103
	addq	$40, %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
.LFE10837:
	.size	_Z18generateLnRnBlocksPyS_S_PiS0_S0_, .-_Z18generateLnRnBlocksPyS_S_PiS0_S0_
	.globl	_Z11reverseLnRnPyS_
	.type	_Z11reverseLnRnPyS_, @function
_Z11reverseLnRnPyS_:
.LFB10838:
	.cfi_startproc
	endbr64
	movq	(%rsi), %rax
	salq	$32, %rax
	orq	(%rdi), %rax
	ret
	.cfi_endproc
.LFE10838:
	.size	_Z11reverseLnRnPyS_, .-_Z11reverseLnRnPyS_
	.globl	_Z21generateCipherMessageyPi
	.type	_Z21generateCipherMessageyPi, @function
_Z21generateCipherMessageyPi:
.LFB10839:
	.cfi_startproc
	endbr64
	movl	$63, %eax
	xorl	%edx, %edx
	movl	$64, %r8d
	movl	$1, %r9d
	movabsq	$-9223372036854775808, %r10
.L109:
	movl	%r8d, %ecx
	movq	%r9, %r11
	subl	(%rsi,%rax,4), %ecx
	salq	%cl, %r11
	testq	%rdi, %r11
	je	.L108
	movq	%r10, %r11
	movl	%eax, %ecx
	shrq	%cl, %r11
	orq	%r11, %rdx
.L108:
	subq	$1, %rax
	jnb	.L109
	movq	%rdx, %rax
	ret
	.cfi_endproc
.LFE10839:
	.size	_Z21generateCipherMessageyPi, .-_Z21generateCipherMessageyPi
	.globl	_Z62__device_stub__Z9cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_PyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_
	.type	_Z62__device_stub__Z9cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_PyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_, @function
_Z62__device_stub__Z9cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_PyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_:
.LFB10865:
	.cfi_startproc
	endbr64
	subq	$328, %rsp
	.cfi_def_cfa_offset 336
	movq	336(%rsp), %rax
	movq	%rdi, 120(%rsp)
	leaq	144(%rsp), %rdi
	movq	%rsi, 112(%rsp)
	leaq	156(%rsp), %rsi
	movq	%rax, 72(%rsp)
	movq	344(%rsp), %rax
	movq	%rdx, 104(%rsp)
	leaq	128(%rsp), %rdx
	movq	%rax, 64(%rsp)
	movq	352(%rsp), %rax
	movq	%rcx, 96(%rsp)
	leaq	136(%rsp), %rcx
	movq	%rax, 56(%rsp)
	movq	360(%rsp), %rax
	movq	%r8, 88(%rsp)
	movq	%rax, 48(%rsp)
	movq	368(%rsp), %rax
	movq	%r9, 80(%rsp)
	movq	%rax, 40(%rsp)
	movq	376(%rsp), %rax
	movq	%rax, 32(%rsp)
	movq	384(%rsp), %rax
	movq	%rax, 24(%rsp)
	movq	392(%rsp), %rax
	movq	%rax, 16(%rsp)
	movq	400(%rsp), %rax
	movq	%rax, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 312(%rsp)
	xorl	%eax, %eax
	leaq	120(%rsp), %rax
	movq	%rax, 192(%rsp)
	leaq	112(%rsp), %rax
	movq	%rax, 200(%rsp)
	leaq	104(%rsp), %rax
	movq	%rax, 208(%rsp)
	leaq	96(%rsp), %rax
	movq	%rax, 216(%rsp)
	leaq	88(%rsp), %rax
	movq	%rax, 224(%rsp)
	leaq	80(%rsp), %rax
	movq	%rax, 232(%rsp)
	leaq	72(%rsp), %rax
	movq	%rax, 240(%rsp)
	leaq	64(%rsp), %rax
	movq	%rax, 248(%rsp)
	leaq	56(%rsp), %rax
	movq	%rax, 256(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 264(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 272(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 280(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 288(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 296(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 304(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 144(%rsp)
	movl	$1, 152(%rsp)
	movq	%rax, 156(%rsp)
	movl	$1, 164(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L114
	pushq	136(%rsp)
	.cfi_def_cfa_offset 344
	leaq	_Z9cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_(%rip), %rdi
	pushq	136(%rsp)
	.cfi_def_cfa_offset 352
	movq	172(%rsp), %rcx
	movl	180(%rsp), %r8d
	movq	160(%rsp), %rsi
	movl	168(%rsp), %edx
	leaq	208(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 344
	popq	%rdx
	.cfi_def_cfa_offset 336
.L114:
	movq	312(%rsp), %rax
	subq	%fs:40, %rax
	je	.L116
	call	__stack_chk_fail@PLT
.L116:
	addq	$328, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE10865:
	.size	_Z62__device_stub__Z9cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_PyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_, .-_Z62__device_stub__Z9cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_PyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_
	.globl	_Z9cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_
	.type	_Z9cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_, @function
_Z9cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_:
.LFB10866:
	.cfi_startproc
	endbr64
	jmp	_Z62__device_stub__Z9cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_PyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_
	.cfi_endproc
.LFE10866:
	.size	_Z9cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_, .-_Z9cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"%i"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB10840:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	subq	$160, %rsp
	.cfi_def_cfa_offset 176
	movq	8(%rsi), %rdi
	leaq	.LC0(%rip), %rsi
	movq	%fs:40, %rax
	movq	%rax, 152(%rsp)
	xorl	%eax, %eax
	leaq	4(%rsp), %rdx
	call	__isoc23_sscanf@PLT
	leaq	8(%rsp), %rdi
	movl	$272, %esi
	call	cudaMalloc@PLT
	leaq	16(%rsp), %rdi
	movl	$272, %esi
	call	cudaMalloc@PLT
	leaq	24(%rsp), %rdi
	movl	$128, %esi
	call	cudaMalloc@PLT
	leaq	32(%rsp), %rdi
	movl	$8000000, %esi
	call	cudaMalloc@PLT
	leaq	40(%rsp), %rdi
	movl	$128, %esi
	call	cudaMalloc@PLT
	leaq	48(%rsp), %rdi
	movl	$224, %esi
	call	cudaMalloc@PLT
	leaq	56(%rsp), %rdi
	movl	$192, %esi
	call	cudaMalloc@PLT
	leaq	64(%rsp), %rdi
	movl	$256, %esi
	call	cudaMalloc@PLT
	leaq	72(%rsp), %rdi
	movl	$256, %esi
	call	cudaMalloc@PLT
	leaq	80(%rsp), %rdi
	movl	$192, %esi
	call	cudaMalloc@PLT
	leaq	88(%rsp), %rdi
	movl	$128, %esi
	call	cudaMalloc@PLT
	leaq	96(%rsp), %rdi
	movl	$2048, %esi
	call	cudaMalloc@PLT
	leaq	104(%rsp), %rdi
	movl	$64, %esi
	call	cudaMalloc@PLT
	leaq	112(%rsp), %rdi
	movl	$64, %esi
	call	cudaMalloc@PLT
	movl	$8000000, %esi
	leaq	120(%rsp), %rdi
	call	cudaMalloc@PLT
	movl	$8000000, %edi
	call	malloc@PLT
	movq	8(%rsp), %rdi
	movl	$1, %ecx
	movl	$272, %edx
	leaq	LnRnBlocks(%rip), %rsi
	movq	%rax, %rbx
	call	cudaMemcpy@PLT
	movq	16(%rsp), %rdi
	movl	$1, %ecx
	movl	$272, %edx
	leaq	CnDnBlocks(%rip), %rsi
	call	cudaMemcpy@PLT
	movq	24(%rsp), %rdi
	movl	$1, %ecx
	movl	$128, %edx
	leaq	keysBlocks(%rip), %rsi
	call	cudaMemcpy@PLT
	movq	32(%rsp), %rdi
	movl	$1, %ecx
	movl	$8000000, %edx
	leaq	allCipherDES(%rip), %rsi
	call	cudaMemcpy@PLT
	movq	40(%rsp), %rdi
	movl	$1, %ecx
	movl	$128, %edx
	leaq	Rotations(%rip), %rsi
	call	cudaMemcpy@PLT
	movq	48(%rsp), %rdi
	movl	$1, %ecx
	movl	$224, %edx
	leaq	PC1(%rip), %rsi
	call	cudaMemcpy@PLT
	movq	56(%rsp), %rdi
	movl	$1, %ecx
	movl	$192, %edx
	leaq	PC2(%rip), %rsi
	call	cudaMemcpy@PLT
	movq	64(%rsp), %rdi
	movl	$1, %ecx
	movl	$256, %edx
	leaq	IniPer(%rip), %rsi
	call	cudaMemcpy@PLT
	movq	72(%rsp), %rdi
	movl	$1, %ecx
	movl	$256, %edx
	leaq	reverseIniPer(%rip), %rsi
	call	cudaMemcpy@PLT
	movq	80(%rsp), %rdi
	movl	$1, %ecx
	movl	$192, %edx
	leaq	Expansion(%rip), %rsi
	call	cudaMemcpy@PLT
	movq	88(%rsp), %rdi
	movl	$1, %ecx
	movl	$128, %edx
	leaq	Pbox(%rip), %rsi
	call	cudaMemcpy@PLT
	movq	96(%rsp), %rdi
	movl	$1, %ecx
	movl	$2048, %edx
	leaq	Sbox(%rip), %rsi
	call	cudaMemcpy@PLT
	movq	104(%rsp), %rdi
	movl	$1, %ecx
	movl	$64, %edx
	leaq	iniKey(%rip), %rsi
	call	cudaMemcpy@PLT
	movq	112(%rsp), %rdi
	movl	$1, %ecx
	leaq	message(%rip), %rsi
	movl	$64, %edx
	call	cudaMemcpy@PLT
	movl	4(%rsp), %esi
	movl	$1000000, %eax
	cmpl	$1000, %esi
	jle	.L125
	movl	$1000000, %ecx
	movl	%ecx, %eax
	cltd
	idivl	%esi
	movl	%ecx, %eax
	subl	%edx, %eax
.L125:
	cltd
	movl	%esi, 140(%rsp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	idivl	%esi
	movabsq	$4294967297, %rdx
	movq	%rdx, 132(%rsp)
	movl	136(%rsp), %esi
	movq	%rdx, 144(%rsp)
	movl	148(%rsp), %ecx
	movq	140(%rsp), %rdx
	movl	%eax, 128(%rsp)
	movq	128(%rsp), %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L122
	pushq	%rax
	.cfi_def_cfa_offset 184
	pushq	128(%rsp)
	.cfi_def_cfa_offset 192
	pushq	128(%rsp)
	.cfi_def_cfa_offset 200
	pushq	128(%rsp)
	.cfi_def_cfa_offset 208
	pushq	128(%rsp)
	.cfi_def_cfa_offset 216
	pushq	128(%rsp)
	.cfi_def_cfa_offset 224
	pushq	128(%rsp)
	.cfi_def_cfa_offset 232
	pushq	128(%rsp)
	.cfi_def_cfa_offset 240
	pushq	128(%rsp)
	.cfi_def_cfa_offset 248
	pushq	128(%rsp)
	.cfi_def_cfa_offset 256
	movq	128(%rsp), %r9
	movq	120(%rsp), %r8
	movq	112(%rsp), %rcx
	movq	104(%rsp), %rdx
	movq	96(%rsp), %rsi
	movq	88(%rsp), %rdi
	call	_Z62__device_stub__Z9cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_PyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_
	addq	$80, %rsp
	.cfi_def_cfa_offset 176
.L122:
	movq	120(%rsp), %rsi
	movl	$2, %ecx
	movl	$8000000, %edx
	movq	%rbx, %rdi
	call	cudaMemcpy@PLT
	movq	%rbx, %rdi
	call	free@PLT
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	movq	32(%rsp), %rdi
	call	cudaFree@PLT
	movq	40(%rsp), %rdi
	call	cudaFree@PLT
	movq	48(%rsp), %rdi
	call	cudaFree@PLT
	movq	56(%rsp), %rdi
	call	cudaFree@PLT
	movq	64(%rsp), %rdi
	call	cudaFree@PLT
	movq	72(%rsp), %rdi
	call	cudaFree@PLT
	movq	80(%rsp), %rdi
	call	cudaFree@PLT
	movq	88(%rsp), %rdi
	call	cudaFree@PLT
	movq	96(%rsp), %rdi
	call	cudaFree@PLT
	movq	104(%rsp), %rdi
	call	cudaFree@PLT
	movq	112(%rsp), %rdi
	call	cudaFree@PLT
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L123
	call	__stack_chk_fail@PLT
.L123:
	addq	$160, %rsp
	.cfi_def_cfa_offset 16
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE10840:
	.size	main, .-main
	.section	.rodata.str1.1
.LC1:
	.string	"_Z9cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB10868:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC1(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z9cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE10868:
	.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	message
	.data
	.align 32
	.type	message, @object
	.size	message, 64
message:
	.quad	1
	.quad	35
	.quad	69
	.quad	103
	.quad	137
	.quad	171
	.quad	205
	.quad	239
	.globl	iniKey
	.align 32
	.type	iniKey, @object
	.size	iniKey, 64
iniKey:
	.quad	19
	.quad	52
	.quad	87
	.quad	121
	.quad	155
	.quad	188
	.quad	223
	.quad	241
	.globl	Sbox
	.align 32
	.type	Sbox, @object
	.size	Sbox, 2048
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	Pbox
	.align 32
	.type	Pbox, @object
	.size	Pbox, 128
Pbox:
	.long	16
	.long	7
	.long	20
	.long	21
	.long	29
	.long	12
	.long	28
	.long	17
	.long	1
	.long	15
	.long	23
	.long	26
	.long	5
	.long	18
	.long	31
	.long	10
	.long	2
	.long	8
	.long	24
	.long	14
	.long	32
	.long	27
	.long	3
	.long	9
	.long	19
	.long	13
	.long	30
	.long	6
	.long	22
	.long	11
	.long	4
	.long	25
	.globl	Expansion
	.align 32
	.type	Expansion, @object
	.size	Expansion, 192
Expansion:
	.long	32
	.long	1
	.long	2
	.long	3
	.long	4
	.long	5
	.long	4
	.long	5
	.long	6
	.long	7
	.long	8
	.long	9
	.long	8
	.long	9
	.long	10
	.long	11
	.long	12
	.long	13
	.long	12
	.long	13
	.long	14
	.long	15
	.long	16
	.long	17
	.long	16
	.long	17
	.long	18
	.long	19
	.long	20
	.long	21
	.long	20
	.long	21
	.long	22
	.long	23
	.long	24
	.long	25
	.long	24
	.long	25
	.long	26
	.long	27
	.long	28
	.long	29
	.long	28
	.long	29
	.long	30
	.long	31
	.long	32
	.long	1
	.globl	reverseIniPer
	.align 32
	.type	reverseIniPer, @object
	.size	reverseIniPer, 256
reverseIniPer:
	.long	40
	.long	8
	.long	48
	.long	16
	.long	56
	.long	24
	.long	64
	.long	32
	.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
	.globl	IniPer
	.align 32
	.type	IniPer, @object
	.size	IniPer, 256
IniPer:
	.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
	.long	64
	.long	56
	.long	48
	.long	40
	.long	32
	.long	24
	.long	16
	.long	8
	.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
	.globl	PC2
	.align 32
	.type	PC2, @object
	.size	PC2, 192
PC2:
	.long	14
	.long	17
	.long	11
	.long	24
	.long	1
	.long	5
	.long	3
	.long	28
	.long	15
	.long	6
	.long	21
	.long	10
	.long	23
	.long	19
	.long	12
	.long	4
	.long	26
	.long	8
	.long	16
	.long	7
	.long	27
	.long	20
	.long	13
	.long	2
	.long	41
	.long	52
	.long	31
	.long	37
	.long	47
	.long	55
	.long	30
	.long	40
	.long	51
	.long	45
	.long	33
	.long	48
	.long	44
	.long	49
	.long	39
	.long	56
	.long	34
	.long	53
	.long	46
	.long	42
	.long	50
	.long	36
	.long	29
	.long	32
	.globl	PC1
	.align 32
	.type	PC1, @object
	.size	PC1, 224
PC1:
	.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
	.long	27
	.long	19
	.long	11
	.long	3
	.long	60
	.long	52
	.long	44
	.long	36
	.long	63
	.long	55
	.long	47
	.long	39
	.long	31
	.long	23
	.long	15
	.long	7
	.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	28
	.long	20
	.long	12
	.long	4
	.globl	Rotations
	.align 32
	.type	Rotations, @object
	.size	Rotations, 128
Rotations:
	.quad	1
	.quad	1
	.quad	2
	.quad	2
	.quad	2
	.quad	2
	.quad	2
	.quad	2
	.quad	1
	.quad	2
	.quad	2
	.quad	2
	.quad	2
	.quad	2
	.quad	2
	.quad	1
	.globl	allCipherDES
	.bss
	.align 32
	.type	allCipherDES, @object
	.size	allCipherDES, 8000000
allCipherDES:
	.zero	8000000
	.globl	keysBlocks
	.align 32
	.type	keysBlocks, @object
	.size	keysBlocks, 128
keysBlocks:
	.zero	128
	.globl	CnDnBlocks
	.align 32
	.type	CnDnBlocks, @object
	.size	CnDnBlocks, 272
CnDnBlocks:
	.zero	272
	.globl	LnRnBlocks
	.align 32
	.type	LnRnBlocks, @object
	.size	LnRnBlocks, 272
LnRnBlocks:
	.zero	272
	.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<bits/stdc++.h>

using namespace std;

typedef unsigned long long ull;
#define MAX 1000000 //10 e 6

ull LnRnBlocks[17*2]; // from l0r0 to l16r16

ull CnDnBlocks[17*2]; //from c0d0 to c16d16

ull keysBlocks[16];  //from key[1] = k0 to key[16] = k15

ull allCipherDES[MAX];

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

int PC1[56] = {
   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,
   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
};

int PC2[48] = {
   14, 17, 11, 24,  1,  5,
    3, 28, 15,  6, 21, 10,
   23, 19, 12,  4, 26,  8,
   16,  7, 27, 20, 13,  2,
   41, 52, 31, 37, 47, 55,
   30, 40, 51, 45, 33, 48,
   44, 49, 39, 56, 34, 53,
   46, 42, 50, 36, 29, 32
};

int IniPer[64] = {
    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
};

int reverseIniPer[64] ={
    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
};

int Expansion[48] ={
    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
};

int Pbox[32] ={
    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
};

int Sbox[512] = { //8*4*16
    
    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
 };

ull iniKey[8] = {
    0x13,0x34,0x57,0x79,0x9B,0xBC,0xDF,0xF1};

ull message[8] = {    
    0x01,0x23,0x45,0x67,0x89,0xAB,0xCD,0xEF};



__host__ __device__ ull generateKeyPlus(ull *d_iniKey, int *d_PC1){
    ull keyPlus=0L;
    for(int i=56-1;i>=0;i--){
        if( d_iniKey[ d_PC1[i]/8 ] & (1 << ( ( 64-d_PC1[i]) % 8 ) ) ){
            keyPlus|=( 1LL<< (55-i)*1L );
        }
    }
    return keyPlus;
}

__host__ __device__ ull* splitKeyPlus(ull keyPlus){
    ull c0=0L, d0=0L;

    for(int i=0;i<28;i++){
        if(keyPlus & (1LL<<i*1L ) ) d0|=(1LL<<i*1L);
        if(keyPlus & (1LL<< i*1L +28L ) ) c0|=(1LL<<i*1L);
    }
    ull new_array[2]= {c0,d0};
    return new_array;
}

__host__ __device__ ull* splitIniPer(ull codeIniPer){
    ull l0=0L, r0=0L;

    for(int i=0;i<32;i++){
        if(codeIniPer & (1LL<<i*1L ) ) r0|=(1LL<<i*1L);
        if(codeIniPer & (1LL<< i*1L +32L ) ) l0|=(1LL<<i*1L);
    }
    ull new_array[2] ={ l0,r0};
    return new_array;
}

__host__ __device__ void generateCnDnBlocks( ull *seedKey,ull *d_CnDnBlocks, ull *d_Rotations){
    d_CnDnBlocks[0] = seedKey[0];
    d_CnDnBlocks[1] = seedKey[1];
    ull cn ,dn, newCn, newDn;
    ull getOnCn, getOnDn;

    for(int i=1;i<=16;i++){
        
        getOnCn = 0L;
        getOnDn = 0L;

        cn = d_CnDnBlocks[(i*2+0)-2];
        dn = d_CnDnBlocks[(i*2+1)-2];

        for(ull j=0;j< d_Rotations[ i-1 ];j++){

            if(  cn & (1 << (27-j) )  ) getOnCn|= 1LL << (d_Rotations[ i-1 ]==1 ? j: 1-j) ;
            if(  dn & (1 << (27-j) )  ) getOnDn|= 1LL << (d_Rotations[ i-1 ]==1 ? j: 1-j);

        }

        newCn = cn << d_Rotations[ i-1 ];
        newDn = dn << d_Rotations[ i-1 ];

        for(ull j=0; j< d_Rotations[ i-1 ] ;j++){
            newCn &= ~(1<< (28+j) );
            newDn &= ~(1<< (28+j) );
        }
        newCn |= ( getOnCn );
        newDn |= ( getOnDn );
        d_CnDnBlocks[ (i*2)+0 ] = newCn;
        d_CnDnBlocks[ (i*2)+1 ] = newDn;
    }

}

__host__ __device__ ull joinCnDn(ull cn, ull dn){ return (cn<<28) | dn; }

__host__ __device__ void generateKeysBlocks(ull *d_CnDnBlocks, int *d_PC2,ull *d_keysBlocks){
    ull cnDn, keyn;

    for(int i=1;i<=16;i++){
        cnDn = joinCnDn( d_CnDnBlocks[i*2+0], d_CnDnBlocks[i*2+1] );
        keyn = 0L;

        for(int j=48-1;j>=0;j--){
            if( cnDn & ( 1LL << (56-d_PC2[j])*1L ) ) {
                 keyn|= ( 1LL<< (47-j)*1L );
            }
        }
        d_keysBlocks[i-1] = keyn;
    }

}

__host__ __device__ ull generateIniPer(int *d_IniPer, ull *d_message){
    ull keyPlus=0L;

    for(int i=64-1;i>=0;i--){
        if( d_message[ (d_IniPer[i]/8) >=8 ? 7: (d_IniPer[i]/8) ]  & (1LL << ( ( 64-d_IniPer[i]) % 8 ) ) ){
            keyPlus|=( 1LL<< (63-i)*1L );
        }
    }
    
    return keyPlus;
}

__host__ __device__ ull expandRn(ull Rn,int *d_Expansion){
    //from a Rn 32 bit to a Kn 48 bit
    ull exRn=0L;
    for(int j=48-1;j>=0;j--){
            if( Rn & ( 1LL << (32-d_Expansion[j])*1L ) ) {
                 exRn|= ( 1LL<< (47-j)*1L );
            }
    }
    return exRn;
}

__host__ __device__ ull xorOperation(ull En, ull Kn){
    return (Kn ^ En);
}

__host__ __device__ ull getSboxNumber(int Bn, int k, int *d_Sbox){

    int row=0,col=0;
    if( Bn & 1<<0 ) row |= ( 1<<0);
    if( Bn & 1<<6-1) row |=( 1<<1 );

    for(int i=1;i<=4;i++){
        if( Bn & 1<<i ) col |=(1<<(i-1));
    }

    return d_Sbox[ (k*4*16)+(row*16)+col ];
}


__host__ __device__ ull generateSboxCombination(ull Bn,int *d_Sbox){

    int Bbox[8];
    ull sbBox[8];
    int number=0, k=7;
    ull snBnChain=0L;
    ull step=28L;

    for(int i=0;i<=48;i++ ){

        if( i%6==0 && i>=6){
            

            Bbox[ k-- ] = number;
            number = 0;
        }
        if( Bn & (1LL<<i*1L) ){
            number |= (1LL<< ( (i%6)*1L ) );
        }
    }

    for(int i=0;i<8;i++){
        sbBox[i] = getSboxNumber( Bbox[i], i,d_Sbox);

       
    }

    
    for(int i=0;i<8;i++){
        snBnChain |= (sbBox[i]<< step);
        step-=4;
    }

    return snBnChain;

}

__host__ __device__ ull generateFalgorithm(ull snBn, int *d_Pbox){

    ull fn=0L;
    for(int j=32-1;j>=0;j--){
            if( snBn & ( 1LL << (32-d_Pbox[j])*1L ) ) {
                 fn|= ( 1LL<< (31-j)*1L );
            }
    }
    return fn;
}

__host__ __device__ void generateLnRnBlocks(ull *L0R0,ull *d_LnRnBlocks, ull *d_keysBlocks,int *d_Expansion, int *d_Sbox,int *d_Pbox){

    d_LnRnBlocks[0] = L0R0[0];
    d_LnRnBlocks[1] = L0R0[1];
    ull fn;

    for(int time=1; time<=16;time++){

        ull Ln = d_LnRnBlocks[ (time*2+0)-2 ];
        ull Rn = d_LnRnBlocks[ (time*2+1)-2 ];
       
        ull snBn = 
            generateSboxCombination( xorOperation( expandRn( Rn, d_Expansion ),d_keysBlocks[ time-1 ] ),d_Sbox );
        
        fn = generateFalgorithm(snBn,d_Pbox);

        d_LnRnBlocks[ (time*2+0) ] = Rn;
        d_LnRnBlocks[ (time*2+1) ] = (Ln ^ fn);

    }

}

__host__ __device__ ull reverseLnRn( ull *LnRn, ull *LnRn1){
    ull Ln = *LnRn;
    ull Rn = *LnRn1;

    return ( Rn<<32L) | Ln;
}

__host__ __device__ ull generateCipherMessage( ull RnLn, int *d_reverseIniPer ){

    ull cipher=0L;
    for(int j=64-1;j>=0;j--){
            if( RnLn & ( 1LL << (64-d_reverseIniPer[j])*1L ) ) {
                 cipher|= ( 1LL<< (63-j)*1L );
            }
    }
    return cipher;
}


__global__ void cipherDES(
    ull *d_LnRnBlocks,
    ull *d_CnDnBlocks,
    ull *d_keysBlocks,
    ull *d_allCipherDES,
    ull *d_Rotations,
    int *d_PC1,
    int *d_PC2,
    int *d_IniPer,
    int *d_reverseIniPer,
    int *d_Expansion,
    int *d_Pbox,
    int *d_Sbox,
    ull *d_iniKey,
    ull *d_message,
    ull *d_result
){
    ull *keyHalves = splitKeyPlus( generateKeyPlus(d_iniKey,d_PC1) );
    generateCnDnBlocks( keyHalves,d_CnDnBlocks,d_Rotations );
    generateKeysBlocks(d_CnDnBlocks,d_PC2,d_keysBlocks);
    ull *iniPerHalves = splitIniPer(generateIniPer(d_IniPer,d_message) ); //got L0 and R0
    

    generateLnRnBlocks( iniPerHalves,d_LnRnBlocks, d_keysBlocks, d_Expansion, d_Sbox, d_Pbox);

    ull revLnRn = reverseLnRn( &d_LnRnBlocks[16*2+0],&d_LnRnBlocks[16*2+1] );
   
    ull cipherMessage = generateCipherMessage( revLnRn,d_reverseIniPer );
    
    int iindex = threadIdx.x + blockIdx.x * blockDim.x;
    d_result[iindex] = cipherMessage;
    // printf("Index %i - Hex Cipher: %llX\n", iindex, cipherMessage);
   
}

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

    int THREADS_BLOCK;
    
    sscanf (argv[1], "%i", &THREADS_BLOCK);
    

    //host and devices copies
    ull *d_LnRnBlocks;//17 size
    ull *d_CnDnBlocks;//17 size
    ull *d_keysBlocks;//16 size
    ull *d_allCipherDES;//10^6 size
    ull *d_Rotations;//16 size
    int *d_PC1;//56 size
    int *d_PC2;//48 size
    int *d_IniPer;//64 size
    int *d_reverseIniPer;//64 size
    int *d_Expansion;//48 size
    int *d_Pbox;//32 size
    int *d_Sbox;//8*8*16 size [8][4][16]
    ull *d_iniKey;//8 size
    ull *d_message;//8 size
    ull *result; //MAX SIZE
    ull *d_result;//MAX SIZE

    //size of host and device copies
    int sd_LnRnBlocks = 17 * 2 * sizeof(ull);
    int sd_CnDnBlocks = 17 * 2 * sizeof(ull);
    int sd_keysBlocks = 16 * sizeof(ull);
    int sd_allCipherDES = 1000000 * sizeof(ull);
    int sd_Rotations = 16 * sizeof(ull);
    int sd_PC1 = 56 * sizeof(int);
    int sd_PC2 = 48 * sizeof(int);
    int sd_IniPer = 64 * sizeof(int);
    int sd_reverseIniPer = 64 * sizeof(int);
    int sd_Expansion = 48 * sizeof(int);
    int sd_Pbox = 32 * sizeof(int);
    int sd_Sbox = 512 * sizeof(int);
    int sd_iniKey = 8 * sizeof(ull);
    int sd_message = 8 * sizeof(ull);
    int sd_result = MAX * sizeof(ull);

    //alloc space for device copies
    hipMalloc( (void **)&d_LnRnBlocks, sd_LnRnBlocks );
    hipMalloc( (void **)&d_CnDnBlocks, sd_CnDnBlocks );
    hipMalloc( (void **)&d_keysBlocks, sd_keysBlocks );
    hipMalloc( (void **)&d_allCipherDES, sd_allCipherDES );
    hipMalloc( (void **)&d_Rotations, sd_Rotations );
    hipMalloc( (void **)&d_PC1, sd_PC1 );
    hipMalloc( (void **)&d_PC2, sd_PC2 );
    hipMalloc( (void **)&d_IniPer, sd_IniPer );
    hipMalloc( (void **)&d_reverseIniPer, sd_reverseIniPer );
    hipMalloc( (void **)&d_Expansion, sd_Expansion );
    hipMalloc( (void **)&d_Pbox, sd_Pbox );
    hipMalloc( (void **)&d_Sbox, sd_Sbox );
    hipMalloc( (void **)&d_iniKey, sd_iniKey );
    hipMalloc( (void **)&d_message, sd_message );
    hipMalloc( (void **)&d_result, sd_result );

    //alloc space for host copies
    result = (ull *) malloc(sd_result);

    //copy inputs to device
    hipMemcpy(d_LnRnBlocks,LnRnBlocks,sd_LnRnBlocks,hipMemcpyHostToDevice);
    hipMemcpy(d_CnDnBlocks,CnDnBlocks,sd_CnDnBlocks,hipMemcpyHostToDevice);
    hipMemcpy(d_keysBlocks,keysBlocks,sd_keysBlocks,hipMemcpyHostToDevice);
    hipMemcpy(d_allCipherDES,allCipherDES,sd_allCipherDES,hipMemcpyHostToDevice);
    hipMemcpy(d_Rotations,Rotations,sd_Rotations,hipMemcpyHostToDevice);
    hipMemcpy(d_PC1,PC1,sd_PC1,hipMemcpyHostToDevice);
    hipMemcpy(d_PC2,PC2,sd_PC2,hipMemcpyHostToDevice);
    hipMemcpy(d_IniPer,IniPer,sd_IniPer,hipMemcpyHostToDevice);
    hipMemcpy(d_reverseIniPer,reverseIniPer,sd_reverseIniPer,hipMemcpyHostToDevice);
    hipMemcpy(d_Expansion,Expansion,sd_Expansion,hipMemcpyHostToDevice);
    hipMemcpy(d_Pbox,Pbox,sd_Pbox,hipMemcpyHostToDevice);
    hipMemcpy(d_Sbox,Sbox,sd_Sbox,hipMemcpyHostToDevice);
    hipMemcpy(d_iniKey,iniKey,sd_iniKey,hipMemcpyHostToDevice);
    hipMemcpy(d_message,message,sd_message,hipMemcpyHostToDevice);
   


    int global_work_size = ceil(MAX/THREADS_BLOCK);
    int local_work_size = THREADS_BLOCK;
    if(THREADS_BLOCK > 1000 ){
         global_work_size = MAX - (MAX%THREADS_BLOCK);
         global_work_size = ceil( global_work_size/THREADS_BLOCK);
         local_work_size = THREADS_BLOCK;
    }
    //launch kernel     ----------- HERES THE MAGIC ---------
    cipherDES<<<global_work_size,local_work_size>>>(
        d_LnRnBlocks,
        d_CnDnBlocks,
        d_keysBlocks,
        d_allCipherDES,
        d_Rotations,
        d_PC1,
        d_PC2,
        d_IniPer,
        d_reverseIniPer,
        d_Expansion,
        d_Pbox,
        d_Sbox,
        d_iniKey,
        d_message,
        d_result
    );


    //copy result back to host
    hipMemcpy(result,d_result,sd_result,hipMemcpyDeviceToHost);

    //print result

    // for(int i=0;i<MAX;i++){
    //     printf("cipher: %i\t%llX\n",i,result[i]);
    //     fflush(stdout);
    // }


    //free host space
    free(result);

    //free cuda space
    hipFree(d_LnRnBlocks);
    hipFree(d_CnDnBlocks);
    hipFree(d_keysBlocks);
    hipFree(d_allCipherDES);
    hipFree(d_Rotations);
    hipFree(d_PC1);
    hipFree(d_PC2);
    hipFree(d_IniPer);
    hipFree(d_reverseIniPer);
    hipFree(d_Expansion);
    hipFree(d_Pbox);
    hipFree(d_Sbox);
    hipFree(d_iniKey);
    hipFree(d_message);

    return 0;
}


// L16 01000011010000100011001000110100    1128411700
// R16 00001010010011001101100110010101    172808597

// revL16R16 0000101001001100110110011001010101000011010000100011001000110100  742207273711055412

// cipher 1000010111101000000100110101010000001111000010101011010000000101     9648983453391827973
	.text
	.file	"cctest.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.globl	_Z15generateKeyPlusPyPi         # -- Begin function _Z15generateKeyPlusPyPi
	.type	_Z15generateKeyPlusPyPi,@function
_Z15generateKeyPlusPyPi:                # @_Z15generateKeyPlusPyPi
	.cfi_startproc
# %bb.0:
	addq	$220, %rsi
	xorl	%ecx, %ecx
	xorl	%eax, %eax
.LBB0_1:                                # =>This Inner Loop Header: Depth=1
	movl	(%rsi), %edx
	leal	7(%rdx), %r8d
	testl	%edx, %edx
	cmovnsl	%edx, %r8d
	sarl	$3, %r8d
	movslq	%r8d, %r8
	movl	(%rdi,%r8,8), %r8d
	negb	%dl
	movzbl	%dl, %edx
	andl	$7, %edx
	movl	$1, %r9d
	shlq	%cl, %r9
	btl	%edx, %r8d
	movl	$0, %edx
	cmovbq	%r9, %rdx
	orq	%rdx, %rax
	incq	%rcx
	addq	$-4, %rsi
	cmpq	$56, %rcx
	jne	.LBB0_1
# %bb.2:
	retq
.Lfunc_end0:
	.size	_Z15generateKeyPlusPyPi, .Lfunc_end0-_Z15generateKeyPlusPyPi
	.cfi_endproc
                                        # -- End function
	.globl	_Z12splitKeyPlusy               # -- Begin function _Z12splitKeyPlusy
	.type	_Z12splitKeyPlusy,@function
_Z12splitKeyPlusy:                      # @_Z12splitKeyPlusy
	.cfi_startproc
# %bb.0:
	leaq	-24(%rsp), %rax
	retq
.Lfunc_end1:
	.size	_Z12splitKeyPlusy, .Lfunc_end1-_Z12splitKeyPlusy
	.cfi_endproc
                                        # -- End function
	.globl	_Z11splitIniPery                # -- Begin function _Z11splitIniPery
	.type	_Z11splitIniPery,@function
_Z11splitIniPery:                       # @_Z11splitIniPery
	.cfi_startproc
# %bb.0:
	leaq	-24(%rsp), %rax
	retq
.Lfunc_end2:
	.size	_Z11splitIniPery, .Lfunc_end2-_Z11splitIniPery
	.cfi_endproc
                                        # -- End function
	.globl	_Z18generateCnDnBlocksPyS_S_    # -- Begin function _Z18generateCnDnBlocksPyS_S_
	.type	_Z18generateCnDnBlocksPyS_S_,@function
_Z18generateCnDnBlocksPyS_S_:           # @_Z18generateCnDnBlocksPyS_S_
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	(%rdi), %rax
	movq	%rax, (%rsi)
	movq	8(%rdi), %rax
	movq	%rax, 8(%rsi)
	movl	$1, %edi
.LBB3_1:                                # =>This Loop Header: Depth=1
                                        #     Child Loop BB3_6 Depth 2
                                        #     Child Loop BB3_8 Depth 2
	movq	%rdi, %rax
	shlq	$4, %rax
	movq	-16(%rsi,%rax), %r8
	movq	-8(%rsi,%rax), %r9
	movq	-8(%rdx,%rdi,8), %rax
	testq	%rax, %rax
	je	.LBB3_2
# %bb.5:                                # %.lr.ph
                                        #   in Loop: Header=BB3_1 Depth=1
	movl	$1, %ebx
	xorl	%r14d, %r14d
	xorl	%r10d, %r10d
	xorl	%r11d, %r11d
.LBB3_6:                                #   Parent Loop BB3_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	leal	26(%rbx), %ecx
	movl	$1, %ebp
                                        # kill: def $cl killed $cl killed $ecx
	shll	%cl, %ebp
	movslq	%ebp, %r15
	movq	%r8, %r12
	andq	%r15, %r12
	cmpq	$1, %rax
	movl	%ebx, %ecx
	cmovel	%r14d, %ecx
	movl	$1, %r13d
                                        # kill: def $cl killed $cl killed $ecx
	shlq	%cl, %r13
	testq	%r12, %r12
	cmovneq	%r13, %r12
	orq	%r12, %r11
	andq	%r9, %r15
	cmovneq	%r13, %r15
	orq	%r15, %r10
	incq	%r14
	decq	%rbx
	cmpq	%r14, %rax
	jne	.LBB3_6
# %bb.7:                                # %.lr.ph70.preheader
                                        #   in Loop: Header=BB3_1 Depth=1
	movl	%eax, %ecx
	shlq	%cl, %r8
	shlq	%cl, %r9
	xorl	%ebx, %ebx
.LBB3_8:                                # %.lr.ph70
                                        #   Parent Loop BB3_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	leal	28(%rbx), %ecx
	movl	$-2, %ebp
                                        # kill: def $cl killed $cl killed $ecx
	roll	%cl, %ebp
	movslq	%ebp, %rcx
	andq	%rcx, %r8
	andq	%rcx, %r9
	incq	%rbx
	cmpq	%rbx, %rax
	jne	.LBB3_8
	jmp	.LBB3_3
.LBB3_2:                                #   in Loop: Header=BB3_1 Depth=1
	xorl	%r10d, %r10d
	xorl	%r11d, %r11d
.LBB3_3:                                # %._crit_edge71
                                        #   in Loop: Header=BB3_1 Depth=1
	orq	%r9, %r10
	orq	%r8, %r11
	leaq	(%rdi,%rdi), %rax
	leaq	1(,%rdi,2), %rcx
	movq	%r11, (%rsi,%rax,8)
	movq	%r10, (%rsi,%rcx,8)
	incq	%rdi
	cmpq	$17, %rdi
	jne	.LBB3_1
# %bb.4:
	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	_Z18generateCnDnBlocksPyS_S_, .Lfunc_end3-_Z18generateCnDnBlocksPyS_S_
	.cfi_endproc
                                        # -- End function
	.globl	_Z8joinCnDnyy                   # -- Begin function _Z8joinCnDnyy
	.type	_Z8joinCnDnyy,@function
_Z8joinCnDnyy:                          # @_Z8joinCnDnyy
	.cfi_startproc
# %bb.0:
	movq	%rdi, %rax
	shlq	$28, %rax
	orq	%rsi, %rax
	retq
.Lfunc_end4:
	.size	_Z8joinCnDnyy, .Lfunc_end4-_Z8joinCnDnyy
	.cfi_endproc
                                        # -- End function
	.globl	_Z18generateKeysBlocksPyPiS_    # -- Begin function _Z18generateKeysBlocksPyPiS_
	.type	_Z18generateKeysBlocksPyPiS_,@function
_Z18generateKeysBlocksPyPiS_:           # @_Z18generateKeysBlocksPyPiS_
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset %rbx, -16
	addq	$188, %rsi
	movl	$1, %eax
.LBB5_1:                                # =>This Loop Header: Depth=1
                                        #     Child Loop BB5_2 Depth 2
	movq	%rax, %rcx
	shlq	$4, %rcx
	movq	(%rdi,%rcx), %r8
	shlq	$28, %r8
	orq	8(%rdi,%rcx), %r8
	movq	%rsi, %r9
	xorl	%ecx, %ecx
	xorl	%r10d, %r10d
.LBB5_2:                                #   Parent Loop BB5_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movb	$56, %r11b
	subb	(%r9), %r11b
	movzbl	%r11b, %ebx
	movl	$1, %r11d
	shlq	%cl, %r11
	btq	%rbx, %r8
	jb	.LBB5_4
# %bb.3:                                #   in Loop: Header=BB5_2 Depth=2
	xorl	%r11d, %r11d
.LBB5_4:                                #   in Loop: Header=BB5_2 Depth=2
	orq	%r11, %r10
	incq	%rcx
	addq	$-4, %r9
	cmpq	$48, %rcx
	jne	.LBB5_2
# %bb.5:                                #   in Loop: Header=BB5_1 Depth=1
	movq	%r10, -8(%rdx,%rax,8)
	incq	%rax
	cmpq	$17, %rax
	jne	.LBB5_1
# %bb.6:
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end5:
	.size	_Z18generateKeysBlocksPyPiS_, .Lfunc_end5-_Z18generateKeysBlocksPyPiS_
	.cfi_endproc
                                        # -- End function
	.globl	_Z14generateIniPerPiPy          # -- Begin function _Z14generateIniPerPiPy
	.type	_Z14generateIniPerPiPy,@function
_Z14generateIniPerPiPy:                 # @_Z14generateIniPerPiPy
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset %rbx, -16
	movl	$63, %ecx
	xorl	%eax, %eax
	movl	$7, %edx
	movabsq	$-9223372036854775808, %r8      # imm = 0x8000000000000000
.LBB6_1:                                # =>This Inner Loop Header: Depth=1
	movl	(%rdi,%rcx,4), %r9d
	leal	7(%r9), %r10d
	testl	%r9d, %r9d
	cmovnsl	%r9d, %r10d
	sarl	$3, %r10d
	movl	$71, %r11d
	subl	%r9d, %r11d
	subl	$64, %r9d
	cmovgel	%edx, %r10d
	negl	%r9d
	movslq	%r10d, %r10
	movq	(%rsi,%r10,8), %rbx
	testl	%r9d, %r9d
	cmovnsl	%r9d, %r11d
	andl	$56, %r11d
	subl	%r11d, %r9d
	movq	%r8, %r10
	shrq	%cl, %r10
	btq	%r9, %rbx
	jb	.LBB6_3
# %bb.2:                                #   in Loop: Header=BB6_1 Depth=1
	xorl	%r10d, %r10d
.LBB6_3:                                #   in Loop: Header=BB6_1 Depth=1
	orq	%r10, %rax
	addq	$-1, %rcx
	jb	.LBB6_1
# %bb.4:
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end6:
	.size	_Z14generateIniPerPiPy, .Lfunc_end6-_Z14generateIniPerPiPy
	.cfi_endproc
                                        # -- End function
	.globl	_Z8expandRnyPi                  # -- Begin function _Z8expandRnyPi
	.type	_Z8expandRnyPi,@function
_Z8expandRnyPi:                         # @_Z8expandRnyPi
	.cfi_startproc
# %bb.0:
	addq	$188, %rsi
	xorl	%ecx, %ecx
	xorl	%eax, %eax
.LBB7_1:                                # =>This Inner Loop Header: Depth=1
	movb	$32, %dl
	subb	(%rsi), %dl
	movzbl	%dl, %r8d
	movl	$1, %edx
	shlq	%cl, %rdx
	btq	%r8, %rdi
	jb	.LBB7_3
# %bb.2:                                #   in Loop: Header=BB7_1 Depth=1
	xorl	%edx, %edx
.LBB7_3:                                #   in Loop: Header=BB7_1 Depth=1
	orq	%rdx, %rax
	incq	%rcx
	addq	$-4, %rsi
	cmpq	$48, %rcx
	jne	.LBB7_1
# %bb.4:
	retq
.Lfunc_end7:
	.size	_Z8expandRnyPi, .Lfunc_end7-_Z8expandRnyPi
	.cfi_endproc
                                        # -- End function
	.globl	_Z12xorOperationyy              # -- Begin function _Z12xorOperationyy
	.type	_Z12xorOperationyy,@function
_Z12xorOperationyy:                     # @_Z12xorOperationyy
	.cfi_startproc
# %bb.0:
	movq	%rdi, %rax
	xorq	%rsi, %rax
	retq
.Lfunc_end8:
	.size	_Z12xorOperationyy, .Lfunc_end8-_Z12xorOperationyy
	.cfi_endproc
                                        # -- End function
	.globl	_Z13getSboxNumberiiPi           # -- Begin function _Z13getSboxNumberiiPi
	.type	_Z13getSboxNumberiiPi,@function
_Z13getSboxNumberiiPi:                  # @_Z13getSboxNumberiiPi
	.cfi_startproc
# %bb.0:
	xorl	%r8d, %r8d
	movl	$1, %r9d
	xorl	%eax, %eax
.LBB9_1:                                # =>This Inner Loop Header: Depth=1
	leal	-1(%r9), %ecx
	movl	$1, %r10d
                                        # kill: def $cl killed $cl killed $ecx
	shll	%cl, %r10d
	btl	%r9d, %edi
	cmovael	%r8d, %r10d
	orl	%r10d, %eax
	incl	%r9d
	cmpl	$5, %r9d
	jne	.LBB9_1
# %bb.2:
	shll	$6, %esi
	movl	%edi, %ecx
	shll	$4, %ecx
	andl	$16, %ecx
	andl	$32, %edi
	orl	%ecx, %edi
	orl	%esi, %edi
	addl	%eax, %edi
	movslq	%edi, %rax
	movslq	(%rdx,%rax,4), %rax
	retq
.Lfunc_end9:
	.size	_Z13getSboxNumberiiPi, .Lfunc_end9-_Z13getSboxNumberiiPi
	.cfi_endproc
                                        # -- End function
	.globl	_Z23generateSboxCombinationyPi  # -- Begin function _Z23generateSboxCombinationyPi
	.type	_Z23generateSboxCombinationyPi,@function
_Z23generateSboxCombinationyPi:         # @_Z23generateSboxCombinationyPi
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	subq	$104, %rsp
	.cfi_def_cfa_offset 128
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	movq	%rsi, %rbx
	xorl	%eax, %eax
	movl	$7, %edx
	movl	$2863311531, %esi               # imm = 0xAAAAAAAB
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
.LBB10_1:                               # =>This Inner Loop Header: Depth=1
	movl	%r8d, %ecx
	imulq	%rsi, %rcx
	shrq	$34, %rcx
	addl	%ecx, %ecx
	leal	(%rcx,%rcx,2), %r10d
	movl	%r8d, %ecx
	subl	%r10d, %ecx
	cmpq	$6, %r8
	jb	.LBB10_4
# %bb.2:                                #   in Loop: Header=BB10_1 Depth=1
	testb	%cl, %cl
	jne	.LBB10_4
# %bb.3:                                #   in Loop: Header=BB10_1 Depth=1
	movslq	%edx, %r10
	decl	%edx
	movl	%r9d, (%rsp,%r10,4)
	xorl	%r9d, %r9d
.LBB10_4:                               #   in Loop: Header=BB10_1 Depth=1
	movl	$1, %r10d
                                        # kill: def $cl killed $cl killed $ecx
	shll	%cl, %r10d
	btq	%r8, %rdi
	cmovael	%eax, %r10d
	orl	%r10d, %r9d
	incq	%r8
	cmpq	$49, %r8
	jne	.LBB10_1
# %bb.5:                                # %.preheader30.preheader
	xorl	%r14d, %r14d
.LBB10_6:                               # %.preheader30
                                        # =>This Inner Loop Header: Depth=1
	movl	(%rsp,%r14,4), %edi
	movl	%r14d, %esi
	movq	%rbx, %rdx
	callq	_Z13getSboxNumberiiPi
	movq	%rax, 32(%rsp,%r14,8)
	incq	%r14
	cmpq	$8, %r14
	jne	.LBB10_6
# %bb.7:                                # %.preheader.preheader
	movl	$28, %ecx
	leaq	32(%rsp), %rdx
	xorl	%eax, %eax
.LBB10_8:                               # %.preheader
                                        # =>This Inner Loop Header: Depth=1
	movq	(%rdx), %rsi
	shlq	%cl, %rsi
	orq	%rsi, %rax
	addq	$-4, %rcx
	addq	$8, %rdx
	cmpq	$-4, %rcx
	jne	.LBB10_8
# %bb.9:
	addq	$104, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end10:
	.size	_Z23generateSboxCombinationyPi, .Lfunc_end10-_Z23generateSboxCombinationyPi
	.cfi_endproc
                                        # -- End function
	.globl	_Z18generateFalgorithmyPi       # -- Begin function _Z18generateFalgorithmyPi
	.type	_Z18generateFalgorithmyPi,@function
_Z18generateFalgorithmyPi:              # @_Z18generateFalgorithmyPi
	.cfi_startproc
# %bb.0:
	addq	$124, %rsi
	xorl	%ecx, %ecx
	xorl	%eax, %eax
.LBB11_1:                               # =>This Inner Loop Header: Depth=1
	movb	$32, %dl
	subb	(%rsi), %dl
	movzbl	%dl, %r8d
	movl	$1, %edx
	shlq	%cl, %rdx
	btq	%r8, %rdi
	jb	.LBB11_3
# %bb.2:                                #   in Loop: Header=BB11_1 Depth=1
	xorl	%edx, %edx
.LBB11_3:                               #   in Loop: Header=BB11_1 Depth=1
	orq	%rdx, %rax
	incq	%rcx
	addq	$-4, %rsi
	cmpq	$32, %rcx
	jne	.LBB11_1
# %bb.4:
	retq
.Lfunc_end11:
	.size	_Z18generateFalgorithmyPi, .Lfunc_end11-_Z18generateFalgorithmyPi
	.cfi_endproc
                                        # -- End function
	.globl	_Z18generateLnRnBlocksPyS_S_PiS0_S0_ # -- Begin function _Z18generateLnRnBlocksPyS_S_PiS0_S0_
	.type	_Z18generateLnRnBlocksPyS_S_PiS0_S0_,@function
_Z18generateLnRnBlocksPyS_S_PiS0_S0_:   # @_Z18generateLnRnBlocksPyS_S_PiS0_S0_
	.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
	movq	%r9, %r15
	movq	%r8, 16(%rsp)                   # 8-byte Spill
	movq	%rdx, (%rsp)                    # 8-byte Spill
	movq	%rsi, %r13
	movq	(%rdi), %rax
	movq	%rax, (%rsi)
	movq	8(%rdi), %rax
	movq	%rax, 8(%rsi)
	addq	$188, %rcx
	movq	%rcx, 8(%rsp)                   # 8-byte Spill
	addq	$124, %r15
	movl	$1, %ebp
.LBB12_1:                               # =>This Loop Header: Depth=1
                                        #     Child Loop BB12_2 Depth 2
                                        #     Child Loop BB12_6 Depth 2
	leaq	(,%rbp,2), %r14
	movq	%rbp, %rax
	shlq	$4, %rax
	movq	-16(%r13,%rax), %r12
	movq	-8(%r13,%rax), %rbx
	movq	8(%rsp), %rax                   # 8-byte Reload
	xorl	%ecx, %ecx
	xorl	%edi, %edi
.LBB12_2:                               #   Parent Loop BB12_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movb	$32, %dl
	subb	(%rax), %dl
	movzbl	%dl, %esi
	movl	$1, %edx
	shlq	%cl, %rdx
	btq	%rsi, %rbx
	jb	.LBB12_4
# %bb.3:                                #   in Loop: Header=BB12_2 Depth=2
	xorl	%edx, %edx
.LBB12_4:                               #   in Loop: Header=BB12_2 Depth=2
	orq	%rdx, %rdi
	incq	%rcx
	addq	$-4, %rax
	cmpq	$48, %rcx
	jne	.LBB12_2
# %bb.5:                                # %_Z8expandRnyPi.exit
                                        #   in Loop: Header=BB12_1 Depth=1
	movq	(%rsp), %rax                    # 8-byte Reload
	xorq	-8(%rax,%rbp,8), %rdi
	movq	16(%rsp), %rsi                  # 8-byte Reload
	callq	_Z23generateSboxCombinationyPi
	movq	%r15, %rsi
	xorl	%ecx, %ecx
	xorl	%edx, %edx
.LBB12_6:                               #   Parent Loop BB12_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movb	$32, %dil
	subb	(%rsi), %dil
	movzbl	%dil, %r8d
	movl	$1, %edi
	shlq	%cl, %rdi
	btq	%r8, %rax
	jb	.LBB12_8
# %bb.7:                                #   in Loop: Header=BB12_6 Depth=2
	xorl	%edi, %edi
.LBB12_8:                               #   in Loop: Header=BB12_6 Depth=2
	orq	%rdi, %rdx
	incq	%rcx
	addq	$-4, %rsi
	cmpq	$32, %rcx
	jne	.LBB12_6
# %bb.9:                                # %_Z18generateFalgorithmyPi.exit
                                        #   in Loop: Header=BB12_1 Depth=1
	movq	%rbx, (%r13,%r14,8)
	xorq	%r12, %rdx
	movq	%rdx, 8(%r13,%r14,8)
	incq	%rbp
	cmpq	$17, %rbp
	jne	.LBB12_1
# %bb.10:
	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
	retq
.Lfunc_end12:
	.size	_Z18generateLnRnBlocksPyS_S_PiS0_S0_, .Lfunc_end12-_Z18generateLnRnBlocksPyS_S_PiS0_S0_
	.cfi_endproc
                                        # -- End function
	.globl	_Z11reverseLnRnPyS_             # -- Begin function _Z11reverseLnRnPyS_
	.type	_Z11reverseLnRnPyS_,@function
_Z11reverseLnRnPyS_:                    # @_Z11reverseLnRnPyS_
	.cfi_startproc
# %bb.0:
	movq	(%rsi), %rax
	shlq	$32, %rax
	orq	(%rdi), %rax
	retq
.Lfunc_end13:
	.size	_Z11reverseLnRnPyS_, .Lfunc_end13-_Z11reverseLnRnPyS_
	.cfi_endproc
                                        # -- End function
	.globl	_Z21generateCipherMessageyPi    # -- Begin function _Z21generateCipherMessageyPi
	.type	_Z21generateCipherMessageyPi,@function
_Z21generateCipherMessageyPi:           # @_Z21generateCipherMessageyPi
	.cfi_startproc
# %bb.0:
	movl	$63, %ecx
	xorl	%eax, %eax
	movabsq	$-9223372036854775808, %rdx     # imm = 0x8000000000000000
.LBB14_1:                               # =>This Inner Loop Header: Depth=1
	movb	$64, %r8b
	subb	(%rsi,%rcx,4), %r8b
	movzbl	%r8b, %r9d
	movq	%rdx, %r8
	shrq	%cl, %r8
	btq	%r9, %rdi
	jb	.LBB14_3
# %bb.2:                                #   in Loop: Header=BB14_1 Depth=1
	xorl	%r8d, %r8d
.LBB14_3:                               #   in Loop: Header=BB14_1 Depth=1
	orq	%r8, %rax
	addq	$-1, %rcx
	jb	.LBB14_1
# %bb.4:
	retq
.Lfunc_end14:
	.size	_Z21generateCipherMessageyPi, .Lfunc_end14-_Z21generateCipherMessageyPi
	.cfi_endproc
                                        # -- End function
	.globl	_Z24__device_stub__cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_ # -- Begin function _Z24__device_stub__cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_
	.type	_Z24__device_stub__cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_,@function
_Z24__device_stub__cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_: # @_Z24__device_stub__cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_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	$224, %rsp
	.cfi_def_cfa_offset 272
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	56(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	48(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	24(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	16(%rsp), %r8
	movq	%r9, (%r8)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	272(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	280(%rsp), %rax
	movq	%rax, 56(%rbx)
	leaq	288(%rsp), %rax
	movq	%rax, 64(%rbx)
	leaq	296(%rsp), %rax
	movq	%rax, 72(%rbx)
	leaq	304(%rsp), %rax
	movq	%rax, 80(%rbx)
	leaq	312(%rsp), %rax
	movq	%rax, 88(%rbx)
	leaq	320(%rsp), %rax
	movq	%rax, 96(%rbx)
	leaq	328(%rsp), %rax
	movq	%rax, 104(%rbx)
	leaq	336(%rsp), %rax
	movq	%rax, 112(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%rsp, %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z9cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$240, %rsp
	.cfi_adjust_cfa_offset -240
	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_end15:
	.size	_Z24__device_stub__cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_, .Lfunc_end15-_Z24__device_stub__cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	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	$136, %rsp
	.cfi_def_cfa_offset 192
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	8(%rsi), %rdi
	leaq	84(%rsp), %rdx
	movl	$.L.str, %esi
	xorl	%eax, %eax
	callq	__isoc23_sscanf
	leaq	128(%rsp), %rbp
	movl	$272, %esi                      # imm = 0x110
	movq	%rbp, %rdi
	callq	hipMalloc
	leaq	120(%rsp), %r13
	movl	$272, %esi                      # imm = 0x110
	movq	%r13, %rdi
	callq	hipMalloc
	leaq	112(%rsp), %r12
	movl	$128, %esi
	movq	%r12, %rdi
	callq	hipMalloc
	leaq	104(%rsp), %r15
	movl	$8000000, %esi                  # imm = 0x7A1200
	movq	%r15, %rdi
	callq	hipMalloc
	leaq	96(%rsp), %rbx
	movl	$128, %esi
	movq	%rbx, %rdi
	callq	hipMalloc
	leaq	72(%rsp), %rdi
	movl	$224, %esi
	callq	hipMalloc
	leaq	64(%rsp), %rdi
	movl	$192, %esi
	callq	hipMalloc
	leaq	56(%rsp), %rdi
	movl	$256, %esi                      # imm = 0x100
	callq	hipMalloc
	leaq	48(%rsp), %rdi
	movl	$256, %esi                      # imm = 0x100
	callq	hipMalloc
	leaq	40(%rsp), %rdi
	movl	$192, %esi
	callq	hipMalloc
	leaq	32(%rsp), %rdi
	movl	$128, %esi
	callq	hipMalloc
	leaq	24(%rsp), %rdi
	movl	$2048, %esi                     # imm = 0x800
	callq	hipMalloc
	leaq	16(%rsp), %rdi
	movl	$64, %esi
	callq	hipMalloc
	leaq	8(%rsp), %rdi
	movl	$64, %esi
	callq	hipMalloc
	leaq	88(%rsp), %rdi
	movl	$8000000, %esi                  # imm = 0x7A1200
	callq	hipMalloc
	movl	$8000000, %edi                  # imm = 0x7A1200
	callq	malloc
	movq	%rax, %r14
	movq	(%rbp), %rdi
	movl	$LnRnBlocks, %esi
	movl	$272, %edx                      # imm = 0x110
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%r13), %rdi
	movl	$CnDnBlocks, %esi
	movl	$272, %edx                      # imm = 0x110
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%r12), %rdi
	movl	$keysBlocks, %esi
	movl	$128, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%r15), %rdi
	movl	$allCipherDES, %esi
	movl	$8000000, %edx                  # imm = 0x7A1200
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%rbx), %rdi
	movl	$Rotations, %esi
	movl	$128, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	72(%rsp), %rax
	movq	(%rax), %rdi
	movl	$PC1, %esi
	movl	$224, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	64(%rsp), %rax
	movq	(%rax), %rdi
	movl	$PC2, %esi
	movl	$192, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	56(%rsp), %rax
	movq	(%rax), %rdi
	movl	$IniPer, %esi
	movl	$256, %edx                      # imm = 0x100
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	48(%rsp), %rax
	movq	(%rax), %rdi
	movl	$reverseIniPer, %esi
	movl	$256, %edx                      # imm = 0x100
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	40(%rsp), %rax
	movq	(%rax), %rdi
	movl	$Expansion, %esi
	movl	$192, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	32(%rsp), %rax
	movq	(%rax), %rdi
	movl	$Pbox, %esi
	movl	$128, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	24(%rsp), %rax
	movq	(%rax), %rdi
	movl	$Sbox, %esi
	movl	$2048, %edx                     # imm = 0x800
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	16(%rsp), %rax
	movq	(%rax), %rdi
	movl	$iniKey, %esi
	movl	$64, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	8(%rsp), %rax
	movq	(%rax), %rdi
	movl	$message, %esi
	movl	$64, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	84(%rsp), %rax
	movl	(%rax), %ecx
	movl	$1000000, %eax                  # imm = 0xF4240
	xorl	%edx, %edx
	idivl	%ecx
	cmpl	$1001, %ecx                     # imm = 0x3E9
	jl	.LBB16_2
# %bb.1:
	movl	$1000000, %eax                  # imm = 0xF4240
	xorl	%edx, %edx
	divl	%ecx
.LBB16_2:
	movl	%eax, %edi
	btsq	$32, %rdi
	btsq	$32, %rcx
	movl	$1, %esi
	movq	%rcx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB16_4
# %bb.3:
	movq	128(%rsp), %rdi
	movq	120(%rsp), %rsi
	movq	112(%rsp), %rdx
	movq	104(%rsp), %rcx
	movq	96(%rsp), %r8
	movq	72(%rsp), %r9
	subq	$8, %rsp
	.cfi_adjust_cfa_offset 8
	pushq	96(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	24(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	40(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	56(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	72(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	88(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	104(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	120(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	136(%rsp)
	.cfi_adjust_cfa_offset 8
	callq	_Z24__device_stub__cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_
	addq	$80, %rsp
	.cfi_adjust_cfa_offset -80
.LBB16_4:
	movq	88(%rsp), %rsi
	movl	$8000000, %edx                  # imm = 0x7A1200
	movq	%r14, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	%r14, %rdi
	callq	free
	movq	128(%rsp), %rdi
	callq	hipFree
	movq	120(%rsp), %rdi
	callq	hipFree
	movq	112(%rsp), %rdi
	callq	hipFree
	movq	104(%rsp), %rdi
	callq	hipFree
	movq	96(%rsp), %rdi
	callq	hipFree
	movq	72(%rsp), %rdi
	callq	hipFree
	movq	64(%rsp), %rdi
	callq	hipFree
	movq	56(%rsp), %rdi
	callq	hipFree
	movq	48(%rsp), %rdi
	callq	hipFree
	movq	40(%rsp), %rdi
	callq	hipFree
	movq	32(%rsp), %rdi
	callq	hipFree
	movq	24(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	movq	8(%rsp), %rdi
	callq	hipFree
	xorl	%eax, %eax
	addq	$136, %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_end16:
	.size	main, .Lfunc_end16-main
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB17_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB17_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z9cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_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_end17:
	.size	__hip_module_ctor, .Lfunc_end17-__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	.LBB18_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
.LBB18_2:
	retq
.Lfunc_end18:
	.size	__hip_module_dtor, .Lfunc_end18-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	LnRnBlocks,@object              # @LnRnBlocks
	.bss
	.globl	LnRnBlocks
	.p2align	4, 0x0
LnRnBlocks:
	.zero	272
	.size	LnRnBlocks, 272

	.type	CnDnBlocks,@object              # @CnDnBlocks
	.globl	CnDnBlocks
	.p2align	4, 0x0
CnDnBlocks:
	.zero	272
	.size	CnDnBlocks, 272

	.type	keysBlocks,@object              # @keysBlocks
	.globl	keysBlocks
	.p2align	4, 0x0
keysBlocks:
	.zero	128
	.size	keysBlocks, 128

	.type	allCipherDES,@object            # @allCipherDES
	.globl	allCipherDES
	.p2align	4, 0x0
allCipherDES:
	.zero	8000000
	.size	allCipherDES, 8000000

	.type	Rotations,@object               # @Rotations
	.data
	.globl	Rotations
	.p2align	4, 0x0
Rotations:
	.quad	1                               # 0x1
	.quad	1                               # 0x1
	.quad	2                               # 0x2
	.quad	2                               # 0x2
	.quad	2                               # 0x2
	.quad	2                               # 0x2
	.quad	2                               # 0x2
	.quad	2                               # 0x2
	.quad	1                               # 0x1
	.quad	2                               # 0x2
	.quad	2                               # 0x2
	.quad	2                               # 0x2
	.quad	2                               # 0x2
	.quad	2                               # 0x2
	.quad	2                               # 0x2
	.quad	1                               # 0x1
	.size	Rotations, 128

	.type	PC1,@object                     # @PC1
	.globl	PC1
	.p2align	4, 0x0
PC1:
	.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
	.long	27                              # 0x1b
	.long	19                              # 0x13
	.long	11                              # 0xb
	.long	3                               # 0x3
	.long	60                              # 0x3c
	.long	52                              # 0x34
	.long	44                              # 0x2c
	.long	36                              # 0x24
	.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	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	28                              # 0x1c
	.long	20                              # 0x14
	.long	12                              # 0xc
	.long	4                               # 0x4
	.size	PC1, 224

	.type	PC2,@object                     # @PC2
	.globl	PC2
	.p2align	4, 0x0
PC2:
	.long	14                              # 0xe
	.long	17                              # 0x11
	.long	11                              # 0xb
	.long	24                              # 0x18
	.long	1                               # 0x1
	.long	5                               # 0x5
	.long	3                               # 0x3
	.long	28                              # 0x1c
	.long	15                              # 0xf
	.long	6                               # 0x6
	.long	21                              # 0x15
	.long	10                              # 0xa
	.long	23                              # 0x17
	.long	19                              # 0x13
	.long	12                              # 0xc
	.long	4                               # 0x4
	.long	26                              # 0x1a
	.long	8                               # 0x8
	.long	16                              # 0x10
	.long	7                               # 0x7
	.long	27                              # 0x1b
	.long	20                              # 0x14
	.long	13                              # 0xd
	.long	2                               # 0x2
	.long	41                              # 0x29
	.long	52                              # 0x34
	.long	31                              # 0x1f
	.long	37                              # 0x25
	.long	47                              # 0x2f
	.long	55                              # 0x37
	.long	30                              # 0x1e
	.long	40                              # 0x28
	.long	51                              # 0x33
	.long	45                              # 0x2d
	.long	33                              # 0x21
	.long	48                              # 0x30
	.long	44                              # 0x2c
	.long	49                              # 0x31
	.long	39                              # 0x27
	.long	56                              # 0x38
	.long	34                              # 0x22
	.long	53                              # 0x35
	.long	46                              # 0x2e
	.long	42                              # 0x2a
	.long	50                              # 0x32
	.long	36                              # 0x24
	.long	29                              # 0x1d
	.long	32                              # 0x20
	.size	PC2, 192

	.type	IniPer,@object                  # @IniPer
	.globl	IniPer
	.p2align	4, 0x0
IniPer:
	.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
	.long	64                              # 0x40
	.long	56                              # 0x38
	.long	48                              # 0x30
	.long	40                              # 0x28
	.long	32                              # 0x20
	.long	24                              # 0x18
	.long	16                              # 0x10
	.long	8                               # 0x8
	.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
	.size	IniPer, 256

	.type	reverseIniPer,@object           # @reverseIniPer
	.globl	reverseIniPer
	.p2align	4, 0x0
reverseIniPer:
	.long	40                              # 0x28
	.long	8                               # 0x8
	.long	48                              # 0x30
	.long	16                              # 0x10
	.long	56                              # 0x38
	.long	24                              # 0x18
	.long	64                              # 0x40
	.long	32                              # 0x20
	.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
	.size	reverseIniPer, 256

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

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

	.type	Sbox,@object                    # @Sbox
	.globl	Sbox
	.p2align	4, 0x0
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	Sbox, 2048

	.type	iniKey,@object                  # @iniKey
	.globl	iniKey
	.p2align	4, 0x0
iniKey:
	.quad	19                              # 0x13
	.quad	52                              # 0x34
	.quad	87                              # 0x57
	.quad	121                             # 0x79
	.quad	155                             # 0x9b
	.quad	188                             # 0xbc
	.quad	223                             # 0xdf
	.quad	241                             # 0xf1
	.size	iniKey, 64

	.type	message,@object                 # @message
	.globl	message
	.p2align	4, 0x0
message:
	.quad	1                               # 0x1
	.quad	35                              # 0x23
	.quad	69                              # 0x45
	.quad	103                             # 0x67
	.quad	137                             # 0x89
	.quad	171                             # 0xab
	.quad	205                             # 0xcd
	.quad	239                             # 0xef
	.size	message, 64

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

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

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z9cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_"
	.size	.L__unnamed_1, 49

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

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

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym _Z24__device_stub__cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym LnRnBlocks
	.addrsig_sym CnDnBlocks
	.addrsig_sym keysBlocks
	.addrsig_sym allCipherDES
	.addrsig_sym Rotations
	.addrsig_sym PC1
	.addrsig_sym PC2
	.addrsig_sym IniPer
	.addrsig_sym reverseIniPer
	.addrsig_sym Expansion
	.addrsig_sym Pbox
	.addrsig_sym Sbox
	.addrsig_sym iniKey
	.addrsig_sym message
	.addrsig_sym _Z9cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z9cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_ ; -- Begin function _Z9cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_
	.globl	_Z9cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_
	.p2align	8
	.type	_Z9cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_,@function
_Z9cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_: ; @_Z9cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_
; %bb.0:                                ; %_Z15generateKeyPlusPyPi.exit.preheader
	s_clause 0x2
	s_load_b64 s[6:7], s[0:1], 0x20
	s_load_b128 s[16:19], s[0:1], 0x0
	s_load_b64 s[2:3], s[0:1], 0x10
	v_mov_b32_e32 v7, 0
	s_mov_b32 s4, 1
	s_waitcnt lgkmcnt(0)
	s_add_u32 s14, s6, -8
	s_addc_u32 s26, s7, -1
	s_mov_b32 s7, 0
.LBB0_1:                                ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_3 Depth 2
                                        ;     Child Loop BB0_10 Depth 2
	s_lshl_b32 s6, s4, 1
	s_mov_b32 s5, s7
	s_lshl_b64 s[8:9], s[6:7], 3
	s_mov_b64 s[12:13], 0
	s_add_u32 s8, s18, s8
	s_addc_u32 s9, s19, s9
	s_lshl_b64 s[10:11], s[4:5], 3
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s10, s14, s10
	s_addc_u32 s11, s26, s11
	s_clause 0x1
	global_load_b64 v[5:6], v7, s[10:11]
	global_load_b128 v[1:4], v7, s[8:9] offset:-16
	s_mov_b64 s[10:11], 0
	s_waitcnt vmcnt(1)
	v_cmp_eq_u64_e32 vcc_lo, 0, v[5:6]
	v_readfirstlane_b32 s20, v5
	v_readfirstlane_b32 s21, v6
	s_cbranch_vccnz .LBB0_8
; %bb.2:                                ; %.lr.ph.i
                                        ;   in Loop: Header=BB0_1 Depth=1
	s_delay_alu instid0(VALU_DEP_1)
	s_cmp_eq_u64 s[20:21], 1
	s_mov_b64 s[22:23], 0
	s_mov_b64 s[24:25], 1
	s_cselect_b32 s5, -1, 0
.LBB0_3:                                ;   Parent Loop BB0_1 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_add_i32 s6, s24, 26
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_lshl_b32 s6, 1, s6
	s_ashr_i32 s27, s6, 31
	s_waitcnt vmcnt(0)
	v_and_b32_e32 v8, s6, v1
	v_and_b32_e32 v9, s27, v2
	s_delay_alu instid0(VALU_DEP_1)
	v_cmp_eq_u64_e32 vcc_lo, 0, v[8:9]
	s_cbranch_vccnz .LBB0_5
; %bb.4:                                ;   in Loop: Header=BB0_3 Depth=2
	s_and_b32 s28, s5, exec_lo
	s_cselect_b32 s28, s22, s24
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_lshl_b64 s[28:29], 1, s28
	s_or_b64 s[12:13], s[28:29], s[12:13]
.LBB0_5:                                ;   in Loop: Header=BB0_3 Depth=2
	v_and_b32_e32 v8, s6, v3
	v_and_b32_e32 v9, s27, v4
	s_delay_alu instid0(VALU_DEP_1)
	v_cmp_eq_u64_e32 vcc_lo, 0, v[8:9]
	s_cbranch_vccnz .LBB0_7
; %bb.6:                                ;   in Loop: Header=BB0_3 Depth=2
	s_and_b32 s6, s5, exec_lo
	s_cselect_b32 s6, s22, s24
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_lshl_b64 s[28:29], 1, s6
	s_or_b64 s[10:11], s[28:29], s[10:11]
.LBB0_7:                                ;   in Loop: Header=BB0_3 Depth=2
	s_add_u32 s22, s22, 1
	s_addc_u32 s23, s23, 0
	s_add_u32 s24, s24, -1
	s_addc_u32 s25, s25, -1
	s_cmp_eq_u64 s[20:21], s[22:23]
	s_cbranch_scc0 .LBB0_3
.LBB0_8:                                ; %Flow143
                                        ;   in Loop: Header=BB0_1 Depth=1
	v_cmp_ne_u64_e32 vcc_lo, 0, v[5:6]
	s_waitcnt vmcnt(0)
	v_lshlrev_b64 v[5:6], s20, v[1:2]
	v_lshlrev_b64 v[1:2], s20, v[3:4]
	s_and_not1_b32 vcc_lo, exec_lo, vcc_lo
	s_cbranch_vccnz .LBB0_11
; %bb.9:                                ; %.lr.ph70.i.preheader
                                        ;   in Loop: Header=BB0_1 Depth=1
	s_mov_b64 s[22:23], 0
.LBB0_10:                               ; %.lr.ph70.i
                                        ;   Parent Loop BB0_1 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_add_i32 s5, s22, 28
	s_lshl_b32 s5, 1, s5
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_not_b32 s5, s5
	s_ashr_i32 s6, s5, 31
	v_and_b32_e32 v5, s5, v5
	v_and_b32_e32 v1, s5, v1
	s_add_u32 s22, s22, 1
	v_and_b32_e32 v6, s6, v6
	v_and_b32_e32 v2, s6, v2
	s_addc_u32 s23, s23, 0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_eq_u64 s[20:21], s[22:23]
	s_cbranch_scc0 .LBB0_10
.LBB0_11:                               ; %Flow141
                                        ;   in Loop: Header=BB0_1 Depth=1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_or_b32_e32 v3, s12, v5
	v_or_b32_e32 v4, s13, v6
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	v_or_b32_e32 v5, s10, v1
	v_or_b32_e32 v6, s11, v2
	s_add_i32 s4, s4, 1
	s_cmp_eq_u32 s4, 17
	global_store_b128 v7, v[3:6], s[8:9]
	s_cbranch_scc0 .LBB0_1
; %bb.12:                               ; %_Z18generateCnDnBlocksPyS_S_.exit
	s_load_b64 s[4:5], s[0:1], 0x30
	s_add_u32 s14, s2, -8
	v_mov_b32_e32 v5, 0
	s_addc_u32 s28, s3, -1
	s_mov_b32 s2, 1
	s_mov_b32 s7, 0
	s_waitcnt lgkmcnt(0)
	s_add_u32 s4, s4, 0xbc
	s_addc_u32 s5, s5, 0
.LBB0_13:                               ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_14 Depth 2
	s_lshl_b32 s6, s2, 1
	s_mov_b64 s[10:11], s[4:5]
	s_lshl_b64 s[8:9], s[6:7], 3
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s8, s18, s8
	s_addc_u32 s9, s19, s9
	global_load_b128 v[1:4], v5, s[8:9]
	s_mov_b64 s[8:9], 0
	s_waitcnt vmcnt(0)
	v_lshlrev_b64 v[6:7], 28, v[1:2]
	v_mov_b32_e32 v1, 0
	v_mov_b32_e32 v2, 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_or_b32_e32 v3, v6, v3
	v_or_b32_e32 v4, v7, v4
.LBB0_14:                               ;   Parent Loop BB0_13 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_load_b32 s3, s[10:11], 0x0
	s_lshl_b64 s[12:13], 1, s8
	s_waitcnt lgkmcnt(0)
	s_sub_i32 s3, 56, s3
	s_add_u32 s8, s8, 1
	v_lshrrev_b64 v[6:7], s3, v[3:4]
	s_addc_u32 s9, s9, 0
	s_add_u32 s10, s10, -4
	s_addc_u32 s11, s11, -1
	s_cmp_eq_u32 s8, 48
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_and_b32_e32 v6, 1, v6
	v_cmp_eq_u32_e32 vcc_lo, 1, v6
	v_cndmask_b32_e64 v6, 0, s13, vcc_lo
	v_cndmask_b32_e64 v7, 0, s12, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_or_b32_e32 v2, v6, v2
	v_or_b32_e32 v1, v7, v1
	s_cbranch_scc0 .LBB0_14
; %bb.15:                               ;   in Loop: Header=BB0_13 Depth=1
	s_mov_b32 s3, s7
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_lshl_b64 s[8:9], s[2:3], 3
	s_add_u32 s8, s14, s8
	s_addc_u32 s9, s28, s9
	s_add_i32 s2, s2, 1
	global_store_b64 v5, v[1:2], s[8:9]
	s_cmp_lg_u32 s2, 17
	s_cbranch_scc1 .LBB0_13
; %bb.16:                               ; %_Z11splitIniPery.exit.preheader
	s_load_b256 s[4:11], s[0:1], 0x40
	v_mov_b32_e32 v30, 0
	s_mov_b32 s12, 1
	s_mov_b32 s3, 0
                                        ; implicit-def: $vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7_vgpr8_vgpr9_vgpr10_vgpr11_vgpr12_vgpr13_vgpr14_vgpr15_vgpr16
	s_waitcnt lgkmcnt(0)
	s_add_u32 s6, s6, 0xbc
	s_addc_u32 s7, s7, 0
	s_add_u32 s8, s8, 0x7c
	s_addc_u32 s9, s9, 0
.LBB0_17:                               ; %_Z11splitIniPery.exit
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_18 Depth 2
                                        ;     Child Loop BB0_20 Depth 2
                                        ;     Child Loop BB0_22 Depth 2
                                        ;     Child Loop BB0_24 Depth 2
                                        ;     Child Loop BB0_26 Depth 2
	s_lshl_b32 s2, s12, 1
	v_mov_b32_e32 v17, 0
	s_lshl_b64 s[18:19], s[2:3], 3
	v_mov_b32_e32 v18, 0
	s_add_u32 s18, s16, s18
	s_addc_u32 s19, s17, s19
	s_mov_b64 s[20:21], 0
	global_load_b128 v[25:28], v30, s[18:19] offset:-16
	s_mov_b64 s[22:23], s[6:7]
.LBB0_18:                               ;   Parent Loop BB0_17 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_load_b32 s2, s[22:23], 0x0
	s_lshl_b64 s[24:25], 1, s20
	s_waitcnt lgkmcnt(0)
	s_sub_i32 s2, 32, s2
	s_add_u32 s20, s20, 1
	s_waitcnt vmcnt(0)
	v_lshrrev_b64 v[19:20], s2, v[27:28]
	s_addc_u32 s21, s21, 0
	s_add_u32 s22, s22, -4
	s_addc_u32 s23, s23, -1
	s_cmp_lg_u32 s20, 48
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_and_b32_e32 v19, 1, v19
	v_cmp_eq_u32_e32 vcc_lo, 1, v19
	v_cndmask_b32_e64 v19, 0, s25, vcc_lo
	v_cndmask_b32_e64 v20, 0, s24, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_or_b32_e32 v18, v19, v18
	v_or_b32_e32 v17, v20, v17
	s_cbranch_scc1 .LBB0_18
; %bb.19:                               ; %_Z8expandRnyPi.exit.i
                                        ;   in Loop: Header=BB0_17 Depth=1
	s_mov_b32 s13, s3
	v_mov_b32_e32 v29, 0
	s_lshl_b64 s[20:21], s[12:13], 3
	s_mov_b32 s13, 7
	s_add_u32 s20, s14, s20
	s_addc_u32 s21, s28, s21
	global_load_b64 v[19:20], v30, s[20:21]
	s_mov_b64 s[20:21], 0
	s_waitcnt vmcnt(0)
	v_xor_b32_e32 v31, v19, v17
	v_xor_b32_e32 v32, v20, v18
                                        ; implicit-def: $vgpr17
.LBB0_20:                               ;   Parent Loop BB0_17 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	v_and_b32_e64 v43, 0xff, s20
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_4)
	v_lshrrev_b64 v[41:42], s20, v[31:32]
	v_dual_mov_b32 v40, v24 :: v_dual_mov_b32 v39, v23
	v_dual_mov_b32 v34, v18 :: v_dual_mov_b32 v33, v17
	v_mul_lo_u16 v42, 0xab, v43
	s_mov_b32 m0, s13
	v_dual_mov_b32 v38, v22 :: v_dual_mov_b32 v37, v21
	v_dual_mov_b32 v36, v20 :: v_dual_mov_b32 v35, v19
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_3) | instid1(VALU_DEP_3)
	v_lshrrev_b16 v42, 10, v42
	v_movreld_b32_e32 v33, v29
	v_and_b32_e32 v41, 1, v41
	s_cmp_gt_u32 s20, 5
	v_mul_lo_u16 v42, v42, 6
	s_cselect_b32 s22, -1, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_eq_u32_e32 vcc_lo, 1, v41
	v_sub_nc_u16 v42, s20, v42
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_and_b32_e32 v43, 0xff, v42
	v_lshlrev_b32_e64 v42, v42, 1
	v_cmp_eq_u16_e64 s2, 0, v43
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v41, 0, v42, vcc_lo
	s_and_b32 vcc_lo, s22, s2
	v_cndmask_b32_e32 v20, v20, v36, vcc_lo
	v_cndmask_b32_e64 v29, v29, 0, vcc_lo
	s_cmp_lg_u32 vcc_lo, 0
	v_dual_cndmask_b32 v24, v24, v40 :: v_dual_cndmask_b32 v23, v23, v39
	v_dual_cndmask_b32 v22, v22, v38 :: v_dual_cndmask_b32 v21, v21, v37
	v_dual_cndmask_b32 v18, v18, v34 :: v_dual_cndmask_b32 v19, v19, v35
	v_cndmask_b32_e32 v17, v17, v33, vcc_lo
	s_subb_u32 s13, s13, 0
	v_or_b32_e32 v29, v29, v41
	s_add_u32 s20, s20, 1
	s_addc_u32 s21, s21, 0
	s_cmp_lg_u32 s20, 49
	s_cbranch_scc1 .LBB0_20
; %bb.21:                               ; %.preheader31.i.i.preheader
                                        ;   in Loop: Header=BB0_17 Depth=1
	s_mov_b64 s[20:21], 0
	s_mov_b32 s2, 0
.LBB0_22:                               ; %.preheader31.i.i
                                        ;   Parent Loop BB0_17 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_mov_b32 m0, s20
	v_movrels_b32_e32 v29, v17
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_lshlrev_b32_e32 v31, 4, v29
	v_and_b32_e32 v32, 32, v29
	v_bfe_u32 v29, v29, 1, 4
	v_and_b32_e32 v31, 16, v31
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_or_b32_e32 v31, v32, v31
	v_or_b32_e32 v29, v31, v29
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v29, s2, v29
	v_lshlrev_b64 v[31:32], 2, v[29:30]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v29, vcc_lo, s10, v31
	v_add_co_ci_u32_e32 v31, vcc_lo, s11, v32, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_readfirstlane_b32 s22, v29
	v_readfirstlane_b32 s23, v31
	s_load_b32 s13, s[22:23], 0x0
	s_lshl_b32 s22, s20, 1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mov_b32 m0, s22
	s_waitcnt lgkmcnt(0)
	v_movreld_b32_e32 v1, s13
	s_ashr_i32 s13, s13, 31
	s_add_u32 s20, s20, 1
	s_addc_u32 s21, s21, 0
	s_add_i32 s2, s2, 64
	s_cmp_lg_u32 s20, 8
	v_movreld_b32_e32 v2, s13
	s_cbranch_scc1 .LBB0_22
; %bb.23:                               ; %.preheader.i.i.preheader
                                        ;   in Loop: Header=BB0_17 Depth=1
	s_mov_b64 s[20:21], 0
	s_mov_b64 s[22:23], 28
	s_mov_b32 s2, 0
.LBB0_24:                               ; %.preheader.i.i
                                        ;   Parent Loop BB0_17 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_lshl_b32 s13, s2, 1
	s_mov_b32 m0, s13
	v_movrels_b32_e32 v17, v1
	v_movrels_b32_e32 v18, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_readfirstlane_b32 s24, v17
	v_readfirstlane_b32 s25, v18
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_lshl_b64 s[24:25], s[24:25], s22
	s_or_b64 s[20:21], s[24:25], s[20:21]
	s_add_u32 s22, s22, -4
	s_addc_u32 s23, s23, -1
	s_add_i32 s2, s2, 1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_lg_u32 s2, 8
	s_cbranch_scc1 .LBB0_24
; %bb.25:                               ; %_Z23generateSboxCombinationyPi.exit.i.preheader
                                        ;   in Loop: Header=BB0_17 Depth=1
	s_mov_b64 s[22:23], 0
	s_mov_b64 s[26:27], s[8:9]
	s_mov_b64 s[24:25], 0
.LBB0_26:                               ; %_Z23generateSboxCombinationyPi.exit.i
                                        ;   Parent Loop BB0_17 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_load_b32 s2, s[26:27], 0x0
	s_lshl_b64 s[30:31], 1, s22
	s_waitcnt lgkmcnt(0)
	s_sub_i32 s2, 32, s2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_lshr_b64 s[34:35], s[20:21], s2
	s_bitcmp1_b32 s34, 0
	s_cselect_b32 s31, s31, 0
	s_cselect_b32 s30, s30, 0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b64 s[24:25], s[30:31], s[24:25]
	s_add_u32 s22, s22, 1
	s_addc_u32 s23, s23, 0
	s_add_u32 s26, s26, -4
	s_addc_u32 s27, s27, -1
	s_cmp_lg_u32 s22, 32
	s_cbranch_scc1 .LBB0_26
; %bb.27:                               ; %_Z18generateFalgorithmyPi.exit.i
                                        ;   in Loop: Header=BB0_17 Depth=1
	v_xor_b32_e32 v19, s24, v25
	v_xor_b32_e32 v20, s25, v26
	v_dual_mov_b32 v17, v27 :: v_dual_mov_b32 v18, v28
	s_add_i32 s12, s12, 1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_lg_u32 s12, 17
	global_store_b128 v30, v[17:20], s[18:19]
	s_cbranch_scc1 .LBB0_17
; %bb.28:                               ; %_Z18generateLnRnBlocksPyS_S_PiS0_S0_.exit
	v_mov_b32_e32 v1, 0
	s_mov_b32 s8, 0
	s_mov_b64 s[6:7], 0
	global_load_b128 v[1:4], v1, s[16:17] offset:256
	s_waitcnt vmcnt(0)
	v_readfirstlane_b32 s3, v2
	v_readfirstlane_b32 s2, v1
	v_readfirstlane_b32 s9, v3
	s_delay_alu instid0(VALU_DEP_1)
	s_or_b64 s[8:9], s[8:9], s[2:3]
	s_add_u32 s4, s4, 0xfc
	s_addc_u32 s5, s5, 0
	s_mov_b64 s[2:3], 0
.LBB0_29:                               ; =>This Inner Loop Header: Depth=1
	s_load_b32 s12, s[4:5], 0x0
	s_lshl_b64 s[10:11], 1, s6
	s_waitcnt lgkmcnt(0)
	s_sub_i32 s12, 64, s12
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_lshr_b64 s[12:13], s[8:9], s12
	s_bitcmp1_b32 s12, 0
	s_cselect_b32 s11, s11, 0
	s_cselect_b32 s10, s10, 0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b64 s[2:3], s[10:11], s[2:3]
	s_add_u32 s6, s6, 1
	s_addc_u32 s7, s7, 0
	s_add_u32 s4, s4, -4
	s_addc_u32 s5, s5, -1
	s_cmp_lg_u32 s6, 64
	s_cbranch_scc1 .LBB0_29
; %bb.30:                               ; %_Z21generateCipherMessageyPi.exit
	s_clause 0x1
	s_load_b32 s4, s[0:1], 0x84
	s_load_b64 s[0:1], s[0:1], 0x70
	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_ashrrev_i32_e32 v2, 31, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[0:1], 3, v[1:2]
	v_dual_mov_b32 v2, s2 :: v_dual_mov_b32 v3, s3
	v_add_co_u32 v0, vcc_lo, s0, v0
	s_delay_alu instid0(VALU_DEP_3)
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	global_store_b64 v[0:1], v[2:3], off
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z9cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 376
		.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 44
		.amdhsa_next_free_sgpr 36
		.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	_Z9cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_, .Lfunc_end0-_Z9cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 1720
; NumSgprs: 38
; NumVgprs: 44
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 4
; VGPRBlocks: 5
; NumSGPRsForWavesPerEU: 38
; NumVGPRsForWavesPerEU: 44
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         32
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         40
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         48
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         56
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         64
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         72
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         80
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         88
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         96
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         104
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         112
        .size:           8
        .value_kind:     global_buffer
      - .offset:         120
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         124
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         128
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         132
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         134
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         136
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         138
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         140
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         142
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         160
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         168
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         176
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         184
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 376
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z9cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_
    .private_segment_fixed_size: 0
    .sgpr_count:     38
    .sgpr_spill_count: 0
    .symbol:         _Z9cipherDESPyS_S_S_S_PiS0_S0_S0_S0_S0_S0_S_S_S_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     44
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
