mirror of
https://github.com/hashcat/hashcat
synced 2024-12-27 05:13:45 +01:00
Merge remote-tracking branch 'upstream/master'
This commit is contained in:
Sein Coray
commit
e39a9284e1
@ -6,6 +6,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_rp.h"
|
||||
#include "inc_rp.cl"
|
||||
|
||||
@ -6,6 +6,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#endif
|
||||
|
||||
|
||||
@ -6,6 +6,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#endif
|
||||
|
||||
KERNEL_FQ void amp (GLOBAL_AS pw_t *pws, GLOBAL_AS pw_t *pws_amp, GLOBAL_AS const kernel_rule_t *rules_buf, GLOBAL_AS const pw_t *combs_buf, CONSTANT_AS bf_t *bfs_buf, const u32 combs_mode, const u64 gid_max)
|
||||
|
||||
@ -5,10 +5,11 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.h"
|
||||
#include "inc_cipher_aes.h"
|
||||
|
||||
CONSTANT_AS u32a te0[256] =
|
||||
CONSTANT_VK u32a te0[256] =
|
||||
{
|
||||
0xc66363a5, 0xf87c7c84, 0xee777799, 0xf67b7b8d,
|
||||
0xfff2f20d, 0xd66b6bbd, 0xde6f6fb1, 0x91c5c554,
|
||||
@ -76,7 +77,7 @@ CONSTANT_AS u32a te0[256] =
|
||||
0x7bb0b0cb, 0xa85454fc, 0x6dbbbbd6, 0x2c16163a,
|
||||
};
|
||||
|
||||
CONSTANT_AS u32a te1[256] =
|
||||
CONSTANT_VK u32a te1[256] =
|
||||
{
|
||||
0xa5c66363, 0x84f87c7c, 0x99ee7777, 0x8df67b7b,
|
||||
0x0dfff2f2, 0xbdd66b6b, 0xb1de6f6f, 0x5491c5c5,
|
||||
@ -144,7 +145,7 @@ CONSTANT_AS u32a te1[256] =
|
||||
0xcb7bb0b0, 0xfca85454, 0xd66dbbbb, 0x3a2c1616,
|
||||
};
|
||||
|
||||
CONSTANT_AS u32a te2[256] =
|
||||
CONSTANT_VK u32a te2[256] =
|
||||
{
|
||||
0x63a5c663, 0x7c84f87c, 0x7799ee77, 0x7b8df67b,
|
||||
0xf20dfff2, 0x6bbdd66b, 0x6fb1de6f, 0xc55491c5,
|
||||
@ -212,7 +213,7 @@ CONSTANT_AS u32a te2[256] =
|
||||
0xb0cb7bb0, 0x54fca854, 0xbbd66dbb, 0x163a2c16,
|
||||
};
|
||||
|
||||
CONSTANT_AS u32a te3[256] =
|
||||
CONSTANT_VK u32a te3[256] =
|
||||
{
|
||||
0x6363a5c6, 0x7c7c84f8, 0x777799ee, 0x7b7b8df6,
|
||||
0xf2f20dff, 0x6b6bbdd6, 0x6f6fb1de, 0xc5c55491,
|
||||
@ -280,7 +281,7 @@ CONSTANT_AS u32a te3[256] =
|
||||
0xb0b0cb7b, 0x5454fca8, 0xbbbbd66d, 0x16163a2c,
|
||||
};
|
||||
|
||||
CONSTANT_AS u32a te4[256] =
|
||||
CONSTANT_VK u32a te4[256] =
|
||||
{
|
||||
0x63636363, 0x7c7c7c7c, 0x77777777, 0x7b7b7b7b,
|
||||
0xf2f2f2f2, 0x6b6b6b6b, 0x6f6f6f6f, 0xc5c5c5c5,
|
||||
@ -348,7 +349,7 @@ CONSTANT_AS u32a te4[256] =
|
||||
0xb0b0b0b0, 0x54545454, 0xbbbbbbbb, 0x16161616,
|
||||
};
|
||||
|
||||
CONSTANT_AS u32a td0[256] =
|
||||
CONSTANT_VK u32a td0[256] =
|
||||
{
|
||||
0x51f4a750, 0x7e416553, 0x1a17a4c3, 0x3a275e96,
|
||||
0x3bab6bcb, 0x1f9d45f1, 0xacfa58ab, 0x4be30393,
|
||||
@ -416,7 +417,7 @@ CONSTANT_AS u32a td0[256] =
|
||||
0x7bcb8461, 0xd532b670, 0x486c5c74, 0xd0b85742,
|
||||
};
|
||||
|
||||
CONSTANT_AS u32a td1[256] =
|
||||
CONSTANT_VK u32a td1[256] =
|
||||
{
|
||||
0x5051f4a7, 0x537e4165, 0xc31a17a4, 0x963a275e,
|
||||
0xcb3bab6b, 0xf11f9d45, 0xabacfa58, 0x934be303,
|
||||
@ -484,7 +485,7 @@ CONSTANT_AS u32a td1[256] =
|
||||
0x617bcb84, 0x70d532b6, 0x74486c5c, 0x42d0b857,
|
||||
};
|
||||
|
||||
CONSTANT_AS u32a td2[256] =
|
||||
CONSTANT_VK u32a td2[256] =
|
||||
{
|
||||
0xa75051f4, 0x65537e41, 0xa4c31a17, 0x5e963a27,
|
||||
0x6bcb3bab, 0x45f11f9d, 0x58abacfa, 0x03934be3,
|
||||
@ -552,7 +553,7 @@ CONSTANT_AS u32a td2[256] =
|
||||
0x84617bcb, 0xb670d532, 0x5c74486c, 0x5742d0b8,
|
||||
};
|
||||
|
||||
CONSTANT_AS u32a td3[256] =
|
||||
CONSTANT_VK u32a td3[256] =
|
||||
{
|
||||
0xf4a75051, 0x4165537e, 0x17a4c31a, 0x275e963a,
|
||||
0xab6bcb3b, 0x9d45f11f, 0xfa58abac, 0xe303934b,
|
||||
@ -620,7 +621,7 @@ CONSTANT_AS u32a td3[256] =
|
||||
0xcb84617b, 0x32b670d5, 0x6c5c7448, 0xb85742d0,
|
||||
};
|
||||
|
||||
CONSTANT_AS u32a td4[256] =
|
||||
CONSTANT_VK u32a td4[256] =
|
||||
{
|
||||
0x52525252, 0x09090909, 0x6a6a6a6a, 0xd5d5d5d5,
|
||||
0x30303030, 0x36363636, 0xa5a5a5a5, 0x38383838,
|
||||
|
||||
@ -17,10 +17,11 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.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,
|
||||
0xe4, 0x85, 0x57, 0x35, 0xea, 0x0c, 0xae, 0x41,
|
||||
|
||||
@ -5,10 +5,11 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.h"
|
||||
#include "inc_cipher_des.h"
|
||||
|
||||
CONSTANT_AS u32a c_SPtrans[8][64] =
|
||||
CONSTANT_VK u32a c_SPtrans[8][64] =
|
||||
{
|
||||
{
|
||||
/* nibble 0 */
|
||||
@ -164,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,
|
||||
@ -322,6 +323,9 @@ DECLSPEC void _des_crypt_encrypt (u32 *out, const u32 *in, const u32 *Kc, const
|
||||
r = hc_rotl32_S (r, 3u);
|
||||
l = hc_rotl32_S (l, 3u);
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (u32 i = 0; i < 16; i += 2)
|
||||
{
|
||||
u32 u;
|
||||
@ -371,6 +375,9 @@ DECLSPEC void _des_crypt_decrypt (u32 *out, const u32 *in, const u32 *Kc, const
|
||||
r = hc_rotl32_S (r, 3u);
|
||||
l = hc_rotl32_S (l, 3u);
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (u32 i = 16; i > 0; i -= 2)
|
||||
{
|
||||
u32 u;
|
||||
@ -426,6 +433,9 @@ DECLSPEC void _des_crypt_keysetup (u32 c, u32 d, u32 *Kc, u32 *Kd, SHM_TYPE u32
|
||||
|
||||
c = c & 0x0fffffff;
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (u32 i = 0; i < 16; i++)
|
||||
{
|
||||
if ((i < 2) || (i == 8) || (i == 15))
|
||||
@ -487,6 +497,9 @@ DECLSPEC void _des_crypt_encrypt_vect (u32x *out, const u32x *in, const u32x *Kc
|
||||
r = hc_rotl32 (r, 3u);
|
||||
l = hc_rotl32 (l, 3u);
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (u32 i = 0; i < 16; i += 2)
|
||||
{
|
||||
u32x u;
|
||||
@ -536,6 +549,9 @@ DECLSPEC void _des_crypt_decrypt_vect (u32x *out, const u32x *in, const u32x *Kc
|
||||
r = hc_rotl32 (r, 3u);
|
||||
l = hc_rotl32 (l, 3u);
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (u32 i = 16; i > 0; i -= 2)
|
||||
{
|
||||
u32x u;
|
||||
@ -591,6 +607,9 @@ DECLSPEC void _des_crypt_keysetup_vect (u32x c, u32x d, u32x *Kc, u32x *Kd, SHM_
|
||||
|
||||
c = c & 0x0fffffff;
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (u32 i = 0; i < 16; i++)
|
||||
{
|
||||
if ((i < 2) || (i == 8) || (i == 15))
|
||||
|
||||
@ -92,13 +92,13 @@
|
||||
#if VECT_SIZE == 1
|
||||
#define DES_BOX(i,n,S) (S)[(n)][(i)]
|
||||
#elif VECT_SIZE == 2
|
||||
#define DES_BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
|
||||
#define DES_BOX(i,n,S) make_u32x ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
|
||||
#elif VECT_SIZE == 4
|
||||
#define DES_BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
|
||||
#define DES_BOX(i,n,S) make_u32x ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
|
||||
#elif VECT_SIZE == 8
|
||||
#define DES_BOX(i,n,S) (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])
|
||||
#define DES_BOX(i,n,S) make_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])
|
||||
#elif VECT_SIZE == 16
|
||||
#define DES_BOX(i,n,S) (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])
|
||||
#define DES_BOX(i,n,S) make_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
|
||||
|
||||
DECLSPEC void _des_crypt_encrypt (u32 *out, const u32 *in, const u32 *Kc, const u32 *Kd, SHM_TYPE u32 (*s_SPtrans)[64]);
|
||||
|
||||
@ -14,10 +14,11 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.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,
|
||||
0xfb, 0xc4, 0xfa, 0xda, 0x23, 0xc5, 0x04, 0x4d,
|
||||
@ -53,7 +54,7 @@ CONSTANT_AS u32a k_sbox[256] =
|
||||
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,
|
||||
0x54, 0xe6, 0x9e, 0x39, 0x55, 0x7e, 0x52, 0x91,
|
||||
@ -91,12 +92,23 @@ CONSTANT_AS u32a k_sbox_inv[256] =
|
||||
|
||||
#define extract_byte(x,n) (((x) >> (8 * (n))) & 0xff)
|
||||
|
||||
#define k_lookup(w,sbox) \
|
||||
for (int i = 0; i < 4; i++) \
|
||||
w[i] = sbox[extract_byte (w[i], 0)] << 0 \
|
||||
| sbox[extract_byte (w[i], 1)] << 8 \
|
||||
| sbox[extract_byte (w[i], 2)] << 16 \
|
||||
| sbox[extract_byte (w[i], 3)] << 24
|
||||
#define k_lookup(w,sbox) \
|
||||
w[0] = sbox[extract_byte (w[0], 0)] << 0 \
|
||||
| sbox[extract_byte (w[0], 1)] << 8 \
|
||||
| sbox[extract_byte (w[0], 2)] << 16 \
|
||||
| sbox[extract_byte (w[0], 3)] << 24; \
|
||||
w[1] = sbox[extract_byte (w[1], 0)] << 0 \
|
||||
| sbox[extract_byte (w[1], 1)] << 8 \
|
||||
| sbox[extract_byte (w[1], 2)] << 16 \
|
||||
| sbox[extract_byte (w[1], 3)] << 24; \
|
||||
w[2] = sbox[extract_byte (w[2], 0)] << 0 \
|
||||
| sbox[extract_byte (w[2], 1)] << 8 \
|
||||
| sbox[extract_byte (w[2], 2)] << 16 \
|
||||
| sbox[extract_byte (w[2], 3)] << 24; \
|
||||
w[3] = sbox[extract_byte (w[3], 0)] << 0 \
|
||||
| sbox[extract_byte (w[3], 1)] << 8 \
|
||||
| sbox[extract_byte (w[3], 2)] << 16 \
|
||||
| sbox[extract_byte (w[3], 3)] << 24;
|
||||
|
||||
#define k_xor(n) \
|
||||
for (int i = (n); i > 0; i /= 2) \
|
||||
|
||||
@ -18,6 +18,7 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.h"
|
||||
#include "inc_cipher_serpent.h"
|
||||
|
||||
|
||||
@ -21,10 +21,11 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.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,
|
||||
@ -76,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,
|
||||
0xE2E22BFB, 0x9E9EFAC8, 0xC9C9EC4A, 0xD4D409D3, 0x18186BE6, 0x1E1E9F6B,
|
||||
@ -392,6 +393,9 @@ DECLSPEC void twofish128_set_key (u32 *sk, u32 *lk, const u32 *ukey)
|
||||
sk[1] = mds_rem (me_key[0], mo_key[0]);
|
||||
sk[0] = mds_rem (me_key[1], mo_key[1]);
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (int i = 0; i < 40; i += 2)
|
||||
{
|
||||
u32 a = 0x01010101 * i;
|
||||
@ -517,6 +521,9 @@ DECLSPEC void twofish256_set_key (u32 *sk, u32 *lk, const u32 *ukey)
|
||||
sk[1] = mds_rem (me_key[2], mo_key[2]);
|
||||
sk[0] = mds_rem (me_key[3], mo_key[3]);
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (int i = 0; i < 40; i += 2)
|
||||
{
|
||||
u32 a = 0x01010101 * i;
|
||||
|
||||
@ -5,6 +5,7 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.h"
|
||||
|
||||
/**
|
||||
@ -17,7 +18,7 @@ DECLSPEC u8 v8a_from_v32_S (const u32 v32)
|
||||
|
||||
v.v32 = v32;
|
||||
|
||||
return v.v8a;
|
||||
return v.v8.a;
|
||||
}
|
||||
|
||||
DECLSPEC u8 v8b_from_v32_S (const u32 v32)
|
||||
@ -26,7 +27,7 @@ DECLSPEC u8 v8b_from_v32_S (const u32 v32)
|
||||
|
||||
v.v32 = v32;
|
||||
|
||||
return v.v8b;
|
||||
return v.v8.b;
|
||||
}
|
||||
|
||||
DECLSPEC u8 v8c_from_v32_S (const u32 v32)
|
||||
@ -35,7 +36,7 @@ DECLSPEC u8 v8c_from_v32_S (const u32 v32)
|
||||
|
||||
v.v32 = v32;
|
||||
|
||||
return v.v8c;
|
||||
return v.v8.c;
|
||||
}
|
||||
|
||||
DECLSPEC u8 v8d_from_v32_S (const u32 v32)
|
||||
@ -44,7 +45,7 @@ DECLSPEC u8 v8d_from_v32_S (const u32 v32)
|
||||
|
||||
v.v32 = v32;
|
||||
|
||||
return v.v8d;
|
||||
return v.v8.d;
|
||||
}
|
||||
|
||||
DECLSPEC u16 v16a_from_v32_S (const u32 v32)
|
||||
@ -53,7 +54,7 @@ DECLSPEC u16 v16a_from_v32_S (const u32 v32)
|
||||
|
||||
v.v32 = v32;
|
||||
|
||||
return v.v16a;
|
||||
return v.v16.a;
|
||||
}
|
||||
|
||||
DECLSPEC u16 v16b_from_v32_S (const u32 v32)
|
||||
@ -62,15 +63,15 @@ DECLSPEC u16 v16b_from_v32_S (const u32 v32)
|
||||
|
||||
v.v32 = v32;
|
||||
|
||||
return v.v16b;
|
||||
return v.v16.b;
|
||||
}
|
||||
|
||||
DECLSPEC u32 v32_from_v16ab_S (const u16 v16a, const u16 v16b)
|
||||
{
|
||||
vconv32_t v;
|
||||
|
||||
v.v16a = v16a;
|
||||
v.v16b = v16b;
|
||||
v.v16.a = v16a;
|
||||
v.v16.b = v16b;
|
||||
|
||||
return v.v32;
|
||||
}
|
||||
@ -81,7 +82,7 @@ DECLSPEC u32 v32a_from_v64_S (const u64 v64)
|
||||
|
||||
v.v64 = v64;
|
||||
|
||||
return v.v32a;
|
||||
return v.v32.a;
|
||||
}
|
||||
|
||||
DECLSPEC u32 v32b_from_v64_S (const u64 v64)
|
||||
@ -90,15 +91,15 @@ DECLSPEC u32 v32b_from_v64_S (const u64 v64)
|
||||
|
||||
v.v64 = v64;
|
||||
|
||||
return v.v32b;
|
||||
return v.v32.b;
|
||||
}
|
||||
|
||||
DECLSPEC u64 v64_from_v32ab_S (const u32 v32a, const u32 v32b)
|
||||
{
|
||||
vconv64_t v;
|
||||
|
||||
v.v32a = v32a;
|
||||
v.v32b = v32b;
|
||||
v.v32.a = v32a;
|
||||
v.v32.b = v32b;
|
||||
|
||||
return v.v64;
|
||||
}
|
||||
@ -304,26 +305,32 @@ DECLSPEC u64x hl32_to_64 (const u32x a, const u32x b)
|
||||
|
||||
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);
|
||||
#else
|
||||
return rotate (a, (u32x) (n));
|
||||
return rotate (a, make_u32x (n));
|
||||
#endif
|
||||
}
|
||||
|
||||
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);
|
||||
#else
|
||||
return rotate (a, (u32x) (32 - n));
|
||||
return rotate (a, make_u32x (32 - n));
|
||||
#endif
|
||||
}
|
||||
|
||||
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);
|
||||
#elif defined IS_CUDA
|
||||
return rotl32_S (a, n);
|
||||
#else
|
||||
return rotate (a, (u32) (n));
|
||||
#endif
|
||||
@ -331,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)
|
||||
{
|
||||
#ifdef _CPU_OPENCL_EMU_H
|
||||
#if defined _CPU_OPENCL_EMU_H
|
||||
return rotr32 (a, n);
|
||||
#elif defined IS_CUDA
|
||||
return rotr32_S (a, n);
|
||||
#else
|
||||
return rotate (a, (u32) (32 - n));
|
||||
#endif
|
||||
@ -340,26 +349,32 @@ DECLSPEC u32 hc_rotr32_S (const u32 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);
|
||||
#else
|
||||
return rotate (a, (u64x) (n));
|
||||
return rotate (a, make_u64x (n));
|
||||
#endif
|
||||
}
|
||||
|
||||
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);
|
||||
#else
|
||||
return rotate (a, (u64x) (64 - n));
|
||||
return rotate (a, make_u64x (64 - n));
|
||||
#endif
|
||||
}
|
||||
|
||||
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);
|
||||
#elif defined IS_CUDA
|
||||
return rotl64_S (a, n);
|
||||
#else
|
||||
return rotate (a, (u64) (n));
|
||||
#endif
|
||||
@ -367,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)
|
||||
{
|
||||
#ifdef _CPU_OPENCL_EMU_H
|
||||
#if defined _CPU_OPENCL_EMU_H
|
||||
return rotr64 (a, n);
|
||||
#elif defined IS_CUDA
|
||||
return rotr64_S (a, n);
|
||||
#else
|
||||
return rotate (a, (u64) (64 - n));
|
||||
#endif
|
||||
@ -454,9 +471,9 @@ DECLSPEC u32x hc_swap32 (const u32x v)
|
||||
#endif
|
||||
|
||||
#else
|
||||
r = bitselect (rotate (v, (u32x) (24)),
|
||||
rotate (v, (u32x) ( 8)),
|
||||
(u32x) (0x00ff00ff));
|
||||
r = bitselect (rotate (v, make_u32x (24)),
|
||||
rotate (v, make_u32x ( 8)),
|
||||
make_u32x (0x00ff00ff));
|
||||
#endif
|
||||
#endif
|
||||
|
||||
@ -672,13 +689,13 @@ DECLSPEC u64x hc_swap64 (const u64x v)
|
||||
#endif
|
||||
|
||||
#else
|
||||
r = bitselect (bitselect (rotate (v, (u64x) (24)),
|
||||
rotate (v, (u64x) ( 8)),
|
||||
(u64x) (0x000000ff000000ff)),
|
||||
bitselect (rotate (v, (u64x) (56)),
|
||||
rotate (v, (u64x) (40)),
|
||||
(u64x) (0x00ff000000ff0000)),
|
||||
(u64x) (0xffff0000ffff0000));
|
||||
r = bitselect (bitselect (rotate (v, make_u64x (24)),
|
||||
rotate (v, make_u64x ( 8)),
|
||||
make_u64x (0x000000ff000000ff)),
|
||||
bitselect (rotate (v, make_u64x (56)),
|
||||
rotate (v, make_u64x (40)),
|
||||
make_u64x (0x00ff000000ff0000)),
|
||||
make_u64x (0xffff0000ffff0000));
|
||||
#endif
|
||||
#endif
|
||||
|
||||
@ -730,7 +747,7 @@ DECLSPEC u64 hc_swap64_S (const u64 v)
|
||||
|
||||
DECLSPEC u32x hc_bfe (const u32x a, const u32x b, const u32x c)
|
||||
{
|
||||
#define BIT(x) ((u32x) (1u) << (x))
|
||||
#define BIT(x) (make_u32x (1u) << (x))
|
||||
#define BIT_MASK(x) (BIT (x) - 1)
|
||||
#define BFE(x,y,z) (((x) >> (y)) & BIT_MASK (z))
|
||||
|
||||
@ -1164,7 +1181,7 @@ DECLSPEC u32 hc_lop_0x96_S (const u32 a, const u32 b, const u32 c)
|
||||
|
||||
DECLSPEC u32x hc_bfe (const u32x a, const u32x b, const u32x c)
|
||||
{
|
||||
#define BIT(x) ((u32x) (1u) << (x))
|
||||
#define BIT(x) (make_u32x (1u) << (x))
|
||||
#define BIT_MASK(x) (BIT (x) - 1)
|
||||
#define BFE(x,y,z) (((x) >> (y)) & BIT_MASK (z))
|
||||
|
||||
@ -1415,8 +1432,8 @@ DECLSPEC int is_valid_hex_8 (const u8 v)
|
||||
{
|
||||
// direct lookup table is slower thanks to CMOV
|
||||
|
||||
if ((v >= '0') && (v <= '9')) return 1;
|
||||
if ((v >= 'a') && (v <= 'f')) return 1;
|
||||
if ((v >= (u8) '0') && (v <= (u8) '9')) return 1;
|
||||
if ((v >= (u8) 'a') && (v <= (u8) 'f')) return 1;
|
||||
|
||||
return 0;
|
||||
}
|
||||
@ -1433,10 +1450,10 @@ DECLSPEC int is_valid_hex_32 (const u32 v)
|
||||
|
||||
DECLSPEC int is_valid_base58_8 (const u8 v)
|
||||
{
|
||||
if (v > 'z') return 0;
|
||||
if (v < '1') return 0;
|
||||
if ((v > '9') && (v < 'A')) return 0;
|
||||
if ((v > 'Z') && (v < 'a')) return 0;
|
||||
if (v > (u8) 'z') return 0;
|
||||
if (v < (u8) '1') return 0;
|
||||
if ((v > (u8) '9') && (v < (u8) 'A')) return 0;
|
||||
if ((v > (u8) 'Z') && (v < (u8) 'a')) return 0;
|
||||
|
||||
return 1;
|
||||
}
|
||||
@ -60860,7 +60877,23 @@ KERNEL_FQ void gpu_memset (GLOBAL_AS uint4 *buf, const u32 value, const u64 gid_
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
buf[gid] = (uint4) (value);
|
||||
uint4 r;
|
||||
|
||||
#if defined IS_NATIVE
|
||||
r = value;
|
||||
#elif defined IS_OPENCL
|
||||
r.s0 = value;
|
||||
r.s1 = value;
|
||||
r.s2 = value;
|
||||
r.s3 = value;
|
||||
#elif defined IS_CUDA
|
||||
r.x = value;
|
||||
r.y = value;
|
||||
r.z = value;
|
||||
r.w = value;
|
||||
#endif
|
||||
|
||||
buf[gid] = r;
|
||||
}
|
||||
|
||||
KERNEL_FQ void gpu_atinit (GLOBAL_AS pw_t *buf, const u64 gid_max)
|
||||
|
||||
@ -26,6 +26,44 @@
|
||||
* - P19: Type of the esalt_bufs structure with additional data, or void.
|
||||
*/
|
||||
|
||||
#ifdef IS_CUDA
|
||||
#define KERN_ATTR(p2,p4,p5,p6,p19) \
|
||||
MAYBE_UNUSED GLOBAL_AS pw_t *pws, \
|
||||
MAYBE_UNUSED p2 const kernel_rule_t *g_rules_buf, \
|
||||
MAYBE_UNUSED GLOBAL_AS const pw_t *combs_buf, \
|
||||
MAYBE_UNUSED p4, \
|
||||
MAYBE_UNUSED GLOBAL_AS p5 *tmps, \
|
||||
MAYBE_UNUSED GLOBAL_AS p6 *hooks, \
|
||||
MAYBE_UNUSED GLOBAL_AS const u32 *bitmaps_buf_s1_a, \
|
||||
MAYBE_UNUSED GLOBAL_AS const u32 *bitmaps_buf_s1_b, \
|
||||
MAYBE_UNUSED GLOBAL_AS const u32 *bitmaps_buf_s1_c, \
|
||||
MAYBE_UNUSED GLOBAL_AS const u32 *bitmaps_buf_s1_d, \
|
||||
MAYBE_UNUSED GLOBAL_AS const u32 *bitmaps_buf_s2_a, \
|
||||
MAYBE_UNUSED GLOBAL_AS const u32 *bitmaps_buf_s2_b, \
|
||||
MAYBE_UNUSED GLOBAL_AS const u32 *bitmaps_buf_s2_c, \
|
||||
MAYBE_UNUSED GLOBAL_AS const u32 *bitmaps_buf_s2_d, \
|
||||
MAYBE_UNUSED GLOBAL_AS plain_t *plains_buf, \
|
||||
MAYBE_UNUSED GLOBAL_AS const digest_t *digests_buf, \
|
||||
MAYBE_UNUSED GLOBAL_AS u32 *hashes_shown, \
|
||||
MAYBE_UNUSED GLOBAL_AS const salt_t *salt_bufs, \
|
||||
MAYBE_UNUSED GLOBAL_AS const p19 *esalt_bufs, \
|
||||
MAYBE_UNUSED GLOBAL_AS u32 *d_return_buf, \
|
||||
MAYBE_UNUSED GLOBAL_AS void *d_extra0_buf, \
|
||||
MAYBE_UNUSED GLOBAL_AS void *d_extra1_buf, \
|
||||
MAYBE_UNUSED GLOBAL_AS void *d_extra2_buf, \
|
||||
MAYBE_UNUSED GLOBAL_AS void *d_extra3_buf, \
|
||||
MAYBE_UNUSED const u32 bitmap_mask, \
|
||||
MAYBE_UNUSED const u32 bitmap_shift1, \
|
||||
MAYBE_UNUSED const u32 bitmap_shift2, \
|
||||
MAYBE_UNUSED const u32 salt_pos, \
|
||||
MAYBE_UNUSED const u32 loop_pos, \
|
||||
MAYBE_UNUSED const u32 loop_cnt, \
|
||||
MAYBE_UNUSED const u32 il_cnt, \
|
||||
MAYBE_UNUSED const u32 digests_cnt, \
|
||||
MAYBE_UNUSED const u32 digests_offset, \
|
||||
MAYBE_UNUSED const u32 combs_mode, \
|
||||
MAYBE_UNUSED const u64 gid_max
|
||||
#else
|
||||
#define KERN_ATTR(p2,p4,p5,p6,p19) \
|
||||
MAYBE_UNUSED GLOBAL_AS pw_t *pws, \
|
||||
MAYBE_UNUSED p2 const kernel_rule_t *rules_buf, \
|
||||
@ -62,7 +100,7 @@
|
||||
MAYBE_UNUSED const u32 digests_offset, \
|
||||
MAYBE_UNUSED const u32 combs_mode, \
|
||||
MAYBE_UNUSED const u64 gid_max
|
||||
|
||||
#endif
|
||||
/*
|
||||
* Shortcut macros for usage in the actual kernels
|
||||
*
|
||||
@ -71,16 +109,29 @@
|
||||
* do not use rules or tmps, etc.
|
||||
*/
|
||||
|
||||
#define KERN_ATTR_BASIC() KERN_ATTR (GLOBAL_AS, GLOBAL_AS const bf_t *bfs_buf, void, void, void)
|
||||
#define KERN_ATTR_BITSLICE() KERN_ATTR (GLOBAL_AS, CONSTANT_AS const bs_word_t *words_buf_r, void, void, void)
|
||||
#define KERN_ATTR_ESALT(e) KERN_ATTR (GLOBAL_AS, GLOBAL_AS const bf_t *bfs_buf, void, void, e)
|
||||
#define KERN_ATTR_RULES() KERN_ATTR (CONSTANT_AS, GLOBAL_AS const bf_t *bfs_buf, void, void, void)
|
||||
#define KERN_ATTR_RULES_ESALT(e) KERN_ATTR (CONSTANT_AS, GLOBAL_AS const bf_t *bfs_buf, void, void, e)
|
||||
#define KERN_ATTR_TMPS(t) KERN_ATTR (GLOBAL_AS, GLOBAL_AS const bf_t *bfs_buf, t, void, void)
|
||||
#define KERN_ATTR_TMPS_ESALT(t,e) KERN_ATTR (GLOBAL_AS, GLOBAL_AS const bf_t *bfs_buf, t, void, e)
|
||||
#define KERN_ATTR_TMPS_HOOKS(t,h) KERN_ATTR (GLOBAL_AS, GLOBAL_AS const bf_t *bfs_buf, t, h, void)
|
||||
#define KERN_ATTR_VECTOR() KERN_ATTR (GLOBAL_AS, CONSTANT_AS const u32x *words_buf_r, void, void, void)
|
||||
#define KERN_ATTR_VECTOR_ESALT(e) KERN_ATTR (GLOBAL_AS, CONSTANT_AS const u32x *words_buf_r, void, void, e)
|
||||
#ifdef IS_CUDA
|
||||
#define KERN_ATTR_BASIC() KERN_ATTR (GLOBAL_AS, GLOBAL_AS const bf_t *g_bfs_buf, void, void, void)
|
||||
#define KERN_ATTR_BITSLICE() KERN_ATTR (GLOBAL_AS, GLOBAL_AS const bs_word_t *g_words_buf_s, void, void, void)
|
||||
#define KERN_ATTR_ESALT(e) KERN_ATTR (GLOBAL_AS, GLOBAL_AS const bf_t *g_bfs_buf, void, void, e)
|
||||
#define KERN_ATTR_RULES() KERN_ATTR (GLOBAL_AS, GLOBAL_AS const bf_t *g_bfs_buf, void, void, void)
|
||||
#define KERN_ATTR_RULES_ESALT(e) KERN_ATTR (GLOBAL_AS, GLOBAL_AS const bf_t *g_bfs_buf, void, void, e)
|
||||
#define KERN_ATTR_TMPS(t) KERN_ATTR (GLOBAL_AS, GLOBAL_AS const bf_t *g_bfs_buf, t, void, void)
|
||||
#define KERN_ATTR_TMPS_ESALT(t,e) KERN_ATTR (GLOBAL_AS, GLOBAL_AS const bf_t *g_bfs_buf, t, void, e)
|
||||
#define KERN_ATTR_TMPS_HOOKS(t,h) KERN_ATTR (GLOBAL_AS, GLOBAL_AS const bf_t *g_bfs_buf, t, h, void)
|
||||
#define KERN_ATTR_VECTOR() KERN_ATTR (GLOBAL_AS, GLOBAL_AS const u32x *g_words_buf_r, void, void, void)
|
||||
#define KERN_ATTR_VECTOR_ESALT(e) KERN_ATTR (GLOBAL_AS, GLOBAL_AS const u32x *g_words_buf_r, void, void, e)
|
||||
#else
|
||||
#define KERN_ATTR_BASIC() KERN_ATTR (GLOBAL_AS, CONSTANT_AS const bf_t *bfs_buf, void, void, void)
|
||||
#define KERN_ATTR_BITSLICE() KERN_ATTR (GLOBAL_AS, CONSTANT_AS const bs_word_t *words_buf_s, void, void, void)
|
||||
#define KERN_ATTR_ESALT(e) KERN_ATTR (GLOBAL_AS, CONSTANT_AS const bf_t *bfs_buf, void, void, e)
|
||||
#define KERN_ATTR_RULES() KERN_ATTR (CONSTANT_AS, GLOBAL_AS const bf_t *bfs_buf, void, void, void)
|
||||
#define KERN_ATTR_RULES_ESALT(e) KERN_ATTR (CONSTANT_AS, GLOBAL_AS const bf_t *bfs_buf, void, void, e)
|
||||
#define KERN_ATTR_TMPS(t) KERN_ATTR (GLOBAL_AS, CONSTANT_AS const bf_t *bfs_buf, t, void, void)
|
||||
#define KERN_ATTR_TMPS_ESALT(t,e) KERN_ATTR (GLOBAL_AS, CONSTANT_AS const bf_t *bfs_buf, t, void, e)
|
||||
#define KERN_ATTR_TMPS_HOOKS(t,h) KERN_ATTR (GLOBAL_AS, CONSTANT_AS const bf_t *bfs_buf, t, h, void)
|
||||
#define KERN_ATTR_VECTOR() KERN_ATTR (GLOBAL_AS, CONSTANT_AS const u32x *words_buf_r, void, void, void)
|
||||
#define KERN_ATTR_VECTOR_ESALT(e) KERN_ATTR (GLOBAL_AS, CONSTANT_AS const u32x *words_buf_r, void, void, e)
|
||||
#endif
|
||||
|
||||
// union based packing
|
||||
|
||||
|
||||
@ -1,3 +1,17 @@
|
||||
/**
|
||||
* Author......: See docs/credits.txt
|
||||
* License.....: MIT
|
||||
*/
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.h"
|
||||
#include "inc_cipher_aes.h"
|
||||
#include "inc_cipher_serpent.h"
|
||||
#include "inc_cipher_twofish.h"
|
||||
#include "inc_diskcryptor_xts.h"
|
||||
|
||||
DECLSPEC void dcrp_xts_mul2 (u32 *in, u32 *out)
|
||||
{
|
||||
const u32 c = in[3] >> 31;
|
||||
|
||||
@ -5,6 +5,7 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.h"
|
||||
#include "inc_hash_md4.h"
|
||||
|
||||
|
||||
@ -5,6 +5,7 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.h"
|
||||
#include "inc_hash_md5.h"
|
||||
|
||||
|
||||
@ -5,6 +5,7 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.h"
|
||||
#include "inc_hash_ripemd160.h"
|
||||
|
||||
|
||||
@ -5,6 +5,7 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.h"
|
||||
#include "inc_hash_sha1.h"
|
||||
|
||||
|
||||
@ -5,10 +5,11 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.h"
|
||||
#include "inc_hash_sha224.h"
|
||||
|
||||
CONSTANT_AS u32a k_sha224[64] =
|
||||
CONSTANT_VK u32a k_sha224[64] =
|
||||
{
|
||||
SHA224C00, SHA224C01, SHA224C02, SHA224C03,
|
||||
SHA224C04, SHA224C05, SHA224C06, SHA224C07,
|
||||
|
||||
@ -5,10 +5,11 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.h"
|
||||
#include "inc_hash_sha256.h"
|
||||
|
||||
CONSTANT_AS u32a k_sha256[64] =
|
||||
CONSTANT_VK u32a k_sha256[64] =
|
||||
{
|
||||
SHA256C00, SHA256C01, SHA256C02, SHA256C03,
|
||||
SHA256C04, SHA256C05, SHA256C06, SHA256C07,
|
||||
|
||||
@ -5,10 +5,11 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.h"
|
||||
#include "inc_hash_sha384.h"
|
||||
|
||||
CONSTANT_AS u64a k_sha384[80] =
|
||||
CONSTANT_VK u64a k_sha384[80] =
|
||||
{
|
||||
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
|
||||
SHA512C04, SHA512C05, SHA512C06, SHA512C07,
|
||||
|
||||
@ -5,10 +5,11 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.h"
|
||||
#include "inc_hash_sha512.h"
|
||||
|
||||
CONSTANT_AS u64a k_sha512[80] =
|
||||
CONSTANT_VK u64a k_sha512[80] =
|
||||
{
|
||||
SHA512C00, SHA512C01, SHA512C02, SHA512C03,
|
||||
SHA512C04, SHA512C05, SHA512C06, SHA512C07,
|
||||
|
||||
@ -5,10 +5,11 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.h"
|
||||
#include "inc_hash_streebog256.h"
|
||||
|
||||
CONSTANT_AS u64a sbob256_sl64[8][256] =
|
||||
CONSTANT_VK u64a sbob256_sl64[8][256] =
|
||||
{
|
||||
{
|
||||
0xd031c397ce553fe6, 0x16ba5b01b006b525, 0xa89bade6296e70c8, 0x6a1f525d77d3435b,
|
||||
@ -540,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,
|
||||
@ -1362,7 +1363,7 @@ DECLSPEC void streebog256_add_vector (u64x *x, const u64x *y)
|
||||
const u64x right = hc_swap64 (y[i]);
|
||||
const u64x sum = left + right + carry;
|
||||
|
||||
carry = (sum < left) ? (u64x) 1 : (u64x) 0;
|
||||
carry = (sum < left) ? make_u64x (1) : make_u64x (0);
|
||||
|
||||
x[i] = hc_swap64 (sum);
|
||||
}
|
||||
@ -1710,7 +1711,7 @@ DECLSPEC void streebog256_final_vector (streebog256_ctx_vector_t *ctx)
|
||||
streebog256_g_vector (ctx->h, ctx->n, m, ctx->s_sbob_sl64);
|
||||
|
||||
u64x sizebuf[8] = { 0 };
|
||||
sizebuf[7] = hc_swap64 ((u64x) (pos << 3));
|
||||
sizebuf[7] = hc_swap64 (make_u64x (pos << 3));
|
||||
|
||||
streebog256_add_vector (ctx->n, sizebuf);
|
||||
|
||||
|
||||
@ -10,17 +10,17 @@
|
||||
#define BOX(S,n,i) ((S)[(n)][(i)])
|
||||
|
||||
#elif VECT_SIZE == 2
|
||||
#define BOX(S,n,i) (u64x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
|
||||
#define BOX(S,n,i) make_u64x ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
|
||||
|
||||
#elif VECT_SIZE == 4
|
||||
#define BOX(S,n,i) (u64x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
|
||||
#define BOX(S,n,i) make_u64x ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
|
||||
|
||||
#elif VECT_SIZE == 8
|
||||
#define BOX(S,n,i) (u64x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], \
|
||||
#define BOX(S,n,i) make_u64x ((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])
|
||||
|
||||
#elif VECT_SIZE == 16
|
||||
#define BOX(S,n,i) (u64x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], \
|
||||
#define BOX(S,n,i) make_u64x ((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])
|
||||
|
||||
@ -5,10 +5,11 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.h"
|
||||
#include "inc_hash_streebog512.h"
|
||||
|
||||
CONSTANT_AS u64a sbob512_sl64[8][256] =
|
||||
CONSTANT_VK u64a sbob512_sl64[8][256] =
|
||||
{
|
||||
{
|
||||
0xd031c397ce553fe6, 0x16ba5b01b006b525, 0xa89bade6296e70c8, 0x6a1f525d77d3435b,
|
||||
@ -540,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,
|
||||
@ -1381,7 +1382,7 @@ DECLSPEC void streebog512_add_vector (u64x *x, const u64x *y)
|
||||
const u64x right = hc_swap64 (y[i]);
|
||||
const u64x sum = left + right + carry;
|
||||
|
||||
carry = (sum < left) ? (u64x) 1 : (u64x) 0;
|
||||
carry = (sum < left) ? make_u64x (1) : make_u64x (0);
|
||||
|
||||
x[i] = hc_swap64 (sum);
|
||||
}
|
||||
@ -1729,7 +1730,7 @@ DECLSPEC void streebog512_final_vector (streebog512_ctx_vector_t *ctx)
|
||||
streebog512_g_vector (ctx->h, ctx->n, m, ctx->s_sbob_sl64);
|
||||
|
||||
u64x sizebuf[8] = { 0 };
|
||||
sizebuf[7] = hc_swap64 ((u64x) (pos << 3));
|
||||
sizebuf[7] = hc_swap64 (make_u64x (pos << 3));
|
||||
|
||||
streebog512_add_vector (ctx->n, sizebuf);
|
||||
|
||||
|
||||
@ -10,17 +10,17 @@
|
||||
#define BOX(S,n,i) ((S)[(n)][(i)])
|
||||
|
||||
#elif VECT_SIZE == 2
|
||||
#define BOX(S,n,i) (u64x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
|
||||
#define BOX(S,n,i) make_u64x ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
|
||||
|
||||
#elif VECT_SIZE == 4
|
||||
#define BOX(S,n,i) (u64x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
|
||||
#define BOX(S,n,i) make_u64x ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
|
||||
|
||||
#elif VECT_SIZE == 8
|
||||
#define BOX(S,n,i) (u64x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], \
|
||||
#define BOX(S,n,i) make_u64x ((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])
|
||||
|
||||
#elif VECT_SIZE == 16
|
||||
#define BOX(S,n,i) (u64x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], \
|
||||
#define BOX(S,n,i) make_u64x ((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])
|
||||
|
||||
@ -5,10 +5,11 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.h"
|
||||
#include "inc_hash_whirlpool.h"
|
||||
|
||||
CONSTANT_AS u32a Ch[8][256] =
|
||||
CONSTANT_VK u32a Ch[8][256] =
|
||||
{
|
||||
{
|
||||
0x18186018, 0x23238c23, 0xc6c63fc6, 0xe8e887e8,
|
||||
@ -540,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,
|
||||
@ -1072,7 +1073,7 @@ CONSTANT_AS u32a Cl[8][256] =
|
||||
},
|
||||
};
|
||||
|
||||
CONSTANT_AS u32a rch[R + 1] =
|
||||
CONSTANT_VK u32a rch[R + 1] =
|
||||
{
|
||||
0x00000000,
|
||||
0x1823c6e8,
|
||||
@ -1087,7 +1088,7 @@ CONSTANT_AS u32a rch[R + 1] =
|
||||
0xca2dbf07,
|
||||
};
|
||||
|
||||
CONSTANT_AS u32a rcl[R + 1] =
|
||||
CONSTANT_VK u32a rcl[R + 1] =
|
||||
{
|
||||
0x00000000,
|
||||
0x87b8014f,
|
||||
|
||||
@ -11,13 +11,13 @@
|
||||
#if VECT_SIZE == 1
|
||||
#define BOX(S,n,i) (S)[(n)][(i)]
|
||||
#elif VECT_SIZE == 2
|
||||
#define BOX(S,n,i) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
|
||||
#define BOX(S,n,i) make_u32x ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
|
||||
#elif VECT_SIZE == 4
|
||||
#define BOX(S,n,i) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
|
||||
#define BOX(S,n,i) make_u32x ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
|
||||
#elif VECT_SIZE == 8
|
||||
#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])
|
||||
#define BOX(S,n,i) make_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])
|
||||
#elif VECT_SIZE == 16
|
||||
#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])
|
||||
#define BOX(S,n,i) make_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
|
||||
|
||||
#define BOX_S(S,n,i) (S)[(n)][(i)]
|
||||
|
||||
@ -5,6 +5,7 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.h"
|
||||
#include "inc_cipher_aes.h"
|
||||
#include "inc_luks_af.h"
|
||||
|
||||
@ -5,6 +5,7 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.h"
|
||||
#include "inc_hash_ripemd160.h"
|
||||
#include "inc_hash_sha1.h"
|
||||
|
||||
@ -5,6 +5,7 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.h"
|
||||
#include "inc_hash_sha256.h"
|
||||
#include "inc_luks_essiv.h"
|
||||
|
||||
@ -5,6 +5,7 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.h"
|
||||
#include "inc_luks_serpent.h"
|
||||
|
||||
|
||||
@ -5,6 +5,7 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.h"
|
||||
#include "inc_luks_twofish.h"
|
||||
|
||||
|
||||
@ -5,6 +5,7 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.h"
|
||||
#include "inc_luks_xts.h"
|
||||
|
||||
|
||||
115
OpenCL/inc_platform.cl
Normal file
115
OpenCL/inc_platform.cl
Normal file
@ -0,0 +1,115 @@
|
||||
/**
|
||||
* Author......: See docs/credits.txt
|
||||
* License.....: MIT
|
||||
*/
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
|
||||
#ifdef IS_NATIVE
|
||||
#define SYNC_THREADS()
|
||||
#endif
|
||||
|
||||
#ifdef IS_CUDA
|
||||
|
||||
#if ATTACK_EXEC == 11
|
||||
|
||||
CONSTANT_VK u32 generic_constant[8192]; // 32k
|
||||
|
||||
#if ATTACK_KERN == 0
|
||||
#define bfs_buf g_bfs_buf
|
||||
#define rules_buf ((const kernel_rule_t *) generic_constant)
|
||||
#define words_buf_s g_words_buf_s
|
||||
#define words_buf_r g_words_buf_r
|
||||
#elif ATTACK_KERN == 1
|
||||
#define bfs_buf g_bfs_buf
|
||||
#define rules_buf g_rules_buf
|
||||
#define words_buf_s g_words_buf_s
|
||||
#define words_buf_r g_words_buf_r
|
||||
#elif ATTACK_KERN == 3
|
||||
#define rules_buf g_rules_buf
|
||||
#define bfs_buf ((const bf_t *) generic_constant)
|
||||
#define words_buf_s ((const bs_word_t *) generic_constant)
|
||||
#define words_buf_r ((const u32x *) generic_constant)
|
||||
#endif
|
||||
|
||||
#endif
|
||||
|
||||
DECLSPEC u32 atomic_dec (u32 *p)
|
||||
{
|
||||
return atomicSub (p, 1);
|
||||
}
|
||||
|
||||
DECLSPEC u32 atomic_inc (u32 *p)
|
||||
{
|
||||
return atomicAdd (p, 1);
|
||||
}
|
||||
|
||||
DECLSPEC u32 atomic_or (u32 *p, u32 val)
|
||||
{
|
||||
return atomicOr (p, val);
|
||||
}
|
||||
|
||||
DECLSPEC size_t get_global_id (const u32 dimindx __attribute__((unused)))
|
||||
{
|
||||
return (blockIdx.x * blockDim.x) + threadIdx.x;
|
||||
}
|
||||
|
||||
DECLSPEC size_t get_local_id (const u32 dimindx __attribute__((unused)))
|
||||
{
|
||||
return threadIdx.x;
|
||||
}
|
||||
|
||||
DECLSPEC size_t get_local_size (const u32 dimindx __attribute__((unused)))
|
||||
{
|
||||
// verify
|
||||
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 ()
|
||||
#endif
|
||||
|
||||
#ifdef IS_OPENCL
|
||||
#define SYNC_THREADS() barrier (CLK_LOCAL_MEM_FENCE)
|
||||
#endif
|
||||
30
OpenCL/inc_platform.h
Normal file
30
OpenCL/inc_platform.h
Normal file
@ -0,0 +1,30 @@
|
||||
/**
|
||||
* Author......: See docs/credits.txt
|
||||
* License.....: MIT
|
||||
*/
|
||||
|
||||
#ifndef _INC_PLATFORM_H
|
||||
#define _INC_PLATFORM_H
|
||||
|
||||
#ifdef IS_CUDA
|
||||
DECLSPEC u32 atomic_dec (u32 *p);
|
||||
DECLSPEC u32 atomic_inc (u32 *p);
|
||||
DECLSPEC u32 atomic_or (u32 *p, u32 val);
|
||||
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_size (const u32 dimindx __attribute__((unused)));
|
||||
|
||||
DECLSPEC u32x rotl32 (const u32x a, const int n);
|
||||
DECLSPEC u32x rotr32 (const u32x a, const int n);
|
||||
DECLSPEC u32 rotl32_S (const u32 a, const int n);
|
||||
DECLSPEC u32 rotr32_S (const u32 a, const int n);
|
||||
DECLSPEC u64x rotl64 (const u64x a, const int n);
|
||||
DECLSPEC u64x rotr64 (const u64x a, const int n);
|
||||
DECLSPEC u64 rotl64_S (const u64 a, const int n);
|
||||
DECLSPEC u64 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))))
|
||||
#endif
|
||||
|
||||
#endif // _INC_PLATFORM_H
|
||||
@ -5,6 +5,7 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.h"
|
||||
#include "inc_rp.h"
|
||||
|
||||
@ -14,7 +15,7 @@
|
||||
|
||||
#ifdef REAL_SHM
|
||||
#define COPY_PW(x) \
|
||||
LOCAL_AS pw_t s_pws[64]; \
|
||||
LOCAL_VK pw_t s_pws[64]; \
|
||||
s_pws[get_local_id (0)] = (x);
|
||||
#else
|
||||
#define COPY_PW(x) \
|
||||
|
||||
@ -5,6 +5,7 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.h"
|
||||
#include "inc_rp_optimized.h"
|
||||
|
||||
@ -2354,7 +2355,7 @@ DECLSPEC u32 apply_rules_optimized (CONSTANT_AS const u32 *cmds, u32 *buf0, u32
|
||||
return out_len;
|
||||
}
|
||||
|
||||
DECLSPEC u32x apply_rules_vect_optimized (const u32 *pw_buf0, const u32 *pw_buf1, const u32 pw_len, CONSTANT_AS const kernel_rule_t *rules_buf, const u32 il_pos, u32x *buf0, u32x *buf1)
|
||||
DECLSPEC u32x apply_rules_vect_optimized (const u32 *pw_buf0, const u32 *pw_buf1, const u32 pw_len, CONSTANT_AS const kernel_rule_t *kernel_rules, const u32 il_pos, u32x *buf0, u32x *buf1)
|
||||
{
|
||||
#if VECT_SIZE == 1
|
||||
|
||||
@ -2367,7 +2368,7 @@ DECLSPEC u32x apply_rules_vect_optimized (const u32 *pw_buf0, const u32 *pw_buf1
|
||||
buf1[2] = pw_buf1[2];
|
||||
buf1[3] = pw_buf1[3];
|
||||
|
||||
return apply_rules_optimized (rules_buf[il_pos].cmds, buf0, buf1, pw_len);
|
||||
return apply_rules_optimized (kernel_rules[il_pos].cmds, buf0, buf1, pw_len);
|
||||
|
||||
#else
|
||||
|
||||
@ -2390,7 +2391,7 @@ DECLSPEC u32x apply_rules_vect_optimized (const u32 *pw_buf0, const u32 *pw_buf1
|
||||
tmp1[2] = pw_buf1[2];
|
||||
tmp1[3] = pw_buf1[3];
|
||||
|
||||
const u32 tmp_len = apply_rules_optimized (rules_buf[il_pos + i].cmds, tmp0, tmp1, pw_len);
|
||||
const u32 tmp_len = apply_rules_optimized (kernel_rules[il_pos + i].cmds, tmp0, tmp1, pw_len);
|
||||
|
||||
switch (i)
|
||||
{
|
||||
|
||||
@ -123,6 +123,6 @@ DECLSPEC u32 toggle_on_register (const u32 in, const u32 r);
|
||||
DECLSPEC u32 rule_op_mangle_title_sep (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
|
||||
DECLSPEC u32 apply_rule_optimized (const u32 name, const u32 p0, const u32 p1, u32 *buf0, u32 *buf1, const u32 in_len);
|
||||
DECLSPEC u32 apply_rules_optimized (CONSTANT_AS const u32 *cmds, u32 *buf0, u32 *buf1, const u32 len);
|
||||
DECLSPEC u32x apply_rules_vect_optimized (const u32 *pw_buf0, const u32 *pw_buf1, const u32 pw_len, CONSTANT_AS const kernel_rule_t *rules_buf, const u32 il_pos, u32x *buf0, u32x *buf1);
|
||||
DECLSPEC u32x apply_rules_vect_optimized (const u32 *pw_buf0, const u32 *pw_buf1, const u32 pw_len, CONSTANT_AS const kernel_rule_t *kernel_rules, const u32 il_pos, u32x *buf0, u32x *buf1);
|
||||
|
||||
#endif // _INC_RP_OPTIMIZED_H
|
||||
|
||||
@ -5,23 +5,24 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.h"
|
||||
#include "inc_simd.h"
|
||||
|
||||
// attack-mode 0
|
||||
|
||||
DECLSPEC u32x ix_create_bft (GLOBAL_AS const bf_t *bfs_buf, const u32 il_pos)
|
||||
DECLSPEC u32x ix_create_bft (CONSTANT_AS const bf_t *arr, const u32 il_pos)
|
||||
{
|
||||
#if VECT_SIZE == 1
|
||||
const u32x ix = (u32x) (bfs_buf[il_pos + 0].i);
|
||||
const u32x ix = make_u32x (arr[il_pos + 0].i);
|
||||
#elif VECT_SIZE == 2
|
||||
const u32x ix = (u32x) (bfs_buf[il_pos + 0].i, bfs_buf[il_pos + 1].i);
|
||||
const u32x ix = make_u32x (arr[il_pos + 0].i, arr[il_pos + 1].i);
|
||||
#elif VECT_SIZE == 4
|
||||
const u32x ix = (u32x) (bfs_buf[il_pos + 0].i, bfs_buf[il_pos + 1].i, bfs_buf[il_pos + 2].i, bfs_buf[il_pos + 3].i);
|
||||
const u32x ix = make_u32x (arr[il_pos + 0].i, arr[il_pos + 1].i, arr[il_pos + 2].i, arr[il_pos + 3].i);
|
||||
#elif VECT_SIZE == 8
|
||||
const u32x ix = (u32x) (bfs_buf[il_pos + 0].i, bfs_buf[il_pos + 1].i, bfs_buf[il_pos + 2].i, bfs_buf[il_pos + 3].i, bfs_buf[il_pos + 4].i, bfs_buf[il_pos + 5].i, bfs_buf[il_pos + 6].i, bfs_buf[il_pos + 7].i);
|
||||
const u32x ix = make_u32x (arr[il_pos + 0].i, arr[il_pos + 1].i, arr[il_pos + 2].i, arr[il_pos + 3].i, arr[il_pos + 4].i, arr[il_pos + 5].i, arr[il_pos + 6].i, arr[il_pos + 7].i);
|
||||
#elif VECT_SIZE == 16
|
||||
const u32x ix = (u32x) (bfs_buf[il_pos + 0].i, bfs_buf[il_pos + 1].i, bfs_buf[il_pos + 2].i, bfs_buf[il_pos + 3].i, bfs_buf[il_pos + 4].i, bfs_buf[il_pos + 5].i, bfs_buf[il_pos + 6].i, bfs_buf[il_pos + 7].i, bfs_buf[il_pos + 8].i, bfs_buf[il_pos + 9].i, bfs_buf[il_pos + 10].i, bfs_buf[il_pos + 11].i, bfs_buf[il_pos + 12].i, bfs_buf[il_pos + 13].i, bfs_buf[il_pos + 14].i, bfs_buf[il_pos + 15].i);
|
||||
const u32x ix = make_u32x (arr[il_pos + 0].i, arr[il_pos + 1].i, arr[il_pos + 2].i, arr[il_pos + 3].i, arr[il_pos + 4].i, arr[il_pos + 5].i, arr[il_pos + 6].i, arr[il_pos + 7].i, arr[il_pos + 8].i, arr[il_pos + 9].i, arr[il_pos + 10].i, arr[il_pos + 11].i, arr[il_pos + 12].i, arr[il_pos + 13].i, arr[il_pos + 14].i, arr[il_pos + 15].i);
|
||||
#endif
|
||||
|
||||
return ix;
|
||||
@ -29,35 +30,35 @@ DECLSPEC u32x ix_create_bft (GLOBAL_AS const bf_t *bfs_buf, const u32 il_pos)
|
||||
|
||||
// attack-mode 1
|
||||
|
||||
DECLSPEC u32x pwlenx_create_combt (GLOBAL_AS const pw_t *combs_buf, const u32 il_pos)
|
||||
DECLSPEC u32x pwlenx_create_combt (GLOBAL_AS const pw_t *arr, const u32 il_pos)
|
||||
{
|
||||
#if VECT_SIZE == 1
|
||||
const u32x pw_lenx = (u32x) (combs_buf[il_pos + 0].pw_len);
|
||||
const u32x pw_lenx = make_u32x (arr[il_pos + 0].pw_len);
|
||||
#elif VECT_SIZE == 2
|
||||
const u32x pw_lenx = (u32x) (combs_buf[il_pos + 0].pw_len, combs_buf[il_pos + 1].pw_len);
|
||||
const u32x pw_lenx = make_u32x (arr[il_pos + 0].pw_len, arr[il_pos + 1].pw_len);
|
||||
#elif VECT_SIZE == 4
|
||||
const u32x pw_lenx = (u32x) (combs_buf[il_pos + 0].pw_len, combs_buf[il_pos + 1].pw_len, combs_buf[il_pos + 2].pw_len, combs_buf[il_pos + 3].pw_len);
|
||||
const u32x pw_lenx = make_u32x (arr[il_pos + 0].pw_len, arr[il_pos + 1].pw_len, arr[il_pos + 2].pw_len, arr[il_pos + 3].pw_len);
|
||||
#elif VECT_SIZE == 8
|
||||
const u32x pw_lenx = (u32x) (combs_buf[il_pos + 0].pw_len, combs_buf[il_pos + 1].pw_len, combs_buf[il_pos + 2].pw_len, combs_buf[il_pos + 3].pw_len, combs_buf[il_pos + 4].pw_len, combs_buf[il_pos + 5].pw_len, combs_buf[il_pos + 6].pw_len, combs_buf[il_pos + 7].pw_len);
|
||||
const u32x pw_lenx = make_u32x (arr[il_pos + 0].pw_len, arr[il_pos + 1].pw_len, arr[il_pos + 2].pw_len, arr[il_pos + 3].pw_len, arr[il_pos + 4].pw_len, arr[il_pos + 5].pw_len, arr[il_pos + 6].pw_len, arr[il_pos + 7].pw_len);
|
||||
#elif VECT_SIZE == 16
|
||||
const u32x pw_lenx = (u32x) (combs_buf[il_pos + 0].pw_len, combs_buf[il_pos + 1].pw_len, combs_buf[il_pos + 2].pw_len, combs_buf[il_pos + 3].pw_len, combs_buf[il_pos + 4].pw_len, combs_buf[il_pos + 5].pw_len, combs_buf[il_pos + 6].pw_len, combs_buf[il_pos + 7].pw_len, combs_buf[il_pos + 8].pw_len, combs_buf[il_pos + 9].pw_len, combs_buf[il_pos + 10].pw_len, combs_buf[il_pos + 11].pw_len, combs_buf[il_pos + 12].pw_len, combs_buf[il_pos + 13].pw_len, combs_buf[il_pos + 14].pw_len, combs_buf[il_pos + 15].pw_len);
|
||||
const u32x pw_lenx = make_u32x (arr[il_pos + 0].pw_len, arr[il_pos + 1].pw_len, arr[il_pos + 2].pw_len, arr[il_pos + 3].pw_len, arr[il_pos + 4].pw_len, arr[il_pos + 5].pw_len, arr[il_pos + 6].pw_len, arr[il_pos + 7].pw_len, arr[il_pos + 8].pw_len, arr[il_pos + 9].pw_len, arr[il_pos + 10].pw_len, arr[il_pos + 11].pw_len, arr[il_pos + 12].pw_len, arr[il_pos + 13].pw_len, arr[il_pos + 14].pw_len, arr[il_pos + 15].pw_len);
|
||||
#endif
|
||||
|
||||
return pw_lenx;
|
||||
}
|
||||
|
||||
DECLSPEC u32x ix_create_combt (GLOBAL_AS const pw_t *combs_buf, const u32 il_pos, const int idx)
|
||||
DECLSPEC u32x ix_create_combt (GLOBAL_AS const pw_t *arr, const u32 il_pos, const int idx)
|
||||
{
|
||||
#if VECT_SIZE == 1
|
||||
const u32x ix = (u32x) (combs_buf[il_pos + 0].i[idx]);
|
||||
const u32x ix = make_u32x (arr[il_pos + 0].i[idx]);
|
||||
#elif VECT_SIZE == 2
|
||||
const u32x ix = (u32x) (combs_buf[il_pos + 0].i[idx], combs_buf[il_pos + 1].i[idx]);
|
||||
const u32x ix = make_u32x (arr[il_pos + 0].i[idx], arr[il_pos + 1].i[idx]);
|
||||
#elif VECT_SIZE == 4
|
||||
const u32x ix = (u32x) (combs_buf[il_pos + 0].i[idx], combs_buf[il_pos + 1].i[idx], combs_buf[il_pos + 2].i[idx], combs_buf[il_pos + 3].i[idx]);
|
||||
const u32x ix = make_u32x (arr[il_pos + 0].i[idx], arr[il_pos + 1].i[idx], arr[il_pos + 2].i[idx], arr[il_pos + 3].i[idx]);
|
||||
#elif VECT_SIZE == 8
|
||||
const u32x ix = (u32x) (combs_buf[il_pos + 0].i[idx], combs_buf[il_pos + 1].i[idx], combs_buf[il_pos + 2].i[idx], combs_buf[il_pos + 3].i[idx], combs_buf[il_pos + 4].i[idx], combs_buf[il_pos + 5].i[idx], combs_buf[il_pos + 6].i[idx], combs_buf[il_pos + 7].i[idx]);
|
||||
const u32x ix = make_u32x (arr[il_pos + 0].i[idx], arr[il_pos + 1].i[idx], arr[il_pos + 2].i[idx], arr[il_pos + 3].i[idx], arr[il_pos + 4].i[idx], arr[il_pos + 5].i[idx], arr[il_pos + 6].i[idx], arr[il_pos + 7].i[idx]);
|
||||
#elif VECT_SIZE == 16
|
||||
const u32x ix = (u32x) (combs_buf[il_pos + 0].i[idx], combs_buf[il_pos + 1].i[idx], combs_buf[il_pos + 2].i[idx], combs_buf[il_pos + 3].i[idx], combs_buf[il_pos + 4].i[idx], combs_buf[il_pos + 5].i[idx], combs_buf[il_pos + 6].i[idx], combs_buf[il_pos + 7].i[idx], combs_buf[il_pos + 8].i[idx], combs_buf[il_pos + 9].i[idx], combs_buf[il_pos + 10].i[idx], combs_buf[il_pos + 11].i[idx], combs_buf[il_pos + 12].i[idx], combs_buf[il_pos + 13].i[idx], combs_buf[il_pos + 14].i[idx], combs_buf[il_pos + 15].i[idx]);
|
||||
const u32x ix = make_u32x (arr[il_pos + 0].i[idx], arr[il_pos + 1].i[idx], arr[il_pos + 2].i[idx], arr[il_pos + 3].i[idx], arr[il_pos + 4].i[idx], arr[il_pos + 5].i[idx], arr[il_pos + 6].i[idx], arr[il_pos + 7].i[idx], arr[il_pos + 8].i[idx], arr[il_pos + 9].i[idx], arr[il_pos + 10].i[idx], arr[il_pos + 11].i[idx], arr[il_pos + 12].i[idx], arr[il_pos + 13].i[idx], arr[il_pos + 14].i[idx], arr[il_pos + 15].i[idx]);
|
||||
#endif
|
||||
|
||||
return ix;
|
||||
|
||||
@ -1050,51 +1050,51 @@
|
||||
#define MATCHES_NONE_VS(a,b) !(MATCHES_ONE_VS ((a), (b)))
|
||||
|
||||
#if VECT_SIZE == 1
|
||||
#define packv(arr,var,gid,idx) (u32x) ((arr)[((gid) * 1) + 0].var[(idx)])
|
||||
#define packv(arr,var,gid,idx) make_u32x ((arr)[((gid) * 1) + 0].var[(idx)])
|
||||
#elif VECT_SIZE == 2
|
||||
#define packv(arr,var,gid,idx) (u32x) ((arr)[((gid) * 2) + 0].var[(idx)], (arr)[((gid) * 2) + 1].var[(idx)])
|
||||
#define packv(arr,var,gid,idx) make_u32x ((arr)[((gid) * 2) + 0].var[(idx)], (arr)[((gid) * 2) + 1].var[(idx)])
|
||||
#elif VECT_SIZE == 4
|
||||
#define packv(arr,var,gid,idx) (u32x) ((arr)[((gid) * 4) + 0].var[(idx)], (arr)[((gid) * 4) + 1].var[(idx)], (arr)[((gid) * 4) + 2].var[(idx)], (arr)[((gid) * 4) + 3].var[(idx)])
|
||||
#define packv(arr,var,gid,idx) make_u32x ((arr)[((gid) * 4) + 0].var[(idx)], (arr)[((gid) * 4) + 1].var[(idx)], (arr)[((gid) * 4) + 2].var[(idx)], (arr)[((gid) * 4) + 3].var[(idx)])
|
||||
#elif VECT_SIZE == 8
|
||||
#define packv(arr,var,gid,idx) (u32x) ((arr)[((gid) * 8) + 0].var[(idx)], (arr)[((gid) * 8) + 1].var[(idx)], (arr)[((gid) * 8) + 2].var[(idx)], (arr)[((gid) * 8) + 3].var[(idx)], (arr)[((gid) * 8) + 4].var[(idx)], (arr)[((gid) * 8) + 5].var[(idx)], (arr)[((gid) * 8) + 6].var[(idx)], (arr)[((gid) * 8) + 7].var[(idx)])
|
||||
#define packv(arr,var,gid,idx) make_u32x ((arr)[((gid) * 8) + 0].var[(idx)], (arr)[((gid) * 8) + 1].var[(idx)], (arr)[((gid) * 8) + 2].var[(idx)], (arr)[((gid) * 8) + 3].var[(idx)], (arr)[((gid) * 8) + 4].var[(idx)], (arr)[((gid) * 8) + 5].var[(idx)], (arr)[((gid) * 8) + 6].var[(idx)], (arr)[((gid) * 8) + 7].var[(idx)])
|
||||
#elif VECT_SIZE == 16
|
||||
#define packv(arr,var,gid,idx) (u32x) ((arr)[((gid) * 16) + 0].var[(idx)], (arr)[((gid) * 16) + 1].var[(idx)], (arr)[((gid) * 16) + 2].var[(idx)], (arr)[((gid) * 16) + 3].var[(idx)], (arr)[((gid) * 16) + 4].var[(idx)], (arr)[((gid) * 16) + 5].var[(idx)], (arr)[((gid) * 16) + 6].var[(idx)], (arr)[((gid) * 16) + 7].var[(idx)], (arr)[((gid) * 16) + 8].var[(idx)], (arr)[((gid) * 16) + 9].var[(idx)], (arr)[((gid) * 16) + 10].var[(idx)], (arr)[((gid) * 16) + 11].var[(idx)], (arr)[((gid) * 16) + 12].var[(idx)], (arr)[((gid) * 16) + 13].var[(idx)], (arr)[((gid) * 16) + 14].var[(idx)], (arr)[((gid) * 16) + 15].var[(idx)])
|
||||
#define packv(arr,var,gid,idx) make_u32x ((arr)[((gid) * 16) + 0].var[(idx)], (arr)[((gid) * 16) + 1].var[(idx)], (arr)[((gid) * 16) + 2].var[(idx)], (arr)[((gid) * 16) + 3].var[(idx)], (arr)[((gid) * 16) + 4].var[(idx)], (arr)[((gid) * 16) + 5].var[(idx)], (arr)[((gid) * 16) + 6].var[(idx)], (arr)[((gid) * 16) + 7].var[(idx)], (arr)[((gid) * 16) + 8].var[(idx)], (arr)[((gid) * 16) + 9].var[(idx)], (arr)[((gid) * 16) + 10].var[(idx)], (arr)[((gid) * 16) + 11].var[(idx)], (arr)[((gid) * 16) + 12].var[(idx)], (arr)[((gid) * 16) + 13].var[(idx)], (arr)[((gid) * 16) + 14].var[(idx)], (arr)[((gid) * 16) + 15].var[(idx)])
|
||||
#endif
|
||||
|
||||
#if VECT_SIZE == 1
|
||||
#define pack64v(arr,var,gid,idx) (u64x) ((arr)[((gid) * 1) + 0].var[(idx)])
|
||||
#define pack64v(arr,var,gid,idx) make_u64x ((arr)[((gid) * 1) + 0].var[(idx)])
|
||||
#elif VECT_SIZE == 2
|
||||
#define pack64v(arr,var,gid,idx) (u64x) ((arr)[((gid) * 2) + 0].var[(idx)], (arr)[((gid) * 2) + 1].var[(idx)])
|
||||
#define pack64v(arr,var,gid,idx) make_u64x ((arr)[((gid) * 2) + 0].var[(idx)], (arr)[((gid) * 2) + 1].var[(idx)])
|
||||
#elif VECT_SIZE == 4
|
||||
#define pack64v(arr,var,gid,idx) (u64x) ((arr)[((gid) * 4) + 0].var[(idx)], (arr)[((gid) * 4) + 1].var[(idx)], (arr)[((gid) * 4) + 2].var[(idx)], (arr)[((gid) * 4) + 3].var[(idx)])
|
||||
#define pack64v(arr,var,gid,idx) make_u64x ((arr)[((gid) * 4) + 0].var[(idx)], (arr)[((gid) * 4) + 1].var[(idx)], (arr)[((gid) * 4) + 2].var[(idx)], (arr)[((gid) * 4) + 3].var[(idx)])
|
||||
#elif VECT_SIZE == 8
|
||||
#define pack64v(arr,var,gid,idx) (u64x) ((arr)[((gid) * 8) + 0].var[(idx)], (arr)[((gid) * 8) + 1].var[(idx)], (arr)[((gid) * 8) + 2].var[(idx)], (arr)[((gid) * 8) + 3].var[(idx)], (arr)[((gid) * 8) + 4].var[(idx)], (arr)[((gid) * 8) + 5].var[(idx)], (arr)[((gid) * 8) + 6].var[(idx)], (arr)[((gid) * 8) + 7].var[(idx)])
|
||||
#define pack64v(arr,var,gid,idx) make_u64x ((arr)[((gid) * 8) + 0].var[(idx)], (arr)[((gid) * 8) + 1].var[(idx)], (arr)[((gid) * 8) + 2].var[(idx)], (arr)[((gid) * 8) + 3].var[(idx)], (arr)[((gid) * 8) + 4].var[(idx)], (arr)[((gid) * 8) + 5].var[(idx)], (arr)[((gid) * 8) + 6].var[(idx)], (arr)[((gid) * 8) + 7].var[(idx)])
|
||||
#elif VECT_SIZE == 16
|
||||
#define pack64v(arr,var,gid,idx) (u64x) ((arr)[((gid) * 16) + 0].var[(idx)], (arr)[((gid) * 16) + 1].var[(idx)], (arr)[((gid) * 16) + 2].var[(idx)], (arr)[((gid) * 16) + 3].var[(idx)], (arr)[((gid) * 16) + 4].var[(idx)], (arr)[((gid) * 16) + 5].var[(idx)], (arr)[((gid) * 16) + 6].var[(idx)], (arr)[((gid) * 16) + 7].var[(idx)], (arr)[((gid) * 16) + 8].var[(idx)], (arr)[((gid) * 16) + 9].var[(idx)], (arr)[((gid) * 16) + 10].var[(idx)], (arr)[((gid) * 16) + 11].var[(idx)], (arr)[((gid) * 16) + 12].var[(idx)], (arr)[((gid) * 16) + 13].var[(idx)], (arr)[((gid) * 16) + 14].var[(idx)], (arr)[((gid) * 16) + 15].var[(idx)])
|
||||
#define pack64v(arr,var,gid,idx) make_u64x ((arr)[((gid) * 16) + 0].var[(idx)], (arr)[((gid) * 16) + 1].var[(idx)], (arr)[((gid) * 16) + 2].var[(idx)], (arr)[((gid) * 16) + 3].var[(idx)], (arr)[((gid) * 16) + 4].var[(idx)], (arr)[((gid) * 16) + 5].var[(idx)], (arr)[((gid) * 16) + 6].var[(idx)], (arr)[((gid) * 16) + 7].var[(idx)], (arr)[((gid) * 16) + 8].var[(idx)], (arr)[((gid) * 16) + 9].var[(idx)], (arr)[((gid) * 16) + 10].var[(idx)], (arr)[((gid) * 16) + 11].var[(idx)], (arr)[((gid) * 16) + 12].var[(idx)], (arr)[((gid) * 16) + 13].var[(idx)], (arr)[((gid) * 16) + 14].var[(idx)], (arr)[((gid) * 16) + 15].var[(idx)])
|
||||
#endif
|
||||
|
||||
#if VECT_SIZE == 1
|
||||
#define packvf(arr,var,gid) (u32x) ((arr)[((gid) * 1) + 0].var)
|
||||
#define packvf(arr,var,gid) make_u32x ((arr)[((gid) * 1) + 0].var)
|
||||
#elif VECT_SIZE == 2
|
||||
#define packvf(arr,var,gid) (u32x) ((arr)[((gid) * 2) + 0].var, (arr)[((gid) * 2) + 1].var)
|
||||
#define packvf(arr,var,gid) make_u32x ((arr)[((gid) * 2) + 0].var, (arr)[((gid) * 2) + 1].var)
|
||||
#elif VECT_SIZE == 4
|
||||
#define packvf(arr,var,gid) (u32x) ((arr)[((gid) * 4) + 0].var, (arr)[((gid) * 4) + 1].var, (arr)[((gid) * 4) + 2].var, (arr)[((gid) * 4) + 3].var)
|
||||
#define packvf(arr,var,gid) make_u32x ((arr)[((gid) * 4) + 0].var, (arr)[((gid) * 4) + 1].var, (arr)[((gid) * 4) + 2].var, (arr)[((gid) * 4) + 3].var)
|
||||
#elif VECT_SIZE == 8
|
||||
#define packvf(arr,var,gid) (u32x) ((arr)[((gid) * 8) + 0].var, (arr)[((gid) * 8) + 1].var, (arr)[((gid) * 8) + 2].var, (arr)[((gid) * 8) + 3].var, (arr)[((gid) * 8) + 4].var, (arr)[((gid) * 8) + 5].var, (arr)[((gid) * 8) + 6].var, (arr)[((gid) * 8) + 7].var)
|
||||
#define packvf(arr,var,gid) make_u32x ((arr)[((gid) * 8) + 0].var, (arr)[((gid) * 8) + 1].var, (arr)[((gid) * 8) + 2].var, (arr)[((gid) * 8) + 3].var, (arr)[((gid) * 8) + 4].var, (arr)[((gid) * 8) + 5].var, (arr)[((gid) * 8) + 6].var, (arr)[((gid) * 8) + 7].var)
|
||||
#elif VECT_SIZE == 16
|
||||
#define packvf(arr,var,gid) (u32x) ((arr)[((gid) * 16) + 0].var, (arr)[((gid) * 16) + 1].var, (arr)[((gid) * 16) + 2].var, (arr)[((gid) * 16) + 3].var, (arr)[((gid) * 16) + 4].var, (arr)[((gid) * 16) + 5].var, (arr)[((gid) * 16) + 6].var, (arr)[((gid) * 16) + 7].var, (arr)[((gid) * 16) + 8].var, (arr)[((gid) * 16) + 9].var, (arr)[((gid) * 16) + 10].var, (arr)[((gid) * 16) + 11].var, (arr)[((gid) * 16) + 12].var, (arr)[((gid) * 16) + 13].var, (arr)[((gid) * 16) + 14].var, (arr)[((gid) * 16) + 15].var)
|
||||
#define packvf(arr,var,gid) make_u32x ((arr)[((gid) * 16) + 0].var, (arr)[((gid) * 16) + 1].var, (arr)[((gid) * 16) + 2].var, (arr)[((gid) * 16) + 3].var, (arr)[((gid) * 16) + 4].var, (arr)[((gid) * 16) + 5].var, (arr)[((gid) * 16) + 6].var, (arr)[((gid) * 16) + 7].var, (arr)[((gid) * 16) + 8].var, (arr)[((gid) * 16) + 9].var, (arr)[((gid) * 16) + 10].var, (arr)[((gid) * 16) + 11].var, (arr)[((gid) * 16) + 12].var, (arr)[((gid) * 16) + 13].var, (arr)[((gid) * 16) + 14].var, (arr)[((gid) * 16) + 15].var)
|
||||
#endif
|
||||
|
||||
#if VECT_SIZE == 1
|
||||
#define pack64vf(arr,var,gid) (u64x) ((arr)[((gid) * 1) + 0].var)
|
||||
#define pack64vf(arr,var,gid) make_u64x ((arr)[((gid) * 1) + 0].var)
|
||||
#elif VECT_SIZE == 2
|
||||
#define pack64vf(arr,var,gid) (u64x) ((arr)[((gid) * 2) + 0].var, (arr)[((gid) * 2) + 1].var)
|
||||
#define pack64vf(arr,var,gid) make_u64x ((arr)[((gid) * 2) + 0].var, (arr)[((gid) * 2) + 1].var)
|
||||
#elif VECT_SIZE == 4
|
||||
#define pack64vf(arr,var,gid) (u64x) ((arr)[((gid) * 4) + 0].var, (arr)[((gid) * 4) + 1].var, (arr)[((gid) * 4) + 2].var, (arr)[((gid) * 4) + 3].var)
|
||||
#define pack64vf(arr,var,gid) make_u64x ((arr)[((gid) * 4) + 0].var, (arr)[((gid) * 4) + 1].var, (arr)[((gid) * 4) + 2].var, (arr)[((gid) * 4) + 3].var)
|
||||
#elif VECT_SIZE == 8
|
||||
#define pack64vf(arr,var,gid) (u64x) ((arr)[((gid) * 8) + 0].var, (arr)[((gid) * 8) + 1].var, (arr)[((gid) * 8) + 2].var, (arr)[((gid) * 8) + 3].var, (arr)[((gid) * 8) + 4].var, (arr)[((gid) * 8) + 5].var, (arr)[((gid) * 8) + 6].var, (arr)[((gid) * 8) + 7].var)
|
||||
#define pack64vf(arr,var,gid) make_u64x ((arr)[((gid) * 8) + 0].var, (arr)[((gid) * 8) + 1].var, (arr)[((gid) * 8) + 2].var, (arr)[((gid) * 8) + 3].var, (arr)[((gid) * 8) + 4].var, (arr)[((gid) * 8) + 5].var, (arr)[((gid) * 8) + 6].var, (arr)[((gid) * 8) + 7].var)
|
||||
#elif VECT_SIZE == 16
|
||||
#define pack64vf(arr,var,gid) (u64x) ((arr)[((gid) * 16) + 0].var, (arr)[((gid) * 16) + 1].var, (arr)[((gid) * 16) + 2].var, (arr)[((gid) * 16) + 3].var, (arr)[((gid) * 16) + 4].var, (arr)[((gid) * 16) + 5].var, (arr)[((gid) * 16) + 6].var, (arr)[((gid) * 16) + 7].var, (arr)[((gid) * 16) + 8].var, (arr)[((gid) * 16) + 9].var, (arr)[((gid) * 16) + 10].var, (arr)[((gid) * 16) + 11].var, (arr)[((gid) * 16) + 12].var, (arr)[((gid) * 16) + 13].var, (arr)[((gid) * 16) + 14].var, (arr)[((gid) * 16) + 15].var)
|
||||
#define pack64vf(arr,var,gid) make_u64x ((arr)[((gid) * 16) + 0].var, (arr)[((gid) * 16) + 1].var, (arr)[((gid) * 16) + 2].var, (arr)[((gid) * 16) + 3].var, (arr)[((gid) * 16) + 4].var, (arr)[((gid) * 16) + 5].var, (arr)[((gid) * 16) + 6].var, (arr)[((gid) * 16) + 7].var, (arr)[((gid) * 16) + 8].var, (arr)[((gid) * 16) + 9].var, (arr)[((gid) * 16) + 10].var, (arr)[((gid) * 16) + 11].var, (arr)[((gid) * 16) + 12].var, (arr)[((gid) * 16) + 13].var, (arr)[((gid) * 16) + 14].var, (arr)[((gid) * 16) + 15].var)
|
||||
#endif
|
||||
|
||||
#if VECT_SIZE == 1
|
||||
@ -1133,8 +1133,8 @@
|
||||
#define unpackv_xor(arr,var,gid,idx,val) (arr)[((gid) * 16) + 0].var[(idx)] ^= val.s0; (arr)[((gid) * 16) + 1].var[(idx)] ^= val.s1; (arr)[((gid) * 16) + 2].var[(idx)] ^= val.s2; (arr)[((gid) * 16) + 3].var[(idx)] ^= val.s3; (arr)[((gid) * 16) + 4].var[(idx)] ^= val.s4; (arr)[((gid) * 16) + 5].var[(idx)] ^= val.s5; (arr)[((gid) * 16) + 6].var[(idx)] ^= val.s6; (arr)[((gid) * 16) + 7].var[(idx)] ^= val.s7; (arr)[((gid) * 16) + 8].var[(idx)] ^= val.s8; (arr)[((gid) * 16) + 9].var[(idx)] ^= val.s9; (arr)[((gid) * 16) + 10].var[(idx)] ^= val.sa; (arr)[((gid) * 16) + 11].var[(idx)] ^= val.sb; (arr)[((gid) * 16) + 12].var[(idx)] ^= val.sc; (arr)[((gid) * 16) + 13].var[(idx)] ^= val.sd; (arr)[((gid) * 16) + 14].var[(idx)] ^= val.se; (arr)[((gid) * 16) + 15].var[(idx)] ^= val.sf;
|
||||
#endif
|
||||
|
||||
DECLSPEC u32x ix_create_bft (GLOBAL_AS const bf_t *bfs_buf, const u32 il_pos);
|
||||
DECLSPEC u32x pwlenx_create_combt (GLOBAL_AS const pw_t *combs_buf, const u32 il_pos);
|
||||
DECLSPEC u32x ix_create_combt (GLOBAL_AS const pw_t *combs_buf, const u32 il_pos, const int idx);
|
||||
DECLSPEC u32x ix_create_bft (CONSTANT_AS const bf_t *arr, const u32 il_pos);
|
||||
DECLSPEC u32x pwlenx_create_combt (GLOBAL_AS const pw_t *arr, const u32 il_pos);
|
||||
DECLSPEC u32x ix_create_combt (GLOBAL_AS const pw_t *arr, const u32 il_pos, const int idx);
|
||||
|
||||
#endif
|
||||
|
||||
@ -5,10 +5,11 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.h"
|
||||
#include "inc_truecrypt_crc32.h"
|
||||
|
||||
CONSTANT_AS u32a crc32tab[0x100] =
|
||||
CONSTANT_VK u32a crc32tab[0x100] =
|
||||
{
|
||||
0x00000000, 0x77073096, 0xee0e612c, 0x990951ba,
|
||||
0x076dc419, 0x706af48f, 0xe963a535, 0x9e6495a3,
|
||||
|
||||
@ -5,6 +5,7 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.h"
|
||||
#include "inc_truecrypt_keyfile.h"
|
||||
|
||||
|
||||
@ -5,6 +5,7 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.h"
|
||||
#include "inc_cipher_aes.h"
|
||||
#include "inc_cipher_serpent.h"
|
||||
|
||||
@ -6,6 +6,14 @@
|
||||
#ifndef _INC_TYPES_H
|
||||
#define _INC_TYPES_H
|
||||
|
||||
#ifdef IS_CUDA
|
||||
//https://docs.nvidia.com/cuda/nvrtc/index.html#integer-size
|
||||
typedef unsigned char uchar;
|
||||
typedef unsigned short ushort;
|
||||
typedef unsigned int uint;
|
||||
typedef unsigned long long ulong;
|
||||
#endif
|
||||
|
||||
#ifdef KERNEL_STATIC
|
||||
typedef uchar u8;
|
||||
typedef ushort u16;
|
||||
@ -43,11 +51,757 @@ typedef u8 u8x;
|
||||
typedef u16 u16x;
|
||||
typedef u32 u32x;
|
||||
typedef u64 u64x;
|
||||
|
||||
#define make_u8x (u8)
|
||||
#define make_u16x (u16)
|
||||
#define make_u32x (u32)
|
||||
#define make_u64x (u64)
|
||||
|
||||
#else
|
||||
#ifdef IS_CUDA
|
||||
|
||||
#if VECT_SIZE == 2
|
||||
|
||||
struct __device_builtin__ __builtin_align__(2) u8x
|
||||
{
|
||||
u8 s0;
|
||||
u8 s1;
|
||||
|
||||
inline __device__ u8x (const u8 a, const u8 b) : s0(a), s1(b) { }
|
||||
inline __device__ u8x (const u8 a) : s0(a), s1(a) { }
|
||||
|
||||
inline __device__ u8x (void) : s0(0), s1(0) { }
|
||||
inline __device__ ~u8x (void) { }
|
||||
};
|
||||
|
||||
struct __device_builtin__ __builtin_align__(4) u16x
|
||||
{
|
||||
u16 s0;
|
||||
u16 s1;
|
||||
|
||||
inline __device__ u16x (const u16 a, const u16 b) : s0(a), s1(b) { }
|
||||
inline __device__ u16x (const u16 a) : s0(a), s1(a) { }
|
||||
|
||||
inline __device__ u16x (void) : s0(0), s1(0) { }
|
||||
inline __device__ ~u16x (void) { }
|
||||
};
|
||||
|
||||
struct __device_builtin__ __builtin_align__(8) u32x
|
||||
{
|
||||
u32 s0;
|
||||
u32 s1;
|
||||
|
||||
inline __device__ u32x (const u32 a, const u32 b) : s0(a), s1(b) { }
|
||||
inline __device__ u32x (const u32 a) : s0(a), s1(a) { }
|
||||
|
||||
inline __device__ u32x (void) : s0(0), s1(0) { }
|
||||
inline __device__ ~u32x (void) { }
|
||||
};
|
||||
|
||||
struct __device_builtin__ __builtin_align__(16) u64x
|
||||
{
|
||||
u64 s0;
|
||||
u64 s1;
|
||||
|
||||
inline __device__ u64x (const u64 a, const u64 b) : s0(a), s1(b) { }
|
||||
inline __device__ u64x (const u64 a) : s0(a), s1(a) { }
|
||||
|
||||
inline __device__ u64x (void) : s0(0), s1(0) { }
|
||||
inline __device__ ~u64x (void) { }
|
||||
};
|
||||
|
||||
inline __device__ bool operator != (const u32x a, const u32 b) { return ((a.s0 != b) && (a.s1 != b)); }
|
||||
inline __device__ bool operator != (const u32x a, const u32x b) { return ((a.s0 != b.s0) && (a.s1 != b.s1)); }
|
||||
|
||||
inline __device__ void operator ^= (u32x &a, const u32 b) { a.s0 ^= b; a.s1 ^= b; }
|
||||
inline __device__ void operator ^= (u32x &a, const u32x b) { a.s0 ^= b.s0; a.s1 ^= b.s1; }
|
||||
|
||||
inline __device__ void operator |= (u32x &a, const u32 b) { a.s0 |= b; a.s1 |= b; }
|
||||
inline __device__ void operator |= (u32x &a, const u32x b) { a.s0 |= b.s0; a.s1 |= b.s1; }
|
||||
|
||||
inline __device__ void operator &= (u32x &a, const u32 b) { a.s0 &= b; a.s1 &= b; }
|
||||
inline __device__ void operator &= (u32x &a, const u32x b) { a.s0 &= b.s0; a.s1 &= b.s1; }
|
||||
|
||||
inline __device__ void operator += (u32x &a, const u32 b) { a.s0 += b; a.s1 += b; }
|
||||
inline __device__ void operator += (u32x &a, const u32x b) { a.s0 += b.s0; a.s1 += b.s1; }
|
||||
|
||||
inline __device__ void operator -= (u32x &a, const u32 b) { a.s0 -= b; a.s1 -= b; }
|
||||
inline __device__ void operator -= (u32x &a, const u32x b) { a.s0 -= b.s0; a.s1 -= b.s1; }
|
||||
|
||||
inline __device__ void operator *= (u32x &a, const u32 b) { a.s0 *= b; a.s1 *= b; }
|
||||
inline __device__ void operator *= (u32x &a, const u32x b) { a.s0 *= b.s0; a.s1 *= b.s1; }
|
||||
|
||||
inline __device__ void operator >>= (u32x &a, const u32 b) { a.s0 >>= b; a.s1 >>= b; }
|
||||
inline __device__ void operator >>= (u32x &a, const u32x b) { a.s0 >>= b.s0; a.s1 >>= b.s1; }
|
||||
|
||||
inline __device__ void operator <<= (u32x &a, const u32 b) { a.s0 <<= b; a.s1 <<= b; }
|
||||
inline __device__ void operator <<= (u32x &a, const u32x b) { a.s0 <<= b.s0; a.s1 <<= b.s1; }
|
||||
|
||||
inline __device__ u32x operator << (const u32x a, const u32 b) { return u32x ((a.s0 << b), (a.s1 << b) ); }
|
||||
inline __device__ u32x operator << (const u32x a, const u32x b) { return u32x ((a.s0 << b.s0), (a.s1 << b.s1)); }
|
||||
|
||||
inline __device__ u32x operator >> (const u32x a, const u32 b) { return u32x ((a.s0 >> b), (a.s1 >> b) ); }
|
||||
inline __device__ u32x operator >> (const u32x a, const u32x b) { return u32x ((a.s0 >> b.s0), (a.s1 >> b.s1)); }
|
||||
|
||||
inline __device__ u32x operator ^ (const u32x a, const u32 b) { return u32x ((a.s0 ^ b), (a.s1 ^ b) ); }
|
||||
inline __device__ u32x operator ^ (const u32x a, const u32x b) { return u32x ((a.s0 ^ b.s0), (a.s1 ^ b.s1)); }
|
||||
|
||||
inline __device__ u32x operator | (const u32x a, const u32 b) { return u32x ((a.s0 | b), (a.s1 | b) ); }
|
||||
inline __device__ u32x operator | (const u32x a, const u32x b) { return u32x ((a.s0 | b.s0), (a.s1 | b.s1)); }
|
||||
|
||||
inline __device__ u32x operator & (const u32x a, const u32 b) { return u32x ((a.s0 & b), (a.s1 & b) ); }
|
||||
inline __device__ u32x operator & (const u32x a, const u32x b) { return u32x ((a.s0 & b.s0), (a.s1 & b.s1)); }
|
||||
|
||||
inline __device__ u32x operator + (const u32x a, const u32 b) { return u32x ((a.s0 + b), (a.s1 + b) ); }
|
||||
inline __device__ u32x operator + (const u32x a, const u32x b) { return u32x ((a.s0 + b.s0), (a.s1 + b.s1)); }
|
||||
|
||||
inline __device__ u32x operator - (const u32x a, const u32 b) { return u32x ((a.s0 - b), (a.s1 - b) ); }
|
||||
inline __device__ u32x operator - (const u32x a, const u32x b) { return u32x ((a.s0 - b.s0), (a.s1 - b.s1)); }
|
||||
|
||||
inline __device__ u32x operator * (const u32x a, const u32 b) { return u32x ((a.s0 * b), (a.s1 * b) ); }
|
||||
inline __device__ u32x operator * (const u32x a, const u32x b) { return u32x ((a.s0 * b.s0), (a.s1 * b.s1)); }
|
||||
|
||||
inline __device__ u32x operator ~ (const u32x a) { return u32x (~a.s0, ~a.s1); }
|
||||
|
||||
inline __device__ bool operator != (const u64x a, const u64 b) { return ((a.s0 != b) && (a.s1 != b)); }
|
||||
inline __device__ bool operator != (const u64x a, const u64x b) { return ((a.s0 != b.s0) && (a.s1 != b.s1)); }
|
||||
|
||||
inline __device__ void operator ^= (u64x &a, const u64 b) { a.s0 ^= b; a.s1 ^= b; }
|
||||
inline __device__ void operator ^= (u64x &a, const u64x b) { a.s0 ^= b.s0; a.s1 ^= b.s1; }
|
||||
|
||||
inline __device__ void operator |= (u64x &a, const u64 b) { a.s0 |= b; a.s1 |= b; }
|
||||
inline __device__ void operator |= (u64x &a, const u64x b) { a.s0 |= b.s0; a.s1 |= b.s1; }
|
||||
|
||||
inline __device__ void operator &= (u64x &a, const u64 b) { a.s0 &= b; a.s1 &= b; }
|
||||
inline __device__ void operator &= (u64x &a, const u64x b) { a.s0 &= b.s0; a.s1 &= b.s1; }
|
||||
|
||||
inline __device__ void operator += (u64x &a, const u64 b) { a.s0 += b; a.s1 += b; }
|
||||
inline __device__ void operator += (u64x &a, const u64x b) { a.s0 += b.s0; a.s1 += b.s1; }
|
||||
|
||||
inline __device__ void operator -= (u64x &a, const u64 b) { a.s0 -= b; a.s1 -= b; }
|
||||
inline __device__ void operator -= (u64x &a, const u64x b) { a.s0 -= b.s0; a.s1 -= b.s1; }
|
||||
|
||||
inline __device__ void operator *= (u64x &a, const u64 b) { a.s0 *= b; a.s1 *= b; }
|
||||
inline __device__ void operator *= (u64x &a, const u64x b) { a.s0 *= b.s0; a.s1 *= b.s1; }
|
||||
|
||||
inline __device__ void operator >>= (u64x &a, const u64 b) { a.s0 >>= b; a.s1 >>= b; }
|
||||
inline __device__ void operator >>= (u64x &a, const u64x b) { a.s0 >>= b.s0; a.s1 >>= b.s1; }
|
||||
|
||||
inline __device__ void operator <<= (u64x &a, const u64 b) { a.s0 <<= b; a.s1 <<= b; }
|
||||
inline __device__ void operator <<= (u64x &a, const u64x b) { a.s0 <<= b.s0; a.s1 <<= b.s1; }
|
||||
|
||||
inline __device__ u64x operator << (const u64x a, const u64 b) { return u64x ((a.s0 << b), (a.s1 << b) ); }
|
||||
inline __device__ u64x operator << (const u64x a, const u64x b) { return u64x ((a.s0 << b.s0), (a.s1 << b.s1)); }
|
||||
|
||||
inline __device__ u64x operator >> (const u64x a, const u64 b) { return u64x ((a.s0 >> b), (a.s1 >> b) ); }
|
||||
inline __device__ u64x operator >> (const u64x a, const u64x b) { return u64x ((a.s0 >> b.s0), (a.s1 >> b.s1)); }
|
||||
|
||||
inline __device__ u64x operator ^ (const u64x a, const u64 b) { return u64x ((a.s0 ^ b), (a.s1 ^ b) ); }
|
||||
inline __device__ u64x operator ^ (const u64x a, const u64x b) { return u64x ((a.s0 ^ b.s0), (a.s1 ^ b.s1)); }
|
||||
|
||||
inline __device__ u64x operator | (const u64x a, const u64 b) { return u64x ((a.s0 | b), (a.s1 | b) ); }
|
||||
inline __device__ u64x operator | (const u64x a, const u64x b) { return u64x ((a.s0 | b.s0), (a.s1 | b.s1)); }
|
||||
|
||||
inline __device__ u64x operator & (const u64x a, const u64 b) { return u64x ((a.s0 & b), (a.s1 & b) ); }
|
||||
inline __device__ u64x operator & (const u64x a, const u64x b) { return u64x ((a.s0 & b.s0), (a.s1 & b.s1)); }
|
||||
|
||||
inline __device__ u64x operator + (const u64x a, const u64 b) { return u64x ((a.s0 + b), (a.s1 + b) ); }
|
||||
inline __device__ u64x operator + (const u64x a, const u64x b) { return u64x ((a.s0 + b.s0), (a.s1 + b.s1)); }
|
||||
|
||||
inline __device__ u64x operator - (const u64x a, const u64 b) { return u64x ((a.s0 - b), (a.s1 - b) ); }
|
||||
inline __device__ u64x operator - (const u64x a, const u64x b) { return u64x ((a.s0 - b.s0), (a.s1 - b.s1)); }
|
||||
|
||||
inline __device__ u64x operator * (const u64x a, const u64 b) { return u64x ((a.s0 * b), (a.s1 * b) ); }
|
||||
inline __device__ u64x operator * (const u64x a, const u64x b) { return u64x ((a.s0 * b.s0), (a.s1 * b.s1)); }
|
||||
|
||||
inline __device__ u64x operator ~ (const u64x a) { return u64x (~a.s0, ~a.s1); }
|
||||
|
||||
#endif
|
||||
|
||||
#if VECT_SIZE == 4
|
||||
|
||||
struct __device_builtin__ __builtin_align__(4) u8x
|
||||
{
|
||||
u8 s0;
|
||||
u8 s1;
|
||||
u8 s2;
|
||||
u8 s3;
|
||||
|
||||
inline __device__ u8x (const u8 a, const u8 b, const u8 c, const u8 d) : s0(a), s1(b), s2(c), s3(d) { }
|
||||
inline __device__ u8x (const u8 a) : s0(a), s1(a), s2(a), s3(a) { }
|
||||
|
||||
inline __device__ u8x (void) : s0(0), s1(0), s2(0), s3(0) { }
|
||||
inline __device__ ~u8x (void) { }
|
||||
};
|
||||
|
||||
struct __device_builtin__ __builtin_align__(8) u16x
|
||||
{
|
||||
u16 s0;
|
||||
u16 s1;
|
||||
u16 s2;
|
||||
u16 s3;
|
||||
|
||||
inline __device__ u16x (const u16 a, const u16 b, const u16 c, const u16 d) : s0(a), s1(b), s2(c), s3(d) { }
|
||||
inline __device__ u16x (const u16 a) : s0(a), s1(a), s2(a), s3(a) { }
|
||||
|
||||
inline __device__ u16x (void) : s0(0), s1(0), s2(0), s3(0) { }
|
||||
inline __device__ ~u16x (void) { }
|
||||
};
|
||||
|
||||
struct __device_builtin__ __builtin_align__(16) u32x
|
||||
{
|
||||
u32 s0;
|
||||
u32 s1;
|
||||
u32 s2;
|
||||
u32 s3;
|
||||
|
||||
inline __device__ u32x (const u32 a, const u32 b, const u32 c, const u32 d) : s0(a), s1(b), s2(c), s3(d) { }
|
||||
inline __device__ u32x (const u32 a) : s0(a), s1(a), s2(a), s3(a) { }
|
||||
|
||||
inline __device__ u32x (void) : s0(0), s1(0), s2(0), s3(0) { }
|
||||
inline __device__ ~u32x (void) { }
|
||||
};
|
||||
|
||||
struct __device_builtin__ __builtin_align__(32) u64x
|
||||
{
|
||||
u64 s0;
|
||||
u64 s1;
|
||||
u64 s2;
|
||||
u64 s3;
|
||||
|
||||
inline __device__ u64x (const u64 a, const u64 b, const u64 c, const u64 d) : s0(a), s1(b), s2(c), s3(d) { }
|
||||
inline __device__ u64x (const u64 a) : s0(a), s1(a), s2(a), s3(a) { }
|
||||
|
||||
inline __device__ u64x (void) : s0(0), s1(0), s2(0), s3(0) { }
|
||||
inline __device__ ~u64x (void) { }
|
||||
};
|
||||
|
||||
inline __device__ bool operator != (const u32x a, const u32 b) { return ((a.s0 != b) && (a.s1 != b) && (a.s2 != b) && (a.s3 != b) ); }
|
||||
inline __device__ bool operator != (const u32x a, const u32x b) { return ((a.s0 != b.s0) && (a.s1 != b.s1) && (a.s2 != b.s2) && (a.s3 != b.s3)); }
|
||||
|
||||
inline __device__ void operator ^= (u32x &a, const u32 b) { a.s0 ^= b; a.s1 ^= b; a.s2 ^= b; a.s3 ^= b; }
|
||||
inline __device__ void operator ^= (u32x &a, const u32x b) { a.s0 ^= b.s0; a.s1 ^= b.s1; a.s2 ^= b.s2; a.s3 ^= b.s3; }
|
||||
|
||||
inline __device__ void operator |= (u32x &a, const u32 b) { a.s0 |= b; a.s1 |= b; a.s2 |= b; a.s3 |= b; }
|
||||
inline __device__ void operator |= (u32x &a, const u32x b) { a.s0 |= b.s0; a.s1 |= b.s1; a.s2 |= b.s2; a.s3 |= b.s3; }
|
||||
|
||||
inline __device__ void operator &= (u32x &a, const u32 b) { a.s0 &= b; a.s1 &= b; a.s2 &= b; a.s3 &= b; }
|
||||
inline __device__ void operator &= (u32x &a, const u32x b) { a.s0 &= b.s0; a.s1 &= b.s1; a.s2 &= b.s2; a.s3 &= b.s3; }
|
||||
|
||||
inline __device__ void operator += (u32x &a, const u32 b) { a.s0 += b; a.s1 += b; a.s2 += b; a.s3 += b; }
|
||||
inline __device__ void operator += (u32x &a, const u32x b) { a.s0 += b.s0; a.s1 += b.s1; a.s2 += b.s2; a.s3 += b.s3; }
|
||||
|
||||
inline __device__ void operator -= (u32x &a, const u32 b) { a.s0 -= b; a.s1 -= b; a.s2 -= b; a.s3 -= b; }
|
||||
inline __device__ void operator -= (u32x &a, const u32x b) { a.s0 -= b.s0; a.s1 -= b.s1; a.s2 -= b.s2; a.s3 -= b.s3; }
|
||||
|
||||
inline __device__ void operator *= (u32x &a, const u32 b) { a.s0 *= b; a.s1 *= b; a.s2 *= b; a.s3 *= b; }
|
||||
inline __device__ void operator *= (u32x &a, const u32x b) { a.s0 *= b.s0; a.s1 *= b.s1; a.s2 *= b.s2; a.s3 *= b.s3; }
|
||||
|
||||
inline __device__ void operator >>= (u32x &a, const u32 b) { a.s0 >>= b; a.s1 >>= b; a.s2 >>= b; a.s3 >>= b; }
|
||||
inline __device__ void operator >>= (u32x &a, const u32x b) { a.s0 >>= b.s0; a.s1 >>= b.s1; a.s2 >>= b.s2; a.s3 >>= b.s3; }
|
||||
|
||||
inline __device__ void operator <<= (u32x &a, const u32 b) { a.s0 <<= b; a.s1 <<= b; a.s2 <<= b; a.s3 <<= b; }
|
||||
inline __device__ void operator <<= (u32x &a, const u32x b) { a.s0 <<= b.s0; a.s1 <<= b.s1; a.s2 <<= b.s2; a.s3 <<= b.s3; }
|
||||
|
||||
inline __device__ u32x operator << (const u32x a, const u32 b) { return u32x ((a.s0 << b), (a.s1 << b) , (a.s2 << b), (a.s3 << b) ); }
|
||||
inline __device__ u32x operator << (const u32x a, const u32x b) { return u32x ((a.s0 << b.s0), (a.s1 << b.s1), (a.s2 << b.s2), (a.s3 << b.s3)); }
|
||||
|
||||
inline __device__ u32x operator >> (const u32x a, const u32 b) { return u32x ((a.s0 >> b), (a.s1 >> b) , (a.s2 >> b), (a.s3 >> b) ); }
|
||||
inline __device__ u32x operator >> (const u32x a, const u32x b) { return u32x ((a.s0 >> b.s0), (a.s1 >> b.s1), (a.s2 >> b.s2), (a.s3 >> b.s3)); }
|
||||
|
||||
inline __device__ u32x operator ^ (const u32x a, const u32 b) { return u32x ((a.s0 ^ b), (a.s1 ^ b) , (a.s2 ^ b), (a.s3 ^ b) ); }
|
||||
inline __device__ u32x operator ^ (const u32x a, const u32x b) { return u32x ((a.s0 ^ b.s0), (a.s1 ^ b.s1), (a.s2 ^ b.s2), (a.s3 ^ b.s3)); }
|
||||
|
||||
inline __device__ u32x operator | (const u32x a, const u32 b) { return u32x ((a.s0 | b), (a.s1 | b) , (a.s2 | b), (a.s3 | b) ); }
|
||||
inline __device__ u32x operator | (const u32x a, const u32x b) { return u32x ((a.s0 | b.s0), (a.s1 | b.s1), (a.s2 | b.s2), (a.s3 | b.s3)); }
|
||||
|
||||
inline __device__ u32x operator & (const u32x a, const u32 b) { return u32x ((a.s0 & b), (a.s1 & b) , (a.s2 & b), (a.s3 & b) ); }
|
||||
inline __device__ u32x operator & (const u32x a, const u32x b) { return u32x ((a.s0 & b.s0), (a.s1 & b.s1), (a.s2 & b.s2), (a.s3 & b.s3)); }
|
||||
|
||||
inline __device__ u32x operator + (const u32x a, const u32 b) { return u32x ((a.s0 + b), (a.s1 + b) , (a.s2 + b), (a.s3 + b) ); }
|
||||
inline __device__ u32x operator + (const u32x a, const u32x b) { return u32x ((a.s0 + b.s0), (a.s1 + b.s1), (a.s2 + b.s2), (a.s3 + b.s3)); }
|
||||
|
||||
inline __device__ u32x operator - (const u32x a, const u32 b) { return u32x ((a.s0 - b), (a.s1 - b) , (a.s2 - b), (a.s3 - b) ); }
|
||||
inline __device__ u32x operator - (const u32x a, const u32x b) { return u32x ((a.s0 - b.s0), (a.s1 - b.s1), (a.s2 - b.s2), (a.s3 - b.s3)); }
|
||||
|
||||
inline __device__ u32x operator * (const u32x a, const u32 b) { return u32x ((a.s0 * b), (a.s1 * b) , (a.s2 * b), (a.s3 * b) ); }
|
||||
inline __device__ u32x operator * (const u32x a, const u32x b) { return u32x ((a.s0 * b.s0), (a.s1 * b.s1), (a.s2 * b.s2), (a.s3 * b.s3)); }
|
||||
|
||||
inline __device__ u32x operator ~ (const u32x a) { return u32x (~a.s0, ~a.s1, ~a.s2, ~a.s3); }
|
||||
|
||||
inline __device__ bool operator != (const u64x a, const u64 b) { return ((a.s0 != b) && (a.s1 != b) && (a.s2 != b) && (a.s3 != 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)); }
|
||||
|
||||
inline __device__ void operator ^= (u64x &a, const u64 b) { a.s0 ^= b; a.s1 ^= b; a.s2 ^= b; a.s3 ^= 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; }
|
||||
|
||||
inline __device__ void operator |= (u64x &a, const u64 b) { a.s0 |= b; a.s1 |= b; a.s2 |= b; a.s3 |= 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; }
|
||||
|
||||
inline __device__ void operator &= (u64x &a, const u64 b) { a.s0 &= b; a.s1 &= b; a.s2 &= b; a.s3 &= 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; }
|
||||
|
||||
inline __device__ void operator += (u64x &a, const u64 b) { a.s0 += b; a.s1 += b; a.s2 += b; a.s3 += 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; }
|
||||
|
||||
inline __device__ void operator -= (u64x &a, const u64 b) { a.s0 -= b; a.s1 -= b; a.s2 -= b; a.s3 -= 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; }
|
||||
|
||||
inline __device__ void operator *= (u64x &a, const u64 b) { a.s0 *= b; a.s1 *= b; a.s2 *= b; a.s3 *= 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; }
|
||||
|
||||
inline __device__ void operator >>= (u64x &a, const u64 b) { a.s0 >>= b; a.s1 >>= b; a.s2 >>= b; a.s3 >>= 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; }
|
||||
|
||||
inline __device__ void operator <<= (u64x &a, const u64 b) { a.s0 <<= b; a.s1 <<= b; a.s2 <<= b; a.s3 <<= 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; }
|
||||
|
||||
inline __device__ u64x operator << (const u64x a, const u64 b) { return u64x ((a.s0 << b), (a.s1 << b) , (a.s2 << b), (a.s3 << 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)); }
|
||||
|
||||
inline __device__ u64x operator >> (const u64x a, const u64 b) { return u64x ((a.s0 >> b), (a.s1 >> b) , (a.s2 >> b), (a.s3 >> 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)); }
|
||||
|
||||
inline __device__ u64x operator ^ (const u64x a, const u64 b) { return u64x ((a.s0 ^ b), (a.s1 ^ b) , (a.s2 ^ b), (a.s3 ^ 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)); }
|
||||
|
||||
inline __device__ u64x operator | (const u64x a, const u64 b) { return u64x ((a.s0 | b), (a.s1 | b) , (a.s2 | b), (a.s3 | 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)); }
|
||||
|
||||
inline __device__ u64x operator & (const u64x a, const u64 b) { return u64x ((a.s0 & b), (a.s1 & b) , (a.s2 & b), (a.s3 & 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)); }
|
||||
|
||||
inline __device__ u64x operator + (const u64x a, const u64 b) { return u64x ((a.s0 + b), (a.s1 + b) , (a.s2 + b), (a.s3 + 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)); }
|
||||
|
||||
inline __device__ u64x operator - (const u64x a, const u64 b) { return u64x ((a.s0 - b), (a.s1 - b) , (a.s2 - b), (a.s3 - 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)); }
|
||||
|
||||
inline __device__ u64x operator * (const u64x a, const u64 b) { return u64x ((a.s0 * b), (a.s1 * b) , (a.s2 * b), (a.s3 * 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)); }
|
||||
|
||||
inline __device__ u64x operator ~ (const u64x a) { return u64x (~a.s0, ~a.s1, ~a.s2, ~a.s3); }
|
||||
|
||||
#endif
|
||||
|
||||
#if VECT_SIZE == 8
|
||||
|
||||
struct __device_builtin__ __builtin_align__(8) u8x
|
||||
{
|
||||
u8 s0;
|
||||
u8 s1;
|
||||
u8 s2;
|
||||
u8 s3;
|
||||
u8 s4;
|
||||
u8 s5;
|
||||
u8 s6;
|
||||
u8 s7;
|
||||
|
||||
inline __device__ u8x (const u8 a, const u8 b, const u8 c, const u8 d, const u8 e, const u8 f, const u8 g, const u8 h) : s0(a), s1(b), s2(c), s3(d), s4(e), s5(f), s6(g), s7(h) { }
|
||||
inline __device__ u8x (const u8 a) : s0(a), s1(a), s2(a), s3(a), s4(a), s5(a), s6(a), s7(a) { }
|
||||
|
||||
inline __device__ u8x (void) : s0(0), s1(0), s2(0), s3(0), s4(0), s5(0), s6(0), s7(0) { }
|
||||
inline __device__ ~u8x (void) { }
|
||||
};
|
||||
|
||||
struct __device_builtin__ __builtin_align__(16) u16x
|
||||
{
|
||||
u16 s0;
|
||||
u16 s1;
|
||||
u16 s2;
|
||||
u16 s3;
|
||||
u16 s4;
|
||||
u16 s5;
|
||||
u16 s6;
|
||||
u16 s7;
|
||||
|
||||
inline __device__ u16x (const u16 a, const u16 b, const u16 c, const u16 d, const u16 e, const u16 f, const u16 g, const u16 h) : s0(a), s1(b), s2(c), s3(d), s4(e), s5(f), s6(g), s7(h) { }
|
||||
inline __device__ u16x (const u16 a) : s0(a), s1(a), s2(a), s3(a), s4(a), s5(a), s6(a), s7(a) { }
|
||||
|
||||
inline __device__ u16x (void) : s0(0), s1(0), s2(0), s3(0), s4(0), s5(0), s6(0), s7(0) { }
|
||||
inline __device__ ~u16x (void) { }
|
||||
};
|
||||
|
||||
struct __device_builtin__ __builtin_align__(32) u32x
|
||||
{
|
||||
u32 s0;
|
||||
u32 s1;
|
||||
u32 s2;
|
||||
u32 s3;
|
||||
u32 s4;
|
||||
u32 s5;
|
||||
u32 s6;
|
||||
u32 s7;
|
||||
|
||||
inline __device__ u32x (const u32 a, const u32 b, const u32 c, const u32 d, const u32 e, const u32 f, const u32 g, const u32 h) : s0(a), s1(b), s2(c), s3(d), s4(e), s5(f), s6(g), s7(h) { }
|
||||
inline __device__ u32x (const u32 a) : s0(a), s1(a), s2(a), s3(a), s4(a), s5(a), s6(a), s7(a) { }
|
||||
|
||||
inline __device__ u32x (void) : s0(0), s1(0), s2(0), s3(0), s4(0), s5(0), s6(0), s7(0) { }
|
||||
inline __device__ ~u32x (void) { }
|
||||
};
|
||||
|
||||
struct __device_builtin__ __builtin_align__(64) u64x
|
||||
{
|
||||
u64 s0;
|
||||
u64 s1;
|
||||
u64 s2;
|
||||
u64 s3;
|
||||
u64 s4;
|
||||
u64 s5;
|
||||
u64 s6;
|
||||
u64 s7;
|
||||
|
||||
inline __device__ u64x (const u64 a, const u64 b, const u64 c, const u64 d, const u64 e, const u64 f, const u64 g, const u64 h) : s0(a), s1(b), s2(c), s3(d), s4(e), s5(f), s6(g), s7(h) { }
|
||||
inline __device__ u64x (const u64 a) : s0(a), s1(a), s2(a), s3(a), s4(a), s5(a), s6(a), s7(a) { }
|
||||
|
||||
inline __device__ u64x (void) : s0(0), s1(0), s2(0), s3(0), s4(0), s5(0), s6(0), s7(0) { }
|
||||
inline __device__ ~u64x (void) { }
|
||||
};
|
||||
|
||||
inline __device__ bool operator != (const u32x a, const u32 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 u32x a, const u32x 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 ^= (u32x &a, const u32 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 ^= (u32x &a, const u32x 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 |= (u32x &a, const u32 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 |= (u32x &a, const u32x 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 &= (u32x &a, const u32 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 &= (u32x &a, const u32x 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 += (u32x &a, const u32 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 += (u32x &a, const u32x 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 -= (u32x &a, const u32 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 -= (u32x &a, const u32x 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 *= (u32x &a, const u32 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 *= (u32x &a, const u32x 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 >>= (u32x &a, const u32 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 >>= (u32x &a, const u32x 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 <<= (u32x &a, const u32 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 <<= (u32x &a, const u32x 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__ u32x operator << (const u32x a, const u32 b) { return u32x ((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__ u32x operator << (const u32x a, const u32x b) { return u32x ((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__ u32x operator >> (const u32x a, const u32 b) { return u32x ((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__ u32x operator >> (const u32x a, const u32x b) { return u32x ((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__ u32x operator ^ (const u32x a, const u32 b) { return u32x ((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__ u32x operator ^ (const u32x a, const u32x b) { return u32x ((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__ u32x operator | (const u32x a, const u32 b) { return u32x ((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__ u32x operator | (const u32x a, const u32x b) { return u32x ((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__ u32x operator & (const u32x a, const u32 b) { return u32x ((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__ u32x operator & (const u32x a, const u32x b) { return u32x ((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__ u32x operator + (const u32x a, const u32 b) { return u32x ((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__ u32x operator + (const u32x a, const u32x b) { return u32x ((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__ u32x operator - (const u32x a, const u32 b) { return u32x ((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__ u32x operator - (const u32x a, const u32x b) { return u32x ((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__ u32x operator * (const u32x a, const u32 b) { return u32x ((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__ u32x operator * (const u32x a, const u32x b) { return u32x ((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__ 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__ 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
|
||||
|
||||
#if VECT_SIZE == 16
|
||||
|
||||
struct __device_builtin__ __builtin_align__(16) u8x
|
||||
{
|
||||
u8 s0;
|
||||
u8 s1;
|
||||
u8 s2;
|
||||
u8 s3;
|
||||
u8 s4;
|
||||
u8 s5;
|
||||
u8 s6;
|
||||
u8 s7;
|
||||
u8 s8;
|
||||
u8 s9;
|
||||
u8 sa;
|
||||
u8 sb;
|
||||
u8 sc;
|
||||
u8 sd;
|
||||
u8 se;
|
||||
u8 sf;
|
||||
|
||||
inline __device__ u8x (const u8 a, const u8 b, const u8 c, const u8 d, const u8 e, const u8 f, const u8 g, const u8 h, const u8 i, const u8 j, const u8 k, const u8 l, const u8 m, const u8 n, const u8 o, const u8 p) : s0(a), s1(b), s2(c), s3(d), s4(e), s5(f), s6(g), s7(h), s8(i), s9(j), sa(k), sb(l), sc(m), sd(n), se(o), sf(p) { }
|
||||
inline __device__ u8x (const u8 a) : s0(a), s1(a), s2(a), s3(a), s4(a), s5(a), s6(a), s7(a), s8(a), s9(a), sa(a), sb(a), sc(a), sd(a), se(a), sf(a) { }
|
||||
|
||||
inline __device__ u8x (void) : s0(0), s1(0), s2(0), s3(0), s4(0), s5(0), s6(0), s7(0), s8(0), s9(0), sa(0), sb(0), sc(0), sd(0), se(0), sf(0) { }
|
||||
inline __device__ ~u8x (void) { }
|
||||
};
|
||||
|
||||
struct __device_builtin__ __builtin_align__(32) u16x
|
||||
{
|
||||
u16 s0;
|
||||
u16 s1;
|
||||
u16 s2;
|
||||
u16 s3;
|
||||
u16 s4;
|
||||
u16 s5;
|
||||
u16 s6;
|
||||
u16 s7;
|
||||
u16 s8;
|
||||
u16 s9;
|
||||
u16 sa;
|
||||
u16 sb;
|
||||
u16 sc;
|
||||
u16 sd;
|
||||
u16 se;
|
||||
u16 sf;
|
||||
|
||||
inline __device__ u16x (const u16 a, const u16 b, const u16 c, const u16 d, const u16 e, const u16 f, const u16 g, const u16 h, const u16 i, const u16 j, const u16 k, const u16 l, const u16 m, const u16 n, const u16 o, const u16 p) : s0(a), s1(b), s2(c), s3(d), s4(e), s5(f), s6(g), s7(h), s8(i), s9(j), sa(k), sb(l), sc(m), sd(n), se(o), sf(p) { }
|
||||
inline __device__ u16x (const u16 a) : s0(a), s1(a), s2(a), s3(a), s4(a), s5(a), s6(a), s7(a), s8(a), s9(a), sa(a), sb(a), sc(a), sd(a), se(a), sf(a) { }
|
||||
|
||||
inline __device__ u16x (void) : s0(0), s1(0), s2(0), s3(0), s4(0), s5(0), s6(0), s7(0), s8(0), s9(0), sa(0), sb(0), sc(0), sd(0), se(0), sf(0){ }
|
||||
inline __device__ ~u16x (void) { }
|
||||
};
|
||||
|
||||
struct __device_builtin__ __builtin_align__(64) u32x
|
||||
{
|
||||
u32 s0;
|
||||
u32 s1;
|
||||
u32 s2;
|
||||
u32 s3;
|
||||
u32 s4;
|
||||
u32 s5;
|
||||
u32 s6;
|
||||
u32 s7;
|
||||
u32 s8;
|
||||
u32 s9;
|
||||
u32 sa;
|
||||
u32 sb;
|
||||
u32 sc;
|
||||
u32 sd;
|
||||
u32 se;
|
||||
u32 sf;
|
||||
|
||||
inline __device__ u32x (const u32 a, const u32 b, const u32 c, const u32 d, const u32 e, const u32 f, const u32 g, const u32 h, const u32 i, const u32 j, const u32 k, const u32 l, const u32 m, const u32 n, const u32 o, const u32 p) : s0(a), s1(b), s2(c), s3(d), s4(e), s5(f), s6(g), s7(h), s8(i), s9(j), sa(k), sb(l), sc(m), sd(n), se(o), sf(p) { }
|
||||
inline __device__ u32x (const u32 a) : s0(a), s1(a), s2(a), s3(a), s4(a), s5(a), s6(a), s7(a), s8(a), s9(a), sa(a), sb(a), sc(a), sd(a), se(a), sf(a) { }
|
||||
|
||||
inline __device__ u32x (void) : s0(0), s1(0), s2(0), s3(0), s4(0), s5(0), s6(0), s7(0), s8(0), s9(0), sa(0), sb(0), sc(0), sd(0), se(0), sf(0){ }
|
||||
inline __device__ ~u32x (void) { }
|
||||
};
|
||||
|
||||
struct __device_builtin__ __builtin_align__(128) u64x
|
||||
{
|
||||
u64 s0;
|
||||
u64 s1;
|
||||
u64 s2;
|
||||
u64 s3;
|
||||
u64 s4;
|
||||
u64 s5;
|
||||
u64 s6;
|
||||
u64 s7;
|
||||
u64 s8;
|
||||
u64 s9;
|
||||
u64 sa;
|
||||
u64 sb;
|
||||
u64 sc;
|
||||
u64 sd;
|
||||
u64 se;
|
||||
u64 sf;
|
||||
|
||||
inline __device__ u64x (const u64 a, const u64 b, const u64 c, const u64 d, const u64 e, const u64 f, const u64 g, const u64 h, const u64 i, const u64 j, const u64 k, const u64 l, const u64 m, const u64 n, const u64 o, const u64 p) : s0(a), s1(b), s2(c), s3(d), s4(e), s5(f), s6(g), s7(h), s8(i), s9(j), sa(k), sb(l), sc(m), sd(n), se(o), sf(p) { }
|
||||
inline __device__ u64x (const u64 a) : s0(a), s1(a), s2(a), s3(a), s4(a), s5(a), s6(a), s7(a), s8(a), s9(a), sa(a), sb(a), sc(a), sd(a), se(a), sf(a) { }
|
||||
|
||||
inline __device__ u64x (void) : s0(0), s1(0), s2(0), s3(0), s4(0), s5(0), s6(0), s7(0), s8(0), s9(0), sa(0), sb(0), sc(0), sd(0), se(0), sf(0) { }
|
||||
inline __device__ ~u64x (void) { }
|
||||
};
|
||||
|
||||
inline __device__ bool operator != (const u32x a, const u32 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) && (a.s8 != b) && (a.s9 != b) && (a.sa != b) && (a.sb != b) && (a.sc != b) && (a.sd != b) && (a.se != b) && (a.sf != b) ); }
|
||||
inline __device__ bool operator != (const u32x a, const u32x 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) && (a.s8 != b.s8) && (a.s9 != b.s9) && (a.sa != b.sa) && (a.sb != b.sb) && (a.sc != b.sc) && (a.sd != b.sd) && (a.se != b.se) && (a.sf != b.sf)); }
|
||||
|
||||
inline __device__ void operator ^= (u32x &a, const u32 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; a.s8 ^= b; a.s9 ^= b; a.sa ^= b; a.sb ^= b; a.sc ^= b; a.sd ^= b; a.se ^= b; a.sf ^= b; }
|
||||
inline __device__ void operator ^= (u32x &a, const u32x 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; a.s8 ^= b.s8; a.s9 ^= b.s9; a.sa ^= b.sa; a.sb ^= b.sb; a.sc ^= b.sc; a.sd ^= b.sd; a.se ^= b.se; a.sf ^= b.sf; }
|
||||
|
||||
inline __device__ void operator |= (u32x &a, const u32 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; a.s8 |= b; a.s9 |= b; a.sa |= b; a.sb |= b; a.sc |= b; a.sd |= b; a.se |= b; a.sf |= b; }
|
||||
inline __device__ void operator |= (u32x &a, const u32x 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; a.s8 |= b.s8; a.s9 |= b.s9; a.sa |= b.sa; a.sb |= b.sb; a.sc |= b.sc; a.sd |= b.sd; a.se |= b.se; a.sf |= b.sf; }
|
||||
|
||||
inline __device__ void operator &= (u32x &a, const u32 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; a.s8 &= b; a.s9 &= b; a.sa &= b; a.sb &= b; a.sc &= b; a.sd &= b; a.se &= b; a.sf &= b; }
|
||||
inline __device__ void operator &= (u32x &a, const u32x 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; a.s8 &= b.s8; a.s9 &= b.s9; a.sa &= b.sa; a.sb &= b.sb; a.sc &= b.sc; a.sd &= b.sd; a.se &= b.se; a.sf &= b.sf; }
|
||||
|
||||
inline __device__ void operator += (u32x &a, const u32 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; a.s8 += b; a.s9 += b; a.sa += b; a.sb += b; a.sc += b; a.sd += b; a.se += b; a.sf += b; }
|
||||
inline __device__ void operator += (u32x &a, const u32x 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; a.s8 += b.s8; a.s9 += b.s9; a.sa += b.sa; a.sb += b.sb; a.sc += b.sc; a.sd += b.sd; a.se += b.se; a.sf += b.sf; }
|
||||
|
||||
inline __device__ void operator -= (u32x &a, const u32 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; a.s8 -= b; a.s9 -= b; a.sa -= b; a.sb -= b; a.sc -= b; a.sd -= b; a.se -= b; a.sf -= b; }
|
||||
inline __device__ void operator -= (u32x &a, const u32x 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; a.s8 -= b.s8; a.s9 -= b.s9; a.sa -= b.sa; a.sb -= b.sb; a.sc -= b.sc; a.sd -= b.sd; a.se -= b.se; a.sf -= b.sf; }
|
||||
|
||||
inline __device__ void operator *= (u32x &a, const u32 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; a.s8 *= b; a.s9 *= b; a.sa *= b; a.sb *= b; a.sc *= b; a.sd *= b; a.se *= b; a.sf *= b; }
|
||||
inline __device__ void operator *= (u32x &a, const u32x 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; a.s8 *= b.s8; a.s9 *= b.s9; a.sa *= b.sa; a.sb *= b.sb; a.sc *= b.sc; a.sd *= b.sd; a.se *= b.se; a.sf *= b.sf; }
|
||||
|
||||
inline __device__ void operator >>= (u32x &a, const u32 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; a.s8 >>= b; a.s9 >>= b; a.sa >>= b; a.sb >>= b; a.sc >>= b; a.sd >>= b; a.se >>= b; a.sf >>= b; }
|
||||
inline __device__ void operator >>= (u32x &a, const u32x 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; a.s8 >>= b.s8; a.s9 >>= b.s9; a.sa >>= b.sa; a.sb >>= b.sb; a.sc >>= b.sc; a.sd >>= b.sd; a.se >>= b.se; a.sf >>= b.sf; }
|
||||
|
||||
inline __device__ void operator <<= (u32x &a, const u32 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; a.s8 <<= b; a.s9 <<= b; a.sa <<= b; a.sb <<= b; a.sc <<= b; a.sd <<= b; a.se <<= b; a.sf <<= b; }
|
||||
inline __device__ void operator <<= (u32x &a, const u32x 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; a.s8 <<= b.s8; a.s9 <<= b.s9; a.sa <<= b.sa; a.sb <<= b.sb; a.sc <<= b.sc; a.sd <<= b.sd; a.se <<= b.se; a.sf <<= b.sf; }
|
||||
|
||||
inline __device__ u32x operator << (const u32x a, const u32 b) { return u32x ((a.s0 << b), (a.s1 << b) , (a.s2 << b), (a.s3 << b) , (a.s4 << b), (a.s5 << b) , (a.s6 << b), (a.s7 << b), (a.s8 << b), (a.s9 << b) , (a.sa << b), (a.sb << b) , (a.sc << b), (a.sd << b) , (a.se << b), (a.sf << b) ); }
|
||||
inline __device__ u32x operator << (const u32x a, const u32x b) { return u32x ((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), (a.s8 << b.s8), (a.s9 << b.s9), (a.sa << b.sa), (a.sb << b.sb), (a.sc << b.sc), (a.sd << b.sd), (a.se << b.se), (a.sf << b.sf)); }
|
||||
|
||||
inline __device__ u32x operator >> (const u32x a, const u32 b) { return u32x ((a.s0 >> b), (a.s1 >> b) , (a.s2 >> b), (a.s3 >> b) , (a.s4 >> b), (a.s5 >> b) , (a.s6 >> b), (a.s7 >> b), (a.s8 >> b), (a.s9 >> b) , (a.sa >> b), (a.sb >> b) , (a.sc >> b), (a.sd >> b) , (a.se >> b), (a.sf >> b) ); }
|
||||
inline __device__ u32x operator >> (const u32x a, const u32x b) { return u32x ((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), (a.s8 >> b.s8), (a.s9 >> b.s9), (a.sa >> b.sa), (a.sb >> b.sb), (a.sc >> b.sc), (a.sd >> b.sd), (a.se >> b.se), (a.sf >> b.sf)); }
|
||||
|
||||
inline __device__ u32x operator ^ (const u32x a, const u32 b) { return u32x ((a.s0 ^ b), (a.s1 ^ b) , (a.s2 ^ b), (a.s3 ^ b) , (a.s4 ^ b), (a.s5 ^ b) , (a.s6 ^ b), (a.s7 ^ b), (a.s8 ^ b), (a.s9 ^ b) , (a.sa ^ b), (a.sb ^ b) , (a.sc ^ b), (a.sd ^ b) , (a.se ^ b), (a.sf ^ b) ); }
|
||||
inline __device__ u32x operator ^ (const u32x a, const u32x b) { return u32x ((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), (a.s8 ^ b.s8), (a.s9 ^ b.s9), (a.sa ^ b.sa), (a.sb ^ b.sb), (a.sc ^ b.sc), (a.sd ^ b.sd), (a.se ^ b.se), (a.sf ^ b.sf)); }
|
||||
|
||||
inline __device__ u32x operator | (const u32x a, const u32 b) { return u32x ((a.s0 | b), (a.s1 | b) , (a.s2 | b), (a.s3 | b) , (a.s4 | b), (a.s5 | b) , (a.s6 | b), (a.s7 | b), (a.s8 | b), (a.s9 | b) , (a.sa | b), (a.sb | b) , (a.sc | b), (a.sd | b) , (a.se | b), (a.sf | b) ); }
|
||||
inline __device__ u32x operator | (const u32x a, const u32x b) { return u32x ((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), (a.s8 | b.s8), (a.s9 | b.s9), (a.sa | b.sa), (a.sb | b.sb), (a.sc | b.sc), (a.sd | b.sd), (a.se | b.se), (a.sf | b.sf)); }
|
||||
|
||||
inline __device__ u32x operator & (const u32x a, const u32 b) { return u32x ((a.s0 & b), (a.s1 & b) , (a.s2 & b), (a.s3 & b) , (a.s4 & b), (a.s5 & b) , (a.s6 & b), (a.s7 & b), (a.s8 & b), (a.s9 & b) , (a.sa & b), (a.sb & b) , (a.sc & b), (a.sd & b) , (a.se & b), (a.sf & b) ); }
|
||||
inline __device__ u32x operator & (const u32x a, const u32x b) { return u32x ((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), (a.s8 & b.s8), (a.s9 & b.s9), (a.sa & b.sa), (a.sb & b.sb), (a.sc & b.sc), (a.sd & b.sd), (a.se & b.se), (a.sf & b.sf)); }
|
||||
|
||||
inline __device__ u32x operator + (const u32x a, const u32 b) { return u32x ((a.s0 + b), (a.s1 + b) , (a.s2 + b), (a.s3 + b) , (a.s4 + b), (a.s5 + b) , (a.s6 + b), (a.s7 + b), (a.s8 + b), (a.s9 + b) , (a.sa + b), (a.sb + b) , (a.sc + b), (a.sd + b) , (a.se + b), (a.sf + b) ); }
|
||||
inline __device__ u32x operator + (const u32x a, const u32x b) { return u32x ((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), (a.s8 + b.s8), (a.s9 + b.s9), (a.sa + b.sa), (a.sb + b.sb), (a.sc + b.sc), (a.sd + b.sd), (a.se + b.se), (a.sf + b.sf)); }
|
||||
|
||||
inline __device__ u32x operator - (const u32x a, const u32 b) { return u32x ((a.s0 - b), (a.s1 - b) , (a.s2 - b), (a.s3 - b) , (a.s4 - b), (a.s5 - b) , (a.s6 - b), (a.s7 - b), (a.s8 - b), (a.s9 - b) , (a.sa - b), (a.sb - b) , (a.sc - b), (a.sd - b) , (a.se - b), (a.sf - b) ); }
|
||||
inline __device__ u32x operator - (const u32x a, const u32x b) { return u32x ((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), (a.s8 - b.s8), (a.s9 - b.s9), (a.sa - b.sa), (a.sb - b.sb), (a.sc - b.sc), (a.sd - b.sd), (a.se - b.se), (a.sf - b.sf)); }
|
||||
|
||||
inline __device__ u32x operator * (const u32x a, const u32 b) { return u32x ((a.s0 * b), (a.s1 * b) , (a.s2 * b), (a.s3 * b) , (a.s4 * b), (a.s5 * b) , (a.s6 * b), (a.s7 * b), (a.s8 * b), (a.s9 * b) , (a.sa * b), (a.sb * b) , (a.sc * b), (a.sd * b) , (a.se * b), (a.sf * b) ); }
|
||||
inline __device__ u32x operator * (const u32x a, const u32x b) { return u32x ((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), (a.s8 * b.s8), (a.s9 * b.s9), (a.sa * b.sa), (a.sb * b.sb), (a.sc * b.sc), (a.sd * b.sd), (a.se * b.se), (a.sf * b.sf)); }
|
||||
|
||||
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, ~a.s8, ~a.s9, ~a.sa, ~a.sb, ~a.sc, ~a.sd, ~a.se, ~a.sf); }
|
||||
|
||||
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) && (a.s8 != b) && (a.s9 != b) && (a.sa != b) && (a.sb != b) && (a.sc != b) && (a.sd != b) && (a.se != b) && (a.sf != 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) && (a.s8 != b.s8) && (a.s9 != b.s9) && (a.sa != b.sa) && (a.sb != b.sb) && (a.sc != b.sc) && (a.sd != b.sd) && (a.se != b.se) && (a.sf != b.sf)); }
|
||||
|
||||
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; a.s8 ^= b; a.s9 ^= b; a.sa ^= b; a.sb ^= b; a.sc ^= b; a.sd ^= b; a.se ^= b; a.sf ^= 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; a.s8 ^= b.s8; a.s9 ^= b.s9; a.sa ^= b.sa; a.sb ^= b.sb; a.sc ^= b.sc; a.sd ^= b.sd; a.se ^= b.se; a.sf ^= b.sf; }
|
||||
|
||||
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; a.s8 |= b; a.s9 |= b; a.sa |= b; a.sb |= b; a.sc |= b; a.sd |= b; a.se |= b; a.sf |= 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; a.s8 |= b.s8; a.s9 |= b.s9; a.sa |= b.sa; a.sb |= b.sb; a.sc |= b.sc; a.sd |= b.sd; a.se |= b.se; a.sf |= b.sf; }
|
||||
|
||||
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; a.s8 &= b; a.s9 &= b; a.sa &= b; a.sb &= b; a.sc &= b; a.sd &= b; a.se &= b; a.sf &= 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; a.s8 &= b.s8; a.s9 &= b.s9; a.sa &= b.sa; a.sb &= b.sb; a.sc &= b.sc; a.sd &= b.sd; a.se &= b.se; a.sf &= b.sf; }
|
||||
|
||||
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; a.s8 += b; a.s9 += b; a.sa += b; a.sb += b; a.sc += b; a.sd += b; a.se += b; a.sf += 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; a.s8 += b.s8; a.s9 += b.s9; a.sa += b.sa; a.sb += b.sb; a.sc += b.sc; a.sd += b.sd; a.se += b.se; a.sf += b.sf; }
|
||||
|
||||
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; a.s8 -= b; a.s9 -= b; a.sa -= b; a.sb -= b; a.sc -= b; a.sd -= b; a.se -= b; a.sf -= 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; a.s8 -= b.s8; a.s9 -= b.s9; a.sa -= b.sa; a.sb -= b.sb; a.sc -= b.sc; a.sd -= b.sd; a.se -= b.se; a.sf -= b.sf; }
|
||||
|
||||
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; a.s8 *= b; a.s9 *= b; a.sa *= b; a.sb *= b; a.sc *= b; a.sd *= b; a.se *= b; a.sf *= 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; a.s8 *= b.s8; a.s9 *= b.s9; a.sa *= b.sa; a.sb *= b.sb; a.sc *= b.sc; a.sd *= b.sd; a.se *= b.se; a.sf *= b.sf; }
|
||||
|
||||
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; a.s8 >>= b; a.s9 >>= b; a.sa >>= b; a.sb >>= b; a.sc >>= b; a.sd >>= b; a.se >>= b; a.sf >>= 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; a.s8 >>= b.s8; a.s9 >>= b.s9; a.sa >>= b.sa; a.sb >>= b.sb; a.sc >>= b.sc; a.sd >>= b.sd; a.se >>= b.se; a.sf >>= b.sf; }
|
||||
|
||||
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; a.s8 <<= b; a.s9 <<= b; a.sa <<= b; a.sb <<= b; a.sc <<= b; a.sd <<= b; a.se <<= b; a.sf <<= 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; a.s8 <<= b.s8; a.s9 <<= b.s9; a.sa <<= b.sa; a.sb <<= b.sb; a.sc <<= b.sc; a.sd <<= b.sd; a.se <<= b.se; a.sf <<= b.sf; }
|
||||
|
||||
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), (a.s8 << b), (a.s9 << b) , (a.sa << b), (a.sb << b) , (a.sc << b), (a.sd << b) , (a.se << b), (a.sf << 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), (a.s8 << b.s8), (a.s9 << b.s9), (a.sa << b.sa), (a.sb << b.sb), (a.sc << b.sc), (a.sd << b.sd), (a.se << b.se), (a.sf << b.sf)); }
|
||||
|
||||
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), (a.s8 >> b), (a.s9 >> b) , (a.sa >> b), (a.sb >> b) , (a.sc >> b), (a.sd >> b) , (a.se >> b), (a.sf >> 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), (a.s8 >> b.s8), (a.s9 >> b.s9), (a.sa >> b.sa), (a.sb >> b.sb), (a.sc >> b.sc), (a.sd >> b.sd), (a.se >> b.se), (a.sf >> b.sf)); }
|
||||
|
||||
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), (a.s8 ^ b), (a.s9 ^ b) , (a.sa ^ b), (a.sb ^ b) , (a.sc ^ b), (a.sd ^ b) , (a.se ^ b), (a.sf ^ 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), (a.s8 ^ b.s8), (a.s9 ^ b.s9), (a.sa ^ b.sa), (a.sb ^ b.sb), (a.sc ^ b.sc), (a.sd ^ b.sd), (a.se ^ b.se), (a.sf ^ b.sf)); }
|
||||
|
||||
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), (a.s8 | b), (a.s9 | b) , (a.sa | b), (a.sb | b) , (a.sc | b), (a.sd | b) , (a.se | b), (a.sf | 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), (a.s8 | b.s8), (a.s9 | b.s9), (a.sa | b.sa), (a.sb | b.sb), (a.sc | b.sc), (a.sd | b.sd), (a.se | b.se), (a.sf | b.sf)); }
|
||||
|
||||
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), (a.s8 & b), (a.s9 & b) , (a.sa & b), (a.sb & b) , (a.sc & b), (a.sd & b) , (a.se & b), (a.sf & 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), (a.s8 & b.s8), (a.s9 & b.s9), (a.sa & b.sa), (a.sb & b.sb), (a.sc & b.sc), (a.sd & b.sd), (a.se & b.se), (a.sf & b.sf)); }
|
||||
|
||||
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), (a.s8 + b), (a.s9 + b) , (a.sa + b), (a.sb + b) , (a.sc + b), (a.sd + b) , (a.se + b), (a.sf + 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), (a.s8 + b.s8), (a.s9 + b.s9), (a.sa + b.sa), (a.sb + b.sb), (a.sc + b.sc), (a.sd + b.sd), (a.se + b.se), (a.sf + b.sf)); }
|
||||
|
||||
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), (a.s8 - b), (a.s9 - b) , (a.sa - b), (a.sb - b) , (a.sc - b), (a.sd - b) , (a.se - b), (a.sf - 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), (a.s8 - b.s8), (a.s9 - b.s9), (a.sa - b.sa), (a.sb - b.sb), (a.sc - b.sc), (a.sd - b.sd), (a.se - b.se), (a.sf - b.sf)); }
|
||||
|
||||
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), (a.s8 * b), (a.s9 * b) , (a.sa * b), (a.sb * b) , (a.sc * b), (a.sd * b) , (a.se * b), (a.sf * 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), (a.s8 * b.s8), (a.s9 * b.s9), (a.sa * b.sa), (a.sb * b.sb), (a.sc * b.sc), (a.sd * b.sd), (a.se * b.se), (a.sf * b.sf)); }
|
||||
|
||||
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, ~a.s8, ~a.s9, ~a.sa, ~a.sb, ~a.sc, ~a.sd, ~a.se, ~a.sf); }
|
||||
|
||||
#endif
|
||||
|
||||
typedef __device_builtin__ struct u8x u8x;
|
||||
typedef __device_builtin__ struct u16x u16x;
|
||||
typedef __device_builtin__ struct u32x u32x;
|
||||
typedef __device_builtin__ struct u64x u64x;
|
||||
|
||||
#define make_u8x u8x
|
||||
#define make_u16x u16x
|
||||
#define make_u32x u32x
|
||||
#define make_u64x u64x
|
||||
|
||||
#else
|
||||
typedef VTYPE(uchar, VECT_SIZE) u8x;
|
||||
typedef VTYPE(ushort, VECT_SIZE) u16x;
|
||||
typedef VTYPE(uint, VECT_SIZE) u32x;
|
||||
typedef VTYPE(ulong, VECT_SIZE) u64x;
|
||||
|
||||
#define make_u8x (u8x)
|
||||
#define make_u16x (u16x)
|
||||
#define make_u32x (u32x)
|
||||
#define make_u64x (u64x)
|
||||
|
||||
#endif
|
||||
#endif
|
||||
|
||||
// unions
|
||||
@ -58,17 +812,19 @@ typedef union vconv32
|
||||
|
||||
struct
|
||||
{
|
||||
u16 v16a;
|
||||
u16 v16b;
|
||||
};
|
||||
u16 a;
|
||||
u16 b;
|
||||
|
||||
} v16;
|
||||
|
||||
struct
|
||||
{
|
||||
u8 v8a;
|
||||
u8 v8b;
|
||||
u8 v8c;
|
||||
u8 v8d;
|
||||
};
|
||||
u8 a;
|
||||
u8 b;
|
||||
u8 c;
|
||||
u8 d;
|
||||
|
||||
} v8;
|
||||
|
||||
} vconv32_t;
|
||||
|
||||
@ -78,29 +834,32 @@ typedef union vconv64
|
||||
|
||||
struct
|
||||
{
|
||||
u32 v32a;
|
||||
u32 v32b;
|
||||
};
|
||||
u32 a;
|
||||
u32 b;
|
||||
|
||||
} v32;
|
||||
|
||||
struct
|
||||
{
|
||||
u16 v16a;
|
||||
u16 v16b;
|
||||
u16 v16c;
|
||||
u16 v16d;
|
||||
};
|
||||
u16 a;
|
||||
u16 b;
|
||||
u16 c;
|
||||
u16 d;
|
||||
|
||||
} v16;
|
||||
|
||||
struct
|
||||
{
|
||||
u8 v8a;
|
||||
u8 v8b;
|
||||
u8 v8c;
|
||||
u8 v8d;
|
||||
u8 v8e;
|
||||
u8 v8f;
|
||||
u8 v8g;
|
||||
u8 v8h;
|
||||
};
|
||||
u8 a;
|
||||
u8 b;
|
||||
u8 c;
|
||||
u8 d;
|
||||
u8 e;
|
||||
u8 f;
|
||||
u8 g;
|
||||
u8 h;
|
||||
|
||||
} v8;
|
||||
|
||||
} vconv64_t;
|
||||
|
||||
|
||||
@ -6,14 +6,33 @@
|
||||
#ifndef _INC_VENDOR_H
|
||||
#define _INC_VENDOR_H
|
||||
|
||||
#ifdef _CPU_OPENCL_EMU_H
|
||||
#if defined _CPU_OPENCL_EMU_H
|
||||
#define IS_NATIVE
|
||||
#elif defined __CUDACC__
|
||||
#define IS_CUDA
|
||||
#else
|
||||
#define IS_OPENCL
|
||||
#endif
|
||||
|
||||
#if defined IS_NATIVE
|
||||
#define CONSTANT_VK
|
||||
#define CONSTANT_AS
|
||||
#define GLOBAL_AS
|
||||
#define LOCAL_VK
|
||||
#define LOCAL_AS
|
||||
#define KERNEL_FQ
|
||||
#else
|
||||
#elif defined IS_CUDA
|
||||
#define CONSTANT_VK __constant__
|
||||
#define CONSTANT_AS
|
||||
#define GLOBAL_AS
|
||||
#define LOCAL_VK __shared__
|
||||
#define LOCAL_AS
|
||||
#define KERNEL_FQ extern "C" __global__
|
||||
#elif defined IS_OPENCL
|
||||
#define CONSTANT_VK __constant
|
||||
#define CONSTANT_AS __constant
|
||||
#define GLOBAL_AS __global
|
||||
#define LOCAL_VK __local
|
||||
#define LOCAL_AS __local
|
||||
#define KERNEL_FQ __kernel
|
||||
#endif
|
||||
|
||||
@ -5,6 +5,7 @@
|
||||
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.h"
|
||||
#include "inc_common.h"
|
||||
#include "inc_cipher_aes.h"
|
||||
#include "inc_cipher_serpent.h"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_rp_optimized.h"
|
||||
#include "inc_rp_optimized.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_rp.h"
|
||||
#include "inc_rp.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_scalar.cl"
|
||||
#include "inc_simd.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_scalar.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
@ -16,7 +17,7 @@
|
||||
#define MD5_STEP_REV(f,a,b,c,d,x,t,s) \
|
||||
{ \
|
||||
a -= b; \
|
||||
a = hc_rotr32_S (a, s); \
|
||||
a = hc_rotr32_S (a, s); \
|
||||
a -= f (b, c, d); \
|
||||
a -= x; \
|
||||
a -= t; \
|
||||
@ -25,7 +26,7 @@
|
||||
#define MD5_STEP_REV1(f,a,b,c,d,x,t,s) \
|
||||
{ \
|
||||
a -= b; \
|
||||
a = hc_rotr32_S (a, s); \
|
||||
a = hc_rotr32_S (a, s); \
|
||||
a -= x; \
|
||||
a -= t; \
|
||||
}
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_rp_optimized.h"
|
||||
#include "inc_rp_optimized.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_rp.h"
|
||||
#include "inc_rp.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_scalar.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_rp_optimized.h"
|
||||
#include "inc_rp_optimized.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_rp.h"
|
||||
#include "inc_rp.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_scalar.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_rp_optimized.h"
|
||||
#include "inc_rp_optimized.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_rp.h"
|
||||
#include "inc_rp.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_scalar.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_rp_optimized.h"
|
||||
#include "inc_rp_optimized.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_rp.h"
|
||||
#include "inc_rp.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_scalar.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_rp_optimized.h"
|
||||
#include "inc_rp_optimized.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_rp.h"
|
||||
#include "inc_rp.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_scalar.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_rp_optimized.h"
|
||||
#include "inc_rp_optimized.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_rp.h"
|
||||
#include "inc_rp.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_scalar.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_rp_optimized.h"
|
||||
#include "inc_rp_optimized.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_rp.h"
|
||||
#include "inc_rp.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
#include "inc_hash_sha1.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_scalar.cl"
|
||||
#include "inc_hash_sha1.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
#include "inc_hash_sha1.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
#include "inc_hash_sha1.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_rp_optimized.h"
|
||||
#include "inc_rp_optimized.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_rp.h"
|
||||
#include "inc_rp.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
#include "inc_hash_sha1.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_scalar.cl"
|
||||
#include "inc_hash_sha1.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
#include "inc_hash_sha1.cl"
|
||||
|
||||
@ -8,6 +8,7 @@
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
#include "inc_hash_sha1.cl"
|
||||
|
||||
Some files were not shown because too many files have changed in this diff Show More
Loading…
Reference in New Issue
Block a user