Align all the __constant buffers to workaround OpenCL JIT compiler errors in NV drivers 378.x

pull/1134/head
jsteube 7 years ago
parent 3fb433de60
commit 6401c58568

@ -3,7 +3,7 @@
* License.....: MIT
*/
__constant u32 te0[256] =
__constant u32a te0[256] =
{
0xc66363a5, 0xf87c7c84, 0xee777799, 0xf67b7b8d,
0xfff2f20d, 0xd66b6bbd, 0xde6f6fb1, 0x91c5c554,
@ -71,7 +71,7 @@ __constant u32 te0[256] =
0x7bb0b0cb, 0xa85454fc, 0x6dbbbbd6, 0x2c16163a,
};
__constant u32 te1[256] =
__constant u32a te1[256] =
{
0xa5c66363, 0x84f87c7c, 0x99ee7777, 0x8df67b7b,
0x0dfff2f2, 0xbdd66b6b, 0xb1de6f6f, 0x5491c5c5,
@ -139,7 +139,7 @@ __constant u32 te1[256] =
0xcb7bb0b0, 0xfca85454, 0xd66dbbbb, 0x3a2c1616,
};
__constant u32 te2[256] =
__constant u32a te2[256] =
{
0x63a5c663, 0x7c84f87c, 0x7799ee77, 0x7b8df67b,
0xf20dfff2, 0x6bbdd66b, 0x6fb1de6f, 0xc55491c5,
@ -207,7 +207,7 @@ __constant u32 te2[256] =
0xb0cb7bb0, 0x54fca854, 0xbbd66dbb, 0x163a2c16,
};
__constant u32 te3[256] =
__constant u32a te3[256] =
{
0x6363a5c6, 0x7c7c84f8, 0x777799ee, 0x7b7b8df6,
0xf2f20dff, 0x6b6bbdd6, 0x6f6fb1de, 0xc5c55491,
@ -275,7 +275,7 @@ __constant u32 te3[256] =
0xb0b0cb7b, 0x5454fca8, 0xbbbbd66d, 0x16163a2c,
};
__constant u32 te4[256] =
__constant u32a te4[256] =
{
0x63636363, 0x7c7c7c7c, 0x77777777, 0x7b7b7b7b,
0xf2f2f2f2, 0x6b6b6b6b, 0x6f6f6f6f, 0xc5c5c5c5,
@ -343,7 +343,7 @@ __constant u32 te4[256] =
0xb0b0b0b0, 0x54545454, 0xbbbbbbbb, 0x16161616,
};
__constant u32 td0[256] =
__constant u32a td0[256] =
{
0x51f4a750, 0x7e416553, 0x1a17a4c3, 0x3a275e96,
0x3bab6bcb, 0x1f9d45f1, 0xacfa58ab, 0x4be30393,
@ -411,7 +411,7 @@ __constant u32 td0[256] =
0x7bcb8461, 0xd532b670, 0x486c5c74, 0xd0b85742,
};
__constant u32 td1[256] =
__constant u32a td1[256] =
{
0x5051f4a7, 0x537e4165, 0xc31a17a4, 0x963a275e,
0xcb3bab6b, 0xf11f9d45, 0xabacfa58, 0x934be303,
@ -479,7 +479,7 @@ __constant u32 td1[256] =
0x617bcb84, 0x70d532b6, 0x74486c5c, 0x42d0b857,
};
__constant u32 td2[256] =
__constant u32a td2[256] =
{
0xa75051f4, 0x65537e41, 0xa4c31a17, 0x5e963a27,
0x6bcb3bab, 0x45f11f9d, 0x58abacfa, 0x03934be3,
@ -547,7 +547,7 @@ __constant u32 td2[256] =
0x84617bcb, 0xb670d532, 0x5c74486c, 0x5742d0b8,
};
__constant u32 td3[256] =
__constant u32a td3[256] =
{
0xf4a75051, 0x4165537e, 0x17a4c31a, 0x275e963a,
0xab6bcb3b, 0x9d45f11f, 0xfa58abac, 0xe303934b,
@ -615,7 +615,7 @@ __constant u32 td3[256] =
0xcb84617b, 0x32b670d5, 0x6c5c7448, 0xb85742d0,
};
__constant u32 td4[256] =
__constant u32a td4[256] =
{
0x52525252, 0x09090909, 0x6a6a6a6a, 0xd5d5d5d5,
0x30303030, 0x36363636, 0xa5a5a5a5, 0x38383838,
@ -683,7 +683,7 @@ __constant u32 td4[256] =
0x55555555, 0x21212121, 0x0c0c0c0c, 0x7d7d7d7d,
};
__constant u32 rcon[] =
__constant u32a rcon[] =
{
0x01000000, 0x02000000, 0x04000000, 0x08000000,
0x10000000, 0x20000000, 0x40000000, 0x80000000,

@ -21,7 +21,7 @@
#define extract_byte(x,n) (((x) >> (8 * (n))) & 0xff)
__constant u32 q_tab[2][256] =
__constant u32a q_tab[2][256] =
{
{
0xA9, 0x67, 0xB3, 0xE8, 0x04, 0xFD, 0xA3, 0x76, 0x9A, 0x92, 0x80, 0x78,
@ -75,7 +75,7 @@ __constant u32 q_tab[2][256] =
#define q(n,x) q_tab[n][x]
__constant u32 m_tab[4][256] =
__constant u32a m_tab[4][256] =
{
{ 0xBCBC3275, 0xECEC21F3, 0x202043C6, 0xB3B3C9F4, 0xDADA03DB, 0x02028B7B,
0xE2E22BFB, 0x9E9EFAC8, 0xC9C9EC4A, 0xD4D409D3, 0x18186BE6, 0x1E1E9F6B,

@ -1,7 +1,7 @@
// basically normal XXX_transform() but with a different name to avoid collisions with function nameing
__constant u32 AF_k_sha256[64] =
__constant u32a AF_k_sha256[64] =
{
SHA256C00, SHA256C01, SHA256C02, SHA256C03,
SHA256C04, SHA256C05, SHA256C06, SHA256C07,
@ -21,7 +21,7 @@ __constant u32 AF_k_sha256[64] =
SHA256C3c, SHA256C3d, SHA256C3e, SHA256C3f,
};
__constant u64 AF_k_sha512[80] =
__constant u64a AF_k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,

@ -1,4 +1,4 @@
__constant u32 ESSIV_k_sha256[64] =
__constant u32a ESSIV_k_sha256[64] =
{
SHA256C00, SHA256C01, SHA256C02, SHA256C03,
SHA256C04, SHA256C05, SHA256C06, SHA256C07,

@ -1,4 +1,4 @@
__constant u32 crc32tab[0x100] =
__constant u32a crc32tab[0x100] =
{
0x00000000, 0x77073096, 0xee0e612c, 0x990951ba,
0x076dc419, 0x706af48f, 0xe963a535, 0x9e6495a3,

@ -16,7 +16,7 @@
#include "inc_rp.cl"
#include "inc_simd.cl"
__constant u32 k_sha256[64] =
__constant u32a k_sha256[64] =
{
SHA256C00, SHA256C01, SHA256C02, SHA256C03,
SHA256C04, SHA256C05, SHA256C06, SHA256C07,

@ -14,7 +14,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u32 k_sha256[64] =
__constant u32a k_sha256[64] =
{
SHA256C00, SHA256C01, SHA256C02, SHA256C03,
SHA256C04, SHA256C05, SHA256C06, SHA256C07,

@ -14,7 +14,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u32 k_sha256[64] =
__constant u32a k_sha256[64] =
{
SHA256C00, SHA256C01, SHA256C02, SHA256C03,
SHA256C04, SHA256C05, SHA256C06, SHA256C07,

@ -16,7 +16,7 @@
#include "inc_rp.cl"
#include "inc_simd.cl"
__constant u32 k_sha256[64] =
__constant u32a k_sha256[64] =
{
SHA256C00, SHA256C01, SHA256C02, SHA256C03,
SHA256C04, SHA256C05, SHA256C06, SHA256C07,

@ -14,7 +14,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u32 k_sha256[64] =
__constant u32a k_sha256[64] =
{
SHA256C00, SHA256C01, SHA256C02, SHA256C03,
SHA256C04, SHA256C05, SHA256C06, SHA256C07,

@ -14,7 +14,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u32 k_sha256[64] =
__constant u32a k_sha256[64] =
{
SHA256C00, SHA256C01, SHA256C02, SHA256C03,
SHA256C04, SHA256C05, SHA256C06, SHA256C07,

@ -36,7 +36,7 @@
a = a ^ tt; \
}
__constant u32 c_SPtrans[8][64] =
__constant u32a c_SPtrans[8][64] =
{
{
0x00820200, 0x00020000, 0x80800000, 0x80820200,
@ -184,7 +184,7 @@ __constant u32 c_SPtrans[8][64] =
},
};
__constant u32 c_skb[8][64] =
__constant u32a c_skb[8][64] =
{
{
0x00000000, 0x00000010, 0x20000000, 0x20000010,

@ -34,7 +34,7 @@
a = a ^ tt; \
}
__constant u32 c_SPtrans[8][64] =
__constant u32a c_SPtrans[8][64] =
{
{
0x00820200, 0x00020000, 0x80800000, 0x80820200,
@ -182,7 +182,7 @@ __constant u32 c_SPtrans[8][64] =
},
};
__constant u32 c_skb[8][64] =
__constant u32a c_skb[8][64] =
{
{
0x00000000, 0x00000010, 0x20000000, 0x20000010,

@ -16,7 +16,7 @@
#include "inc_rp.cl"
#include "inc_simd.cl"
__constant u64 k_sha512[80] =
__constant u64a k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,

@ -14,7 +14,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u64 k_sha512[80] =
__constant u64a k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,

@ -14,7 +14,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u64 k_sha512[80] =
__constant u64a k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,

@ -16,7 +16,7 @@
#include "inc_rp.cl"
#include "inc_simd.cl"
__constant u64 k_sha512[80] =
__constant u64a k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,

@ -14,7 +14,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u64 k_sha512[80] =
__constant u64a k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,

@ -14,7 +14,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u64 k_sha512[80] =
__constant u64a k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,

@ -16,7 +16,7 @@
#include "inc_rp.cl"
#include "inc_simd.cl"
__constant u64 k_sha512[80] =
__constant u64a k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,

@ -14,7 +14,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u64 k_sha512[80] =
__constant u64a k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,

@ -14,7 +14,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u64 k_sha512[80] =
__constant u64a k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,

@ -16,7 +16,7 @@
#include "inc_rp.cl"
#include "inc_simd.cl"
__constant u64 k_sha512[80] =
__constant u64a k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,

@ -14,7 +14,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u64 k_sha512[80] =
__constant u64a k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,

@ -14,7 +14,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u64 k_sha512[80] =
__constant u64a k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,

@ -16,7 +16,7 @@
#include "inc_rp.cl"
#include "inc_simd.cl"
__constant u64 k_sha512[80] =
__constant u64a k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,

@ -14,7 +14,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u64 k_sha512[80] =
__constant u64a k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,

@ -14,7 +14,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u64 k_sha512[80] =
__constant u64a k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,

@ -16,7 +16,7 @@
#include "inc_rp.cl"
#include "inc_simd.cl"
__constant u64 k_sha512[80] =
__constant u64a k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,

@ -14,7 +14,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u64 k_sha512[80] =
__constant u64a k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,

@ -14,7 +14,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u64 k_sha512[80] =
__constant u64a k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,

@ -16,7 +16,7 @@
#include "inc_rp.cl"
#include "inc_simd.cl"
__constant u64 k_sha512[80] =
__constant u64a k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,

@ -14,7 +14,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u64 k_sha512[80] =
__constant u64a k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,

@ -14,7 +14,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u64 k_sha512[80] =
__constant u64a k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,

@ -34,7 +34,7 @@ typedef struct
} sha512_ctx_t;
__constant u64 k_sha512[80] =
__constant u64a k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,

@ -36,7 +36,7 @@
a = a ^ tt; \
}
__constant u32 c_SPtrans[8][64] =
__constant u32a c_SPtrans[8][64] =
{
{
0x02080800, 0x00080000, 0x02000002, 0x02080802,
@ -184,7 +184,7 @@ __constant u32 c_SPtrans[8][64] =
}
};
__constant u32 c_skb[8][64] =
__constant u32a c_skb[8][64] =
{
{
0x00000000, 0x00000010, 0x20000000, 0x20000010,

@ -34,7 +34,7 @@
a = a ^ tt; \
}
__constant u32 c_SPtrans[8][64] =
__constant u32a c_SPtrans[8][64] =
{
{
0x02080800, 0x00080000, 0x02000002, 0x02080802,
@ -182,7 +182,7 @@ __constant u32 c_SPtrans[8][64] =
}
};
__constant u32 c_skb[8][64] =
__constant u32a c_skb[8][64] =
{
{
0x00000000, 0x00000010, 0x20000000, 0x20000010,

@ -55,7 +55,7 @@
PERM_OP (l, r, tt, 4, 0x0f0f0f0f); \
}
__constant u32 c_SPtrans[8][64] =
__constant u32a c_SPtrans[8][64] =
{
{
0x02080800, 0x00080000, 0x02000002, 0x02080802,
@ -203,7 +203,7 @@ __constant u32 c_SPtrans[8][64] =
}
};
__constant u32 c_skb[8][64] =
__constant u32a c_skb[8][64] =
{
{
0x00000000, 0x00000010, 0x20000000, 0x20000010,

@ -53,7 +53,7 @@
PERM_OP (l, r, tt, 4, 0x0f0f0f0f); \
}
__constant u32 c_SPtrans[8][64] =
__constant u32a c_SPtrans[8][64] =
{
{
0x02080800, 0x00080000, 0x02000002, 0x02080802,
@ -201,7 +201,7 @@ __constant u32 c_SPtrans[8][64] =
}
};
__constant u32 c_skb[8][64] =
__constant u32a c_skb[8][64] =
{
{
0x00000000, 0x00000010, 0x20000000, 0x20000010,

@ -52,7 +52,7 @@
PERM_OP (l, r, tt, 4, 0x0f0f0f0f); \
}
__constant u32 c_SPtrans[8][64] =
__constant u32a c_SPtrans[8][64] =
{
{
0x02080800, 0x00080000, 0x02000002, 0x02080802,
@ -200,7 +200,7 @@ __constant u32 c_SPtrans[8][64] =
}
};
__constant u32 c_skb[8][64] =
__constant u32a c_skb[8][64] =
{
{
0x00000000, 0x00000010, 0x20000000, 0x20000010,

@ -16,7 +16,7 @@
// http://www.schneier.com/code/constants.txt
__constant u32 c_sbox0[256] =
__constant u32a c_sbox0[256] =
{
0xd1310ba6, 0x98dfb5ac, 0x2ffd72db, 0xd01adfb7,
0xb8e1afed, 0x6a267e96, 0xba7c9045, 0xf12c7f99,
@ -84,7 +84,7 @@ __constant u32 c_sbox0[256] =
0x53b02d5d, 0xa99f8fa1, 0x08ba4799, 0x6e85076a
};
__constant u32 c_sbox1[256] =
__constant u32a c_sbox1[256] =
{
0x4b7a70e9, 0xb5b32944, 0xdb75092e, 0xc4192623,
0xad6ea6b0, 0x49a7df7d, 0x9cee60b8, 0x8fedb266,
@ -152,7 +152,7 @@ __constant u32 c_sbox1[256] =
0x153e21e7, 0x8fb03d4a, 0xe6e39f2b, 0xdb83adf7
};
__constant u32 c_sbox2[256] =
__constant u32a c_sbox2[256] =
{
0xe93d5a68, 0x948140f7, 0xf64c261c, 0x94692934,
0x411520f7, 0x7602d4f7, 0xbcf46b2e, 0xd4a20068,
@ -220,7 +220,7 @@ __constant u32 c_sbox2[256] =
0xd79a3234, 0x92638212, 0x670efa8e, 0x406000e0
};
__constant u32 c_sbox3[256] =
__constant u32a c_sbox3[256] =
{
0x3a39ce37, 0xd3faf5cf, 0xabc27737, 0x5ac52d1b,
0x5cb0679e, 0x4fa33742, 0xd3822740, 0x99bc9bbe,

@ -16,7 +16,7 @@
#include "inc_rp.cl"
#include "inc_simd.cl"
__constant u64 keccakf_rndc[24] =
__constant u64a keccakf_rndc[24] =
{
0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
0x8000000080008000, 0x000000000000808b, 0x0000000080000001,

@ -14,7 +14,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u64 keccakf_rndc[24] =
__constant u64a keccakf_rndc[24] =
{
0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
0x8000000080008000, 0x000000000000808b, 0x0000000080000001,

@ -14,7 +14,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u64 keccakf_rndc[24] =
__constant u64a keccakf_rndc[24] =
{
0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
0x8000000080008000, 0x000000000000808b, 0x0000000080000001,

@ -14,7 +14,7 @@
#define COMPARE_S "inc_comp_single.cl"
#define COMPARE_M "inc_comp_multi.cl"
__constant u32 k_sha256[64] =
__constant u32a k_sha256[64] =
{
SHA256C00, SHA256C01, SHA256C02, SHA256C03,
SHA256C04, SHA256C05, SHA256C06, SHA256C07,

@ -36,7 +36,7 @@
a = a ^ tt; \
}
__constant u32 c_SPtrans[8][64] =
__constant u32a c_SPtrans[8][64] =
{
{
0x02080800, 0x00080000, 0x02000002, 0x02080802,
@ -184,7 +184,7 @@ __constant u32 c_SPtrans[8][64] =
}
};
__constant u32 c_skb[8][64] =
__constant u32a c_skb[8][64] =
{
{
0x00000000, 0x00000010, 0x20000000, 0x20000010,

@ -34,7 +34,7 @@
a = a ^ tt; \
}
__constant u32 c_SPtrans[8][64] =
__constant u32a c_SPtrans[8][64] =
{
{
0x02080800, 0x00080000, 0x02000002, 0x02080802,
@ -182,7 +182,7 @@ __constant u32 c_SPtrans[8][64] =
}
};
__constant u32 c_skb[8][64] =
__constant u32a c_skb[8][64] =
{
{
0x00000000, 0x00000010, 0x20000000, 0x20000010,

@ -34,7 +34,7 @@
a = a ^ tt; \
}
__constant u32 c_SPtrans[8][64] =
__constant u32a c_SPtrans[8][64] =
{
{
0x02080800, 0x00080000, 0x02000002, 0x02080802,
@ -182,7 +182,7 @@ __constant u32 c_SPtrans[8][64] =
}
};
__constant u32 c_skb[8][64] =
__constant u32a c_skb[8][64] =
{
{
0x00000000, 0x00000010, 0x20000000, 0x20000010,

@ -14,7 +14,7 @@
#define COMPARE_S "inc_comp_single.cl"
#define COMPARE_M "inc_comp_multi.cl"
__constant u32 c_pc_dec[1024] =
__constant u32a c_pc_dec[1024] =
{
0x00000030,
0x00000031,
@ -1042,7 +1042,7 @@ __constant u32 c_pc_dec[1024] =
0x33323031,
};
__constant u32 c_pc_len[1024] =
__constant u32a c_pc_len[1024] =
{
1,
1,

@ -30,7 +30,7 @@
#define BOX(S,n,i) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], (S)[(n)][(i).s4], (S)[(n)][(i).s5], (S)[(n)][(i).s6], (S)[(n)][(i).s7], (S)[(n)][(i).s8], (S)[(n)][(i).s9], (S)[(n)][(i).sa], (S)[(n)][(i).sb], (S)[(n)][(i).sc], (S)[(n)][(i).sd], (S)[(n)][(i).se], (S)[(n)][(i).sf])
#endif
__constant u32 Ch[8][256] =
__constant u32a Ch[8][256] =
{
{
0x18186018, 0x23238c23, 0xc6c63fc6, 0xe8e887e8,
@ -562,7 +562,7 @@ __constant u32 Ch[8][256] =
}
};
__constant u32 Cl[8][256] =
__constant u32a Cl[8][256] =
{
{
0xc07830d8, 0x05af4626, 0x7ef991b8, 0x136fcdfb,
@ -1094,7 +1094,7 @@ __constant u32 Cl[8][256] =
},
};
__constant u32 rch[R + 1] =
__constant u32a rch[R + 1] =
{
0x00000000,
0x1823c6e8,
@ -1109,7 +1109,7 @@ __constant u32 rch[R + 1] =
0xca2dbf07,
};
__constant u32 rcl[R + 1] =
__constant u32a rcl[R + 1] =
{
0x00000000,
0x87b8014f,

@ -28,7 +28,7 @@
#define BOX(S,n,i) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], (S)[(n)][(i).s4], (S)[(n)][(i).s5], (S)[(n)][(i).s6], (S)[(n)][(i).s7], (S)[(n)][(i).s8], (S)[(n)][(i).s9], (S)[(n)][(i).sa], (S)[(n)][(i).sb], (S)[(n)][(i).sc], (S)[(n)][(i).sd], (S)[(n)][(i).se], (S)[(n)][(i).sf])
#endif
__constant u32 Ch[8][256] =
__constant u32a Ch[8][256] =
{
{
0x18186018, 0x23238c23, 0xc6c63fc6, 0xe8e887e8,
@ -560,7 +560,7 @@ __constant u32 Ch[8][256] =
}
};
__constant u32 Cl[8][256] =
__constant u32a Cl[8][256] =
{
{
0xc07830d8, 0x05af4626, 0x7ef991b8, 0x136fcdfb,
@ -1092,7 +1092,7 @@ __constant u32 Cl[8][256] =
},
};
__constant u32 rch[R + 1] =
__constant u32a rch[R + 1] =
{
0x00000000,
0x1823c6e8,
@ -1107,7 +1107,7 @@ __constant u32 rch[R + 1] =
0xca2dbf07,
};
__constant u32 rcl[R + 1] =
__constant u32a rcl[R + 1] =
{
0x00000000,
0x87b8014f,

@ -28,7 +28,7 @@
#define BOX(S,n,i) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], (S)[(n)][(i).s4], (S)[(n)][(i).s5], (S)[(n)][(i).s6], (S)[(n)][(i).s7], (S)[(n)][(i).s8], (S)[(n)][(i).s9], (S)[(n)][(i).sa], (S)[(n)][(i).sb], (S)[(n)][(i).sc], (S)[(n)][(i).sd], (S)[(n)][(i).se], (S)[(n)][(i).sf])
#endif
__constant u32 Ch[8][256] =
__constant u32a Ch[8][256] =
{
{
0x18186018, 0x23238c23, 0xc6c63fc6, 0xe8e887e8,
@ -560,7 +560,7 @@ __constant u32 Ch[8][256] =
}
};
__constant u32 Cl[8][256] =
__constant u32a Cl[8][256] =
{
{
0xc07830d8, 0x05af4626, 0x7ef991b8, 0x136fcdfb,
@ -1092,7 +1092,7 @@ __constant u32 Cl[8][256] =
},
};
__constant u32 rch[R + 1] =
__constant u32a rch[R + 1] =
{
0x00000000,
0x1823c6e8,
@ -1107,7 +1107,7 @@ __constant u32 rch[R + 1] =
0xca2dbf07,
};
__constant u32 rcl[R + 1] =
__constant u32a rcl[R + 1] =
{
0x00000000,
0x87b8014f,

@ -652,17 +652,17 @@ __kernel void m06211_comp (__global pw_t *pws, __global const kernel_rule_t *rul
#else
__constant u32 *s_td0 = td0;
__constant u32 *s_td1 = td1;
__constant u32 *s_td2 = td2;
__constant u32 *s_td3 = td3;
__constant u32 *s_td4 = td4;
__constant u32 *s_te0 = te0;
__constant u32 *s_te1 = te1;
__constant u32 *s_te2 = te2;
__constant u32 *s_te3 = te3;
__constant u32 *s_te4 = te4;
__constant u32a *s_td0 = td0;
__constant u32a *s_td1 = td1;
__constant u32a *s_td2 = td2;
__constant u32a *s_td3 = td3;
__constant u32a *s_td4 = td4;
__constant u32a *s_te0 = te0;
__constant u32a *s_te1 = te1;
__constant u32a *s_te2 = te2;
__constant u32a *s_te3 = te3;
__constant u32a *s_te4 = te4;
#endif

@ -652,17 +652,17 @@ __kernel void m06212_comp (__global pw_t *pws, __global const kernel_rule_t *rul
#else
__constant u32 *s_td0 = td0;
__constant u32 *s_td1 = td1;
__constant u32 *s_td2 = td2;
__constant u32 *s_td3 = td3;
__constant u32 *s_td4 = td4;
__constant u32 *s_te0 = te0;
__constant u32 *s_te1 = te1;
__constant u32 *s_te2 = te2;
__constant u32 *s_te3 = te3;
__constant u32 *s_te4 = te4;
__constant u32a *s_td0 = td0;
__constant u32a *s_td1 = td1;
__constant u32a *s_td2 = td2;
__constant u32a *s_td3 = td3;
__constant u32a *s_td4 = td4;
__constant u32a *s_te0 = te0;
__constant u32a *s_te1 = te1;
__constant u32a *s_te2 = te2;
__constant u32a *s_te3 = te3;
__constant u32a *s_te4 = te4;
#endif

@ -652,17 +652,17 @@ __kernel void m06213_comp (__global pw_t *pws, __global const kernel_rule_t *rul
#else
__constant u32 *s_td0 = td0;
__constant u32 *s_td1 = td1;
__constant u32 *s_td2 = td2;
__constant u32 *s_td3 = td3;
__constant u32 *s_td4 = td4;
__constant u32 *s_te0 = te0;
__constant u32 *s_te1 = te1;
__constant u32 *s_te2 = te2;
__constant u32 *s_te3 = te3;
__constant u32 *s_te4 = te4;
__constant u32a *s_td0 = td0;
__constant u32a *s_td1 = td1;
__constant u32a *s_td2 = td2;
__constant u32a *s_td3 = td3;
__constant u32a *s_td4 = td4;
__constant u32a *s_te0 = te0;
__constant u32a *s_te1 = te1;
__constant u32a *s_te2 = te2;
__constant u32a *s_te3 = te3;
__constant u32a *s_te4 = te4;
#endif

@ -18,7 +18,7 @@
#include "inc_truecrypt_crc32.cl"
#include "inc_truecrypt_xts.cl"
__constant u64 k_sha512[80] =
__constant u64a k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,
@ -560,17 +560,17 @@ __kernel void m06221_comp (__global pw_t *pws, __global const kernel_rule_t *rul
#else
__constant u32 *s_td0 = td0;
__constant u32 *s_td1 = td1;
__constant u32 *s_td2 = td2;
__constant u32 *s_td3 = td3;
__constant u32 *s_td4 = td4;
__constant u32 *s_te0 = te0;
__constant u32 *s_te1 = te1;
__constant u32 *s_te2 = te2;
__constant u32 *s_te3 = te3;
__constant u32 *s_te4 = te4;
__constant u32a *s_td0 = td0;
__constant u32a *s_td1 = td1;
__constant u32a *s_td2 = td2;
__constant u32a *s_td3 = td3;
__constant u32a *s_td4 = td4;
__constant u32a *s_te0 = te0;
__constant u32a *s_te1 = te1;
__constant u32a *s_te2 = te2;
__constant u32a *s_te3 = te3;
__constant u32a *s_te4 = te4;
#endif

@ -18,7 +18,7 @@
#include "inc_truecrypt_crc32.cl"
#include "inc_truecrypt_xts.cl"
__constant u64 k_sha512[80] =
__constant u64a k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,
@ -560,17 +560,17 @@ __kernel void m06222_comp (__global pw_t *pws, __global const kernel_rule_t *rul
#else
__constant u32 *s_td0 = td0;
__constant u32 *s_td1 = td1;
__constant u32 *s_td2 = td2;
__constant u32 *s_td3 = td3;
__constant u32 *s_td4 = td4;
__constant u32 *s_te0 = te0;
__constant u32 *s_te1 = te1;
__constant u32 *s_te2 = te2;
__constant u32 *s_te3 = te3;
__constant u32 *s_te4 = te4;
__constant u32a *s_td0 = td0;
__constant u32a *s_td1 = td1;
__constant u32a *s_td2 = td2;
__constant u32a *s_td3 = td3;
__constant u32a *s_td4 = td4;
__constant u32a *s_te0 = te0;
__constant u32a *s_te1 = te1;
__constant u32a *s_te2 = te2;
__constant u32a *s_te3 = te3;
__constant u32a *s_te4 = te4;
#endif

@ -18,7 +18,7 @@
#include "inc_truecrypt_crc32.cl"
#include "inc_truecrypt_xts.cl"
__constant u64 k_sha512[80] =
__constant u64a k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,
@ -301,17 +301,17 @@ __kernel void m06223_init (__global pw_t *pws, __global const kernel_rule_t *rul
#else
__constant u32 *s_td0 = td0;
__constant u32 *s_td1 = td1;
__constant u32 *s_td2 = td2;
__constant u32 *s_td3 = td3;
__constant u32 *s_td4 = td4;
__constant u32a *s_td0 = td0;
__constant u32a *s_td1 = td1;
__constant u32a *s_td2 = td2;
__constant u32a *s_td3 = td3;
__constant u32a *s_td4 = td4;
__constant u32 *s_te0 = te0;
__constant u32 *s_te1 = te1;
__constant u32 *s_te2 = te2;
__constant u32 *s_te3 = te3;
__constant u32 *s_te4 = te4;
__constant u32a *s_te0 = te0;
__constant u32a *s_te1 = te1;
__constant u32a *s_te2 = te2;
__constant u32a *s_te3 = te3;
__constant u32a *s_te4 = te4;
#endif
@ -609,17 +609,17 @@ __kernel void m06223_comp (__global pw_t *pws, __global const kernel_rule_t *rul
#else
__constant u32 *s_td0 = td0;
__constant u32 *s_td1 = td1;
__constant u32 *s_td2 = td2;
__constant u32 *s_td3 = td3;
__constant u32 *s_td4 = td4;
__constant u32 *s_te0 = te0;
__constant u32 *s_te1 = te1;
__constant u32 *s_te2 = te2;
__constant u32 *s_te3 = te3;
__constant u32 *s_te4 = te4;
__constant u32a *s_td0 = td0;
__constant u32a *s_td1 = td1;
__constant u32a *s_td2 = td2;
__constant u32a *s_td3 = td3;
__constant u32a *s_td4 = td4;
__constant u32a *s_te0 = te0;
__constant u32a *s_te1 = te1;
__constant u32a *s_te2 = te2;
__constant u32a *s_te3 = te3;
__constant u32a *s_te4 = te4;
#endif

@ -20,7 +20,7 @@
#define R 10
__constant u32 Ch[8][256] =
__constant u32a Ch[8][256] =
{
{
0x18186018, 0x23238c23, 0xc6c63fc6, 0xe8e887e8,
@ -552,7 +552,7 @@ __constant u32 Ch[8][256] =
}
};
__constant u32 Cl[8][256] =
__constant u32a Cl[8][256] =
{
{
0xc07830d8, 0x05af4626, 0x7ef991b8, 0x136fcdfb,
@ -2195,17 +2195,17 @@ __kernel void m06231_comp (__global pw_t *pws, __global const kernel_rule_t *rul
#else
__constant u32 *s_td0 = td0;
__constant u32 *s_td1 = td1;
__constant u32 *s_td2 = td2;
__constant u32 *s_td3 = td3;
__constant u32 *s_td4 = td4;
__constant u32 *s_te0 = te0;
__constant u32 *s_te1 = te1;
__constant u32 *s_te2 = te2;
__constant u32 *s_te3 = te3;
__constant u32 *s_te4 = te4;
__constant u32a *s_td0 = td0;
__constant u32a *s_td1 = td1;
__constant u32a *s_td2 = td2;
__constant u32a *s_td3 = td3;
__constant u32a *s_td4 = td4;
__constant u32a *s_te0 = te0;
__constant u32a *s_te1 = te1;
__constant u32a *s_te2 = te2;
__constant u32a *s_te3 = te3;
__constant u32a *s_te4 = te4;
#endif

@ -20,7 +20,7 @@
#define R 10
__constant u32 Ch[8][256] =
__constant u32a Ch[8][256] =
{
{
0x18186018, 0x23238c23, 0xc6c63fc6, 0xe8e887e8,
@ -552,7 +552,7 @@ __constant u32 Ch[8][256] =
}
};
__constant u32 Cl[8][256] =
__constant u32a Cl[8][256] =
{
{
0xc07830d8, 0x05af4626, 0x7ef991b8, 0x136fcdfb,
@ -1964,17 +1964,17 @@ __kernel void m06232_comp (__global pw_t *pws, __global const kernel_rule_t *rul
#else
__constant u32 *s_td0 = td0;
__constant u32 *s_td1 = td1;
__constant u32 *s_td2 = td2;
__constant u32 *s_td3 = td3;
__constant u32 *s_td4 = td4;
__constant u32 *s_te0 = te0;
__constant u32 *s_te1 = te1;
__constant u32 *s_te2 = te2;
__constant u32 *s_te3 = te3;
__constant u32 *s_te4 = te4;
__constant u32a *s_td0 = td0;
__constant u32a *s_td1 = td1;
__constant u32a *s_td2 = td2;
__constant u32a *s_td3 = td3;
__constant u32a *s_td4 = td4;
__constant u32a *s_te0 = te0;
__constant u32a *s_te1 = te1;
__constant u32a *s_te2 = te2;
__constant u32a *s_te3 = te3;
__constant u32a *s_te4 = te4;
#endif

@ -20,7 +20,7 @@
#define R 10
__constant u32 Ch[8][256] =
__constant u32a Ch[8][256] =
{
{
0x18186018, 0x23238c23, 0xc6c63fc6, 0xe8e887e8,
@ -552,7 +552,7 @@ __constant u32 Ch[8][256] =
}
};
__constant u32 Cl[8][256] =
__constant u32a Cl[8][256] =
{
{
0xc07830d8, 0x05af4626, 0x7ef991b8, 0x136fcdfb,
@ -1964,17 +1964,17 @@ __kernel void m06233_comp (__global pw_t *pws, __global const kernel_rule_t *rul
#else
__constant u32 *s_td0 = td0;
__constant u32 *s_td1 = td1;
__constant u32 *s_td2 = td2;
__constant u32 *s_td3 = td3;
__constant u32 *s_td4 = td4;
__constant u32 *s_te0 = te0;
__constant u32 *s_te1 = te1;
__constant u32 *s_te2 = te2;
__constant u32 *s_te3 = te3;
__constant u32 *s_te4 = te4;
__constant u32a *s_td0 = td0;
__constant u32a *s_td1 = td1;
__constant u32a *s_td2 = td2;
__constant u32a *s_td3 = td3;
__constant u32a *s_td4 = td4;
__constant u32a *s_te0 = te0;
__constant u32a *s_te1 = te1;
__constant u32a *s_te2 = te2;
__constant u32a *s_te3 = te3;
__constant u32a *s_te4 = te4;
#endif

@ -14,7 +14,7 @@
#define COMPARE_S "inc_comp_single.cl"
#define COMPARE_M "inc_comp_multi.cl"
__constant u32 k_sha256[64] =
__constant u32a k_sha256[64] =
{
SHA256C00, SHA256C01, SHA256C02, SHA256C03,
SHA256C04, SHA256C05, SHA256C06, SHA256C07,

@ -14,7 +14,7 @@
#define COMPARE_S "inc_comp_single.cl"
#define COMPARE_M "inc_comp_multi.cl"
__constant u64 k_sha512[80] =
__constant u64a k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,

@ -14,7 +14,7 @@
#define COMPARE_S "inc_comp_single.cl"
#define COMPARE_M "inc_comp_multi.cl"
__constant u32 te0[256] =
__constant u32a te0[256] =
{
0xc66363a5, 0xf87c7c84, 0xee777799, 0xf67b7b8d,
0xfff2f20d, 0xd66b6bbd, 0xde6f6fb1, 0x91c5c554,
@ -82,7 +82,7 @@ __constant u32 te0[256] =
0x7bb0b0cb, 0xa85454fc, 0x6dbbbbd6, 0x2c16163a,
};
__constant u32 te1[256] =
__constant u32a te1[256] =
{
0xa5c66363, 0x84f87c7c, 0x99ee7777, 0x8df67b7b,
0x0dfff2f2, 0xbdd66b6b, 0xb1de6f6f, 0x5491c5c5,
@ -150,7 +150,7 @@ __constant u32 te1[256] =
0xcb7bb0b0, 0xfca85454, 0xd66dbbbb, 0x3a2c1616,
};
__constant u32 te2[256] =
__constant u32a te2[256] =
{
0x63a5c663, 0x7c84f87c, 0x7799ee77, 0x7b8df67b,
0xf20dfff2, 0x6bbdd66b, 0x6fb1de6f, 0xc55491c5,
@ -218,7 +218,7 @@ __constant u32 te2[256] =
0xb0cb7bb0, 0x54fca854, 0xbbd66dbb, 0x163a2c16,
};
__constant u32 te3[256] =
__constant u32a te3[256] =
{
0x6363a5c6, 0x7c7c84f8, 0x777799ee, 0x7b7b8df6,
0xf2f20dff, 0x6b6bbdd6, 0x6f6fb1de, 0xc5c55491,
@ -286,7 +286,7 @@ __constant u32 te3[256] =
0xb0b0cb7b, 0x5454fca8, 0xbbbbd66d, 0x16163a2c,
};
__constant u32 te4[256] =
__constant u32a te4[256] =
{
0x63636363, 0x7c7c7c7c, 0x77777777, 0x7b7b7b7b,
0xf2f2f2f2, 0x6b6b6b6b, 0x6f6f6f6f, 0xc5c5c5c5,
@ -354,7 +354,7 @@ __constant u32 te4[256] =
0xb0b0b0b0, 0x54545454, 0xbbbbbbbb, 0x16161616,
};
__constant u32 td0[256] =
__constant u32a td0[256] =
{
0x51f4a750, 0x7e416553, 0x1a17a4c3, 0x3a275e96,
0x3bab6bcb, 0x1f9d45f1, 0xacfa58ab, 0x4be30393,
@ -422,7 +422,7 @@ __constant u32 td0[256] =
0x7bcb8461, 0xd532b670, 0x486c5c74, 0xd0b85742,
};
__constant u32 td1[256] =
__constant u32a td1[256] =
{
0x5051f4a7, 0x537e4165, 0xc31a17a4, 0x963a275e,
0xcb3bab6b, 0xf11f9d45, 0xabacfa58, 0x934be303,
@ -490,7 +490,7 @@ __constant u32 td1[256] =
0x617bcb84, 0x70d532b6, 0x74486c5c, 0x42d0b857,
};
__constant u32 td2[256] =
__constant u32a td2[256] =
{
0xa75051f4, 0x65537e41, 0xa4c31a17, 0x5e963a27,
0x6bcb3bab, 0x45f11f9d, 0x58abacfa, 0x03934be3,
@ -558,7 +558,7 @@ __constant u32 td2[256] =
0x84617bcb, 0xb670d532, 0x5c74486c, 0x5742d0b8,
};
__constant u32 td3[256] =
__constant u32a td3[256] =
{
0xf4a75051, 0x4165537e, 0x17a4c31a, 0x275e963a,
0xab6bcb3b, 0x9d45f11f, 0xfa58abac, 0xe303934b,
@ -626,7 +626,7 @@ __constant u32 td3[256] =
0xcb84617b, 0x32b670d5, 0x6c5c7448, 0xb85742d0,
};
__constant u32 td4[256] =
__constant u32a td4[256] =
{
0x52525252, 0x09090909, 0x6a6a6a6a, 0xd5d5d5d5,
0x30303030, 0x36363636, 0xa5a5a5a5, 0x38383838,
@ -694,7 +694,7 @@ __constant u32 td4[256] =
0x55555555, 0x21212121, 0x0c0c0c0c, 0x7d7d7d7d,
};
__constant u32 rcon[] =
__constant u32a rcon[] =
{
0x01000000, 0x02000000, 0x04000000, 0x08000000,
0x10000000, 0x20000000, 0x40000000, 0x80000000,

@ -14,7 +14,7 @@
#define COMPARE_S "inc_comp_single.cl"
#define COMPARE_M "inc_comp_multi.cl"
__constant u32 te0[256] =
__constant u32a te0[256] =
{
0xc66363a5, 0xf87c7c84, 0xee777799, 0xf67b7b8d,
0xfff2f20d, 0xd66b6bbd, 0xde6f6fb1, 0x91c5c554,
@ -82,7 +82,7 @@ __constant u32 te0[256] =
0x7bb0b0cb, 0xa85454fc, 0x6dbbbbd6, 0x2c16163a,
};
__constant u32 te1[256] =
__constant u32a te1[256] =
{
0xa5c66363, 0x84f87c7c, 0x99ee7777, 0x8df67b7b,
0x0dfff2f2, 0xbdd66b6b, 0xb1de6f6f, 0x5491c5c5,
@ -150,7 +150,7 @@ __constant u32 te1[256] =
0xcb7bb0b0, 0xfca85454, 0xd66dbbbb, 0x3a2c1616,
};
__constant u32 te2[256] =
__constant u32a te2[256] =
{
0x63a5c663, 0x7c84f87c, 0x7799ee77, 0x7b8df67b,
0xf20dfff2, 0x6bbdd66b, 0x6fb1de6f, 0xc55491c5,
@ -218,7 +218,7 @@ __constant u32 te2[256] =
0xb0cb7bb0, 0x54fca854, 0xbbd66dbb, 0x163a2c16,
};
__constant u32 te3[256] =
__constant u32a te3[256] =
{
0x6363a5c6, 0x7c7c84f8, 0x777799ee, 0x7b7b8df6,
0xf2f20dff, 0x6b6bbdd6, 0x6f6fb1de, 0xc5c55491,
@ -286,7 +286,7 @@ __constant u32 te3[256] =
0xb0b0cb7b, 0x5454fca8, 0xbbbbd66d, 0x16163a2c,
};
__constant u32 te4[256] =
__constant u32a te4[256] =
{
0x63636363, 0x7c7c7c7c, 0x77777777, 0x7b7b7b7b,
0xf2f2f2f2, 0x6b6b6b6b, 0x6f6f6f6f, 0xc5c5c5c5,
@ -354,7 +354,7 @@ __constant u32 te4[256] =
0xb0b0b0b0, 0x54545454, 0xbbbbbbbb, 0x16161616,
};
__constant u32 td0[256] =
__constant u32a td0[256] =
{
0x51f4a750, 0x7e416553, 0x1a17a4c3, 0x3a275e96,
0x3bab6bcb, 0x1f9d45f1, 0xacfa58ab, 0x4be30393,
@ -422,7 +422,7 @@ __constant u32 td0[256] =
0x7bcb8461, 0xd532b670, 0x486c5c74, 0xd0b85742,
};
__constant u32 td1[256] =
__constant u32a td1[256] =
{
0x5051f4a7, 0x537e4165, 0xc31a17a4, 0x963a275e,
0xcb3bab6b, 0xf11f9d45, 0xabacfa58, 0x934be303,
@ -490,7 +490,7 @@ __constant u32 td1[256] =
0x617bcb84, 0x70d532b6, 0x74486c5c, 0x42d0b857,
};
__constant u32 td2[256] =
__constant u32a td2[256] =
{
0xa75051f4, 0x65537e41, 0xa4c31a17, 0x5e963a27,
0x6bcb3bab, 0x45f11f9d, 0x58abacfa, 0x03934be3,
@ -558,7 +558,7 @@ __constant u32 td2[256] =
0x84617bcb, 0xb670d532, 0x5c74486c, 0x5742d0b8,
};
__constant u32 td3[256] =
__constant u32a td3[256] =
{
0xf4a75051, 0x4165537e, 0x17a4c31a, 0x275e963a,
0xab6bcb3b, 0x9d45f11f, 0xfa58abac, 0xe303934b,
@ -626,7 +626,7 @@ __constant u32 td3[256] =
0xcb84617b, 0x32b670d5, 0x6c5c7448, 0xb85742d0,
};
__constant u32 td4[256] =
__constant u32a td4[256] =
{
0x52525252, 0x09090909, 0x6a6a6a6a, 0xd5d5d5d5,
0x30303030, 0x36363636, 0xa5a5a5a5, 0x38383838,
@ -694,7 +694,7 @@ __constant u32 td4[256] =
0x55555555, 0x21212121, 0x0c0c0c0c, 0x7d7d7d7d,
};
__constant u32 rcon[] =
__constant u32a rcon[] =
{
0x01000000, 0x02000000, 0x04000000, 0x08000000,
0x10000000, 0x20000000, 0x40000000, 0x80000000,
@ -977,7 +977,7 @@ void AES256_encrypt (const u32 *in, u32 *out, const u32 *rek, __local u32 *s_te0
^ rek[59];
}
__constant u32 k_sha256[64] =
__constant u32a k_sha256[64] =
{
SHA256C00, SHA256C01, SHA256C02, SHA256C03,
SHA256C04, SHA256C05, SHA256C06, SHA256C07,

@ -16,7 +16,7 @@
#include "inc_rp.cl"
#include "inc_simd.cl"
__constant u32 c_tables[4][256] =
__constant u32a c_tables[4][256] =
{
{
0x00072000, 0x00075000, 0x00074800, 0x00071000,

@ -14,7 +14,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u32 c_tables[4][256] =
__constant u32a c_tables[4][256] =
{
{
0x00072000, 0x00075000, 0x00074800, 0x00071000,

@ -14,7 +14,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u32 c_tables[4][256] =
__constant u32a c_tables[4][256] =
{
{
0x00072000, 0x00075000, 0x00074800, 0x00071000,

@ -17,7 +17,7 @@
#define COMPARE_S "inc_comp_single.cl"
#define COMPARE_M "inc_comp_multi.cl"
__constant u64 k_sha512[80] =
__constant u64a k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,

@ -14,7 +14,7 @@
#define COMPARE_S "inc_comp_single.cl"
#define COMPARE_M "inc_comp_multi.cl"
__constant u32 k_sha256[64] =
__constant u32a k_sha256[64] =
{
SHA256C00, SHA256C01, SHA256C02, SHA256C03,
SHA256C04, SHA256C05, SHA256C06, SHA256C07,

@ -29,7 +29,7 @@
(a)[((n)/4)+1] = x >> 32; \
}
__constant u32 sapb_trans_tbl[256] =
__constant u32a sapb_trans_tbl[256] =
{
// first value hack for 0 byte as part of an optimization
0x00, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
@ -50,7 +50,7 @@ __constant u32 sapb_trans_tbl[256] =
0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff
};
__constant u32 bcodeArray[48] =
__constant u32a bcodeArray[48] =
{
0x14, 0x77, 0xf3, 0xd4, 0xbb, 0x71, 0x23, 0xd0, 0x03, 0xff, 0x47, 0x93, 0x55, 0xaa, 0x66, 0x91,
0xf2, 0x88, 0x6b, 0x99, 0xbf, 0xcb, 0x32, 0x1a, 0x19, 0xd9, 0xa7, 0x82, 0x22, 0x49, 0xa2, 0x51,

@ -27,7 +27,7 @@
(a)[((n)/4)+1] = x >> 32; \
}
__constant u32 sapb_trans_tbl[256] =
__constant u32a sapb_trans_tbl[256] =
{
// first value hack for 0 byte as part of an optimization
0x00, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
@ -48,7 +48,7 @@ __constant u32 sapb_trans_tbl[256] =
0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff
};
__constant u32 bcodeArray[48] =
__constant u32a bcodeArray[48] =
{
0x14, 0x77, 0xf3, 0xd4, 0xbb, 0x71, 0x23, 0xd0, 0x03, 0xff, 0x47, 0x93, 0x55, 0xaa, 0x66, 0x91,
0xf2, 0x88, 0x6b, 0x99, 0xbf, 0xcb, 0x32, 0x1a, 0x19, 0xd9, 0xa7, 0x82, 0x22, 0x49, 0xa2, 0x51,

@ -27,7 +27,7 @@
(a)[((n)/4)+1] = x >> 32; \
}
__constant u32 sapb_trans_tbl[256] =
__constant u32a sapb_trans_tbl[256] =
{
// first value hack for 0 byte as part of an optimization
0x00, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
@ -48,7 +48,7 @@ __constant u32 sapb_trans_tbl[256] =
0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff
};
__constant u32 bcodeArray[48] =
__constant u32a bcodeArray[48] =
{
0x14, 0x77, 0xf3, 0xd4, 0xbb, 0x71, 0x23, 0xd0, 0x03, 0xff, 0x47, 0x93, 0x55, 0xaa, 0x66, 0x91,
0xf2, 0x88, 0x6b, 0x99, 0xbf, 0xcb, 0x32, 0x1a, 0x19, 0xd9, 0xa7, 0x82, 0x22, 0x49, 0xa2, 0x51,

@ -17,7 +17,7 @@
#include "inc_rp.cl"
#include "inc_simd.cl"
__constant u32 theMagicArray[64] =
__constant u32a theMagicArray[64] =
{
0x91ac5114, 0x9f675443, 0x24e73be0, 0x28747bc2, 0x863313eb, 0x5a4fcb5c, 0x080a7337, 0x0e5d1c2f,
0x338fe6e5, 0xf89baedd, 0x16f24b8d, 0x2ce1d4dc, 0xb0cbdf9d, 0xd4706d17, 0xf94d423f, 0x9b1b1194,
@ -29,7 +29,7 @@ __constant u32 theMagicArray[64] =
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
};
u32 GETSHIFTEDINT_CONST (__constant u32 *a, const int n)
u32 GETSHIFTEDINT_CONST (__constant u32a *a, const int n)
{
const int d = n / 4;
const int m = n & 3;

@ -15,7 +15,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u32 theMagicArray[64] =
__constant u32a theMagicArray[64] =
{
0x91ac5114, 0x9f675443, 0x24e73be0, 0x28747bc2, 0x863313eb, 0x5a4fcb5c, 0x080a7337, 0x0e5d1c2f,
0x338fe6e5, 0xf89baedd, 0x16f24b8d, 0x2ce1d4dc, 0xb0cbdf9d, 0xd4706d17, 0xf94d423f, 0x9b1b1194,
@ -27,7 +27,7 @@ __constant u32 theMagicArray[64] =
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
};
u32 GETSHIFTEDINT_CONST (__constant u32 *a, const int n)
u32 GETSHIFTEDINT_CONST (__constant u32a *a, const int n)
{
const int d = n / 4;
const int m = n & 3;

@ -15,7 +15,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u32 theMagicArray[64] =
__constant u32a theMagicArray[64] =
{
0x91ac5114, 0x9f675443, 0x24e73be0, 0x28747bc2, 0x863313eb, 0x5a4fcb5c, 0x080a7337, 0x0e5d1c2f,
0x338fe6e5, 0xf89baedd, 0x16f24b8d, 0x2ce1d4dc, 0xb0cbdf9d, 0xd4706d17, 0xf94d423f, 0x9b1b1194,
@ -27,7 +27,7 @@ __constant u32 theMagicArray[64] =
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
};
u32 GETSHIFTEDINT_CONST (__constant u32 *a, const int n)
u32 GETSHIFTEDINT_CONST (__constant u32a *a, const int n)
{
const int d = n / 4;
const int m = n & 3;

@ -14,7 +14,7 @@
#define COMPARE_S "inc_comp_single.cl"
#define COMPARE_M "inc_comp_multi.cl"
__constant u64 k_sha512[80] =
__constant u64a k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,

@ -16,7 +16,7 @@
#include "inc_rp.cl"
#include "inc_simd.cl"
__constant u32 k_sha256[64] =
__constant u32a k_sha256[64] =
{
SHA256C00, SHA256C01, SHA256C02, SHA256C03,
SHA256C04, SHA256C05, SHA256C06, SHA256C07,

@ -14,7 +14,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u32 k_sha256[64] =
__constant u32a k_sha256[64] =
{
SHA256C00, SHA256C01, SHA256C02, SHA256C03,
SHA256C04, SHA256C05, SHA256C06, SHA256C07,

@ -14,7 +14,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u32 k_sha256[64] =
__constant u32a k_sha256[64] =
{
SHA256C00, SHA256C01, SHA256C02, SHA256C03,
SHA256C04, SHA256C05, SHA256C06, SHA256C07,

@ -14,7 +14,7 @@
#define COMPARE_S "inc_comp_single.cl"
#define COMPARE_M "inc_comp_multi.cl"
__constant u32 k_sha256[64] =
__constant u32a k_sha256[64] =
{
SHA256C00, SHA256C01, SHA256C02, SHA256C03,
SHA256C04, SHA256C05, SHA256C06, SHA256C07,
@ -34,7 +34,7 @@ __constant u32 k_sha256[64] =
SHA256C3c, SHA256C3d, SHA256C3e, SHA256C3f,
};
__constant u64 k_sha512[80] =
__constant u64a k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,

@ -54,7 +54,7 @@
PERM_OP (l, r, tt, 4, 0x0f0f0f0f); \
}
__constant u32 c_ascii_to_ebcdic_pc[256] =
__constant u32a c_ascii_to_ebcdic_pc[256] =
{
// little hack, can't crack 0-bytes in password, but who cares
// 0xab, 0xa8, 0xae, 0xad, 0xc4, 0xf1, 0xf7, 0xf4, 0x86, 0xa1, 0xe0, 0xbc, 0xb3, 0xb0, 0xb6, 0xb5,
@ -76,7 +76,7 @@ __constant u32 c_ascii_to_ebcdic_pc[256] =
0x13, 0x10, 0x16, 0x15, 0x7f, 0x7c, 0x73, 0x70, 0x76, 0x75, 0x5e, 0x5d, 0x52, 0x51, 0x57, 0x54,
};
__constant u32 c_SPtrans[8][64] =
__constant u32a c_SPtrans[8][64] =
{
{
0x02080800, 0x00080000, 0x02000002, 0x02080802,
@ -224,7 +224,7 @@ __constant u32 c_SPtrans[8][64] =
}
};
__constant u32 c_skb[8][64] =
__constant u32a c_skb[8][64] =
{
{
0x00000000, 0x00000010, 0x20000000, 0x20000010,

@ -52,7 +52,7 @@
PERM_OP (l, r, tt, 4, 0x0f0f0f0f); \
}
__constant u32 c_ascii_to_ebcdic_pc[256] =
__constant u32a c_ascii_to_ebcdic_pc[256] =
{
// little hack, can't crack 0-bytes in password, but who cares
// 0xab, 0xa8, 0xae, 0xad, 0xc4, 0xf1, 0xf7, 0xf4, 0x86, 0xa1, 0xe0, 0xbc, 0xb3, 0xb0, 0xb6, 0xb5,
@ -74,7 +74,7 @@ __constant u32 c_ascii_to_ebcdic_pc[256] =
0x13, 0x10, 0x16, 0x15, 0x7f, 0x7c, 0x73, 0x70, 0x76, 0x75, 0x5e, 0x5d, 0x52, 0x51, 0x57, 0x54,
};
__constant u32 c_SPtrans[8][64] =
__constant u32a c_SPtrans[8][64] =
{
{
0x02080800, 0x00080000, 0x02000002, 0x02080802,
@ -222,7 +222,7 @@ __constant u32 c_SPtrans[8][64] =
}
};
__constant u32 c_skb[8][64] =
__constant u32a c_skb[8][64] =
{
{
0x00000000, 0x00000010, 0x20000000, 0x20000010,

@ -52,7 +52,7 @@
PERM_OP (l, r, tt, 4, 0x0f0f0f0f); \
}
__constant u32 c_ascii_to_ebcdic_pc[256] =
__constant u32a c_ascii_to_ebcdic_pc[256] =
{
// little hack, can't crack 0-bytes in password, but who cares
// 0xab, 0xa8, 0xae, 0xad, 0xc4, 0xf1, 0xf7, 0xf4, 0x86, 0xa1, 0xe0, 0xbc, 0xb3, 0xb0, 0xb6, 0xb5,
@ -74,7 +74,7 @@ __constant u32 c_ascii_to_ebcdic_pc[256] =
0x13, 0x10, 0x16, 0x15, 0x7f, 0x7c, 0x73, 0x70, 0x76, 0x75, 0x5e, 0x5d, 0x52, 0x51, 0x57, 0x54,
};
__constant u32 c_SPtrans[8][64] =
__constant u32a c_SPtrans[8][64] =
{
{
0x02080800, 0x00080000, 0x02000002, 0x02080802,
@ -222,7 +222,7 @@ __constant u32 c_SPtrans[8][64] =
}
};
__constant u32 c_skb[8][64] =
__constant u32a c_skb[8][64] =
{
{
0x00000000, 0x00000010, 0x20000000, 0x20000010,

@ -17,7 +17,7 @@
#include "inc_rp.cl"
#include "inc_simd.cl"
__constant u32 lotus_magic_table[256] =
__constant u32a lotus_magic_table[256] =
{
0xbd, 0x56, 0xea, 0xf2, 0xa2, 0xf1, 0xac, 0x2a,
0xb0, 0x93, 0xd1, 0x9c, 0x1b, 0x33, 0xfd, 0xd0,

@ -15,7 +15,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u32 lotus_magic_table[256] =
__constant u32a lotus_magic_table[256] =
{
0xbd, 0x56, 0xea, 0xf2, 0xa2, 0xf1, 0xac, 0x2a,
0xb0, 0x93, 0xd1, 0x9c, 0x1b, 0x33, 0xfd, 0xd0,

@ -14,7 +14,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u32 lotus_magic_table[256] =
__constant u32a lotus_magic_table[256] =
{
0xbd, 0x56, 0xea, 0xf2, 0xa2, 0xf1, 0xac, 0x2a,
0xb0, 0x93, 0xd1, 0x9c, 0x1b, 0x33, 0xfd, 0xd0,

@ -17,7 +17,7 @@
#include "inc_rp.cl"
#include "inc_simd.cl"
__constant u32 lotus_magic_table[256] =
__constant u32a lotus_magic_table[256] =
{
0xbd, 0x56, 0xea, 0xf2, 0xa2, 0xf1, 0xac, 0x2a,
0xb0, 0x93, 0xd1, 0x9c, 0x1b, 0x33, 0xfd, 0xd0,

@ -15,7 +15,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u32 lotus_magic_table[256] =
__constant u32a lotus_magic_table[256] =
{
0xbd, 0x56, 0xea, 0xf2, 0xa2, 0xf1, 0xac, 0x2a,
0xb0, 0x93, 0xd1, 0x9c, 0x1b, 0x33, 0xfd, 0xd0,

@ -14,7 +14,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u32 lotus_magic_table[256] =
__constant u32a lotus_magic_table[256] =
{
0xbd, 0x56, 0xea, 0xf2, 0xa2, 0xf1, 0xac, 0x2a,
0xb0, 0x93, 0xd1, 0x9c, 0x1b, 0x33, 0xfd, 0xd0,

@ -11,7 +11,7 @@
#include "inc_types.cl"
#include "inc_common.cl"
__constant u32 te0[256] =
__constant u32a te0[256] =
{
0xc66363a5, 0xf87c7c84, 0xee777799, 0xf67b7b8d,
0xfff2f20d, 0xd66b6bbd, 0xde6f6fb1, 0x91c5c554,
@ -79,7 +79,7 @@ __constant u32 te0[256] =
0x7bb0b0cb, 0xa85454fc, 0x6dbbbbd6, 0x2c16163a,
};
__constant u32 te1[256] =
__constant u32a te1[256] =
{
0xa5c66363, 0x84f87c7c, 0x99ee7777, 0x8df67b7b,
0x0dfff2f2, 0xbdd66b6b, 0xb1de6f6f, 0x5491c5c5,
@ -147,7 +147,7 @@ __constant u32 te1[256] =
0xcb7bb0b0, 0xfca85454, 0xd66dbbbb, 0x3a2c1616,
};
__constant u32 te2[256] =
__constant u32a te2[256] =
{
0x63a5c663, 0x7c84f87c, 0x7799ee77, 0x7b8df67b,
0xf20dfff2, 0x6bbdd66b, 0x6fb1de6f, 0xc55491c5,
@ -215,7 +215,7 @@ __constant u32 te2[256] =
0xb0cb7bb0, 0x54fca854, 0xbbd66dbb, 0x163a2c16,
};
__constant u32 te3[256] =
__constant u32a te3[256] =
{
0x6363a5c6, 0x7c7c84f8, 0x777799ee, 0x7b7b8df6,
0xf2f20dff, 0x6b6bbdd6, 0x6f6fb1de, 0xc5c55491,
@ -283,7 +283,7 @@ __constant u32 te3[256] =
0xb0b0cb7b, 0x5454fca8, 0xbbbbd66d, 0x16163a2c,
};
__constant u32 te4[256] =
__constant u32a te4[256] =
{
0x63636363, 0x7c7c7c7c, 0x77777777, 0x7b7b7b7b,
0xf2f2f2f2, 0x6b6b6b6b, 0x6f6f6f6f, 0xc5c5c5c5,
@ -351,7 +351,7 @@ __constant u32 te4[256] =
0xb0b0b0b0, 0x54545454, 0xbbbbbbbb, 0x16161616,
};
__constant u32 td0[256] =
__constant u32a td0[256] =
{
0x51f4a750, 0x7e416553, 0x1a17a4c3, 0x3a275e96,
0x3bab6bcb, 0x1f9d45f1, 0xacfa58ab, 0x4be30393,
@ -419,7 +419,7 @@ __constant u32 td0[256] =
0x7bcb8461, 0xd532b670, 0x486c5c74, 0xd0b85742,
};
__constant u32 td1[256] =
__constant u32a td1[256] =
{
0x5051f4a7, 0x537e4165, 0xc31a17a4, 0x963a275e,
0xcb3bab6b, 0xf11f9d45, 0xabacfa58, 0x934be303,
@ -487,7 +487,7 @@ __constant u32 td1[256] =
0x617bcb84, 0x70d532b6, 0x74486c5c, 0x42d0b857,
};
__constant u32 td2[256] =
__constant u32a td2[256] =
{
0xa75051f4, 0x65537e41, 0xa4c31a17, 0x5e963a27,
0x6bcb3bab, 0x45f11f9d, 0x58abacfa, 0x03934be3,
@ -555,7 +555,7 @@ __constant u32 td2[256] =
0x84617bcb, 0xb670d532, 0x5c74486c, 0x5742d0b8,
};
__constant u32 td3[256] =
__constant u32a td3[256] =
{
0xf4a75051, 0x4165537e, 0x17a4c31a, 0x275e963a,
0xab6bcb3b, 0x9d45f11f, 0xfa58abac, 0xe303934b,
@ -623,7 +623,7 @@ __constant u32 td3[256] =
0xcb84617b, 0x32b670d5, 0x6c5c7448, 0xb85742d0,
};
__constant u32 td4[256] =
__constant u32a td4[256] =
{
0x52525252, 0x09090909, 0x6a6a6a6a, 0xd5d5d5d5,
0x30303030, 0x36363636, 0xa5a5a5a5, 0x38383838,
@ -691,7 +691,7 @@ __constant u32 td4[256] =
0x55555555, 0x21212121, 0x0c0c0c0c, 0x7d7d7d7d,
};
__constant u32 rcon[] =
__constant u32a rcon[] =
{
0x01000000, 0x02000000, 0x04000000, 0x08000000,
0x10000000, 0x20000000, 0x40000000, 0x80000000,
@ -1114,7 +1114,7 @@ void AES256_encrypt (const u32 *in, u32 *out, const u32 *rek, __local u32 *s_te0
^ rek[59];
}
__constant u32 k_sha256[64] =
__constant u32a k_sha256[64] =
{
SHA256C00, SHA256C01, SHA256C02, SHA256C03,
SHA256C04, SHA256C05, SHA256C06, SHA256C07,

@ -14,7 +14,7 @@
#define COMPARE_S "inc_comp_single.cl"
#define COMPARE_M "inc_comp_multi.cl"
__constant u32 k_sha256[64] =
__constant u32a k_sha256[64] =
{
SHA256C00, SHA256C01, SHA256C02, SHA256C03,
SHA256C04, SHA256C05, SHA256C06, SHA256C07,

@ -16,7 +16,7 @@
// http://www.schneier.com/code/constants.txt
__constant u32 c_sbox0[256] =
__constant u32a c_sbox0[256] =
{
0xd1310ba6, 0x98dfb5ac, 0x2ffd72db, 0xd01adfb7,
0xb8e1afed, 0x6a267e96, 0xba7c9045, 0xf12c7f99,
@ -84,7 +84,7 @@ __constant u32 c_sbox0[256] =
0x53b02d5d, 0xa99f8fa1, 0x08ba4799, 0x6e85076a
};
__constant u32 c_sbox1[256] =
__constant u32a c_sbox1[256] =
{
0x4b7a70e9, 0xb5b32944, 0xdb75092e, 0xc4192623,
0xad6ea6b0, 0x49a7df7d, 0x9cee60b8, 0x8fedb266,
@ -152,7 +152,7 @@ __constant u32 c_sbox1[256] =
0x153e21e7, 0x8fb03d4a, 0xe6e39f2b, 0xdb83adf7
};
__constant u32 c_sbox2[256] =
__constant u32a c_sbox2[256] =
{
0xe93d5a68, 0x948140f7, 0xf64c261c, 0x94692934,
0x411520f7, 0x7602d4f7, 0xbcf46b2e, 0xd4a20068,
@ -220,7 +220,7 @@ __constant u32 c_sbox2[256] =
0xd79a3234, 0x92638212, 0x670efa8e, 0x406000e0
};
__constant u32 c_sbox3[256] =
__constant u32a c_sbox3[256] =
{
0x3a39ce37, 0xd3faf5cf, 0xabc27737, 0x5ac52d1b,
0x5cb0679e, 0x4fa33742, 0xd3822740, 0x99bc9bbe,
@ -288,7 +288,7 @@ __constant u32 c_sbox3[256] =
0xb74e6132, 0xce77e25b, 0x578fdfe3, 0x3ac372e6
};
__constant u32 c_pbox[18] =
__constant u32a c_pbox[18] =
{
0x243f6a88, 0x85a308d3, 0x13198a2e, 0x03707344,
0xa4093822, 0x299f31d0, 0x082efa98, 0xec4e6c89,

@ -14,7 +14,7 @@
#define COMPARE_S "inc_comp_single.cl"
#define COMPARE_M "inc_comp_multi.cl"
__constant u32 te0[256] =
__constant u32a te0[256] =
{
0xc66363a5, 0xf87c7c84, 0xee777799, 0xf67b7b8d,
0xfff2f20d, 0xd66b6bbd, 0xde6f6fb1, 0x91c5c554,
@ -82,7 +82,7 @@ __constant u32 te0[256] =
0x7bb0b0cb, 0xa85454fc, 0x6dbbbbd6, 0x2c16163a,
};
__constant u32 te1[256] =
__constant u32a te1[256] =
{
0xa5c66363, 0x84f87c7c, 0x99ee7777, 0x8df67b7b,
0x0dfff2f2, 0xbdd66b6b, 0xb1de6f6f, 0x5491c5c5,
@ -150,7 +150,7 @@ __constant u32 te1[256] =
0xcb7bb0b0, 0xfca85454, 0xd66dbbbb, 0x3a2c1616,
};
__constant u32 te2[256] =
__constant u32a te2[256] =
{
0x63a5c663, 0x7c84f87c, 0x7799ee77, 0x7b8df67b,
0xf20dfff2, 0x6bbdd66b, 0x6fb1de6f, 0xc55491c5,
@ -218,7 +218,7 @@ __constant u32 te2[256] =
0xb0cb7bb0, 0x54fca854, 0xbbd66dbb, 0x163a2c16,
};
__constant u32 te3[256] =
__constant u32a te3[256] =
{
0x6363a5c6, 0x7c7c84f8, 0x777799ee, 0x7b7b8df6,
0xf2f20dff, 0x6b6bbdd6, 0x6f6fb1de, 0xc5c55491,
@ -286,7 +286,7 @@ __constant u32 te3[256] =
0xb0b0cb7b, 0x5454fca8, 0xbbbbd66d, 0x16163a2c,
};
__constant u32 te4[256] =
__constant u32a te4[256] =
{
0x63636363, 0x7c7c7c7c, 0x77777777, 0x7b7b7b7b,
0xf2f2f2f2, 0x6b6b6b6b, 0x6f6f6f6f, 0xc5c5c5c5,
@ -354,7 +354,7 @@ __constant u32 te4[256] =
0xb0b0b0b0, 0x54545454, 0xbbbbbbbb, 0x16161616,
};
__constant u32 td0[256] =
__constant u32a td0[256] =
{
0x51f4a750, 0x7e416553, 0x1a17a4c3, 0x3a275e96,
0x3bab6bcb, 0x1f9d45f1, 0xacfa58ab, 0x4be30393,
@ -422,7 +422,7 @@ __constant u32 td0[256] =
0x7bcb8461, 0xd532b670, 0x486c5c74, 0xd0b85742,
};
__constant u32 td1[256] =
__constant u32a td1[256] =
{
0x5051f4a7, 0x537e4165, 0xc31a17a4, 0x963a275e,
0xcb3bab6b, 0xf11f9d45, 0xabacfa58, 0x934be303,
@ -490,7 +490,7 @@ __constant u32 td1[256] =
0x617bcb84, 0x70d532b6, 0x74486c5c, 0x42d0b857,
};
__constant u32 td2[256] =
__constant u32a td2[256] =
{
0xa75051f4, 0x65537e41, 0xa4c31a17, 0x5e963a27,
0x6bcb3bab, 0x45f11f9d, 0x58abacfa, 0x03934be3,
@ -558,7 +558,7 @@ __constant u32 td2[256] =
0x84617bcb, 0xb670d532, 0x5c74486c, 0x5742d0b8,
};
__constant u32 td3[256] =
__constant u32a td3[256] =
{
0xf4a75051, 0x4165537e, 0x17a4c31a, 0x275e963a,
0xab6bcb3b, 0x9d45f11f, 0xfa58abac, 0xe303934b,
@ -626,7 +626,7 @@ __constant u32 td3[256] =
0xcb84617b, 0x32b670d5, 0x6c5c7448, 0xb85742d0,
};
__constant u32 td4[256] =
__constant u32a td4[256] =
{
0x52525252, 0x09090909, 0x6a6a6a6a, 0xd5d5d5d5,
0x30303030, 0x36363636, 0xa5a5a5a5, 0x38383838,
@ -694,7 +694,7 @@ __constant u32 td4[256] =
0x55555555, 0x21212121, 0x0c0c0c0c, 0x7d7d7d7d,
};
__constant u32 rcon[] =
__constant u32a rcon[] =
{
0x01000000, 0x02000000, 0x04000000, 0x08000000,
0x10000000, 0x20000000, 0x40000000, 0x80000000,

@ -14,7 +14,7 @@
#define COMPARE_S "inc_comp_single.cl"
#define COMPARE_M "inc_comp_multi.cl"
__constant u32 te0[256] =
__constant u32a te0[256] =
{
0xc66363a5, 0xf87c7c84, 0xee777799, 0xf67b7b8d,
0xfff2f20d, 0xd66b6bbd, 0xde6f6fb1, 0x91c5c554,
@ -82,7 +82,7 @@ __constant u32 te0[256] =
0x7bb0b0cb, 0xa85454fc, 0x6dbbbbd6, 0x2c16163a,
};
__constant u32 te1[256] =
__constant u32a te1[256] =
{
0xa5c66363, 0x84f87c7c, 0x99ee7777, 0x8df67b7b,
0x0dfff2f2, 0xbdd66b6b, 0xb1de6f6f, 0x5491c5c5,
@ -150,7 +150,7 @@ __constant u32 te1[256] =
0xcb7bb0b0, 0xfca85454, 0xd66dbbbb, 0x3a2c1616,
};
__constant u32 te2[256] =
__constant u32a te2[256] =
{
0x63a5c663, 0x7c84f87c, 0x7799ee77, 0x7b8df67b,
0xf20dfff2, 0x6bbdd66b, 0x6fb1de6f, 0xc55491c5,
@ -218,7 +218,7 @@ __constant u32 te2[256] =
0xb0cb7bb0, 0x54fca854, 0xbbd66dbb, 0x163a2c16,
};
__constant u32 te3[256] =
__constant u32a te3[256] =
{
0x6363a5c6, 0x7c7c84f8, 0x777799ee, 0x7b7b8df6,
0xf2f20dff, 0x6b6bbdd6, 0x6f6fb1de, 0xc5c55491,
@ -286,7 +286,7 @@ __constant u32 te3[256] =
0xb0b0cb7b, 0x5454fca8, 0xbbbbd66d, 0x16163a2c,
};
__constant u32 te4[256] =
__constant u32a te4[256] =
{
0x63636363, 0x7c7c7c7c, 0x77777777, 0x7b7b7b7b,
0xf2f2f2f2, 0x6b6b6b6b, 0x6f6f6f6f, 0xc5c5c5c5,
@ -354,7 +354,7 @@ __constant u32 te4[256] =
0xb0b0b0b0, 0x54545454, 0xbbbbbbbb, 0x16161616,
};
__constant u32 td0[256] =
__constant u32a td0[256] =
{
0x51f4a750, 0x7e416553, 0x1a17a4c3, 0x3a275e96,
0x3bab6bcb, 0x1f9d45f1, 0xacfa58ab, 0x4be30393,
@ -422,7 +422,7 @@ __constant u32 td0[256] =
0x7bcb8461, 0xd532b670, 0x486c5c74, 0xd0b85742,
};
__constant u32 td1[256] =
__constant u32a td1[256] =
{
0x5051f4a7, 0x537e4165, 0xc31a17a4, 0x963a275e,
0xcb3bab6b, 0xf11f9d45, 0xabacfa58, 0x934be303,
@ -490,7 +490,7 @@ __constant u32 td1[256] =
0x617bcb84, 0x70d532b6, 0x74486c5c, 0x42d0b857,
};
__constant u32 td2[256] =
__constant u32a td2[256] =
{
0xa75051f4, 0x65537e41, 0xa4c31a17, 0x5e963a27,
0x6bcb3bab, 0x45f11f9d, 0x58abacfa, 0x03934be3,
@ -558,7 +558,7 @@ __constant u32 td2[256] =
0x84617bcb, 0xb670d532, 0x5c74486c, 0x5742d0b8,
};
__constant u32 td3[256] =
__constant u32a td3[256] =
{
0xf4a75051, 0x4165537e, 0x17a4c31a, 0x275e963a,
0xab6bcb3b, 0x9d45f11f, 0xfa58abac, 0xe303934b,
@ -626,7 +626,7 @@ __constant u32 td3[256] =
0xcb84617b, 0x32b670d5, 0x6c5c7448, 0xb85742d0,
};
__constant u32 td4[256] =
__constant u32a td4[256] =
{
0x52525252, 0x09090909, 0x6a6a6a6a, 0xd5d5d5d5,
0x30303030, 0x36363636, 0xa5a5a5a5, 0x38383838,
@ -694,7 +694,7 @@ __constant u32 td4[256] =
0x55555555, 0x21212121, 0x0c0c0c0c, 0x7d7d7d7d,
};
__constant u32 rcon[] =
__constant u32a rcon[] =
{
0x01000000, 0x02000000, 0x04000000, 0x08000000,
0x10000000, 0x20000000, 0x40000000, 0x80000000,

@ -14,7 +14,7 @@
#define COMPARE_S "inc_comp_single.cl"
#define COMPARE_M "inc_comp_multi.cl"
__constant u32 te0[256] =
__constant u32a te0[256] =
{
0xc66363a5, 0xf87c7c84, 0xee777799, 0xf67b7b8d,
0xfff2f20d, 0xd66b6bbd, 0xde6f6fb1, 0x91c5c554,
@ -82,7 +82,7 @@ __constant u32 te0[256] =
0x7bb0b0cb, 0xa85454fc, 0x6dbbbbd6, 0x2c16163a,
};
__constant u32 te1[256] =
__constant u32a te1[256] =
{
0xa5c66363, 0x84f87c7c, 0x99ee7777, 0x8df67b7b,
0x0dfff2f2, 0xbdd66b6b, 0xb1de6f6f, 0x5491c5c5,
@ -150,7 +150,7 @@ __constant u32 te1[256] =
0xcb7bb0b0, 0xfca85454, 0xd66dbbbb, 0x3a2c1616,
};
__constant u32 te2[256] =
__constant u32a te2[256] =
{
0x63a5c663, 0x7c84f87c, 0x7799ee77, 0x7b8df67b,
0xf20dfff2, 0x6bbdd66b, 0x6fb1de6f, 0xc55491c5,
@ -218,7 +218,7 @@ __constant u32 te2[256] =
0xb0cb7bb0, 0x54fca854, 0xbbd66dbb, 0x163a2c16,
};
__constant u32 te3[256] =
__constant u32a te3[256] =
{
0x6363a5c6, 0x7c7c84f8, 0x777799ee, 0x7b7b8df6,
0xf2f20dff, 0x6b6bbdd6, 0x6f6fb1de, 0xc5c55491,
@ -286,7 +286,7 @@ __constant u32 te3[256] =
0xb0b0cb7b, 0x5454fca8, 0xbbbbd66d, 0x16163a2c,
};
__constant u32 te4[256] =
__constant u32a te4[256] =
{
0x63636363, 0x7c7c7c7c, 0x77777777, 0x7b7b7b7b,
0xf2f2f2f2, 0x6b6b6b6b, 0x6f6f6f6f, 0xc5c5c5c5,
@ -354,7 +354,7 @@ __constant u32 te4[256] =
0xb0b0b0b0, 0x54545454, 0xbbbbbbbb, 0x16161616,
};
__constant u32 td0[256] =
__constant u32a td0[256] =
{
0x51f4a750, 0x7e416553, 0x1a17a4c3, 0x3a275e96,
0x3bab6bcb, 0x1f9d45f1, 0xacfa58ab, 0x4be30393,
@ -422,7 +422,7 @@ __constant u32 td0[256] =
0x7bcb8461, 0xd532b670, 0x486c5c74, 0xd0b85742,
};
__constant u32 td1[256] =
__constant u32a td1[256] =
{
0x5051f4a7, 0x537e4165, 0xc31a17a4, 0x963a275e,
0xcb3bab6b, 0xf11f9d45, 0xabacfa58, 0x934be303,
@ -490,7 +490,7 @@ __constant u32 td1[256] =
0x617bcb84, 0x70d532b6, 0x74486c5c, 0x42d0b857,
};
__constant u32 td2[256] =
__constant u32a td2[256] =
{
0xa75051f4, 0x65537e41, 0xa4c31a17, 0x5e963a27,
0x6bcb3bab, 0x45f11f9d, 0x58abacfa, 0x03934be3,
@ -558,7 +558,7 @@ __constant u32 td2[256] =
0x84617bcb, 0xb670d532, 0x5c74486c, 0x5742d0b8,
};
__constant u32 td3[256] =
__constant u32a td3[256] =
{
0xf4a75051, 0x4165537e, 0x17a4c31a, 0x275e963a,
0xab6bcb3b, 0x9d45f11f, 0xfa58abac, 0xe303934b,
@ -626,7 +626,7 @@ __constant u32 td3[256] =
0xcb84617b, 0x32b670d5, 0x6c5c7448, 0xb85742d0,
};
__constant u32 td4[256] =
__constant u32a td4[256] =
{
0x52525252, 0x09090909, 0x6a6a6a6a, 0xd5d5d5d5,
0x30303030, 0x36363636, 0xa5a5a5a5, 0x38383838,
@ -694,7 +694,7 @@ __constant u32 td4[256] =
0x55555555, 0x21212121, 0x0c0c0c0c, 0x7d7d7d7d,
};
__constant u32 rcon[] =
__constant u32a rcon[] =
{
0x01000000, 0x02000000, 0x04000000, 0x08000000,
0x10000000, 0x20000000, 0x40000000, 0x80000000,
@ -971,7 +971,7 @@ void AES256_encrypt (const u32 *in, u32 *out, const u32 *rek, __local u32 *s_te0
^ rek[59];
}
__constant u64 k_sha512[80] =
__constant u64a k_sha512[80] =
{
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07,

@ -17,7 +17,7 @@
#include "inc_rp.cl"
#include "inc_simd.cl"
__constant u32 padding[8] =
__constant u32a padding[8] =
{
0x5e4ebf28,
0x418a754e,
@ -84,7 +84,7 @@ void rc4_init_16 (__local RC4_KEY *rc4_key, const u32 data[4])
j += rc4_key->S[255] + d0; swap (rc4_key, 255, j);
}
u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, __constant u32 *in, u32 out[4])
u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, __constant u32a *in, u32 out[4])
{
#ifdef _unroll
#pragma unroll

@ -15,7 +15,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u32 padding[8] =
__constant u32a padding[8] =
{
0x5e4ebf28,
0x418a754e,
@ -82,7 +82,7 @@ void rc4_init_16 (__local RC4_KEY *rc4_key, const u32 data[4])
j += rc4_key->S[255] + d0; swap (rc4_key, 255, j);
}
u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, __constant u32 *in, u32 out[4])
u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, __constant u32a *in, u32 out[4])
{
#ifdef _unroll
#pragma unroll

@ -15,7 +15,7 @@
#include "inc_common.cl"
#include "inc_simd.cl"
__constant u32 padding[8] =
__constant u32a padding[8] =
{
0x5e4ebf28,
0x418a754e,
@ -82,7 +82,7 @@ void rc4_init_16 (__local RC4_KEY *rc4_key, const u32 data[4])
j += rc4_key->S[255] + d0; swap (rc4_key, 255, j);
}
u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, __constant u32 *in, u32 out[4])
u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, __constant u32a *in, u32 out[4])
{
#ifdef _unroll
#pragma unroll

@ -17,7 +17,7 @@
#include "inc_rp.cl"
#include "inc_simd.cl"
__constant u32 padding[8] =
__constant u32a padding[8] =
{
0x5e4ebf28,
0x418a754e,
@ -84,7 +84,7 @@ void rc4_init_16 (__local RC4_KEY *rc4_key, const u32 data[4])
j += rc4_key->S[255] + d0; swap (rc4_key, 255, j);
}
u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, __constant u32 *in, u32 out[4])
u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, __constant u32a *in, u32 out[4])
{
#ifdef _unroll
#pragma unroll

Some files were not shown because too many files have changed in this diff Show More

Loading…
Cancel
Save