arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2025-05-22 16:58:52 +00:00
Code Issues Releases Wiki Activity

Rename CONSTANT_AS to CONSTANT_VK

Browse Source
This commit is contained in:
Jens Steube 2019-05-06 14:34:16 +02:00
parent 64c495dfa5
commit d0bd33c9d1
111 changed files with 295 additions and 184 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

20
OpenCL/inc_cipher_aes.cl
View File

@ -9,7 +9,7 @@
#include "inc_common.h" #include "inc_common.h"
#include "inc_cipher_aes.h" #include "inc_cipher_aes.h"
CONSTANT_AS u32a te0[256] = CONSTANT_VK u32a te0[256] =
{ {
0xc66363a5, 0xf87c7c84, 0xee777799, 0xf67b7b8d, 0xc66363a5, 0xf87c7c84, 0xee777799, 0xf67b7b8d,
0xfff2f20d, 0xd66b6bbd, 0xde6f6fb1, 0x91c5c554, 0xfff2f20d, 0xd66b6bbd, 0xde6f6fb1, 0x91c5c554,
@ -77,7 +77,7 @@ CONSTANT_AS u32a te0[256] =
0x7bb0b0cb, 0xa85454fc, 0x6dbbbbd6, 0x2c16163a, 0x7bb0b0cb, 0xa85454fc, 0x6dbbbbd6, 0x2c16163a,
}; };
CONSTANT_AS u32a te1[256] = CONSTANT_VK u32a te1[256] =
{ {
0xa5c66363, 0x84f87c7c, 0x99ee7777, 0x8df67b7b, 0xa5c66363, 0x84f87c7c, 0x99ee7777, 0x8df67b7b,
0x0dfff2f2, 0xbdd66b6b, 0xb1de6f6f, 0x5491c5c5, 0x0dfff2f2, 0xbdd66b6b, 0xb1de6f6f, 0x5491c5c5,
@ -145,7 +145,7 @@ CONSTANT_AS u32a te1[256] =
0xcb7bb0b0, 0xfca85454, 0xd66dbbbb, 0x3a2c1616, 0xcb7bb0b0, 0xfca85454, 0xd66dbbbb, 0x3a2c1616,
}; };
CONSTANT_AS u32a te2[256] = CONSTANT_VK u32a te2[256] =
{ {
0x63a5c663, 0x7c84f87c, 0x7799ee77, 0x7b8df67b, 0x63a5c663, 0x7c84f87c, 0x7799ee77, 0x7b8df67b,
0xf20dfff2, 0x6bbdd66b, 0x6fb1de6f, 0xc55491c5, 0xf20dfff2, 0x6bbdd66b, 0x6fb1de6f, 0xc55491c5,
@ -213,7 +213,7 @@ CONSTANT_AS u32a te2[256] =
0xb0cb7bb0, 0x54fca854, 0xbbd66dbb, 0x163a2c16, 0xb0cb7bb0, 0x54fca854, 0xbbd66dbb, 0x163a2c16,
}; };
CONSTANT_AS u32a te3[256] = CONSTANT_VK u32a te3[256] =
{ {
0x6363a5c6, 0x7c7c84f8, 0x777799ee, 0x7b7b8df6, 0x6363a5c6, 0x7c7c84f8, 0x777799ee, 0x7b7b8df6,
0xf2f20dff, 0x6b6bbdd6, 0x6f6fb1de, 0xc5c55491, 0xf2f20dff, 0x6b6bbdd6, 0x6f6fb1de, 0xc5c55491,
@ -281,7 +281,7 @@ CONSTANT_AS u32a te3[256] =
0xb0b0cb7b, 0x5454fca8, 0xbbbbd66d, 0x16163a2c, 0xb0b0cb7b, 0x5454fca8, 0xbbbbd66d, 0x16163a2c,
}; };
CONSTANT_AS u32a te4[256] = CONSTANT_VK u32a te4[256] =
{ {
0x63636363, 0x7c7c7c7c, 0x77777777, 0x7b7b7b7b, 0x63636363, 0x7c7c7c7c, 0x77777777, 0x7b7b7b7b,
0xf2f2f2f2, 0x6b6b6b6b, 0x6f6f6f6f, 0xc5c5c5c5, 0xf2f2f2f2, 0x6b6b6b6b, 0x6f6f6f6f, 0xc5c5c5c5,
@ -349,7 +349,7 @@ CONSTANT_AS u32a te4[256] =
0xb0b0b0b0, 0x54545454, 0xbbbbbbbb, 0x16161616, 0xb0b0b0b0, 0x54545454, 0xbbbbbbbb, 0x16161616,
}; };
CONSTANT_AS u32a td0[256] = CONSTANT_VK u32a td0[256] =
{ {
0x51f4a750, 0x7e416553, 0x1a17a4c3, 0x3a275e96, 0x51f4a750, 0x7e416553, 0x1a17a4c3, 0x3a275e96,
0x3bab6bcb, 0x1f9d45f1, 0xacfa58ab, 0x4be30393, 0x3bab6bcb, 0x1f9d45f1, 0xacfa58ab, 0x4be30393,
@ -417,7 +417,7 @@ CONSTANT_AS u32a td0[256] =
0x7bcb8461, 0xd532b670, 0x486c5c74, 0xd0b85742, 0x7bcb8461, 0xd532b670, 0x486c5c74, 0xd0b85742,
}; };
CONSTANT_AS u32a td1[256] = CONSTANT_VK u32a td1[256] =
{ {
0x5051f4a7, 0x537e4165, 0xc31a17a4, 0x963a275e, 0x5051f4a7, 0x537e4165, 0xc31a17a4, 0x963a275e,
0xcb3bab6b, 0xf11f9d45, 0xabacfa58, 0x934be303, 0xcb3bab6b, 0xf11f9d45, 0xabacfa58, 0x934be303,
@ -485,7 +485,7 @@ CONSTANT_AS u32a td1[256] =
0x617bcb84, 0x70d532b6, 0x74486c5c, 0x42d0b857, 0x617bcb84, 0x70d532b6, 0x74486c5c, 0x42d0b857,
}; };
CONSTANT_AS u32a td2[256] = CONSTANT_VK u32a td2[256] =
{ {
0xa75051f4, 0x65537e41, 0xa4c31a17, 0x5e963a27, 0xa75051f4, 0x65537e41, 0xa4c31a17, 0x5e963a27,
0x6bcb3bab, 0x45f11f9d, 0x58abacfa, 0x03934be3, 0x6bcb3bab, 0x45f11f9d, 0x58abacfa, 0x03934be3,
@ -553,7 +553,7 @@ CONSTANT_AS u32a td2[256] =
0x84617bcb, 0xb670d532, 0x5c74486c, 0x5742d0b8, 0x84617bcb, 0xb670d532, 0x5c74486c, 0x5742d0b8,
}; };
CONSTANT_AS u32a td3[256] = CONSTANT_VK u32a td3[256] =
{ {
0xf4a75051, 0x4165537e, 0x17a4c31a, 0x275e963a, 0xf4a75051, 0x4165537e, 0x17a4c31a, 0x275e963a,
0xab6bcb3b, 0x9d45f11f, 0xfa58abac, 0xe303934b, 0xab6bcb3b, 0x9d45f11f, 0xfa58abac, 0xe303934b,
@ -621,7 +621,7 @@ CONSTANT_AS u32a td3[256] =
0xcb84617b, 0x32b670d5, 0x6c5c7448, 0xb85742d0, 0xcb84617b, 0x32b670d5, 0x6c5c7448, 0xb85742d0,
}; };
CONSTANT_AS u32a td4[256] = CONSTANT_VK u32a td4[256] =
{ {
0x52525252, 0x09090909, 0x6a6a6a6a, 0xd5d5d5d5, 0x52525252, 0x09090909, 0x6a6a6a6a, 0xd5d5d5d5,
0x30303030, 0x36363636, 0xa5a5a5a5, 0x38383838, 0x30303030, 0x36363636, 0xa5a5a5a5, 0x38383838,

2
OpenCL/inc_cipher_camellia.cl
View File

@ -21,7 +21,7 @@
#include "inc_common.h" #include "inc_common.h"
#include "inc_cipher_camellia.h" #include "inc_cipher_camellia.h"
CONSTANT_AS u32a c_sbox[256] = CONSTANT_VK u32a c_sbox[256] =
{ {
0x70, 0x82, 0x2c, 0xec, 0xb3, 0x27, 0xc0, 0xe5, 0x70, 0x82, 0x2c, 0xec, 0xb3, 0x27, 0xc0, 0xe5,
0xe4, 0x85, 0x57, 0x35, 0xea, 0x0c, 0xae, 0x41, 0xe4, 0x85, 0x57, 0x35, 0xea, 0x0c, 0xae, 0x41,

4
OpenCL/inc_cipher_des.cl
View File

@ -9,7 +9,7 @@
#include "inc_common.h" #include "inc_common.h"
#include "inc_cipher_des.h" #include "inc_cipher_des.h"
CONSTANT_AS u32a c_SPtrans[8][64] = CONSTANT_VK u32a c_SPtrans[8][64] =
{ {
{ {
/* nibble 0 */ /* nibble 0 */
@ -165,7 +165,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
}, },
}; };
CONSTANT_AS u32a c_skb[8][64] = CONSTANT_VK u32a c_skb[8][64] =
{ {
{ {
0x00000000, 0x00000010, 0x20000000, 0x20000010, 0x00000000, 0x00000010, 0x20000000, 0x20000010,

4
OpenCL/inc_cipher_kuznyechik.cl
View File

@ -18,7 +18,7 @@
#include "inc_common.h" #include "inc_common.h"
#include "inc_cipher_kuznyechik.h" #include "inc_cipher_kuznyechik.h"
CONSTANT_AS u32a k_sbox[256] = CONSTANT_VK u32a k_sbox[256] =
{ {
0xfc, 0xee, 0xdd, 0x11, 0xcf, 0x6e, 0x31, 0x16, 0xfc, 0xee, 0xdd, 0x11, 0xcf, 0x6e, 0x31, 0x16,
0xfb, 0xc4, 0xfa, 0xda, 0x23, 0xc5, 0x04, 0x4d, 0xfb, 0xc4, 0xfa, 0xda, 0x23, 0xc5, 0x04, 0x4d,
@ -54,7 +54,7 @@ CONSTANT_AS u32a k_sbox[256] =
0xd1, 0x66, 0xaf, 0xc2, 0x39, 0x4b, 0x63, 0xb6 0xd1, 0x66, 0xaf, 0xc2, 0x39, 0x4b, 0x63, 0xb6
}; };
CONSTANT_AS u32a k_sbox_inv[256] = CONSTANT_VK u32a k_sbox_inv[256] =
{ {
0xa5, 0x2d, 0x32, 0x8f, 0x0e, 0x30, 0x38, 0xc0, 0xa5, 0x2d, 0x32, 0x8f, 0x0e, 0x30, 0x38, 0xc0,
0x54, 0xe6, 0x9e, 0x39, 0x55, 0x7e, 0x52, 0x91, 0x54, 0xe6, 0x9e, 0x39, 0x55, 0x7e, 0x52, 0x91,

4
OpenCL/inc_cipher_twofish.cl
View File

@ -25,7 +25,7 @@
#include "inc_common.h" #include "inc_common.h"
#include "inc_cipher_twofish.h" #include "inc_cipher_twofish.h"
CONSTANT_AS u32a q_tab[2][256] = CONSTANT_VK u32a q_tab[2][256] =
{ {
{ {
0xA9, 0x67, 0xB3, 0xE8, 0x04, 0xFD, 0xA3, 0x76, 0x9A, 0x92, 0x80, 0x78, 0xA9, 0x67, 0xB3, 0xE8, 0x04, 0xFD, 0xA3, 0x76, 0x9A, 0x92, 0x80, 0x78,
@ -77,7 +77,7 @@ CONSTANT_AS u32a q_tab[2][256] =
} }
}; };
CONSTANT_AS u32a m_tab[4][256] = CONSTANT_VK u32a m_tab[4][256] =
{ {
{ 0xBCBC3275, 0xECEC21F3, 0x202043C6, 0xB3B3C9F4, 0xDADA03DB, 0x02028B7B, { 0xBCBC3275, 0xECEC21F3, 0x202043C6, 0xB3B3C9F4, 0xDADA03DB, 0x02028B7B,
0xE2E22BFB, 0x9E9EFAC8, 0xC9C9EC4A, 0xD4D409D3, 0x18186BE6, 0x1E1E9F6B, 0xE2E22BFB, 0x9E9EFAC8, 0xC9C9EC4A, 0xD4D409D3, 0x18186BE6, 0x1E1E9F6B,

32
OpenCL/inc_common.cl
View File

@ -305,7 +305,9 @@ DECLSPEC u64x hl32_to_64 (const u32x a, const u32x b)
DECLSPEC u32x hc_rotl32 (const u32x a, const int n) DECLSPEC u32x hc_rotl32 (const u32x a, const int n)
{ {
#ifdef _CPU_OPENCL_EMU_H #if defined _CPU_OPENCL_EMU_H
return rotl32 (a, n);
#elif defined IS_CUDA
return rotl32 (a, n); return rotl32 (a, n);
#else #else
return rotate (a, (u32x) (n)); return rotate (a, (u32x) (n));
@ -314,7 +316,9 @@ DECLSPEC u32x hc_rotl32 (const u32x a, const int n)
DECLSPEC u32x hc_rotr32 (const u32x a, const int n) DECLSPEC u32x hc_rotr32 (const u32x a, const int n)
{ {
#ifdef _CPU_OPENCL_EMU_H #if defined _CPU_OPENCL_EMU_H
return rotr32 (a, n);
#elif defined IS_CUDA
return rotr32 (a, n); return rotr32 (a, n);
#else #else
return rotate (a, (u32x) (32 - n)); return rotate (a, (u32x) (32 - n));
@ -323,8 +327,10 @@ DECLSPEC u32x hc_rotr32 (const u32x a, const int n)
DECLSPEC u32 hc_rotl32_S (const u32 a, const int n) DECLSPEC u32 hc_rotl32_S (const u32 a, const int n)
{ {
#ifdef _CPU_OPENCL_EMU_H #if defined _CPU_OPENCL_EMU_H
return rotl32 (a, n); return rotl32 (a, n);
#elif defined IS_CUDA
return rotl32_S (a, n);
#else #else
return rotate (a, (u32) (n)); return rotate (a, (u32) (n));
#endif #endif
@ -332,8 +338,10 @@ DECLSPEC u32 hc_rotl32_S (const u32 a, const int n)
DECLSPEC u32 hc_rotr32_S (const u32 a, const int n) DECLSPEC u32 hc_rotr32_S (const u32 a, const int n)
{ {
#ifdef _CPU_OPENCL_EMU_H #if defined _CPU_OPENCL_EMU_H
return rotr32 (a, n); return rotr32 (a, n);
#elif defined IS_CUDA
return rotr32_S (a, n);
#else #else
return rotate (a, (u32) (32 - n)); return rotate (a, (u32) (32 - n));
#endif #endif
@ -341,7 +349,9 @@ DECLSPEC u32 hc_rotr32_S (const u32 a, const int n)
DECLSPEC u64x hc_rotl64 (const u64x a, const int n) DECLSPEC u64x hc_rotl64 (const u64x a, const int n)
{ {
#ifdef _CPU_OPENCL_EMU_H #if defined _CPU_OPENCL_EMU_H
return rotl64 (a, n);
#elif defined IS_CUDA
return rotl64 (a, n); return rotl64 (a, n);
#else #else
return rotate (a, (u64x) (n)); return rotate (a, (u64x) (n));
@ -350,7 +360,9 @@ DECLSPEC u64x hc_rotl64 (const u64x a, const int n)
DECLSPEC u64x hc_rotr64 (const u64x a, const int n) DECLSPEC u64x hc_rotr64 (const u64x a, const int n)
{ {
#ifdef _CPU_OPENCL_EMU_H #if defined _CPU_OPENCL_EMU_H
return rotr64 (a, n);
#elif defined IS_CUDA
return rotr64 (a, n); return rotr64 (a, n);
#else #else
return rotate (a, (u64x) (64 - n)); return rotate (a, (u64x) (64 - n));
@ -359,8 +371,10 @@ DECLSPEC u64x hc_rotr64 (const u64x a, const int n)
DECLSPEC u64 hc_rotl64_S (const u64 a, const int n) DECLSPEC u64 hc_rotl64_S (const u64 a, const int n)
{ {
#ifdef _CPU_OPENCL_EMU_H #if defined _CPU_OPENCL_EMU_H
return rotl64 (a, n); return rotl64 (a, n);
#elif defined IS_CUDA
return rotl64_S (a, n);
#else #else
return rotate (a, (u64) (n)); return rotate (a, (u64) (n));
#endif #endif
@ -368,8 +382,10 @@ DECLSPEC u64 hc_rotl64_S (const u64 a, const int n)
DECLSPEC u64 hc_rotr64_S (const u64 a, const int n) DECLSPEC u64 hc_rotr64_S (const u64 a, const int n)
{ {
#ifdef _CPU_OPENCL_EMU_H #if defined _CPU_OPENCL_EMU_H
return rotr64 (a, n); return rotr64 (a, n);
#elif defined IS_CUDA
return rotr64_S (a, n);
#else #else
return rotate (a, (u64) (64 - n)); return rotate (a, (u64) (64 - n));
#endif #endif

2
OpenCL/inc_hash_sha224.cl
View File

@ -9,7 +9,7 @@
#include "inc_common.h" #include "inc_common.h"
#include "inc_hash_sha224.h" #include "inc_hash_sha224.h"
CONSTANT_AS u32a k_sha224[64] = CONSTANT_VK u32a k_sha224[64] =
{ {
SHA224C00, SHA224C01, SHA224C02, SHA224C03, SHA224C00, SHA224C01, SHA224C02, SHA224C03,
SHA224C04, SHA224C05, SHA224C06, SHA224C07, SHA224C04, SHA224C05, SHA224C06, SHA224C07,

2
OpenCL/inc_hash_sha256.cl
View File

@ -9,7 +9,7 @@
#include "inc_common.h" #include "inc_common.h"
#include "inc_hash_sha256.h" #include "inc_hash_sha256.h"
CONSTANT_AS u32a k_sha256[64] = CONSTANT_VK u32a k_sha256[64] =
{ {
SHA256C00, SHA256C01, SHA256C02, SHA256C03, SHA256C00, SHA256C01, SHA256C02, SHA256C03,
SHA256C04, SHA256C05, SHA256C06, SHA256C07, SHA256C04, SHA256C05, SHA256C06, SHA256C07,

2
OpenCL/inc_hash_sha384.cl
View File

@ -9,7 +9,7 @@
#include "inc_common.h" #include "inc_common.h"
#include "inc_hash_sha384.h" #include "inc_hash_sha384.h"
CONSTANT_AS u64a k_sha384[80] = CONSTANT_VK u64a k_sha384[80] =
{ {
SHA512C00, SHA512C01, SHA512C02, SHA512C03, SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07, SHA512C04, SHA512C05, SHA512C06, SHA512C07,

2
OpenCL/inc_hash_sha512.cl
View File

@ -9,7 +9,7 @@
#include "inc_common.h" #include "inc_common.h"
#include "inc_hash_sha512.h" #include "inc_hash_sha512.h"
CONSTANT_AS u64a k_sha512[80] = CONSTANT_VK u64a k_sha512[80] =
{ {
SHA512C00, SHA512C01, SHA512C02, SHA512C03, SHA512C00, SHA512C01, SHA512C02, SHA512C03,
SHA512C04, SHA512C05, SHA512C06, SHA512C07, SHA512C04, SHA512C05, SHA512C06, SHA512C07,

4
OpenCL/inc_hash_streebog256.cl
View File

@ -9,7 +9,7 @@
#include "inc_common.h" #include "inc_common.h"
#include "inc_hash_streebog256.h" #include "inc_hash_streebog256.h"
CONSTANT_AS u64a sbob256_sl64[8][256] = CONSTANT_VK u64a sbob256_sl64[8][256] =
{ {
{ {
0xd031c397ce553fe6, 0x16ba5b01b006b525, 0xa89bade6296e70c8, 0x6a1f525d77d3435b, 0xd031c397ce553fe6, 0x16ba5b01b006b525, 0xa89bade6296e70c8, 0x6a1f525d77d3435b,
@ -541,7 +541,7 @@ CONSTANT_AS u64a sbob256_sl64[8][256] =
}, },
}; };
CONSTANT_AS u64a sbob256_rc64[12][8] = CONSTANT_VK u64a sbob256_rc64[12][8] =
{ {
{ {
0xe9daca1eda5b08b1, 0x1f7c65c0812fcbeb, 0x16d0452e43766a2f, 0xfcc485758db84e71, 0xe9daca1eda5b08b1, 0x1f7c65c0812fcbeb, 0x16d0452e43766a2f, 0xfcc485758db84e71,

4
OpenCL/inc_hash_streebog512.cl
View File

@ -9,7 +9,7 @@
#include "inc_common.h" #include "inc_common.h"
#include "inc_hash_streebog512.h" #include "inc_hash_streebog512.h"
CONSTANT_AS u64a sbob512_sl64[8][256] = CONSTANT_VK u64a sbob512_sl64[8][256] =
{ {
{ {
0xd031c397ce553fe6, 0x16ba5b01b006b525, 0xa89bade6296e70c8, 0x6a1f525d77d3435b, 0xd031c397ce553fe6, 0x16ba5b01b006b525, 0xa89bade6296e70c8, 0x6a1f525d77d3435b,
@ -541,7 +541,7 @@ CONSTANT_AS u64a sbob512_sl64[8][256] =
}, },
}; };
CONSTANT_AS u64a sbob512_rc64[12][8] = CONSTANT_VK u64a sbob512_rc64[12][8] =
{ {
{ {
0xe9daca1eda5b08b1, 0x1f7c65c0812fcbeb, 0x16d0452e43766a2f, 0xfcc485758db84e71, 0xe9daca1eda5b08b1, 0x1f7c65c0812fcbeb, 0x16d0452e43766a2f, 0xfcc485758db84e71,

8
OpenCL/inc_hash_whirlpool.cl
View File

@ -9,7 +9,7 @@
#include "inc_common.h" #include "inc_common.h"
#include "inc_hash_whirlpool.h" #include "inc_hash_whirlpool.h"
CONSTANT_AS u32a Ch[8][256] = CONSTANT_VK u32a Ch[8][256] =
{ {
{ {
0x18186018, 0x23238c23, 0xc6c63fc6, 0xe8e887e8, 0x18186018, 0x23238c23, 0xc6c63fc6, 0xe8e887e8,
@ -541,7 +541,7 @@ CONSTANT_AS u32a Ch[8][256] =
} }
}; };
CONSTANT_AS u32a Cl[8][256] = CONSTANT_VK u32a Cl[8][256] =
{ {
{ {
0xc07830d8, 0x05af4626, 0x7ef991b8, 0x136fcdfb, 0xc07830d8, 0x05af4626, 0x7ef991b8, 0x136fcdfb,
@ -1073,7 +1073,7 @@ CONSTANT_AS u32a Cl[8][256] =
}, },
}; };
CONSTANT_AS u32a rch[R + 1] = CONSTANT_VK u32a rch[R + 1] =
{ {
0x00000000, 0x00000000,
0x1823c6e8, 0x1823c6e8,
@ -1088,7 +1088,7 @@ CONSTANT_AS u32a rch[R + 1] =
0xca2dbf07, 0xca2dbf07,
}; };
CONSTANT_AS u32a rcl[R + 1] = CONSTANT_VK u32a rcl[R + 1] =
{ {
0x00000000, 0x00000000,
0x87b8014f, 0x87b8014f,

40
OpenCL/inc_platform.cl
View File

@ -44,6 +44,46 @@ DECLSPEC size_t get_local_size (const u32 dimindx __attribute__((unused)))
return blockDim.x; return blockDim.x;
} }
DECLSPEC u32x rotl32 (const u32x a, const int n)
{
return ((a << n) | ((a >> (32 - n))));
}
DECLSPEC u32x rotr32 (const u32x a, const int n)
{
return ((a >> n) | ((a << (32 - n))));
}
DECLSPEC u32 rotl32_S (const u32 a, const int n)
{
return ((a << n) | ((a >> (32 - n))));
}
DECLSPEC u32 rotr32_S (const u32 a, const int n)
{
return ((a >> n) | ((a << (32 - n))));
}
DECLSPEC u64x rotl64 (const u64x a, const int n)
{
return ((a << n) | ((a >> (64 - n))));
}
DECLSPEC u64x rotr64 (const u64x a, const int n)
{
return ((a >> n) | ((a << (64 - n))));
}
DECLSPEC u64 rotl64_S (const u64 a, const int n)
{
return ((a << n) | ((a >> (64 - n))));
}
DECLSPEC u64 rotr64_S (const u64 a, const int n)
{
return ((a >> n) | ((a << (64 - n))));
}
#define SYNC_THREADS() __syncthreads () #define SYNC_THREADS() __syncthreads ()
#endif #endif

11
OpenCL/inc_platform.h
View File

@ -14,7 +14,16 @@ DECLSPEC size_t get_global_id (const u32 dimindx __attribute__((unused)));
DECLSPEC size_t get_local_id (const u32 dimindx __attribute__((unused))); DECLSPEC size_t get_local_id (const u32 dimindx __attribute__((unused)));
DECLSPEC size_t get_local_size (const u32 dimindx __attribute__((unused))); DECLSPEC size_t get_local_size (const u32 dimindx __attribute__((unused)));
#define rotate(a,n) (((a) << (n)) | ((a) >> (32 - (n)))) DECLSPEC u32x hc_rotl32 (const u32x a, const int n);
DECLSPEC u32x hc_rotr32 (const u32x a, const int n);
DECLSPEC u32 hc_rotl32_S (const u32 a, const int n);
DECLSPEC u32 hc_rotr32_S (const u32 a, const int n);
DECLSPEC u64x hc_rotl64 (const u64x a, const int n);
DECLSPEC u64x hc_rotr64 (const u64x a, const int n);
DECLSPEC u64 hc_rotl64_S (const u64 a, const int n);
DECLSPEC u64 hc_rotr64_S (const u64 a, const int n);
//#define rotate(a,n) (((a) << (n)) | ((a) >> (32 - (n))))
#define bitselect(a,b,c) ((a) ^ ((c) & ((b) ^ (a)))) #define bitselect(a,b,c) ((a) ^ ((c) & ((b) ^ (a))))
#endif #endif

2
OpenCL/inc_truecrypt_crc32.cl
View File

@ -9,7 +9,7 @@
#include "inc_common.h" #include "inc_common.h"
#include "inc_truecrypt_crc32.h" #include "inc_truecrypt_crc32.h"
CONSTANT_AS u32a crc32tab[0x100] = CONSTANT_VK u32a crc32tab[0x100] =
{ {
0x00000000, 0x77073096, 0xee0e612c, 0x990951ba, 0x00000000, 0x77073096, 0xee0e612c, 0x990951ba,
0x076dc419, 0x706af48f, 0xe963a535, 0x9e6495a3, 0x076dc419, 0x706af48f, 0xe963a535, 0x9e6495a3,

45
OpenCL/inc_types.h
View File

@ -368,7 +368,6 @@ inline __device__ u64x operator ~ (const u64x a) { return u64x (~a.s0, ~a.s1, ~
#if VECT_SIZE == 8 #if VECT_SIZE == 8
class u8x class u8x
{ {
private: private:
@ -497,6 +496,50 @@ inline __device__ u32x operator * (const u32x a, const u32x b) { return u32x ((
inline __device__ u32x operator ~ (const u32x a) { return u32x (~a.s0, ~a.s1, ~a.s2, ~a.s3, ~a.s4, ~a.s5, ~a.s6, ~a.s7); } inline __device__ u32x operator ~ (const u32x a) { return u32x (~a.s0, ~a.s1, ~a.s2, ~a.s3, ~a.s4, ~a.s5, ~a.s6, ~a.s7); }
inline __device__ bool operator != (const u64x a, const u64 b) { return ((a.s0 != b) && (a.s1 != b) && (a.s2 != b) && (a.s3 != b) && (a.s4 != b) && (a.s5 != b) && (a.s6 != b) && (a.s7 != b) ); }
inline __device__ bool operator != (const u64x a, const u64x b) { return ((a.s0 != b.s0) && (a.s1 != b.s1) && (a.s2 != b.s2) && (a.s3 != b.s3) && (a.s4 != b.s4) && (a.s5 != b.s5) && (a.s6 != b.s6) && (a.s7 != b.s7)); }
inline __device__ void operator ^= (u64x &a, const u64 b) { a.s0 ^= b; a.s1 ^= b; a.s2 ^= b; a.s3 ^= b; a.s4 ^= b; a.s5 ^= b; a.s6 ^= b; a.s7 ^= b; }
inline __device__ void operator ^= (u64x &a, const u64x b) { a.s0 ^= b.s0; a.s1 ^= b.s1; a.s2 ^= b.s2; a.s3 ^= b.s3; a.s4 ^= b.s4; a.s5 ^= b.s5; a.s6 ^= b.s6; a.s7 ^= b.s7; }
inline __device__ void operator |= (u64x &a, const u64 b) { a.s0 |= b; a.s1 |= b; a.s2 |= b; a.s3 |= b; a.s4 |= b; a.s5 |= b; a.s6 |= b; a.s7 |= b; }
inline __device__ void operator |= (u64x &a, const u64x b) { a.s0 |= b.s0; a.s1 |= b.s1; a.s2 |= b.s2; a.s3 |= b.s3; a.s4 |= b.s4; a.s5 |= b.s5; a.s6 |= b.s6; a.s7 |= b.s7; }
inline __device__ void operator &= (u64x &a, const u64 b) { a.s0 &= b; a.s1 &= b; a.s2 &= b; a.s3 &= b; a.s4 &= b; a.s5 &= b; a.s6 &= b; a.s7 &= b; }
inline __device__ void operator &= (u64x &a, const u64x b) { a.s0 &= b.s0; a.s1 &= b.s1; a.s2 &= b.s2; a.s3 &= b.s3; a.s4 &= b.s4; a.s5 &= b.s5; a.s6 &= b.s6; a.s7 &= b.s7; }
inline __device__ void operator += (u64x &a, const u64 b) { a.s0 += b; a.s1 += b; a.s2 += b; a.s3 += b; a.s4 += b; a.s5 += b; a.s6 += b; a.s7 += b; }
inline __device__ void operator += (u64x &a, const u64x b) { a.s0 += b.s0; a.s1 += b.s1; a.s2 += b.s2; a.s3 += b.s3; a.s4 += b.s4; a.s5 += b.s5; a.s6 += b.s6; a.s7 += b.s7; }
inline __device__ void operator -= (u64x &a, const u64 b) { a.s0 -= b; a.s1 -= b; a.s2 -= b; a.s3 -= b; a.s4 -= b; a.s5 -= b; a.s6 -= b; a.s7 -= b; }
inline __device__ void operator -= (u64x &a, const u64x b) { a.s0 -= b.s0; a.s1 -= b.s1; a.s2 -= b.s2; a.s3 -= b.s3; a.s4 -= b.s4; a.s5 -= b.s5; a.s6 -= b.s6; a.s7 -= b.s7; }
inline __device__ u64x operator << (const u64x a, const u64 b) { return u64x ((a.s0 << b), (a.s1 << b) , (a.s2 << b), (a.s3 << b) , (a.s4 << b), (a.s5 << b) , (a.s6 << b), (a.s7 << b) ); }
inline __device__ u64x operator << (const u64x a, const u64x b) { return u64x ((a.s0 << b.s0), (a.s1 << b.s1), (a.s2 << b.s2), (a.s3 << b.s3), (a.s4 << b.s4), (a.s5 << b.s5), (a.s6 << b.s6), (a.s7 << b.s7)); }
inline __device__ u64x operator >> (const u64x a, const u64 b) { return u64x ((a.s0 >> b), (a.s1 >> b) , (a.s2 >> b), (a.s3 >> b) , (a.s4 >> b), (a.s5 >> b) , (a.s6 >> b), (a.s7 >> b) ); }
inline __device__ u64x operator >> (const u64x a, const u64x b) { return u64x ((a.s0 >> b.s0), (a.s1 >> b.s1), (a.s2 >> b.s2), (a.s3 >> b.s3), (a.s4 >> b.s4), (a.s5 >> b.s5), (a.s6 >> b.s6), (a.s7 >> b.s7)); }
inline __device__ u64x operator ^ (const u64x a, const u64 b) { return u64x ((a.s0 ^ b), (a.s1 ^ b) , (a.s2 ^ b), (a.s3 ^ b) , (a.s4 ^ b), (a.s5 ^ b) , (a.s6 ^ b), (a.s7 ^ b) ); }
inline __device__ u64x operator ^ (const u64x a, const u64x b) { return u64x ((a.s0 ^ b.s0), (a.s1 ^ b.s1), (a.s2 ^ b.s2), (a.s3 ^ b.s3), (a.s4 ^ b.s4), (a.s5 ^ b.s5), (a.s6 ^ b.s6), (a.s7 ^ b.s7)); }
inline __device__ u64x operator | (const u64x a, const u64 b) { return u64x ((a.s0 | b), (a.s1 | b) , (a.s2 | b), (a.s3 | b) , (a.s4 | b), (a.s5 | b) , (a.s6 | b), (a.s7 | b) ); }
inline __device__ u64x operator | (const u64x a, const u64x b) { return u64x ((a.s0 | b.s0), (a.s1 | b.s1), (a.s2 | b.s2), (a.s3 | b.s3), (a.s4 | b.s4), (a.s5 | b.s5), (a.s6 | b.s6), (a.s7 | b.s7)); }
inline __device__ u64x operator & (const u64x a, const u64 b) { return u64x ((a.s0 & b), (a.s1 & b) , (a.s2 & b), (a.s3 & b) , (a.s4 & b), (a.s5 & b) , (a.s6 & b), (a.s7 & b) ); }
inline __device__ u64x operator & (const u64x a, const u64x b) { return u64x ((a.s0 & b.s0), (a.s1 & b.s1), (a.s2 & b.s2), (a.s3 & b.s3), (a.s4 & b.s4), (a.s5 & b.s5), (a.s6 & b.s6), (a.s7 & b.s7)); }
inline __device__ u64x operator + (const u64x a, const u64 b) { return u64x ((a.s0 + b), (a.s1 + b) , (a.s2 + b), (a.s3 + b) , (a.s4 + b), (a.s5 + b) , (a.s6 + b), (a.s7 + b) ); }
inline __device__ u64x operator + (const u64x a, const u64x b) { return u64x ((a.s0 + b.s0), (a.s1 + b.s1), (a.s2 + b.s2), (a.s3 + b.s3), (a.s4 + b.s4), (a.s5 + b.s5), (a.s6 + b.s6), (a.s7 + b.s7)); }
inline __device__ u64x operator - (const u64x a, const u64 b) { return u64x ((a.s0 - b), (a.s1 - b) , (a.s2 - b), (a.s3 - b) , (a.s4 - b), (a.s5 - b) , (a.s6 - b), (a.s7 - b) ); }
inline __device__ u64x operator - (const u64x a, const u64x b) { return u64x ((a.s0 - b.s0), (a.s1 - b.s1), (a.s2 - b.s2), (a.s3 - b.s3), (a.s4 - b.s4), (a.s5 - b.s5), (a.s6 - b.s6), (a.s7 - b.s7)); }
inline __device__ u64x operator * (const u64x a, const u64 b) { return u64x ((a.s0 * b), (a.s1 * b) , (a.s2 * b), (a.s3 * b) , (a.s4 * b), (a.s5 * b) , (a.s6 * b), (a.s7 * b) ); }
inline __device__ u64x operator * (const u64x a, const u64x b) { return u64x ((a.s0 * b.s0), (a.s1 * b.s1), (a.s2 * b.s2), (a.s3 * b.s3), (a.s4 * b.s4), (a.s5 * b.s5), (a.s6 * b.s6), (a.s7 * b.s7)); }
inline __device__ u64x operator ~ (const u64x a) { return u64x (~a.s0, ~a.s1, ~a.s2, ~a.s3, ~a.s4, ~a.s5, ~a.s6, ~a.s7); }
#endif #endif
#if VECT_SIZE == 16 #if VECT_SIZE == 16

3
OpenCL/inc_vendor.h
View File

@ -15,16 +15,19 @@
#endif #endif
#if defined IS_NATIVE #if defined IS_NATIVE
#define CONSTANT_VK
#define CONSTANT_AS #define CONSTANT_AS
#define GLOBAL_AS #define GLOBAL_AS
#define LOCAL_AS #define LOCAL_AS
#define KERNEL_FQ #define KERNEL_FQ
#elif defined IS_CUDA #elif defined IS_CUDA
#define CONSTANT_VK __constant__
#define CONSTANT_AS #define CONSTANT_AS
#define GLOBAL_AS #define GLOBAL_AS
#define LOCAL_AS #define LOCAL_AS
#define KERNEL_FQ extern "C" __global__ #define KERNEL_FQ extern "C" __global__
#elif defined IS_OPENCL #elif defined IS_OPENCL
#define CONSTANT_VK __constant
#define CONSTANT_AS __constant #define CONSTANT_AS __constant
#define GLOBAL_AS __global #define GLOBAL_AS __global
#define LOCAL_AS __local #define LOCAL_AS __local

4
OpenCL/m01500_a0-pure.cl
View File

@ -35,7 +35,7 @@
a = a ^ tt; \ a = a ^ tt; \
} }
CONSTANT_AS u32a c_SPtrans[8][64] = CONSTANT_VK u32a c_SPtrans[8][64] =
{ {
{ {
0x00820200, 0x00020000, 0x80800000, 0x80820200, 0x00820200, 0x00020000, 0x80800000, 0x80820200,
@ -183,7 +183,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
}, },
}; };
CONSTANT_AS u32a c_skb[8][64] = CONSTANT_VK u32a c_skb[8][64] =
{ {
{ {
0x00000000, 0x00000010, 0x20000000, 0x20000010, 0x00000000, 0x00000010, 0x20000000, 0x20000010,

4
OpenCL/m01500_a1-pure.cl
View File

@ -33,7 +33,7 @@
a = a ^ tt; \ a = a ^ tt; \
} }
CONSTANT_AS u32a c_SPtrans[8][64] = CONSTANT_VK u32a c_SPtrans[8][64] =
{ {
{ {
0x00820200, 0x00020000, 0x80800000, 0x80820200, 0x00820200, 0x00020000, 0x80800000, 0x80820200,
@ -181,7 +181,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
}, },
}; };
CONSTANT_AS u32a c_skb[8][64] = CONSTANT_VK u32a c_skb[8][64] =
{ {
{ {
0x00000000, 0x00000010, 0x20000000, 0x20000010, 0x00000000, 0x00000010, 0x20000000, 0x20000010,

4
OpenCL/m03000_a0-pure.cl
View File

@ -35,7 +35,7 @@
a = a ^ tt; \ a = a ^ tt; \
} }
CONSTANT_AS u32a c_SPtrans[8][64] = CONSTANT_VK u32a c_SPtrans[8][64] =
{ {
{ {
0x02080800, 0x00080000, 0x02000002, 0x02080802, 0x02080800, 0x00080000, 0x02000002, 0x02080802,
@ -183,7 +183,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
} }
}; };
CONSTANT_AS u32a c_skb[8][64] = CONSTANT_VK u32a c_skb[8][64] =
{ {
{ {
0x00000000, 0x00000010, 0x20000000, 0x20000010, 0x00000000, 0x00000010, 0x20000000, 0x20000010,

4
OpenCL/m03000_a1-pure.cl
View File

@ -33,7 +33,7 @@
a = a ^ tt; \ a = a ^ tt; \
} }
CONSTANT_AS u32a c_SPtrans[8][64] = CONSTANT_VK u32a c_SPtrans[8][64] =
{ {
{ {
0x02080800, 0x00080000, 0x02000002, 0x02080802, 0x02080800, 0x00080000, 0x02000002, 0x02080802,
@ -181,7 +181,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
} }
}; };
CONSTANT_AS u32a c_skb[8][64] = CONSTANT_VK u32a c_skb[8][64] =
{ {
{ {
0x00000000, 0x00000010, 0x20000000, 0x20000010, 0x00000000, 0x00000010, 0x20000000, 0x20000010,

8
OpenCL/m03200-pure.cl
View File

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

4
OpenCL/m05500_a0-optimized.cl
View File

@ -48,7 +48,7 @@ typedef struct netntlm
a = a ^ tt; \ a = a ^ tt; \
} }
CONSTANT_AS u32a c_SPtrans[8][64] = CONSTANT_VK u32a c_SPtrans[8][64] =
{ {
{ {
0x02080800, 0x00080000, 0x02000002, 0x02080802, 0x02080800, 0x00080000, 0x02000002, 0x02080802,
@ -196,7 +196,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
} }
}; };
CONSTANT_AS u32a c_skb[8][64] = CONSTANT_VK u32a c_skb[8][64] =
{ {
{ {
0x00000000, 0x00000010, 0x20000000, 0x20000010, 0x00000000, 0x00000010, 0x20000000, 0x20000010,

4
OpenCL/m05500_a0-pure.cl
View File

@ -48,7 +48,7 @@ typedef struct netntlm
a = a ^ tt; \ a = a ^ tt; \
} }
CONSTANT_AS u32a c_SPtrans[8][64] = CONSTANT_VK u32a c_SPtrans[8][64] =
{ {
{ {
0x02080800, 0x00080000, 0x02000002, 0x02080802, 0x02080800, 0x00080000, 0x02000002, 0x02080802,
@ -196,7 +196,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
} }
}; };
CONSTANT_AS u32a c_skb[8][64] = CONSTANT_VK u32a c_skb[8][64] =
{ {
{ {
0x00000000, 0x00000010, 0x20000000, 0x20000010, 0x00000000, 0x00000010, 0x20000000, 0x20000010,

4
OpenCL/m05500_a1-optimized.cl
View File

@ -46,7 +46,7 @@ typedef struct netntlm
a = a ^ tt; \ a = a ^ tt; \
} }
CONSTANT_AS u32a c_SPtrans[8][64] = CONSTANT_VK u32a c_SPtrans[8][64] =
{ {
{ {
0x02080800, 0x00080000, 0x02000002, 0x02080802, 0x02080800, 0x00080000, 0x02000002, 0x02080802,
@ -194,7 +194,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
} }
}; };
CONSTANT_AS u32a c_skb[8][64] = CONSTANT_VK u32a c_skb[8][64] =
{ {
{ {
0x00000000, 0x00000010, 0x20000000, 0x20000010, 0x00000000, 0x00000010, 0x20000000, 0x20000010,

4
OpenCL/m05500_a1-pure.cl
View File

@ -46,7 +46,7 @@ typedef struct netntlm
a = a ^ tt; \ a = a ^ tt; \
} }
CONSTANT_AS u32a c_SPtrans[8][64] = CONSTANT_VK u32a c_SPtrans[8][64] =
{ {
{ {
0x02080800, 0x00080000, 0x02000002, 0x02080802, 0x02080800, 0x00080000, 0x02000002, 0x02080802,
@ -194,7 +194,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
} }
}; };
CONSTANT_AS u32a c_skb[8][64] = CONSTANT_VK u32a c_skb[8][64] =
{ {
{ {
0x00000000, 0x00000010, 0x20000000, 0x20000010, 0x00000000, 0x00000010, 0x20000000, 0x20000010,

4
OpenCL/m05500_a3-optimized.cl
View File

@ -46,7 +46,7 @@ typedef struct netntlm
a = a ^ tt; \ a = a ^ tt; \
} }
CONSTANT_AS u32a c_SPtrans[8][64] = CONSTANT_VK u32a c_SPtrans[8][64] =
{ {
{ {
0x02080800, 0x00080000, 0x02000002, 0x02080802, 0x02080800, 0x00080000, 0x02000002, 0x02080802,
@ -194,7 +194,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
} }
}; };
CONSTANT_AS u32a c_skb[8][64] = CONSTANT_VK u32a c_skb[8][64] =
{ {
{ {
0x00000000, 0x00000010, 0x20000000, 0x20000010, 0x00000000, 0x00000010, 0x20000000, 0x20000010,

4
OpenCL/m05500_a3-pure.cl
View File

@ -46,7 +46,7 @@ typedef struct netntlm
a = a ^ tt; \ a = a ^ tt; \
} }
CONSTANT_AS u32a c_SPtrans[8][64] = CONSTANT_VK u32a c_SPtrans[8][64] =
{ {
{ {
0x02080800, 0x00080000, 0x02000002, 0x02080802, 0x02080800, 0x00080000, 0x02000002, 0x02080802,
@ -194,7 +194,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
} }
}; };
CONSTANT_AS u32a c_skb[8][64] = CONSTANT_VK u32a c_skb[8][64] =
{ {
{ {
0x00000000, 0x00000010, 0x20000000, 0x20000010, 0x00000000, 0x00000010, 0x20000000, 0x20000010,

4
OpenCL/m05800-optimized.cl
View File

@ -20,7 +20,7 @@ typedef struct androidpin_tmp
} androidpin_tmp_t; } androidpin_tmp_t;
CONSTANT_AS u32a c_pc_dec[1024] = CONSTANT_VK u32a c_pc_dec[1024] =
{ {
0x00000030, 0x00000030,
0x00000031, 0x00000031,
@ -1048,7 +1048,7 @@ CONSTANT_AS u32a c_pc_dec[1024] =
0x33323031, 0x33323031,
}; };
CONSTANT_AS u32a c_pc_len[1024] = CONSTANT_VK u32a c_pc_len[1024] =
{ {
1, 1,
1, 1,

4
OpenCL/m05800-pure.cl
View File

@ -20,7 +20,7 @@ typedef struct androidpin_tmp
} androidpin_tmp_t; } androidpin_tmp_t;
CONSTANT_AS u32a c_pc_dec[1024] = CONSTANT_VK u32a c_pc_dec[1024] =
{ {
0x00000030, 0x00000030,
0x00000031, 0x00000031,
@ -1048,7 +1048,7 @@ CONSTANT_AS u32a c_pc_dec[1024] =
0x33323031, 0x33323031,
}; };
CONSTANT_AS u32a c_pc_len[1024] = CONSTANT_VK u32a c_pc_len[1024] =
{ {
1, 1,
1, 1,

2
OpenCL/m06900_a0-optimized.cl
View File

@ -15,7 +15,7 @@
#include "inc_simd.cl" #include "inc_simd.cl"
#endif #endif
CONSTANT_AS u32a c_tables[4][256] = CONSTANT_VK u32a c_tables[4][256] =
{ {
{ {
0x00072000, 0x00075000, 0x00074800, 0x00071000, 0x00072000, 0x00075000, 0x00074800, 0x00071000,

2
OpenCL/m06900_a1-optimized.cl
View File

@ -13,7 +13,7 @@
#include "inc_simd.cl" #include "inc_simd.cl"
#endif #endif
CONSTANT_AS u32a c_tables[4][256] = CONSTANT_VK u32a c_tables[4][256] =
{ {
{ {
0x00072000, 0x00075000, 0x00074800, 0x00071000, 0x00072000, 0x00075000, 0x00074800, 0x00071000,

2
OpenCL/m06900_a3-optimized.cl
View File

@ -13,7 +13,7 @@
#include "inc_simd.cl" #include "inc_simd.cl"
#endif #endif
CONSTANT_AS u32a c_tables[4][256] = CONSTANT_VK u32a c_tables[4][256] =
{ {
{ {
0x00072000, 0x00075000, 0x00074800, 0x00071000, 0x00072000, 0x00075000, 0x00074800, 0x00071000,

4
OpenCL/m07700_a0-optimized.cl
View File

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

4
OpenCL/m07700_a1-optimized.cl
View File

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

4
OpenCL/m07700_a3-optimized.cl
View File

@ -18,7 +18,7 @@
#define GETCHAR(a,p) (((a)[(p) / 4] >> (((p) & 3) * 8)) & 0xff) #define GETCHAR(a,p) (((a)[(p) / 4] >> (((p) & 3) * 8)) & 0xff)
#define PUTCHAR(a,p,c) ((a)[(p) / 4] = (((a)[(p) / 4] & ~(0xff << (((p) & 3) * 8))) | ((c) << (((p) & 3) * 8)))) #define PUTCHAR(a,p,c) ((a)[(p) / 4] = (((a)[(p) / 4] & ~(0xff << (((p) & 3) * 8))) | ((c) << (((p) & 3) * 8))))
CONSTANT_AS u32a sapb_trans_tbl[256] = CONSTANT_VK u32a sapb_trans_tbl[256] =
{ {
// first value hack for 0 byte as part of an optimization // 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, 0x00, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
@ -39,7 +39,7 @@ CONSTANT_AS u32a sapb_trans_tbl[256] =
0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff
}; };
CONSTANT_AS u32a bcodeArray[48] = CONSTANT_VK u32a bcodeArray[48] =
{ {
0x14, 0x77, 0xf3, 0xd4, 0xbb, 0x71, 0x23, 0xd0, 0x03, 0xff, 0x47, 0x93, 0x55, 0xaa, 0x66, 0x91, 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, 0xf2, 0x88, 0x6b, 0x99, 0xbf, 0xcb, 0x32, 0x1a, 0x19, 0xd9, 0xa7, 0x82, 0x22, 0x49, 0xa2, 0x51,

4
OpenCL/m07701_a0-optimized.cl
View File

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

4
OpenCL/m07701_a1-optimized.cl
View File

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

4
OpenCL/m07701_a3-optimized.cl
View File

@ -18,7 +18,7 @@
#define GETCHAR(a,p) (((a)[(p) / 4] >> (((p) & 3) * 8)) & 0xff) #define GETCHAR(a,p) (((a)[(p) / 4] >> (((p) & 3) * 8)) & 0xff)
#define PUTCHAR(a,p,c) ((a)[(p) / 4] = (((a)[(p) / 4] & ~(0xff << (((p) & 3) * 8))) | ((c) << (((p) & 3) * 8)))) #define PUTCHAR(a,p,c) ((a)[(p) / 4] = (((a)[(p) / 4] & ~(0xff << (((p) & 3) * 8))) | ((c) << (((p) & 3) * 8))))
CONSTANT_AS u32a sapb_trans_tbl[256] = CONSTANT_VK u32a sapb_trans_tbl[256] =
{ {
// first value hack for 0 byte as part of an optimization // 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, 0x00, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
@ -39,7 +39,7 @@ CONSTANT_AS u32a sapb_trans_tbl[256] =
0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff
}; };
CONSTANT_AS u32a bcodeArray[48] = CONSTANT_VK u32a bcodeArray[48] =
{ {
0x14, 0x77, 0xf3, 0xd4, 0xbb, 0x71, 0x23, 0xd0, 0x03, 0xff, 0x47, 0x93, 0x55, 0xaa, 0x66, 0x91, 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, 0xf2, 0x88, 0x6b, 0x99, 0xbf, 0xcb, 0x32, 0x1a, 0x19, 0xd9, 0xa7, 0x82, 0x22, 0x49, 0xa2, 0x51,

2
OpenCL/m07800_a0-optimized.cl
View File

@ -17,7 +17,7 @@
#include "inc_hash_sha1.cl" #include "inc_hash_sha1.cl"
#endif #endif
CONSTANT_AS u32a theMagicArray[64] = CONSTANT_VK u32a theMagicArray[64] =
{ {
0x91ac5114, 0x9f675443, 0x24e73be0, 0x28747bc2, 0x863313eb, 0x5a4fcb5c, 0x080a7337, 0x0e5d1c2f, 0x91ac5114, 0x9f675443, 0x24e73be0, 0x28747bc2, 0x863313eb, 0x5a4fcb5c, 0x080a7337, 0x0e5d1c2f,
0x338fe6e5, 0xf89baedd, 0x16f24b8d, 0x2ce1d4dc, 0xb0cbdf9d, 0xd4706d17, 0xf94d423f, 0x9b1b1194, 0x338fe6e5, 0xf89baedd, 0x16f24b8d, 0x2ce1d4dc, 0xb0cbdf9d, 0xd4706d17, 0xf94d423f, 0x9b1b1194,

2
OpenCL/m07800_a1-optimized.cl
View File

@ -15,7 +15,7 @@
#include "inc_hash_sha1.cl" #include "inc_hash_sha1.cl"
#endif #endif
CONSTANT_AS u32a theMagicArray[64] = CONSTANT_VK u32a theMagicArray[64] =
{ {
0x91ac5114, 0x9f675443, 0x24e73be0, 0x28747bc2, 0x863313eb, 0x5a4fcb5c, 0x080a7337, 0x0e5d1c2f, 0x91ac5114, 0x9f675443, 0x24e73be0, 0x28747bc2, 0x863313eb, 0x5a4fcb5c, 0x080a7337, 0x0e5d1c2f,
0x338fe6e5, 0xf89baedd, 0x16f24b8d, 0x2ce1d4dc, 0xb0cbdf9d, 0xd4706d17, 0xf94d423f, 0x9b1b1194, 0x338fe6e5, 0xf89baedd, 0x16f24b8d, 0x2ce1d4dc, 0xb0cbdf9d, 0xd4706d17, 0xf94d423f, 0x9b1b1194,

2
OpenCL/m07800_a3-optimized.cl
View File

@ -15,7 +15,7 @@
#include "inc_hash_sha1.cl" #include "inc_hash_sha1.cl"
#endif #endif
CONSTANT_AS u32a theMagicArray[64] = CONSTANT_VK u32a theMagicArray[64] =
{ {
0x91ac5114, 0x9f675443, 0x24e73be0, 0x28747bc2, 0x863313eb, 0x5a4fcb5c, 0x080a7337, 0x0e5d1c2f, 0x91ac5114, 0x9f675443, 0x24e73be0, 0x28747bc2, 0x863313eb, 0x5a4fcb5c, 0x080a7337, 0x0e5d1c2f,
0x338fe6e5, 0xf89baedd, 0x16f24b8d, 0x2ce1d4dc, 0xb0cbdf9d, 0xd4706d17, 0xf94d423f, 0x9b1b1194, 0x338fe6e5, 0xf89baedd, 0x16f24b8d, 0x2ce1d4dc, 0xb0cbdf9d, 0xd4706d17, 0xf94d423f, 0x9b1b1194,

2
OpenCL/m07801_a0-optimized.cl
View File

@ -17,7 +17,7 @@
#include "inc_hash_sha1.cl" #include "inc_hash_sha1.cl"
#endif #endif
CONSTANT_AS u32a theMagicArray[64] = CONSTANT_VK u32a theMagicArray[64] =
{ {
0x91ac5114, 0x9f675443, 0x24e73be0, 0x28747bc2, 0x863313eb, 0x5a4fcb5c, 0x080a7337, 0x0e5d1c2f, 0x91ac5114, 0x9f675443, 0x24e73be0, 0x28747bc2, 0x863313eb, 0x5a4fcb5c, 0x080a7337, 0x0e5d1c2f,
0x338fe6e5, 0xf89baedd, 0x16f24b8d, 0x2ce1d4dc, 0xb0cbdf9d, 0xd4706d17, 0xf94d423f, 0x9b1b1194, 0x338fe6e5, 0xf89baedd, 0x16f24b8d, 0x2ce1d4dc, 0xb0cbdf9d, 0xd4706d17, 0xf94d423f, 0x9b1b1194,

2
OpenCL/m07801_a1-optimized.cl
View File

@ -15,7 +15,7 @@
#include "inc_hash_sha1.cl" #include "inc_hash_sha1.cl"
#endif #endif
CONSTANT_AS u32a theMagicArray[64] = CONSTANT_VK u32a theMagicArray[64] =
{ {
0x91ac5114, 0x9f675443, 0x24e73be0, 0x28747bc2, 0x863313eb, 0x5a4fcb5c, 0x080a7337, 0x0e5d1c2f, 0x91ac5114, 0x9f675443, 0x24e73be0, 0x28747bc2, 0x863313eb, 0x5a4fcb5c, 0x080a7337, 0x0e5d1c2f,
0x338fe6e5, 0xf89baedd, 0x16f24b8d, 0x2ce1d4dc, 0xb0cbdf9d, 0xd4706d17, 0xf94d423f, 0x9b1b1194, 0x338fe6e5, 0xf89baedd, 0x16f24b8d, 0x2ce1d4dc, 0xb0cbdf9d, 0xd4706d17, 0xf94d423f, 0x9b1b1194,

2
OpenCL/m07801_a3-optimized.cl
View File

@ -15,7 +15,7 @@
#include "inc_hash_sha1.cl" #include "inc_hash_sha1.cl"
#endif #endif
CONSTANT_AS u32a theMagicArray[64] = CONSTANT_VK u32a theMagicArray[64] =
{ {
0x91ac5114, 0x9f675443, 0x24e73be0, 0x28747bc2, 0x863313eb, 0x5a4fcb5c, 0x080a7337, 0x0e5d1c2f, 0x91ac5114, 0x9f675443, 0x24e73be0, 0x28747bc2, 0x863313eb, 0x5a4fcb5c, 0x080a7337, 0x0e5d1c2f,
0x338fe6e5, 0xf89baedd, 0x16f24b8d, 0x2ce1d4dc, 0xb0cbdf9d, 0xd4706d17, 0xf94d423f, 0x9b1b1194, 0x338fe6e5, 0xf89baedd, 0x16f24b8d, 0x2ce1d4dc, 0xb0cbdf9d, 0xd4706d17, 0xf94d423f, 0x9b1b1194,

6
OpenCL/m08500_a0-pure.cl
View File

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

6
OpenCL/m08500_a1-pure.cl
View File

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

6
OpenCL/m08500_a3-pure.cl
View File

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

2
OpenCL/m08600_a0-pure.cl
View File

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

2
OpenCL/m08600_a1-pure.cl
View File

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

2
OpenCL/m08600_a3-pure.cl
View File

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

2
OpenCL/m08700_a0-optimized.cl
View File

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

2
OpenCL/m08700_a1-optimized.cl
View File

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

2
OpenCL/m08700_a3-optimized.cl
View File

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

10
OpenCL/m09000-pure.cl
View File

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

4
OpenCL/m09100-pure.cl
View File

@ -27,7 +27,7 @@ typedef struct lotus8_tmp
} lotus8_tmp_t; } lotus8_tmp_t;
CONSTANT_AS u32a lotus64_table[64] = CONSTANT_VK u32a lotus64_table[64] =
{ {
'0', '1', '2', '3', '4', '5', '6', '7', '0', '1', '2', '3', '4', '5', '6', '7',
'8', '9', 'A', 'B', 'C', 'D', 'E', 'F', '8', '9', 'A', 'B', 'C', 'D', 'E', 'F',
@ -39,7 +39,7 @@ CONSTANT_AS u32a lotus64_table[64] =
'u', 'v', 'w', 'x', 'y', 'z', '+', '/', 'u', 'v', 'w', 'x', 'y', 'z', '+', '/',
}; };
CONSTANT_AS u32a lotus_magic_table[256] = CONSTANT_VK u32a lotus_magic_table[256] =
{ {
0xbd, 0x56, 0xea, 0xf2, 0xa2, 0xf1, 0xac, 0x2a, 0xbd, 0x56, 0xea, 0xf2, 0xa2, 0xf1, 0xac, 0x2a,
0xb0, 0x93, 0xd1, 0x9c, 0x1b, 0x33, 0xfd, 0xd0, 0xb0, 0x93, 0xd1, 0x9c, 0x1b, 0x33, 0xfd, 0xd0,

2
OpenCL/m10400_a0-optimized.cl
View File

@ -17,7 +17,7 @@
#include "inc_hash_md5.cl" #include "inc_hash_md5.cl"
#endif #endif
CONSTANT_AS u32a padding[8] = CONSTANT_VK u32a padding[8] =
{ {
0x5e4ebf28, 0x5e4ebf28,
0x418a754e, 0x418a754e,

2
OpenCL/m10400_a1-optimized.cl
View File

@ -15,7 +15,7 @@
#include "inc_hash_md5.cl" #include "inc_hash_md5.cl"
#endif #endif
CONSTANT_AS u32a padding[8] = CONSTANT_VK u32a padding[8] =
{ {
0x5e4ebf28, 0x5e4ebf28,
0x418a754e, 0x418a754e,

2
OpenCL/m10400_a3-optimized.cl
View File

@ -15,7 +15,7 @@
#include "inc_hash_md5.cl" #include "inc_hash_md5.cl"
#endif #endif
CONSTANT_AS u32a padding[8] = CONSTANT_VK u32a padding[8] =
{ {
0x5e4ebf28, 0x5e4ebf28,
0x418a754e, 0x418a754e,

2
OpenCL/m10410_a0-optimized.cl
View File

@ -17,7 +17,7 @@
#include "inc_hash_md5.cl" #include "inc_hash_md5.cl"
#endif #endif
CONSTANT_AS u32a padding[8] = CONSTANT_VK u32a padding[8] =
{ {
0x5e4ebf28, 0x5e4ebf28,
0x418a754e, 0x418a754e,

2
OpenCL/m10410_a1-optimized.cl
View File

@ -15,7 +15,7 @@
#include "inc_hash_md5.cl" #include "inc_hash_md5.cl"
#endif #endif
CONSTANT_AS u32a padding[8] = CONSTANT_VK u32a padding[8] =
{ {
0x5e4ebf28, 0x5e4ebf28,
0x418a754e, 0x418a754e,

2
OpenCL/m10410_a3-optimized.cl
View File

@ -15,7 +15,7 @@
#include "inc_hash_md5.cl" #include "inc_hash_md5.cl"
#endif #endif
CONSTANT_AS u32a padding[8] = CONSTANT_VK u32a padding[8] =
{ {
0x5e4ebf28, 0x5e4ebf28,
0x418a754e, 0x418a754e,

2
OpenCL/m10420_a0-optimized.cl
View File

@ -16,7 +16,7 @@
#include "inc_hash_md5.cl" #include "inc_hash_md5.cl"
#endif #endif
CONSTANT_AS u32a padding[8] = CONSTANT_VK u32a padding[8] =
{ {
0x5e4ebf28, 0x5e4ebf28,
0x418a754e, 0x418a754e,

2
OpenCL/m10420_a1-optimized.cl
View File

@ -14,7 +14,7 @@
#include "inc_hash_md5.cl" #include "inc_hash_md5.cl"
#endif #endif
CONSTANT_AS u32a padding[8] = CONSTANT_VK u32a padding[8] =
{ {
0x5e4ebf28, 0x5e4ebf28,
0x418a754e, 0x418a754e,

2
OpenCL/m10420_a3-optimized.cl
View File

@ -14,7 +14,7 @@
#include "inc_hash_md5.cl" #include "inc_hash_md5.cl"
#endif #endif
CONSTANT_AS u32a padding[8] = CONSTANT_VK u32a padding[8] =
{ {
0x5e4ebf28, 0x5e4ebf28,
0x418a754e, 0x418a754e,

2
OpenCL/m10500-pure.cl
View File

@ -14,7 +14,7 @@
#define COMPARE_S "inc_comp_single.cl" #define COMPARE_S "inc_comp_single.cl"
#define COMPARE_M "inc_comp_multi.cl" #define COMPARE_M "inc_comp_multi.cl"
CONSTANT_AS u32a padding[8] = CONSTANT_VK u32a padding[8] =
{ {
0x5e4ebf28, 0x5e4ebf28,
0x418a754e, 0x418a754e,

2
OpenCL/m11500_a0-optimized.cl
View File

@ -16,7 +16,7 @@
#include "inc_simd.cl" #include "inc_simd.cl"
#endif #endif
CONSTANT_AS u32a crc32tab[0x100] = CONSTANT_VK u32a crc32tab[0x100] =
{ {
0x00000000, 0x77073096, 0xee0e612c, 0x990951ba, 0x00000000, 0x77073096, 0xee0e612c, 0x990951ba,
0x076dc419, 0x706af48f, 0xe963a535, 0x9e6495a3, 0x076dc419, 0x706af48f, 0xe963a535, 0x9e6495a3,

2
OpenCL/m11500_a1-optimized.cl
View File

@ -14,7 +14,7 @@
#include "inc_simd.cl" #include "inc_simd.cl"
#endif #endif
CONSTANT_AS u32a crc32tab[0x100] = CONSTANT_VK u32a crc32tab[0x100] =
{ {
0x00000000, 0x77073096, 0xee0e612c, 0x990951ba, 0x00000000, 0x77073096, 0xee0e612c, 0x990951ba,
0x076dc419, 0x706af48f, 0xe963a535, 0x9e6495a3, 0x076dc419, 0x706af48f, 0xe963a535, 0x9e6495a3,

2
OpenCL/m11500_a3-optimized.cl
View File

@ -14,7 +14,7 @@
#include "inc_simd.cl" #include "inc_simd.cl"
#endif #endif
CONSTANT_AS u32a crc32tab[0x100] = CONSTANT_VK u32a crc32tab[0x100] =
{ {
0x00000000, 0x77073096, 0xee0e612c, 0x990951ba, 0x00000000, 0x77073096, 0xee0e612c, 0x990951ba,
0x076dc419, 0x706af48f, 0xe963a535, 0x9e6495a3, 0x076dc419, 0x706af48f, 0xe963a535, 0x9e6495a3,

4
OpenCL/m12400-pure.cl
View File

@ -60,7 +60,7 @@ typedef struct bsdicrypt_tmp
PERM_OP (l, r, tt, 4, 0x0f0f0f0f); \ PERM_OP (l, r, tt, 4, 0x0f0f0f0f); \
} }
CONSTANT_AS u32a c_SPtrans[8][64] = CONSTANT_VK u32a c_SPtrans[8][64] =
{ {
{ {
0x00820200, 0x00020000, 0x80800000, 0x80820200, 0x00820200, 0x00020000, 0x80800000, 0x80820200,
@ -208,7 +208,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
}, },
}; };
CONSTANT_AS u32a c_skb[8][64] = CONSTANT_VK u32a c_skb[8][64] =
{ {
{ {
0x00000000, 0x00000010, 0x20000000, 0x20000010, 0x00000000, 0x00000010, 0x20000000, 0x20000010,

4
OpenCL/m14000_a0-pure.cl
View File

@ -53,7 +53,7 @@
PERM_OP (l, r, tt, 4, 0x0f0f0f0f); \ PERM_OP (l, r, tt, 4, 0x0f0f0f0f); \
} }
CONSTANT_AS u32a c_SPtrans[8][64] = CONSTANT_VK u32a c_SPtrans[8][64] =
{ {
{ {
/* nibble 0 */ /* nibble 0 */
@ -209,7 +209,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
}, },
}; };
CONSTANT_AS u32a c_skb[8][64] = CONSTANT_VK u32a c_skb[8][64] =
{ {
{ {
0x00000000, 0x00000010, 0x20000000, 0x20000010, 0x00000000, 0x00000010, 0x20000000, 0x20000010,

4
OpenCL/m14000_a1-pure.cl
View File

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

4
OpenCL/m14100_a0-pure.cl
View File

@ -53,7 +53,7 @@
PERM_OP (l, r, tt, 4, 0x0f0f0f0f); \ PERM_OP (l, r, tt, 4, 0x0f0f0f0f); \
} }
CONSTANT_AS u32a c_SPtrans[8][64] = CONSTANT_VK u32a c_SPtrans[8][64] =
{ {
{ {
/* nibble 0 */ /* nibble 0 */
@ -209,7 +209,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
}, },
}; };
CONSTANT_AS u32a c_skb[8][64] = CONSTANT_VK u32a c_skb[8][64] =
{ {
{ {
0x00000000, 0x00000010, 0x20000000, 0x20000010, 0x00000000, 0x00000010, 0x20000000, 0x20000010,

4
OpenCL/m14100_a1-pure.cl
View File

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

4
OpenCL/m14100_a3-pure.cl
View File

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

2
OpenCL/m14900_a0-optimized.cl
View File

@ -16,7 +16,7 @@
#include "inc_simd.cl" #include "inc_simd.cl"
#endif #endif
CONSTANT_AS u8a c_ftable[256] = CONSTANT_VK u8a c_ftable[256] =
{ {
0xa3, 0xd7, 0x09, 0x83, 0xf8, 0x48, 0xf6, 0xf4, 0xa3, 0xd7, 0x09, 0x83, 0xf8, 0x48, 0xf6, 0xf4,
0xb3, 0x21, 0x15, 0x78, 0x99, 0xb1, 0xaf, 0xf9, 0xb3, 0x21, 0x15, 0x78, 0x99, 0xb1, 0xaf, 0xf9,

2
OpenCL/m14900_a1-optimized.cl
View File

@ -14,7 +14,7 @@
#include "inc_simd.cl" #include "inc_simd.cl"
#endif #endif
CONSTANT_AS u8a c_ftable[256] = CONSTANT_VK u8a c_ftable[256] =
{ {
0xa3, 0xd7, 0x09, 0x83, 0xf8, 0x48, 0xf6, 0xf4, 0xa3, 0xd7, 0x09, 0x83, 0xf8, 0x48, 0xf6, 0xf4,
0xb3, 0x21, 0x15, 0x78, 0x99, 0xb1, 0xaf, 0xf9, 0xb3, 0x21, 0x15, 0x78, 0x99, 0xb1, 0xaf, 0xf9,

2
OpenCL/m14900_a3-optimized.cl
View File

@ -14,7 +14,7 @@
#include "inc_simd.cl" #include "inc_simd.cl"
#endif #endif
CONSTANT_AS u8a c_ftable[256] = CONSTANT_VK u8a c_ftable[256] =
{ {
0xa3, 0xd7, 0x09, 0x83, 0xf8, 0x48, 0xf6, 0xf4, 0xa3, 0xd7, 0x09, 0x83, 0xf8, 0x48, 0xf6, 0xf4,
0xb3, 0x21, 0x15, 0x78, 0x99, 0xb1, 0xaf, 0xf9, 0xb3, 0x21, 0x15, 0x78, 0x99, 0xb1, 0xaf, 0xf9,

2
OpenCL/m15600-pure.cl
View File

@ -34,7 +34,7 @@ typedef struct ethereum_pbkdf2
#define COMPARE_S "inc_comp_single.cl" #define COMPARE_S "inc_comp_single.cl"
#define COMPARE_M "inc_comp_multi.cl" #define COMPARE_M "inc_comp_multi.cl"
CONSTANT_AS u64a keccakf_rndc[24] = CONSTANT_VK u64a keccakf_rndc[24] =
{ {
0x0000000000000001, 0x0000000000008082, 0x800000000000808a, 0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
0x8000000080008000, 0x000000000000808b, 0x0000000080000001, 0x8000000080008000, 0x000000000000808b, 0x0000000080000001,

2
OpenCL/m15700-pure.cl
View File

@ -261,7 +261,7 @@ DECLSPEC void scrypt_smix (uint4 *X, uint4 *T, GLOBAL_AS uint4 *V0, GLOBAL_AS ui
st[4 + s] ^= ~bc0 & bc1; \ st[4 + s] ^= ~bc0 & bc1; \
} }
CONSTANT_AS u64a keccakf_rndc[24] = CONSTANT_VK u64a keccakf_rndc[24] =
{ {
0x0000000000000001, 0x0000000000008082, 0x800000000000808a, 0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
0x8000000080008000, 0x000000000000808b, 0x0000000080000001, 0x8000000080008000, 0x000000000000808b, 0x0000000080000001,

6
OpenCL/m16000_a0-pure.cl
View File

@ -35,7 +35,7 @@
a = a ^ tt; \ a = a ^ tt; \
} }
CONSTANT_AS u32a c_SPtrans[8][64] = CONSTANT_VK u32a c_SPtrans[8][64] =
{ {
{ {
0x00820200, 0x00020000, 0x80800000, 0x80820200, 0x00820200, 0x00020000, 0x80800000, 0x80820200,
@ -183,7 +183,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
}, },
}; };
CONSTANT_AS u32a c_skb[8][64] = CONSTANT_VK u32a c_skb[8][64] =
{ {
{ {
0x00000000, 0x00000010, 0x20000000, 0x20000010, 0x00000000, 0x00000010, 0x20000000, 0x20000010,
@ -331,7 +331,7 @@ CONSTANT_AS u32a c_skb[8][64] =
}, },
}; };
CONSTANT_AS u32a c_tripcode_salt[128] = CONSTANT_VK u32a c_tripcode_salt[128] =
{ {
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,

6
OpenCL/m16000_a1-pure.cl
View File

@ -33,7 +33,7 @@
a = a ^ tt; \ a = a ^ tt; \
} }
CONSTANT_AS u32a c_SPtrans[8][64] = CONSTANT_VK u32a c_SPtrans[8][64] =
{ {
{ {
0x00820200, 0x00020000, 0x80800000, 0x80820200, 0x00820200, 0x00020000, 0x80800000, 0x80820200,
@ -181,7 +181,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
}, },
}; };
CONSTANT_AS u32a c_skb[8][64] = CONSTANT_VK u32a c_skb[8][64] =
{ {
{ {
0x00000000, 0x00000010, 0x20000000, 0x20000010, 0x00000000, 0x00000010, 0x20000000, 0x20000010,
@ -329,7 +329,7 @@ CONSTANT_AS u32a c_skb[8][64] =
}, },
}; };
CONSTANT_AS u32a c_tripcode_salt[128] = CONSTANT_VK u32a c_tripcode_salt[128] =
{ {
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,

6
OpenCL/m16000_a3-pure.cl
View File

@ -33,7 +33,7 @@
a = a ^ tt; \ a = a ^ tt; \
} }
CONSTANT_AS u32a c_SPtrans[8][64] = CONSTANT_VK u32a c_SPtrans[8][64] =
{ {
{ {
0x00820200, 0x00020000, 0x80800000, 0x80820200, 0x00820200, 0x00020000, 0x80800000, 0x80820200,
@ -181,7 +181,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
}, },
}; };
CONSTANT_AS u32a c_skb[8][64] = CONSTANT_VK u32a c_skb[8][64] =
{ {
{ {
0x00000000, 0x00000010, 0x20000000, 0x20000010, 0x00000000, 0x00000010, 0x20000000, 0x20000010,
@ -329,7 +329,7 @@ CONSTANT_AS u32a c_skb[8][64] =
}, },
}; };
CONSTANT_AS u32a c_tripcode_salt[128] = CONSTANT_VK u32a c_tripcode_salt[128] =
{ {
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,

2
OpenCL/m16300-pure.cl
View File

@ -36,7 +36,7 @@ typedef struct pbkdf2_sha256_tmp
} pbkdf2_sha256_tmp_t; } pbkdf2_sha256_tmp_t;
CONSTANT_AS u64a keccakf_rndc[24] = CONSTANT_VK u64a keccakf_rndc[24] =
{ {
0x0000000000000001, 0x0000000000008082, 0x800000000000808a, 0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
0x8000000080008000, 0x000000000000808b, 0x0000000080000001, 0x8000000080008000, 0x000000000000808b, 0x0000000080000001,

2
OpenCL/m17300_a0-optimized.cl
View File

@ -15,7 +15,7 @@
#include "inc_simd.cl" #include "inc_simd.cl"
#endif #endif
CONSTANT_AS u64a keccakf_rndc[24] = CONSTANT_VK u64a keccakf_rndc[24] =
{ {
0x0000000000000001, 0x0000000000008082, 0x800000000000808a, 0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
0x8000000080008000, 0x000000000000808b, 0x0000000080000001, 0x8000000080008000, 0x000000000000808b, 0x0000000080000001,

2
OpenCL/m17300_a1-optimized.cl
View File

@ -13,7 +13,7 @@
#include "inc_simd.cl" #include "inc_simd.cl"
#endif #endif
CONSTANT_AS u64a keccakf_rndc[24] = CONSTANT_VK u64a keccakf_rndc[24] =
{ {
0x0000000000000001, 0x0000000000008082, 0x800000000000808a, 0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
0x8000000080008000, 0x000000000000808b, 0x0000000080000001, 0x8000000080008000, 0x000000000000808b, 0x0000000080000001,

2
OpenCL/m17300_a3-optimized.cl
View File

@ -13,7 +13,7 @@
#include "inc_simd.cl" #include "inc_simd.cl"
#endif #endif
CONSTANT_AS u64a keccakf_rndc[24] = CONSTANT_VK u64a keccakf_rndc[24] =
{ {
0x0000000000000001, 0x0000000000008082, 0x800000000000808a, 0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
0x8000000080008000, 0x000000000000808b, 0x0000000080000001, 0x8000000080008000, 0x000000000000808b, 0x0000000080000001,

2
OpenCL/m17400_a0-optimized.cl
View File

@ -15,7 +15,7 @@
#include "inc_simd.cl" #include "inc_simd.cl"
#endif #endif
CONSTANT_AS u64a keccakf_rndc[24] = CONSTANT_VK u64a keccakf_rndc[24] =
{ {
0x0000000000000001, 0x0000000000008082, 0x800000000000808a, 0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
0x8000000080008000, 0x000000000000808b, 0x0000000080000001, 0x8000000080008000, 0x000000000000808b, 0x0000000080000001,

2
OpenCL/m17400_a1-optimized.cl
View File

@ -13,7 +13,7 @@
#include "inc_simd.cl" #include "inc_simd.cl"
#endif #endif
CONSTANT_AS u64a keccakf_rndc[24] = CONSTANT_VK u64a keccakf_rndc[24] =
{ {
0x0000000000000001, 0x0000000000008082, 0x800000000000808a, 0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
0x8000000080008000, 0x000000000000808b, 0x0000000080000001, 0x8000000080008000, 0x000000000000808b, 0x0000000080000001,

2
OpenCL/m17400_a3-optimized.cl
View File

@ -13,7 +13,7 @@
#include "inc_simd.cl" #include "inc_simd.cl"
#endif #endif
CONSTANT_AS u64a keccakf_rndc[24] = CONSTANT_VK u64a keccakf_rndc[24] =
{ {
0x0000000000000001, 0x0000000000008082, 0x800000000000808a, 0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
0x8000000080008000, 0x000000000000808b, 0x0000000080000001, 0x8000000080008000, 0x000000000000808b, 0x0000000080000001,

2
OpenCL/m17500_a0-optimized.cl
View File

@ -15,7 +15,7 @@
#include "inc_simd.cl" #include "inc_simd.cl"
#endif #endif
CONSTANT_AS u64a keccakf_rndc[24] = CONSTANT_VK u64a keccakf_rndc[24] =
{ {
0x0000000000000001, 0x0000000000008082, 0x800000000000808a, 0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
0x8000000080008000, 0x000000000000808b, 0x0000000080000001, 0x8000000080008000, 0x000000000000808b, 0x0000000080000001,

2
OpenCL/m17500_a1-optimized.cl
View File

@ -13,7 +13,7 @@
#include "inc_simd.cl" #include "inc_simd.cl"
#endif #endif
CONSTANT_AS u64a keccakf_rndc[24] = CONSTANT_VK u64a keccakf_rndc[24] =
{ {
0x0000000000000001, 0x0000000000008082, 0x800000000000808a, 0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
0x8000000080008000, 0x000000000000808b, 0x0000000080000001, 0x8000000080008000, 0x000000000000808b, 0x0000000080000001,

2
OpenCL/m17500_a3-optimized.cl
View File

@ -13,7 +13,7 @@
#include "inc_simd.cl" #include "inc_simd.cl"
#endif #endif
CONSTANT_AS u64a keccakf_rndc[24] = CONSTANT_VK u64a keccakf_rndc[24] =
{ {
0x0000000000000001, 0x0000000000008082, 0x800000000000808a, 0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
0x8000000080008000, 0x000000000000808b, 0x0000000080000001, 0x8000000080008000, 0x000000000000808b, 0x0000000080000001,

2
OpenCL/m17600_a0-optimized.cl
View File

@ -15,7 +15,7 @@
#include "inc_simd.cl" #include "inc_simd.cl"
#endif #endif
CONSTANT_AS u64a keccakf_rndc[24] = CONSTANT_VK u64a keccakf_rndc[24] =
{ {
0x0000000000000001, 0x0000000000008082, 0x800000000000808a, 0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
0x8000000080008000, 0x000000000000808b, 0x0000000080000001, 0x8000000080008000, 0x000000000000808b, 0x0000000080000001,

2
OpenCL/m17600_a1-optimized.cl
View File

@ -13,7 +13,7 @@
#include "inc_simd.cl" #include "inc_simd.cl"
#endif #endif
CONSTANT_AS u64a keccakf_rndc[24] = CONSTANT_VK u64a keccakf_rndc[24] =
{ {
0x0000000000000001, 0x0000000000008082, 0x800000000000808a, 0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
0x8000000080008000, 0x000000000000808b, 0x0000000080000001, 0x8000000080008000, 0x000000000000808b, 0x0000000080000001,

2
OpenCL/m17600_a3-optimized.cl
View File

@ -13,7 +13,7 @@
#include "inc_simd.cl" #include "inc_simd.cl"
#endif #endif
CONSTANT_AS u64a keccakf_rndc[24] = CONSTANT_VK u64a keccakf_rndc[24] =
{ {
0x0000000000000001, 0x0000000000008082, 0x800000000000808a, 0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
0x8000000080008000, 0x000000000000808b, 0x0000000080000001, 0x8000000080008000, 0x000000000000808b, 0x0000000080000001,

2
OpenCL/m17700_a0-optimized.cl
View File

@ -15,7 +15,7 @@
#include "inc_simd.cl" #include "inc_simd.cl"
#endif #endif
CONSTANT_AS u64a keccakf_rndc[24] = CONSTANT_VK u64a keccakf_rndc[24] =
{ {
0x0000000000000001, 0x0000000000008082, 0x800000000000808a, 0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
0x8000000080008000, 0x000000000000808b, 0x0000000080000001, 0x8000000080008000, 0x000000000000808b, 0x0000000080000001,

2
OpenCL/m17700_a1-optimized.cl
View File

@ -13,7 +13,7 @@
#include "inc_simd.cl" #include "inc_simd.cl"
#endif #endif
CONSTANT_AS u64a keccakf_rndc[24] = CONSTANT_VK u64a keccakf_rndc[24] =
{ {
0x0000000000000001, 0x0000000000008082, 0x800000000000808a, 0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
0x8000000080008000, 0x000000000000808b, 0x0000000080000001, 0x8000000080008000, 0x000000000000808b, 0x0000000080000001,

2
OpenCL/m17700_a3-optimized.cl
View File

@ -13,7 +13,7 @@
#include "inc_simd.cl" #include "inc_simd.cl"
#endif #endif
CONSTANT_AS u64a keccakf_rndc[24] = CONSTANT_VK u64a keccakf_rndc[24] =
{ {
0x0000000000000001, 0x0000000000008082, 0x800000000000808a, 0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
0x8000000080008000, 0x000000000000808b, 0x0000000080000001, 0x8000000080008000, 0x000000000000808b, 0x0000000080000001,

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