unsigned long permute_pc1(unsigned long src,int tid){
unsigned long dst = 0;
int srcPos = 0;
unsigned char permutation[] = {
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
};
if(tid<56)
{
srcPos = 64 - permutation[tid];
dst = (src>>srcPos & 0x01)<<(55-tid);
}
return dst;
}
__kernel void encrypt(__global unsigned long *message, __global unsigned long *key, __global unsigned long *encrypted)
{
unsigned int tid = get_global_id(0);
//permutarea initiala -----------------------------------------------------------
unsigned long bit = permute_pc1(key[0],tid);
__local unsigned long result[64],perm_key;
result[tid] = 0;
result[tid] = result[tid] | bit;
barrier(CLK_LOCAL_MEM_FENCE);
if (tid == 0) {
for (int i=0;i<56;i++) {
perm_key = perm_key | result[i];
}
encrypted[0] = perm_key;
}
/*atomic_or()*/
//end of permutarea initiala ----------------------------------------------------
//expandarea propriu-zisa -------------------------------------------------------
__local unsigned long cd[16];
if(tid==0)
{
unsigned int c[17],d[17];
c[0] = (unsigned int)(perm_key>>28);
d[0] = (unsigned int)(perm_key & 0x0fffffff);
for(int i=1;i<=16;i++)
{
if((i==1)||(i==2)||(i==9)||(i==16))
{
c[i] = ((c[i-1]<<1)& 0x0FFFFFFF)|(c[i-1]>>27);
d[i] = ((d[i-1]<<1)& 0x0FFFFFFF)|(d[i-1]>>27);
}
else
{
c[i] = ((c[i-1]<<2)& 0x0FFFFFFF)|(c[i-1]>>26);
d[i] = ((d[i-1]<<2)& 0x0FFFFFFF)|(d[i-1]>>26);
}
cd[i-1] = ((unsigned long)c[i] << 28)|d[i];
}
}
//end of expandarea propriu-zisa -------------------------------------------------
}