From 6401c585689b504406b2e518c67db19dc0a511b4 Mon Sep 17 00:00:00 2001 From: jsteube Date: Sun, 26 Feb 2017 15:04:53 +0100 Subject: [PATCH] Align all the __constant buffers to workaround OpenCL JIT compiler errors in NV drivers 378.x --- OpenCL/inc_cipher_aes.cl | 22 +++++++++--------- OpenCL/inc_cipher_twofish.cl | 4 ++-- OpenCL/inc_luks_af.cl | 4 ++-- OpenCL/inc_luks_essiv.cl | 2 +- OpenCL/inc_truecrypt_crc32.cl | 2 +- OpenCL/m01450_a0.cl | 2 +- OpenCL/m01450_a1.cl | 2 +- OpenCL/m01450_a3.cl | 2 +- OpenCL/m01460_a0.cl | 2 +- OpenCL/m01460_a1.cl | 2 +- OpenCL/m01460_a3.cl | 2 +- OpenCL/m01500_a0.cl | 4 ++-- OpenCL/m01500_a1.cl | 4 ++-- OpenCL/m01700_a0.cl | 2 +- OpenCL/m01700_a1.cl | 2 +- OpenCL/m01700_a3.cl | 2 +- OpenCL/m01710_a0.cl | 2 +- OpenCL/m01710_a1.cl | 2 +- OpenCL/m01710_a3.cl | 2 +- OpenCL/m01720_a0.cl | 2 +- OpenCL/m01720_a1.cl | 2 +- OpenCL/m01720_a3.cl | 2 +- OpenCL/m01730_a0.cl | 2 +- OpenCL/m01730_a1.cl | 2 +- OpenCL/m01730_a3.cl | 2 +- OpenCL/m01740_a0.cl | 2 +- OpenCL/m01740_a1.cl | 2 +- OpenCL/m01740_a3.cl | 2 +- OpenCL/m01750_a0.cl | 2 +- OpenCL/m01750_a1.cl | 2 +- OpenCL/m01750_a3.cl | 2 +- OpenCL/m01760_a0.cl | 2 +- OpenCL/m01760_a1.cl | 2 +- OpenCL/m01760_a3.cl | 2 +- OpenCL/m01800.cl | 2 +- OpenCL/m03000_a0.cl | 4 ++-- OpenCL/m03000_a1.cl | 4 ++-- OpenCL/m03100_a0.cl | 4 ++-- OpenCL/m03100_a1.cl | 4 ++-- OpenCL/m03100_a3.cl | 4 ++-- OpenCL/m03200.cl | 8 +++---- OpenCL/m05000_a0.cl | 2 +- OpenCL/m05000_a1.cl | 2 +- OpenCL/m05000_a3.cl | 2 +- OpenCL/m05200.cl | 2 +- OpenCL/m05500_a0.cl | 4 ++-- OpenCL/m05500_a1.cl | 4 ++-- OpenCL/m05500_a3.cl | 4 ++-- OpenCL/m05800.cl | 4 ++-- OpenCL/m06100_a0.cl | 8 +++---- OpenCL/m06100_a1.cl | 8 +++---- OpenCL/m06100_a3.cl | 8 +++---- OpenCL/m06211.cl | 22 +++++++++--------- OpenCL/m06212.cl | 22 +++++++++--------- OpenCL/m06213.cl | 22 +++++++++--------- OpenCL/m06221.cl | 24 +++++++++---------- OpenCL/m06222.cl | 24 +++++++++---------- OpenCL/m06223.cl | 44 +++++++++++++++++------------------ OpenCL/m06231.cl | 26 ++++++++++----------- OpenCL/m06232.cl | 26 ++++++++++----------- OpenCL/m06233.cl | 26 ++++++++++----------- OpenCL/m06400.cl | 2 +- OpenCL/m06500.cl | 2 +- OpenCL/m06600.cl | 22 +++++++++--------- OpenCL/m06800.cl | 24 +++++++++---------- OpenCL/m06900_a0.cl | 2 +- OpenCL/m06900_a1.cl | 2 +- OpenCL/m06900_a3.cl | 2 +- OpenCL/m07100.cl | 2 +- OpenCL/m07400.cl | 2 +- OpenCL/m07700_a0.cl | 4 ++-- OpenCL/m07700_a1.cl | 4 ++-- OpenCL/m07700_a3.cl | 4 ++-- OpenCL/m07800_a0.cl | 4 ++-- OpenCL/m07800_a1.cl | 4 ++-- OpenCL/m07800_a3.cl | 4 ++-- OpenCL/m07900.cl | 2 +- OpenCL/m08000_a0.cl | 2 +- OpenCL/m08000_a1.cl | 2 +- OpenCL/m08000_a3.cl | 2 +- OpenCL/m08200.cl | 4 ++-- OpenCL/m08500_a0.cl | 6 ++--- OpenCL/m08500_a1.cl | 6 ++--- OpenCL/m08500_a3.cl | 6 ++--- OpenCL/m08600_a0.cl | 2 +- OpenCL/m08600_a1.cl | 2 +- OpenCL/m08600_a3.cl | 2 +- OpenCL/m08700_a0.cl | 2 +- OpenCL/m08700_a1.cl | 2 +- OpenCL/m08700_a3.cl | 2 +- OpenCL/m08800.cl | 24 +++++++++---------- OpenCL/m08900.cl | 2 +- OpenCL/m09000.cl | 10 ++++---- OpenCL/m09400.cl | 22 +++++++++--------- OpenCL/m09500.cl | 22 +++++++++--------- OpenCL/m09600.cl | 24 +++++++++---------- OpenCL/m10400_a0.cl | 4 ++-- OpenCL/m10400_a1.cl | 4 ++-- OpenCL/m10400_a3.cl | 4 ++-- OpenCL/m10410_a0.cl | 4 ++-- OpenCL/m10410_a1.cl | 4 ++-- OpenCL/m10410_a3.cl | 4 ++-- OpenCL/m10420_a0.cl | 2 +- OpenCL/m10420_a1.cl | 2 +- OpenCL/m10420_a3.cl | 2 +- OpenCL/m10500.cl | 2 +- OpenCL/m10700.cl | 18 +++++++------- OpenCL/m10800_a0.cl | 2 +- OpenCL/m10800_a1.cl | 2 +- OpenCL/m10800_a3.cl | 2 +- OpenCL/m10900.cl | 2 +- OpenCL/m11300.cl | 24 +++++++++---------- OpenCL/m11500_a0.cl | 2 +- OpenCL/m11500_a1.cl | 2 +- OpenCL/m11500_a3.cl | 2 +- OpenCL/m11600.cl | 2 +- OpenCL/m11700_a0.cl | 4 ++-- OpenCL/m11700_a1.cl | 4 ++-- OpenCL/m11700_a3.cl | 4 ++-- OpenCL/m11800_a0.cl | 4 ++-- OpenCL/m11800_a1.cl | 4 ++-- OpenCL/m11800_a3.cl | 4 ++-- OpenCL/m12200.cl | 2 +- OpenCL/m12300.cl | 2 +- OpenCL/m12400.cl | 4 ++-- OpenCL/m12500.cl | 22 +++++++++--------- OpenCL/m12700.cl | 22 +++++++++--------- OpenCL/m12800.cl | 2 +- OpenCL/m12900.cl | 2 +- OpenCL/m13000.cl | 2 +- OpenCL/m13200.cl | 22 +++++++++--------- OpenCL/m13400.cl | 34 +++++++++++++-------------- OpenCL/m13751.cl | 24 +++++++++---------- OpenCL/m13752.cl | 24 +++++++++---------- OpenCL/m13753.cl | 24 +++++++++---------- OpenCL/m13800_a0.cl | 2 +- OpenCL/m13800_a1.cl | 2 +- OpenCL/m13800_a3.cl | 2 +- OpenCL/m14000_a0.cl | 4 ++-- OpenCL/m14000_a1.cl | 4 ++-- OpenCL/m14100_a0.cl | 4 ++-- OpenCL/m14100_a1.cl | 4 ++-- OpenCL/m14100_a3.cl | 4 ++-- OpenCL/m14611.cl | 22 +++++++++--------- OpenCL/m14621.cl | 24 +++++++++---------- OpenCL/m14622.cl | 2 +- OpenCL/m14623.cl | 2 +- OpenCL/m14631.cl | 24 +++++++++---------- OpenCL/m14632.cl | 2 +- OpenCL/m14633.cl | 2 +- OpenCL/m14641.cl | 22 +++++++++--------- OpenCL/m14700.cl | 22 +++++++++--------- OpenCL/m14800.cl | 24 +++++++++---------- OpenCL/m14900_a0.cl | 2 +- OpenCL/m14900_a1.cl | 2 +- OpenCL/m14900_a3.cl | 2 +- OpenCL/m15000_a0.cl | 2 +- OpenCL/m15000_a1.cl | 2 +- OpenCL/m15000_a3.cl | 2 +- 159 files changed, 562 insertions(+), 562 deletions(-) diff --git a/OpenCL/inc_cipher_aes.cl b/OpenCL/inc_cipher_aes.cl index 575ad9a5d..4add5775c 100644 --- a/OpenCL/inc_cipher_aes.cl +++ b/OpenCL/inc_cipher_aes.cl @@ -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, diff --git a/OpenCL/inc_cipher_twofish.cl b/OpenCL/inc_cipher_twofish.cl index 9d90ce155..a81611cc8 100644 --- a/OpenCL/inc_cipher_twofish.cl +++ b/OpenCL/inc_cipher_twofish.cl @@ -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, diff --git a/OpenCL/inc_luks_af.cl b/OpenCL/inc_luks_af.cl index 57e7ef040..47059bca1 100644 --- a/OpenCL/inc_luks_af.cl +++ b/OpenCL/inc_luks_af.cl @@ -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, diff --git a/OpenCL/inc_luks_essiv.cl b/OpenCL/inc_luks_essiv.cl index 49c9cb7da..5011ffed4 100644 --- a/OpenCL/inc_luks_essiv.cl +++ b/OpenCL/inc_luks_essiv.cl @@ -1,4 +1,4 @@ -__constant u32 ESSIV_k_sha256[64] = +__constant u32a ESSIV_k_sha256[64] = { SHA256C00, SHA256C01, SHA256C02, SHA256C03, SHA256C04, SHA256C05, SHA256C06, SHA256C07, diff --git a/OpenCL/inc_truecrypt_crc32.cl b/OpenCL/inc_truecrypt_crc32.cl index 72936b328..f6b5a5eaa 100644 --- a/OpenCL/inc_truecrypt_crc32.cl +++ b/OpenCL/inc_truecrypt_crc32.cl @@ -1,4 +1,4 @@ -__constant u32 crc32tab[0x100] = +__constant u32a crc32tab[0x100] = { 0x00000000, 0x77073096, 0xee0e612c, 0x990951ba, 0x076dc419, 0x706af48f, 0xe963a535, 0x9e6495a3, diff --git a/OpenCL/m01450_a0.cl b/OpenCL/m01450_a0.cl index c0d15b320..f3f2e6b1d 100644 --- a/OpenCL/m01450_a0.cl +++ b/OpenCL/m01450_a0.cl @@ -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, diff --git a/OpenCL/m01450_a1.cl b/OpenCL/m01450_a1.cl index 122454029..861afe9d8 100644 --- a/OpenCL/m01450_a1.cl +++ b/OpenCL/m01450_a1.cl @@ -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, diff --git a/OpenCL/m01450_a3.cl b/OpenCL/m01450_a3.cl index 1e2fcad05..f7e0dd050 100644 --- a/OpenCL/m01450_a3.cl +++ b/OpenCL/m01450_a3.cl @@ -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, diff --git a/OpenCL/m01460_a0.cl b/OpenCL/m01460_a0.cl index c076781be..186946265 100644 --- a/OpenCL/m01460_a0.cl +++ b/OpenCL/m01460_a0.cl @@ -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, diff --git a/OpenCL/m01460_a1.cl b/OpenCL/m01460_a1.cl index 1c45f96b6..3dffec37d 100644 --- a/OpenCL/m01460_a1.cl +++ b/OpenCL/m01460_a1.cl @@ -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, diff --git a/OpenCL/m01460_a3.cl b/OpenCL/m01460_a3.cl index b2686ce79..995824075 100644 --- a/OpenCL/m01460_a3.cl +++ b/OpenCL/m01460_a3.cl @@ -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, diff --git a/OpenCL/m01500_a0.cl b/OpenCL/m01500_a0.cl index 4e7e3a493..5dd853585 100644 --- a/OpenCL/m01500_a0.cl +++ b/OpenCL/m01500_a0.cl @@ -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, diff --git a/OpenCL/m01500_a1.cl b/OpenCL/m01500_a1.cl index 9ef45dd89..eb09c9284 100644 --- a/OpenCL/m01500_a1.cl +++ b/OpenCL/m01500_a1.cl @@ -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, diff --git a/OpenCL/m01700_a0.cl b/OpenCL/m01700_a0.cl index 598dc675e..3f1f30ae4 100644 --- a/OpenCL/m01700_a0.cl +++ b/OpenCL/m01700_a0.cl @@ -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, diff --git a/OpenCL/m01700_a1.cl b/OpenCL/m01700_a1.cl index 39932b71d..2d5454d58 100644 --- a/OpenCL/m01700_a1.cl +++ b/OpenCL/m01700_a1.cl @@ -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, diff --git a/OpenCL/m01700_a3.cl b/OpenCL/m01700_a3.cl index e23865ca6..cfcac6702 100644 --- a/OpenCL/m01700_a3.cl +++ b/OpenCL/m01700_a3.cl @@ -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, diff --git a/OpenCL/m01710_a0.cl b/OpenCL/m01710_a0.cl index 05ecc1c72..87b8e0da4 100644 --- a/OpenCL/m01710_a0.cl +++ b/OpenCL/m01710_a0.cl @@ -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, diff --git a/OpenCL/m01710_a1.cl b/OpenCL/m01710_a1.cl index f334d0c10..b2fa62dbf 100644 --- a/OpenCL/m01710_a1.cl +++ b/OpenCL/m01710_a1.cl @@ -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, diff --git a/OpenCL/m01710_a3.cl b/OpenCL/m01710_a3.cl index 3bf3f407f..63500b4a2 100644 --- a/OpenCL/m01710_a3.cl +++ b/OpenCL/m01710_a3.cl @@ -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, diff --git a/OpenCL/m01720_a0.cl b/OpenCL/m01720_a0.cl index c79353ef4..75a72ff2d 100644 --- a/OpenCL/m01720_a0.cl +++ b/OpenCL/m01720_a0.cl @@ -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, diff --git a/OpenCL/m01720_a1.cl b/OpenCL/m01720_a1.cl index ba8917d0b..e3f7875b2 100644 --- a/OpenCL/m01720_a1.cl +++ b/OpenCL/m01720_a1.cl @@ -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, diff --git a/OpenCL/m01720_a3.cl b/OpenCL/m01720_a3.cl index 12fe0b26d..a107e5bc2 100644 --- a/OpenCL/m01720_a3.cl +++ b/OpenCL/m01720_a3.cl @@ -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, diff --git a/OpenCL/m01730_a0.cl b/OpenCL/m01730_a0.cl index c57ce012a..2d373fab1 100644 --- a/OpenCL/m01730_a0.cl +++ b/OpenCL/m01730_a0.cl @@ -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, diff --git a/OpenCL/m01730_a1.cl b/OpenCL/m01730_a1.cl index fab4abb58..8ec1cc1e1 100644 --- a/OpenCL/m01730_a1.cl +++ b/OpenCL/m01730_a1.cl @@ -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, diff --git a/OpenCL/m01730_a3.cl b/OpenCL/m01730_a3.cl index a0f606c2b..b3ec8aa85 100644 --- a/OpenCL/m01730_a3.cl +++ b/OpenCL/m01730_a3.cl @@ -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, diff --git a/OpenCL/m01740_a0.cl b/OpenCL/m01740_a0.cl index 6c9ed6391..09726036f 100644 --- a/OpenCL/m01740_a0.cl +++ b/OpenCL/m01740_a0.cl @@ -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, diff --git a/OpenCL/m01740_a1.cl b/OpenCL/m01740_a1.cl index c3576a4a7..01e71bdc5 100644 --- a/OpenCL/m01740_a1.cl +++ b/OpenCL/m01740_a1.cl @@ -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, diff --git a/OpenCL/m01740_a3.cl b/OpenCL/m01740_a3.cl index f5cfd9f62..49198aee8 100644 --- a/OpenCL/m01740_a3.cl +++ b/OpenCL/m01740_a3.cl @@ -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, diff --git a/OpenCL/m01750_a0.cl b/OpenCL/m01750_a0.cl index e91cd1a37..837a88801 100644 --- a/OpenCL/m01750_a0.cl +++ b/OpenCL/m01750_a0.cl @@ -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, diff --git a/OpenCL/m01750_a1.cl b/OpenCL/m01750_a1.cl index e4df9d028..319e4e8ad 100644 --- a/OpenCL/m01750_a1.cl +++ b/OpenCL/m01750_a1.cl @@ -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, diff --git a/OpenCL/m01750_a3.cl b/OpenCL/m01750_a3.cl index bc70b88a8..99016b135 100644 --- a/OpenCL/m01750_a3.cl +++ b/OpenCL/m01750_a3.cl @@ -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, diff --git a/OpenCL/m01760_a0.cl b/OpenCL/m01760_a0.cl index f703c9584..b391849c9 100644 --- a/OpenCL/m01760_a0.cl +++ b/OpenCL/m01760_a0.cl @@ -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, diff --git a/OpenCL/m01760_a1.cl b/OpenCL/m01760_a1.cl index 9c31f27d4..548dac5a0 100644 --- a/OpenCL/m01760_a1.cl +++ b/OpenCL/m01760_a1.cl @@ -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, diff --git a/OpenCL/m01760_a3.cl b/OpenCL/m01760_a3.cl index dcfb4d3dc..b3e5d387e 100644 --- a/OpenCL/m01760_a3.cl +++ b/OpenCL/m01760_a3.cl @@ -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, diff --git a/OpenCL/m01800.cl b/OpenCL/m01800.cl index 3c640b450..776453791 100644 --- a/OpenCL/m01800.cl +++ b/OpenCL/m01800.cl @@ -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, diff --git a/OpenCL/m03000_a0.cl b/OpenCL/m03000_a0.cl index e10eea313..d5b646a3b 100644 --- a/OpenCL/m03000_a0.cl +++ b/OpenCL/m03000_a0.cl @@ -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, diff --git a/OpenCL/m03000_a1.cl b/OpenCL/m03000_a1.cl index affab99ea..7bacb25ad 100644 --- a/OpenCL/m03000_a1.cl +++ b/OpenCL/m03000_a1.cl @@ -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, diff --git a/OpenCL/m03100_a0.cl b/OpenCL/m03100_a0.cl index 926f6aa39..12e03e71d 100644 --- a/OpenCL/m03100_a0.cl +++ b/OpenCL/m03100_a0.cl @@ -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, diff --git a/OpenCL/m03100_a1.cl b/OpenCL/m03100_a1.cl index ba512cf64..316531476 100644 --- a/OpenCL/m03100_a1.cl +++ b/OpenCL/m03100_a1.cl @@ -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, diff --git a/OpenCL/m03100_a3.cl b/OpenCL/m03100_a3.cl index f27da7108..e7124d368 100644 --- a/OpenCL/m03100_a3.cl +++ b/OpenCL/m03100_a3.cl @@ -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, diff --git a/OpenCL/m03200.cl b/OpenCL/m03200.cl index a7fc3b4a9..eadd183e8 100644 --- a/OpenCL/m03200.cl +++ b/OpenCL/m03200.cl @@ -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, diff --git a/OpenCL/m05000_a0.cl b/OpenCL/m05000_a0.cl index a2bcfbb94..86be1830c 100644 --- a/OpenCL/m05000_a0.cl +++ b/OpenCL/m05000_a0.cl @@ -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, diff --git a/OpenCL/m05000_a1.cl b/OpenCL/m05000_a1.cl index 5aa555499..ee0d51ee4 100644 --- a/OpenCL/m05000_a1.cl +++ b/OpenCL/m05000_a1.cl @@ -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, diff --git a/OpenCL/m05000_a3.cl b/OpenCL/m05000_a3.cl index 5d4b38559..d5e567c02 100644 --- a/OpenCL/m05000_a3.cl +++ b/OpenCL/m05000_a3.cl @@ -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, diff --git a/OpenCL/m05200.cl b/OpenCL/m05200.cl index a1e021ee3..ef47f6411 100644 --- a/OpenCL/m05200.cl +++ b/OpenCL/m05200.cl @@ -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, diff --git a/OpenCL/m05500_a0.cl b/OpenCL/m05500_a0.cl index e036e7b41..fb9a47683 100644 --- a/OpenCL/m05500_a0.cl +++ b/OpenCL/m05500_a0.cl @@ -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, diff --git a/OpenCL/m05500_a1.cl b/OpenCL/m05500_a1.cl index 8b61636ac..06914fb38 100644 --- a/OpenCL/m05500_a1.cl +++ b/OpenCL/m05500_a1.cl @@ -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, diff --git a/OpenCL/m05500_a3.cl b/OpenCL/m05500_a3.cl index 065490595..eb210585a 100644 --- a/OpenCL/m05500_a3.cl +++ b/OpenCL/m05500_a3.cl @@ -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, diff --git a/OpenCL/m05800.cl b/OpenCL/m05800.cl index 343d5d491..c855220dd 100644 --- a/OpenCL/m05800.cl +++ b/OpenCL/m05800.cl @@ -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, diff --git a/OpenCL/m06100_a0.cl b/OpenCL/m06100_a0.cl index adb7cbad5..6ac57662c 100644 --- a/OpenCL/m06100_a0.cl +++ b/OpenCL/m06100_a0.cl @@ -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, diff --git a/OpenCL/m06100_a1.cl b/OpenCL/m06100_a1.cl index 878058f38..d4b5dc647 100644 --- a/OpenCL/m06100_a1.cl +++ b/OpenCL/m06100_a1.cl @@ -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, diff --git a/OpenCL/m06100_a3.cl b/OpenCL/m06100_a3.cl index 5cb6e10e6..ba54cfc06 100644 --- a/OpenCL/m06100_a3.cl +++ b/OpenCL/m06100_a3.cl @@ -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, diff --git a/OpenCL/m06211.cl b/OpenCL/m06211.cl index aeaa6c76b..e2c185e5d 100644 --- a/OpenCL/m06211.cl +++ b/OpenCL/m06211.cl @@ -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 diff --git a/OpenCL/m06212.cl b/OpenCL/m06212.cl index ad797a439..67af580c4 100644 --- a/OpenCL/m06212.cl +++ b/OpenCL/m06212.cl @@ -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 diff --git a/OpenCL/m06213.cl b/OpenCL/m06213.cl index 6b345796b..3ddc8d673 100644 --- a/OpenCL/m06213.cl +++ b/OpenCL/m06213.cl @@ -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 diff --git a/OpenCL/m06221.cl b/OpenCL/m06221.cl index dedca8b34..bd3e26ea6 100644 --- a/OpenCL/m06221.cl +++ b/OpenCL/m06221.cl @@ -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 diff --git a/OpenCL/m06222.cl b/OpenCL/m06222.cl index 797bd191b..dc75a04b0 100644 --- a/OpenCL/m06222.cl +++ b/OpenCL/m06222.cl @@ -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 diff --git a/OpenCL/m06223.cl b/OpenCL/m06223.cl index 02adac994..8a4c72e0b 100644 --- a/OpenCL/m06223.cl +++ b/OpenCL/m06223.cl @@ -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 diff --git a/OpenCL/m06231.cl b/OpenCL/m06231.cl index a5530aa8a..76ae4da95 100644 --- a/OpenCL/m06231.cl +++ b/OpenCL/m06231.cl @@ -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 diff --git a/OpenCL/m06232.cl b/OpenCL/m06232.cl index 38a6acc6e..ece775e11 100644 --- a/OpenCL/m06232.cl +++ b/OpenCL/m06232.cl @@ -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 diff --git a/OpenCL/m06233.cl b/OpenCL/m06233.cl index 666d35f86..24779d488 100644 --- a/OpenCL/m06233.cl +++ b/OpenCL/m06233.cl @@ -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 diff --git a/OpenCL/m06400.cl b/OpenCL/m06400.cl index 40e79e86a..fda5bee49 100644 --- a/OpenCL/m06400.cl +++ b/OpenCL/m06400.cl @@ -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, diff --git a/OpenCL/m06500.cl b/OpenCL/m06500.cl index 9ec6cc61a..eef0653dc 100644 --- a/OpenCL/m06500.cl +++ b/OpenCL/m06500.cl @@ -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, diff --git a/OpenCL/m06600.cl b/OpenCL/m06600.cl index d53eabba4..0b5859dd0 100644 --- a/OpenCL/m06600.cl +++ b/OpenCL/m06600.cl @@ -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, diff --git a/OpenCL/m06800.cl b/OpenCL/m06800.cl index 951ce06c4..8ab2faf44 100644 --- a/OpenCL/m06800.cl +++ b/OpenCL/m06800.cl @@ -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, diff --git a/OpenCL/m06900_a0.cl b/OpenCL/m06900_a0.cl index 18e93f2fe..b9bffa474 100644 --- a/OpenCL/m06900_a0.cl +++ b/OpenCL/m06900_a0.cl @@ -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, diff --git a/OpenCL/m06900_a1.cl b/OpenCL/m06900_a1.cl index 33f3586b0..4c93475ac 100644 --- a/OpenCL/m06900_a1.cl +++ b/OpenCL/m06900_a1.cl @@ -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, diff --git a/OpenCL/m06900_a3.cl b/OpenCL/m06900_a3.cl index dc24f052f..5170a78bd 100644 --- a/OpenCL/m06900_a3.cl +++ b/OpenCL/m06900_a3.cl @@ -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, diff --git a/OpenCL/m07100.cl b/OpenCL/m07100.cl index db7e23e51..d7e7df41c 100644 --- a/OpenCL/m07100.cl +++ b/OpenCL/m07100.cl @@ -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, diff --git a/OpenCL/m07400.cl b/OpenCL/m07400.cl index 454761700..543276a49 100644 --- a/OpenCL/m07400.cl +++ b/OpenCL/m07400.cl @@ -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, diff --git a/OpenCL/m07700_a0.cl b/OpenCL/m07700_a0.cl index 762c18f4a..8b3da6771 100644 --- a/OpenCL/m07700_a0.cl +++ b/OpenCL/m07700_a0.cl @@ -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, diff --git a/OpenCL/m07700_a1.cl b/OpenCL/m07700_a1.cl index 5e02856d2..10c1f001d 100644 --- a/OpenCL/m07700_a1.cl +++ b/OpenCL/m07700_a1.cl @@ -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, diff --git a/OpenCL/m07700_a3.cl b/OpenCL/m07700_a3.cl index bab0f66bd..c051e5fcd 100644 --- a/OpenCL/m07700_a3.cl +++ b/OpenCL/m07700_a3.cl @@ -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, diff --git a/OpenCL/m07800_a0.cl b/OpenCL/m07800_a0.cl index 939fd8895..83706c9ed 100644 --- a/OpenCL/m07800_a0.cl +++ b/OpenCL/m07800_a0.cl @@ -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; diff --git a/OpenCL/m07800_a1.cl b/OpenCL/m07800_a1.cl index 4b597d490..49272c380 100644 --- a/OpenCL/m07800_a1.cl +++ b/OpenCL/m07800_a1.cl @@ -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; diff --git a/OpenCL/m07800_a3.cl b/OpenCL/m07800_a3.cl index eca037a94..6b3ac671c 100644 --- a/OpenCL/m07800_a3.cl +++ b/OpenCL/m07800_a3.cl @@ -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; diff --git a/OpenCL/m07900.cl b/OpenCL/m07900.cl index d45d154ff..b35bdf25f 100644 --- a/OpenCL/m07900.cl +++ b/OpenCL/m07900.cl @@ -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, diff --git a/OpenCL/m08000_a0.cl b/OpenCL/m08000_a0.cl index 904cb4bb2..175c0b487 100644 --- a/OpenCL/m08000_a0.cl +++ b/OpenCL/m08000_a0.cl @@ -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, diff --git a/OpenCL/m08000_a1.cl b/OpenCL/m08000_a1.cl index 80dcde3ca..dd7205d15 100644 --- a/OpenCL/m08000_a1.cl +++ b/OpenCL/m08000_a1.cl @@ -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, diff --git a/OpenCL/m08000_a3.cl b/OpenCL/m08000_a3.cl index 22a406028..6f29689d8 100644 --- a/OpenCL/m08000_a3.cl +++ b/OpenCL/m08000_a3.cl @@ -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, diff --git a/OpenCL/m08200.cl b/OpenCL/m08200.cl index 940722382..9a94ff1ab 100644 --- a/OpenCL/m08200.cl +++ b/OpenCL/m08200.cl @@ -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, diff --git a/OpenCL/m08500_a0.cl b/OpenCL/m08500_a0.cl index 77ccb3dcc..e646b4b49 100644 --- a/OpenCL/m08500_a0.cl +++ b/OpenCL/m08500_a0.cl @@ -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, diff --git a/OpenCL/m08500_a1.cl b/OpenCL/m08500_a1.cl index 6ed49c241..755c2f8e0 100644 --- a/OpenCL/m08500_a1.cl +++ b/OpenCL/m08500_a1.cl @@ -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, diff --git a/OpenCL/m08500_a3.cl b/OpenCL/m08500_a3.cl index 9f21a31a5..cd43b8d07 100644 --- a/OpenCL/m08500_a3.cl +++ b/OpenCL/m08500_a3.cl @@ -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, diff --git a/OpenCL/m08600_a0.cl b/OpenCL/m08600_a0.cl index bdc8356ab..a99187d04 100644 --- a/OpenCL/m08600_a0.cl +++ b/OpenCL/m08600_a0.cl @@ -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, diff --git a/OpenCL/m08600_a1.cl b/OpenCL/m08600_a1.cl index ddcefdc9f..8b901c806 100644 --- a/OpenCL/m08600_a1.cl +++ b/OpenCL/m08600_a1.cl @@ -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, diff --git a/OpenCL/m08600_a3.cl b/OpenCL/m08600_a3.cl index 0e0e712e2..f1bdf9b6b 100644 --- a/OpenCL/m08600_a3.cl +++ b/OpenCL/m08600_a3.cl @@ -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, diff --git a/OpenCL/m08700_a0.cl b/OpenCL/m08700_a0.cl index 41894e8a4..c50088470 100644 --- a/OpenCL/m08700_a0.cl +++ b/OpenCL/m08700_a0.cl @@ -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, diff --git a/OpenCL/m08700_a1.cl b/OpenCL/m08700_a1.cl index 797392640..044e58256 100644 --- a/OpenCL/m08700_a1.cl +++ b/OpenCL/m08700_a1.cl @@ -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, diff --git a/OpenCL/m08700_a3.cl b/OpenCL/m08700_a3.cl index 1ea5cc0bc..557d024dc 100644 --- a/OpenCL/m08700_a3.cl +++ b/OpenCL/m08700_a3.cl @@ -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, diff --git a/OpenCL/m08800.cl b/OpenCL/m08800.cl index 9a0a491e9..644961ae2 100644 --- a/OpenCL/m08800.cl +++ b/OpenCL/m08800.cl @@ -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, diff --git a/OpenCL/m08900.cl b/OpenCL/m08900.cl index 1e30408f1..740db2646 100644 --- a/OpenCL/m08900.cl +++ b/OpenCL/m08900.cl @@ -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, diff --git a/OpenCL/m09000.cl b/OpenCL/m09000.cl index 6c5863f68..80afc73d7 100644 --- a/OpenCL/m09000.cl +++ b/OpenCL/m09000.cl @@ -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, diff --git a/OpenCL/m09400.cl b/OpenCL/m09400.cl index 8c1eccc2c..12d4db5ac 100644 --- a/OpenCL/m09400.cl +++ b/OpenCL/m09400.cl @@ -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, diff --git a/OpenCL/m09500.cl b/OpenCL/m09500.cl index b749a8076..08bb6e36c 100644 --- a/OpenCL/m09500.cl +++ b/OpenCL/m09500.cl @@ -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, diff --git a/OpenCL/m09600.cl b/OpenCL/m09600.cl index 24be89f8a..04938c66e 100644 --- a/OpenCL/m09600.cl +++ b/OpenCL/m09600.cl @@ -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, diff --git a/OpenCL/m10400_a0.cl b/OpenCL/m10400_a0.cl index e67c5fec2..f03b08f7b 100644 --- a/OpenCL/m10400_a0.cl +++ b/OpenCL/m10400_a0.cl @@ -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 diff --git a/OpenCL/m10400_a1.cl b/OpenCL/m10400_a1.cl index 91a6fa3a5..e016b3cf7 100644 --- a/OpenCL/m10400_a1.cl +++ b/OpenCL/m10400_a1.cl @@ -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 diff --git a/OpenCL/m10400_a3.cl b/OpenCL/m10400_a3.cl index b79520401..976fa29ee 100644 --- a/OpenCL/m10400_a3.cl +++ b/OpenCL/m10400_a3.cl @@ -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 diff --git a/OpenCL/m10410_a0.cl b/OpenCL/m10410_a0.cl index 10ad3bfd2..31b9b6ff5 100644 --- a/OpenCL/m10410_a0.cl +++ b/OpenCL/m10410_a0.cl @@ -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 diff --git a/OpenCL/m10410_a1.cl b/OpenCL/m10410_a1.cl index f2d48b89e..3f11c89ee 100644 --- a/OpenCL/m10410_a1.cl +++ b/OpenCL/m10410_a1.cl @@ -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 diff --git a/OpenCL/m10410_a3.cl b/OpenCL/m10410_a3.cl index 306137fa5..d3980fcfb 100644 --- a/OpenCL/m10410_a3.cl +++ b/OpenCL/m10410_a3.cl @@ -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 diff --git a/OpenCL/m10420_a0.cl b/OpenCL/m10420_a0.cl index 69f2fb1a2..55201e753 100644 --- a/OpenCL/m10420_a0.cl +++ b/OpenCL/m10420_a0.cl @@ -16,7 +16,7 @@ #include "inc_rp.cl" #include "inc_simd.cl" -__constant u32 padding[8] = +__constant u32a padding[8] = { 0x5e4ebf28, 0x418a754e, diff --git a/OpenCL/m10420_a1.cl b/OpenCL/m10420_a1.cl index c2d6e1020..dbfa292d7 100644 --- a/OpenCL/m10420_a1.cl +++ b/OpenCL/m10420_a1.cl @@ -14,7 +14,7 @@ #include "inc_common.cl" #include "inc_simd.cl" -__constant u32 padding[8] = +__constant u32a padding[8] = { 0x5e4ebf28, 0x418a754e, diff --git a/OpenCL/m10420_a3.cl b/OpenCL/m10420_a3.cl index 901df8f8e..b014d07e3 100644 --- a/OpenCL/m10420_a3.cl +++ b/OpenCL/m10420_a3.cl @@ -14,7 +14,7 @@ #include "inc_common.cl" #include "inc_simd.cl" -__constant u32 padding[8] = +__constant u32a padding[8] = { 0x5e4ebf28, 0x418a754e, diff --git a/OpenCL/m10500.cl b/OpenCL/m10500.cl index 7580a99f8..8045091e8 100644 --- a/OpenCL/m10500.cl +++ b/OpenCL/m10500.cl @@ -14,7 +14,7 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" -__constant u32 padding[8] = +__constant u32a padding[8] = { 0x5e4ebf28, 0x418a754e, diff --git a/OpenCL/m10700.cl b/OpenCL/m10700.cl index 4e5f328e5..4e86bf9da 100644 --- a/OpenCL/m10700.cl +++ b/OpenCL/m10700.cl @@ -34,7 +34,7 @@ typedef struct } ctx_t; -__constant u32 k_sha256[64] = +__constant u32a k_sha256[64] = { SHA256C00, SHA256C01, SHA256C02, SHA256C03, SHA256C04, SHA256C05, SHA256C06, SHA256C07, @@ -142,7 +142,7 @@ void sha256_transform (const u32 *w0, const u32 *w1, const u32 *w2, const u32 *w digest[7] += h; } -__constant u64 k_sha384[80] = +__constant u64a k_sha384[80] = { SHA512C00, SHA512C01, SHA512C02, SHA512C03, SHA512C04, SHA512C05, SHA512C06, SHA512C07, @@ -254,7 +254,7 @@ void sha384_transform (const u64 *w0, const u64 *w1, const u64 *w2, const u64 *w digest[7] += h; } -__constant u64 k_sha512[80] = +__constant u64a k_sha512[80] = { SHA512C00, SHA512C01, SHA512C02, SHA512C03, SHA512C04, SHA512C05, SHA512C06, SHA512C07, @@ -366,7 +366,7 @@ void sha512_transform (const u64 *w0, const u64 *w1, const u64 *w2, const u64 *w digest[7] += h; } -__constant u32 te0[256] = +__constant u32a te0[256] = { 0xc66363a5, 0xf87c7c84, 0xee777799, 0xf67b7b8d, 0xfff2f20d, 0xd66b6bbd, 0xde6f6fb1, 0x91c5c554, @@ -434,7 +434,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, @@ -502,7 +502,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, @@ -570,7 +570,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, @@ -638,7 +638,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, @@ -706,7 +706,7 @@ __constant u32 te4[256] = 0xb0b0b0b0, 0x54545454, 0xbbbbbbbb, 0x16161616, }; -__constant u32 rcon[] = +__constant u32a rcon[] = { 0x01000000, 0x02000000, 0x04000000, 0x08000000, 0x10000000, 0x20000000, 0x40000000, 0x80000000, diff --git a/OpenCL/m10800_a0.cl b/OpenCL/m10800_a0.cl index 29d4306a6..9acaea6f7 100644 --- a/OpenCL/m10800_a0.cl +++ b/OpenCL/m10800_a0.cl @@ -16,7 +16,7 @@ #include "inc_rp.cl" #include "inc_simd.cl" -__constant u64 k_sha384[80] = +__constant u64a k_sha384[80] = { SHA512C00, SHA512C01, SHA512C02, SHA512C03, SHA512C04, SHA512C05, SHA512C06, SHA512C07, diff --git a/OpenCL/m10800_a1.cl b/OpenCL/m10800_a1.cl index b8eeda2cc..0ba493699 100644 --- a/OpenCL/m10800_a1.cl +++ b/OpenCL/m10800_a1.cl @@ -14,7 +14,7 @@ #include "inc_common.cl" #include "inc_simd.cl" -__constant u64 k_sha384[80] = +__constant u64a k_sha384[80] = { SHA512C00, SHA512C01, SHA512C02, SHA512C03, SHA512C04, SHA512C05, SHA512C06, SHA512C07, diff --git a/OpenCL/m10800_a3.cl b/OpenCL/m10800_a3.cl index 1e686da0e..e40f292c7 100644 --- a/OpenCL/m10800_a3.cl +++ b/OpenCL/m10800_a3.cl @@ -14,7 +14,7 @@ #include "inc_common.cl" #include "inc_simd.cl" -__constant u64 k_sha384[80] = +__constant u64a k_sha384[80] = { SHA512C00, SHA512C01, SHA512C02, SHA512C03, SHA512C04, SHA512C05, SHA512C06, SHA512C07, diff --git a/OpenCL/m10900.cl b/OpenCL/m10900.cl index e6ccfcdfa..1c61b2698 100644 --- a/OpenCL/m10900.cl +++ b/OpenCL/m10900.cl @@ -17,7 +17,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, diff --git a/OpenCL/m11300.cl b/OpenCL/m11300.cl index f42dd66cd..90373f268 100644 --- a/OpenCL/m11300.cl +++ b/OpenCL/m11300.cl @@ -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, @@ -884,7 +884,7 @@ void AES256_decrypt (const u32 *in, u32 *out, const u32 *rdk, __local u32 *s_td0 ^ rdk[59]; } -__constant u64 k_sha512[80] = +__constant u64a k_sha512[80] = { SHA512C00, SHA512C01, SHA512C02, SHA512C03, SHA512C04, SHA512C05, SHA512C06, SHA512C07, diff --git a/OpenCL/m11500_a0.cl b/OpenCL/m11500_a0.cl index 9f461cf79..a94b221e0 100644 --- a/OpenCL/m11500_a0.cl +++ b/OpenCL/m11500_a0.cl @@ -17,7 +17,7 @@ #include "inc_rp.cl" #include "inc_simd.cl" -__constant u32 crc32tab[0x100] = +__constant u32a crc32tab[0x100] = { 0x00000000, 0x77073096, 0xee0e612c, 0x990951ba, 0x076dc419, 0x706af48f, 0xe963a535, 0x9e6495a3, diff --git a/OpenCL/m11500_a1.cl b/OpenCL/m11500_a1.cl index d89ebeca9..8e0a512f3 100644 --- a/OpenCL/m11500_a1.cl +++ b/OpenCL/m11500_a1.cl @@ -15,7 +15,7 @@ #include "inc_common.cl" #include "inc_simd.cl" -__constant u32 crc32tab[0x100] = +__constant u32a crc32tab[0x100] = { 0x00000000, 0x77073096, 0xee0e612c, 0x990951ba, 0x076dc419, 0x706af48f, 0xe963a535, 0x9e6495a3, diff --git a/OpenCL/m11500_a3.cl b/OpenCL/m11500_a3.cl index 749e22efa..cf658076c 100644 --- a/OpenCL/m11500_a3.cl +++ b/OpenCL/m11500_a3.cl @@ -15,7 +15,7 @@ #include "inc_common.cl" #include "inc_simd.cl" -__constant u32 crc32tab[0x100] = +__constant u32a crc32tab[0x100] = { 0x00000000, 0x77073096, 0xee0e612c, 0x990951ba, 0x076dc419, 0x706af48f, 0xe963a535, 0x9e6495a3, diff --git a/OpenCL/m11600.cl b/OpenCL/m11600.cl index 00006d356..e34679e61 100644 --- a/OpenCL/m11600.cl +++ b/OpenCL/m11600.cl @@ -11,7 +11,7 @@ #include "inc_types.cl" #include "inc_common.cl" -__constant u32 k_sha256[64] = +__constant u32a k_sha256[64] = { SHA256C00, SHA256C01, SHA256C02, SHA256C03, SHA256C04, SHA256C05, SHA256C06, SHA256C07, diff --git a/OpenCL/m11700_a0.cl b/OpenCL/m11700_a0.cl index 1f139d9f6..16c39ca79 100644 --- a/OpenCL/m11700_a0.cl +++ b/OpenCL/m11700_a0.cl @@ -43,7 +43,7 @@ // constants -__constant u64 sbob_sl64[8][256] = +__constant u64a sbob_sl64[8][256] = { { 0xd031c397ce553fe6, @@ -2111,7 +2111,7 @@ __constant u64 sbob_sl64[8][256] = }, }; -__constant u64 sbob_rc64[12][8] = +__constant u64a sbob_rc64[12][8] = { { 0xe9daca1eda5b08b1, diff --git a/OpenCL/m11700_a1.cl b/OpenCL/m11700_a1.cl index a8202286f..6743ca8e8 100644 --- a/OpenCL/m11700_a1.cl +++ b/OpenCL/m11700_a1.cl @@ -41,7 +41,7 @@ // constants -__constant u64 sbob_sl64[8][256] = +__constant u64a sbob_sl64[8][256] = { { 0xd031c397ce553fe6, @@ -2109,7 +2109,7 @@ __constant u64 sbob_sl64[8][256] = }, }; -__constant u64 sbob_rc64[12][8] = +__constant u64a sbob_rc64[12][8] = { { 0xe9daca1eda5b08b1, diff --git a/OpenCL/m11700_a3.cl b/OpenCL/m11700_a3.cl index 5f5f30e8a..0fa4b2514 100644 --- a/OpenCL/m11700_a3.cl +++ b/OpenCL/m11700_a3.cl @@ -41,7 +41,7 @@ // constants -__constant u64 sbob_sl64[8][256] = +__constant u64a sbob_sl64[8][256] = { { 0xd031c397ce553fe6, @@ -2109,7 +2109,7 @@ __constant u64 sbob_sl64[8][256] = }, }; -__constant u64 sbob_rc64[12][8] = +__constant u64a sbob_rc64[12][8] = { { 0xe9daca1eda5b08b1, diff --git a/OpenCL/m11800_a0.cl b/OpenCL/m11800_a0.cl index 23eb856ca..565f2d547 100644 --- a/OpenCL/m11800_a0.cl +++ b/OpenCL/m11800_a0.cl @@ -43,7 +43,7 @@ // constants -__constant u64 sbob_sl64[8][256] = +__constant u64a sbob_sl64[8][256] = { { 0xd031c397ce553fe6, @@ -2111,7 +2111,7 @@ __constant u64 sbob_sl64[8][256] = }, }; -__constant u64 sbob_rc64[12][8] = +__constant u64a sbob_rc64[12][8] = { { 0xe9daca1eda5b08b1, diff --git a/OpenCL/m11800_a1.cl b/OpenCL/m11800_a1.cl index 548a06c3f..df6afae9c 100644 --- a/OpenCL/m11800_a1.cl +++ b/OpenCL/m11800_a1.cl @@ -41,7 +41,7 @@ // constants -__constant u64 sbob_sl64[8][256] = +__constant u64a sbob_sl64[8][256] = { { 0xd031c397ce553fe6, @@ -2109,7 +2109,7 @@ __constant u64 sbob_sl64[8][256] = }, }; -__constant u64 sbob_rc64[12][8] = +__constant u64a sbob_rc64[12][8] = { { 0xe9daca1eda5b08b1, diff --git a/OpenCL/m11800_a3.cl b/OpenCL/m11800_a3.cl index 8d98b648b..5eab14e1f 100644 --- a/OpenCL/m11800_a3.cl +++ b/OpenCL/m11800_a3.cl @@ -41,7 +41,7 @@ // constants -__constant u64 sbob_sl64[8][256] = +__constant u64a sbob_sl64[8][256] = { { 0xd031c397ce553fe6, @@ -2109,7 +2109,7 @@ __constant u64 sbob_sl64[8][256] = }, }; -__constant u64 sbob_rc64[12][8] = +__constant u64a sbob_rc64[12][8] = { { 0xe9daca1eda5b08b1, diff --git a/OpenCL/m12200.cl b/OpenCL/m12200.cl index 99e49b9ef..d71de6c0f 100644 --- a/OpenCL/m12200.cl +++ b/OpenCL/m12200.cl @@ -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, diff --git a/OpenCL/m12300.cl b/OpenCL/m12300.cl index d284e9390..8e88b4216 100644 --- a/OpenCL/m12300.cl +++ b/OpenCL/m12300.cl @@ -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, diff --git a/OpenCL/m12400.cl b/OpenCL/m12400.cl index 5a7973777..525aac7fb 100644 --- a/OpenCL/m12400.cl +++ b/OpenCL/m12400.cl @@ -52,7 +52,7 @@ PERM_OP (l, r, tt, 4, 0x0f0f0f0f); \ } -__constant u32 c_SPtrans[8][64] = +__constant u32a c_SPtrans[8][64] = { { 0x00820200, 0x00020000, 0x80800000, 0x80820200, @@ -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, diff --git a/OpenCL/m12500.cl b/OpenCL/m12500.cl index 14d09a9a4..d58309337 100644 --- a/OpenCL/m12500.cl +++ b/OpenCL/m12500.cl @@ -25,7 +25,7 @@ #define MIN(a,b) (((a) < (b)) ? (a) : (b)) -__constant u32 te0[256] = +__constant u32a te0[256] = { 0xc66363a5, 0xf87c7c84, 0xee777799, 0xf67b7b8d, 0xfff2f20d, 0xd66b6bbd, 0xde6f6fb1, 0x91c5c554, @@ -93,7 +93,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, @@ -161,7 +161,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, @@ -229,7 +229,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, @@ -297,7 +297,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, @@ -365,7 +365,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, @@ -433,7 +433,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, @@ -501,7 +501,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, @@ -569,7 +569,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, @@ -637,7 +637,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, @@ -705,7 +705,7 @@ __constant u32 td4[256] = 0x55555555, 0x21212121, 0x0c0c0c0c, 0x7d7d7d7d, }; -__constant u32 rcon[] = +__constant u32a rcon[] = { 0x01000000, 0x02000000, 0x04000000, 0x08000000, 0x10000000, 0x20000000, 0x40000000, 0x80000000, diff --git a/OpenCL/m12700.cl b/OpenCL/m12700.cl index 61d38b51e..16a4e4e33 100644 --- a/OpenCL/m12700.cl +++ b/OpenCL/m12700.cl @@ -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, diff --git a/OpenCL/m12800.cl b/OpenCL/m12800.cl index 71af18f66..2b1a5f98b 100644 --- a/OpenCL/m12800.cl +++ b/OpenCL/m12800.cl @@ -17,7 +17,7 @@ #define uint_to_hex_lower8(i) l_bin2asc[(i)] -__constant u32 k_sha256[64] = +__constant u32a k_sha256[64] = { SHA256C00, SHA256C01, SHA256C02, SHA256C03, SHA256C04, SHA256C05, SHA256C06, SHA256C07, diff --git a/OpenCL/m12900.cl b/OpenCL/m12900.cl index 8a200542c..a6b8f5b2c 100644 --- a/OpenCL/m12900.cl +++ b/OpenCL/m12900.cl @@ -15,7 +15,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, diff --git a/OpenCL/m13000.cl b/OpenCL/m13000.cl index 5b53a6b64..cb03ff474 100644 --- a/OpenCL/m13000.cl +++ b/OpenCL/m13000.cl @@ -15,7 +15,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, diff --git a/OpenCL/m13200.cl b/OpenCL/m13200.cl index 586c08b65..4ff4bb058 100644 --- a/OpenCL/m13200.cl +++ b/OpenCL/m13200.cl @@ -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, diff --git a/OpenCL/m13400.cl b/OpenCL/m13400.cl index fb1700188..cd0d43c59 100644 --- a/OpenCL/m13400.cl +++ b/OpenCL/m13400.cl @@ -199,7 +199,7 @@ void AES256_encrypt (const u32 *in, u32 *out, const u32 *rek, SHM_TYPE u32 *s_te ^ rek[59]; } -__constant u32 k_sha256[64] = +__constant u32a k_sha256[64] = { SHA256C00, SHA256C01, SHA256C02, SHA256C03, SHA256C04, SHA256C05, SHA256C06, SHA256C07, @@ -516,11 +516,11 @@ __kernel void m13400_loop (__global pw_t *pws, __global const kernel_rule_t *rul #else - __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 @@ -616,17 +616,17 @@ __kernel void m13400_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 diff --git a/OpenCL/m13751.cl b/OpenCL/m13751.cl index 5e52fec2f..c007a1498 100644 --- a/OpenCL/m13751.cl +++ b/OpenCL/m13751.cl @@ -18,7 +18,7 @@ #include "inc_truecrypt_crc32.cl" #include "inc_truecrypt_xts.cl" -__constant u32 k_sha256[64] = +__constant u32a k_sha256[64] = { SHA256C00, SHA256C01, SHA256C02, SHA256C03, SHA256C04, SHA256C05, SHA256C06, SHA256C07, @@ -623,17 +623,17 @@ __kernel void m13751_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 diff --git a/OpenCL/m13752.cl b/OpenCL/m13752.cl index b0d1a87b8..5f9e7ebb0 100644 --- a/OpenCL/m13752.cl +++ b/OpenCL/m13752.cl @@ -18,7 +18,7 @@ #include "inc_truecrypt_crc32.cl" #include "inc_truecrypt_xts.cl" -__constant u32 k_sha256[64] = +__constant u32a k_sha256[64] = { SHA256C00, SHA256C01, SHA256C02, SHA256C03, SHA256C04, SHA256C05, SHA256C06, SHA256C07, @@ -623,17 +623,17 @@ __kernel void m13752_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 diff --git a/OpenCL/m13753.cl b/OpenCL/m13753.cl index f312d1dfa..6b30b81ae 100644 --- a/OpenCL/m13753.cl +++ b/OpenCL/m13753.cl @@ -18,7 +18,7 @@ #include "inc_truecrypt_crc32.cl" #include "inc_truecrypt_xts.cl" -__constant u32 k_sha256[64] = +__constant u32a k_sha256[64] = { SHA256C00, SHA256C01, SHA256C02, SHA256C03, SHA256C04, SHA256C05, SHA256C06, SHA256C07, @@ -623,17 +623,17 @@ __kernel void m13753_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 diff --git a/OpenCL/m13800_a0.cl b/OpenCL/m13800_a0.cl index 7aa322862..8a9515453 100644 --- a/OpenCL/m13800_a0.cl +++ b/OpenCL/m13800_a0.cl @@ -17,7 +17,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, diff --git a/OpenCL/m13800_a1.cl b/OpenCL/m13800_a1.cl index e33ed1daf..a618b4af9 100644 --- a/OpenCL/m13800_a1.cl +++ b/OpenCL/m13800_a1.cl @@ -15,7 +15,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, diff --git a/OpenCL/m13800_a3.cl b/OpenCL/m13800_a3.cl index 11456ec74..e0c4a0e06 100644 --- a/OpenCL/m13800_a3.cl +++ b/OpenCL/m13800_a3.cl @@ -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, diff --git a/OpenCL/m14000_a0.cl b/OpenCL/m14000_a0.cl index ef3dbb034..f6689c75f 100644 --- a/OpenCL/m14000_a0.cl +++ b/OpenCL/m14000_a0.cl @@ -54,7 +54,7 @@ PERM_OP (l, r, tt, 4, 0x0f0f0f0f); \ } -__constant u32 c_SPtrans[8][64] = +__constant u32a c_SPtrans[8][64] = { { /* nibble 0 */ @@ -210,7 +210,7 @@ __constant u32 c_SPtrans[8][64] = }, }; -__constant u32 c_skb[8][64] = +__constant u32a c_skb[8][64] = { { 0x00000000, 0x00000010, 0x20000000, 0x20000010, diff --git a/OpenCL/m14000_a1.cl b/OpenCL/m14000_a1.cl index e40882c69..3c3a15987 100644 --- a/OpenCL/m14000_a1.cl +++ b/OpenCL/m14000_a1.cl @@ -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, diff --git a/OpenCL/m14100_a0.cl b/OpenCL/m14100_a0.cl index af2b662fc..5d36a0c97 100644 --- a/OpenCL/m14100_a0.cl +++ b/OpenCL/m14100_a0.cl @@ -54,7 +54,7 @@ PERM_OP (l, r, tt, 4, 0x0f0f0f0f); \ } -__constant u32 c_SPtrans[8][64] = +__constant u32a c_SPtrans[8][64] = { { /* nibble 0 */ @@ -210,7 +210,7 @@ __constant u32 c_SPtrans[8][64] = }, }; -__constant u32 c_skb[8][64] = +__constant u32a c_skb[8][64] = { { 0x00000000, 0x00000010, 0x20000000, 0x20000010, diff --git a/OpenCL/m14100_a1.cl b/OpenCL/m14100_a1.cl index cdac30d63..7ff4bd1b1 100644 --- a/OpenCL/m14100_a1.cl +++ b/OpenCL/m14100_a1.cl @@ -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, diff --git a/OpenCL/m14100_a3.cl b/OpenCL/m14100_a3.cl index d0a9ce803..4ee44e149 100644 --- a/OpenCL/m14100_a3.cl +++ b/OpenCL/m14100_a3.cl @@ -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, diff --git a/OpenCL/m14611.cl b/OpenCL/m14611.cl index 56442970b..2eb7e8a7d 100644 --- a/OpenCL/m14611.cl +++ b/OpenCL/m14611.cl @@ -671,17 +671,17 @@ __kernel void m14611_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 diff --git a/OpenCL/m14621.cl b/OpenCL/m14621.cl index 630567b52..2069158b6 100644 --- a/OpenCL/m14621.cl +++ b/OpenCL/m14621.cl @@ -27,7 +27,7 @@ #define MAX_ENTROPY 7.0 -__constant u32 k_sha256[64] = +__constant u32a k_sha256[64] = { SHA256C00, SHA256C01, SHA256C02, SHA256C03, SHA256C04, SHA256C05, SHA256C06, SHA256C07, @@ -662,17 +662,17 @@ __kernel void m14621_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 diff --git a/OpenCL/m14622.cl b/OpenCL/m14622.cl index 250f7142c..2d44fb420 100644 --- a/OpenCL/m14622.cl +++ b/OpenCL/m14622.cl @@ -27,7 +27,7 @@ #define MAX_ENTROPY 7.0 -__constant u32 k_sha256[64] = +__constant u32a k_sha256[64] = { SHA256C00, SHA256C01, SHA256C02, SHA256C03, SHA256C04, SHA256C05, SHA256C06, SHA256C07, diff --git a/OpenCL/m14623.cl b/OpenCL/m14623.cl index 1aa93e2ff..746b8761c 100644 --- a/OpenCL/m14623.cl +++ b/OpenCL/m14623.cl @@ -27,7 +27,7 @@ #define MAX_ENTROPY 7.0 -__constant u32 k_sha256[64] = +__constant u32a k_sha256[64] = { SHA256C00, SHA256C01, SHA256C02, SHA256C03, SHA256C04, SHA256C05, SHA256C06, SHA256C07, diff --git a/OpenCL/m14631.cl b/OpenCL/m14631.cl index 812f5052c..35c9d3885 100644 --- a/OpenCL/m14631.cl +++ b/OpenCL/m14631.cl @@ -27,7 +27,7 @@ #define MAX_ENTROPY 7.0 -__constant u64 k_sha512[80] = +__constant u64a k_sha512[80] = { SHA512C00, SHA512C01, SHA512C02, SHA512C03, SHA512C04, SHA512C05, SHA512C06, SHA512C07, @@ -638,17 +638,17 @@ __kernel void m14631_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 diff --git a/OpenCL/m14632.cl b/OpenCL/m14632.cl index d8e4202f9..3ec7f015c 100644 --- a/OpenCL/m14632.cl +++ b/OpenCL/m14632.cl @@ -27,7 +27,7 @@ #define MAX_ENTROPY 7.0 -__constant u64 k_sha512[80] = +__constant u64a k_sha512[80] = { SHA512C00, SHA512C01, SHA512C02, SHA512C03, SHA512C04, SHA512C05, SHA512C06, SHA512C07, diff --git a/OpenCL/m14633.cl b/OpenCL/m14633.cl index 7f3dfe0b6..c5f5b89ae 100644 --- a/OpenCL/m14633.cl +++ b/OpenCL/m14633.cl @@ -27,7 +27,7 @@ #define MAX_ENTROPY 7.0 -__constant u64 k_sha512[80] = +__constant u64a k_sha512[80] = { SHA512C00, SHA512C01, SHA512C02, SHA512C03, SHA512C04, SHA512C05, SHA512C06, SHA512C07, diff --git a/OpenCL/m14641.cl b/OpenCL/m14641.cl index 6fa65e82f..1ea1486c9 100644 --- a/OpenCL/m14641.cl +++ b/OpenCL/m14641.cl @@ -826,17 +826,17 @@ __kernel void m14641_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 diff --git a/OpenCL/m14700.cl b/OpenCL/m14700.cl index b802c5cf8..9f285b08c 100644 --- a/OpenCL/m14700.cl +++ b/OpenCL/m14700.cl @@ -14,7 +14,7 @@ #include "inc_common.cl" #include "inc_simd.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, diff --git a/OpenCL/m14800.cl b/OpenCL/m14800.cl index f52fa0dbe..f0796daa5 100644 --- a/OpenCL/m14800.cl +++ b/OpenCL/m14800.cl @@ -12,7 +12,7 @@ #include "inc_common.cl" #include "inc_simd.cl" -__constant u32 te0[256] = +__constant u32a te0[256] = { 0xc66363a5, 0xf87c7c84, 0xee777799, 0xf67b7b8d, 0xfff2f20d, 0xd66b6bbd, 0xde6f6fb1, 0x91c5c554, @@ -80,7 +80,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, @@ -148,7 +148,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, @@ -216,7 +216,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, @@ -284,7 +284,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, @@ -352,7 +352,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, @@ -420,7 +420,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, @@ -488,7 +488,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, @@ -556,7 +556,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, @@ -624,7 +624,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, @@ -692,7 +692,7 @@ __constant u32 td4[256] = 0x55555555, 0x21212121, 0x0c0c0c0c, 0x7d7d7d7d, }; -__constant u32 rcon[] = +__constant u32a rcon[] = { 0x01000000, 0x02000000, 0x04000000, 0x08000000, 0x10000000, 0x20000000, 0x40000000, 0x80000000, @@ -969,7 +969,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, diff --git a/OpenCL/m14900_a0.cl b/OpenCL/m14900_a0.cl index 387457ba9..b365a606c 100644 --- a/OpenCL/m14900_a0.cl +++ b/OpenCL/m14900_a0.cl @@ -15,7 +15,7 @@ #include "inc_rp.cl" #include "inc_simd.cl" -__constant u8 c_ftable[256] = +__constant u8a c_ftable[256] = { 0xa3, 0xd7, 0x09, 0x83, 0xf8, 0x48, 0xf6, 0xf4, 0xb3, 0x21, 0x15, 0x78, 0x99, 0xb1, 0xaf, 0xf9, diff --git a/OpenCL/m14900_a1.cl b/OpenCL/m14900_a1.cl index f1db939d1..121e43d73 100644 --- a/OpenCL/m14900_a1.cl +++ b/OpenCL/m14900_a1.cl @@ -13,7 +13,7 @@ #include "inc_common.cl" #include "inc_simd.cl" -__constant u8 c_ftable[256] = +__constant u8a c_ftable[256] = { 0xa3, 0xd7, 0x09, 0x83, 0xf8, 0x48, 0xf6, 0xf4, 0xb3, 0x21, 0x15, 0x78, 0x99, 0xb1, 0xaf, 0xf9, diff --git a/OpenCL/m14900_a3.cl b/OpenCL/m14900_a3.cl index 0ea5858f1..6eaf22ee4 100644 --- a/OpenCL/m14900_a3.cl +++ b/OpenCL/m14900_a3.cl @@ -13,7 +13,7 @@ #include "inc_common.cl" #include "inc_simd.cl" -__constant u8 c_ftable[256] = +__constant u8a c_ftable[256] = { 0xa3, 0xd7, 0x09, 0x83, 0xf8, 0x48, 0xf6, 0xf4, 0xb3, 0x21, 0x15, 0x78, 0x99, 0xb1, 0xaf, 0xf9, diff --git a/OpenCL/m15000_a0.cl b/OpenCL/m15000_a0.cl index b3fd4c095..d1f16ce8d 100644 --- a/OpenCL/m15000_a0.cl +++ b/OpenCL/m15000_a0.cl @@ -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, diff --git a/OpenCL/m15000_a1.cl b/OpenCL/m15000_a1.cl index 89017d370..7f4b4de7b 100644 --- a/OpenCL/m15000_a1.cl +++ b/OpenCL/m15000_a1.cl @@ -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, diff --git a/OpenCL/m15000_a3.cl b/OpenCL/m15000_a3.cl index 1d87c202c..8530fd738 100644 --- a/OpenCL/m15000_a3.cl +++ b/OpenCL/m15000_a3.cl @@ -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,