diff --git a/run/opencl/7z_kernel.cl b/run/opencl/7z_kernel.cl index d9fdec92f7..077e6fb5a7 100644 --- a/run/opencl/7z_kernel.cl +++ b/run/opencl/7z_kernel.cl @@ -46,7 +46,7 @@ } \ } -inline void sha256_zerofinal(uint *W, uint *output, const uint tot_len) +INLINE void sha256_zerofinal(uint *W, uint *output, const uint tot_len) { uint len = ((tot_len & 63) >> 2) + 1; diff --git a/run/opencl/axcrypt_kernel.cl b/run/opencl/axcrypt_kernel.cl index b4ad69511f..411403170c 100644 --- a/run/opencl/axcrypt_kernel.cl +++ b/run/opencl/axcrypt_kernel.cl @@ -43,7 +43,7 @@ typedef struct { (cp)[2] ^= (uchar)((value) >> 16), \ (cp)[3] ^= (uchar)((value) >> 24 ) ) -inline int axcrypt_decrypt(__global const axcrypt_password *inbuffer, uint gid, __constant axcrypt_salt *cur_salt, __global axcrypt_out *output, __local aes_local_t *lt) +INLINE int axcrypt_decrypt(__global const axcrypt_password *inbuffer, uint gid, __constant axcrypt_salt *cur_salt, __global axcrypt_out *output, __local aes_local_t *lt) { uchar password[PLAINTEXT_LENGTH]; uchar keyfile[4096]; diff --git a/run/opencl/bitlocker_kernel.cl b/run/opencl/bitlocker_kernel.cl index e903c387e7..4986b4e601 100644 --- a/run/opencl/bitlocker_kernel.cl +++ b/run/opencl/bitlocker_kernel.cl @@ -24,7 +24,7 @@ #define TS2 Te2 #define TS3 Te3 -inline unsigned int OPT3_XOR(unsigned int a, unsigned int b, unsigned int c) +INLINE unsigned int OPT3_XOR(unsigned int a, unsigned int b, unsigned int c) { #if HAVE_LUT3 return lut3(a, b, c, 0x96); @@ -33,7 +33,7 @@ inline unsigned int OPT3_XOR(unsigned int a, unsigned int b, unsigned int c) #endif } -inline unsigned int OPT3_XORAND(unsigned int a, unsigned int b, unsigned int c) +INLINE unsigned int OPT3_XORAND(unsigned int a, unsigned int b, unsigned int c) { #if HAVE_LUT3 return lut3(a, b, c, 0xb8); @@ -42,7 +42,7 @@ inline unsigned int OPT3_XORAND(unsigned int a, unsigned int b, unsigned int c) #endif } -inline unsigned int OPT3_ANDOR(unsigned int a, unsigned int b, unsigned int c) +INLINE unsigned int OPT3_ANDOR(unsigned int a, unsigned int b, unsigned int c) { #if HAVE_LUT3 return lut3(a, b, c, 0xe8); diff --git a/run/opencl/cryptmd5_kernel.cl b/run/opencl/cryptmd5_kernel.cl index d4b03057c5..bd23c44b0b 100644 --- a/run/opencl/cryptmd5_kernel.cl +++ b/run/opencl/cryptmd5_kernel.cl @@ -140,7 +140,7 @@ __constant uchar g[] = 1*8, 7*8, 3*8, 5*8, 2*8, 7*8, 1*8, 7*8, 3*8, 5*8, 3*8, 6*8, 1*8, 7*8, 3*8, 5*8, 3*8, 7*8 }; #ifdef BUF_UPDATE_SWITCH -inline void buf_update(uint *buf, uint a, uint b, uint c, uint d, uint offset) +INLINE void buf_update(uint *buf, uint a, uint b, uint c, uint d, uint offset) { uint i = offset >> 2; switch (offset & 3) { @@ -198,7 +198,7 @@ inline void buf_update(uint *buf, uint a, uint b, uint c, uint d, uint offset) } } #else -inline void buf_update(uint *buf, uint a, uint b, uint c, uint d, uint offset) +INLINE void buf_update(uint *buf, uint a, uint b, uint c, uint d, uint offset) { uint i = offset >> 2; uint j = offset & 3; @@ -228,7 +228,7 @@ inline void buf_update(uint *buf, uint a, uint b, uint c, uint d, uint offset) } #endif -inline void ctx_update(md5_ctx *ctx, uchar *string, uint len, +INLINE void ctx_update(md5_ctx *ctx, uchar *string, uint len, uchar *ctx_buflen) { uint i; @@ -239,7 +239,7 @@ inline void ctx_update(md5_ctx *ctx, uchar *string, uint len, *ctx_buflen += len; } -inline void ctx_update_prefix(md5_ctx *ctx, uchar prefix, uchar *ctx_buflen) +INLINE void ctx_update_prefix(md5_ctx *ctx, uchar prefix, uchar *ctx_buflen) { uint i; @@ -257,7 +257,7 @@ inline void ctx_update_prefix(md5_ctx *ctx, uchar prefix, uchar *ctx_buflen) // else if (prefix == '\0') do nothing. for {smd5} } -inline void init_ctx(md5_ctx *ctx, uchar *ctx_buflen) +INLINE void init_ctx(md5_ctx *ctx, uchar *ctx_buflen) { #if __OS_X__ @@ -280,7 +280,7 @@ inline void init_ctx(md5_ctx *ctx, uchar *ctx_buflen) *ctx_buflen = 0; } -inline void md5_digest(uint *x, uint *y, uint *z, uint *zmem, uint zmem_offset, uint len, uint unify) +INLINE void md5_digest(uint *x, uint *y, uint *z, uint *zmem, uint zmem_offset, uint len, uint unify) { uint a; uint b = 0xefcdab89; diff --git a/run/opencl/cryptosafe_kernel.cl b/run/opencl/cryptosafe_kernel.cl index 6ff9ce49f5..bbb178d0ed 100644 --- a/run/opencl/cryptosafe_kernel.cl +++ b/run/opencl/cryptosafe_kernel.cl @@ -15,7 +15,7 @@ typedef struct { unsigned char ciphertext[16]; } salt_t; -inline void prepare(__global const uchar *pwbuf, __global const uint *index, uchar *password) +INLINE void prepare(__global const uchar *pwbuf, __global const uint *index, uchar *password) { uint i; uint gid = get_global_id(0); diff --git a/run/opencl/ed25519-donna/ed25519-donna-impl-base.h b/run/opencl/ed25519-donna/ed25519-donna-impl-base.h index 440430d123..1dee53925d 100644 --- a/run/opencl/ed25519-donna/ed25519-donna-impl-base.h +++ b/run/opencl/ed25519-donna/ed25519-donna-impl-base.h @@ -122,7 +122,7 @@ ge25519_scalarmult_base_choose_niels(ge25519_niels *t, uint32_t pos, signed char /* computes [s]basepoint */ static void #if gpu_nvidia(DEVICE_INFO) && DEV_VER_MAJOR > 525 && DEV_VER_MAJOR < 560 -__attribute__((noinline)) +NOINLINE #endif ge25519_scalarmult_base_niels(ge25519 *r, const bignum256modm s) { signed char b[64]; diff --git a/run/opencl/gpg_kernel.cl b/run/opencl/gpg_kernel.cl index be2db62515..df50dec67e 100644 --- a/run/opencl/gpg_kernel.cl +++ b/run/opencl/gpg_kernel.cl @@ -46,7 +46,7 @@ typedef struct { #define LEAN #endif -inline void S2KItSaltedSHA1Generator(__global const uchar *password, +INLINE void S2KItSaltedSHA1Generator(__global const uchar *password, uint password_length, __constant uchar *salt, uint _count, @@ -161,7 +161,7 @@ __kernel void gpg(__global const gpg_password *inbuffer, /* SHA-256 based S2K */ -inline void S2KItSaltedSHA256Generator(__global const uchar *ipassword, +INLINE void S2KItSaltedSHA256Generator(__global const uchar *ipassword, uint password_length, __constant uchar *isalt, uint count, // iterations @@ -227,7 +227,7 @@ __kernel void gpg_sha256(__global const gpg_password *inbuffer, /* SHA-512 based S2K */ -inline void S2KItSaltedSHA512Generator(__global const uchar *ipassword, +INLINE void S2KItSaltedSHA512Generator(__global const uchar *ipassword, uint password_length, __constant uchar *isalt, uint count, // iterations diff --git a/run/opencl/keepass_kernel.cl b/run/opencl/keepass_kernel.cl index 63b1e779ab..28558b49da 100644 --- a/run/opencl/keepass_kernel.cl +++ b/run/opencl/keepass_kernel.cl @@ -86,7 +86,7 @@ typedef struct { uint8_t hash[32]; } keepass_state; -inline void calc_hmac_base_key(const void *master_seed, const void *final_key, void *result) +INLINE void calc_hmac_base_key(const void *master_seed, const void *final_key, void *result) { const uint8_t one_le[1] = "\x01"; SHA512_CTX ctx; @@ -98,7 +98,7 @@ inline void calc_hmac_base_key(const void *master_seed, const void *final_key, v SHA512_Final(result, &ctx); } -inline void calc_hmac_key(const void *block_index, const void *base_key, void *result) +INLINE void calc_hmac_key(const void *block_index, const void *base_key, void *result) { SHA512_CTX ctx; diff --git a/run/opencl/keyring_kernel.cl b/run/opencl/keyring_kernel.cl index 86feacb23a..9b0ab312b7 100644 --- a/run/opencl/keyring_kernel.cl +++ b/run/opencl/keyring_kernel.cl @@ -48,7 +48,7 @@ typedef struct { uchar ct[LINE_BUFFER_SIZE / 2]; /* after hex conversion */ } keyring_salt; -inline int verify_decrypted_buffer(uchar *buffer, int len) +INLINE int verify_decrypted_buffer(uchar *buffer, int len) { uchar digest[16]; MD5_CTX ctx; diff --git a/run/opencl/krb5pa-md5_kernel.cl b/run/opencl/krb5pa-md5_kernel.cl index a45ff98cd8..8fa5c2e6c8 100644 --- a/run/opencl/krb5pa-md5_kernel.cl +++ b/run/opencl/krb5pa-md5_kernel.cl @@ -26,7 +26,7 @@ #ifdef UTF_8 -inline +INLINE void prepare(const __global uint *key, uint length, MAYBE_VOLATILE uint *nt_buffer) { @@ -97,7 +97,7 @@ void prepare(const __global uint *key, uint length, #else -inline +INLINE void prepare(const __global uint *key, uint length, uint *nt_buffer) { uint i, nt_index, keychars; @@ -118,7 +118,7 @@ void prepare(const __global uint *key, uint length, uint *nt_buffer) #define asciidigit(n) ((n) >= '0' && (n) <= '9') -inline +INLINE void krb5pa_md5_final(const uint *K, MAYBE_CONSTANT uint *salts, #ifdef RC4_USE_LOCAL @@ -257,7 +257,7 @@ void krb5pa_md5_final(const uint *K, md5_block(uint, block, K2); /* md5_update(ihash, 16), md5_final() */ } -inline +INLINE void cmp_final(uint gid, uint iter, uint *hash, @@ -305,7 +305,7 @@ void cmp_final(uint gid, } } -inline +INLINE void cmp(uint gid, uint iter, uint *hash, diff --git a/run/opencl/krb5tgs_kernel.cl b/run/opencl/krb5tgs_kernel.cl index dd77d0f922..ad5619ad55 100644 --- a/run/opencl/krb5tgs_kernel.cl +++ b/run/opencl/krb5tgs_kernel.cl @@ -40,7 +40,7 @@ typedef struct { #ifdef UTF_8 -inline +INLINE void prepare_utf16(__global const uchar *source, __global const uint *index, nt_buffer_t *nt_buffer) @@ -116,7 +116,7 @@ void prepare_utf16(__global const uchar *source, #else -inline +INLINE void prepare_utf16(__global const uchar *password, __global const uint *index, nt_buffer_t *nt_buffer) diff --git a/run/opencl/lotus5_kernel.cl b/run/opencl/lotus5_kernel.cl index f05d86d3c9..731bee3ff2 100644 --- a/run/opencl/lotus5_kernel.cl +++ b/run/opencl/lotus5_kernel.cl @@ -51,7 +51,7 @@ __constant uint magic_table[256] = { 0x29, 0x39, 0xb9, 0xe9, 0x4c, 0xff, 0x43, 0xab, }; -inline void +INLINE void lotus_transform_password (unsigned int *i1, unsigned int *o1, MAYBE_LOCAL unsigned int *lotus_magic_table) { @@ -71,7 +71,7 @@ lotus_transform_password (unsigned int *i1, unsigned int *o1, } /* The mixing function: perturbs the first three rows of the matrix */ -inline void +INLINE void lotus_mix (unsigned int *m1, MAYBE_LOCAL unsigned int *lotus_magic_table) { int i, j, k; diff --git a/run/opencl/md4_kernel.cl b/run/opencl/md4_kernel.cl index 52ea69b513..4da8c74b9c 100644 --- a/run/opencl/md4_kernel.cl +++ b/run/opencl/md4_kernel.cl @@ -56,7 +56,7 @@ /* This handles an input of 0xffffffffU correctly */ #define BITMAP_SHIFT ((BITMAP_MASK >> 5) + 1) -inline void md4_encrypt(uint *hash, uint *W, uint len) +INLINE void md4_encrypt(uint *hash, uint *W, uint len) { hash[0] = 0x67452301; hash[1] = 0xefcdab89; @@ -118,7 +118,7 @@ inline void md4_encrypt(uint *hash, uint *W, uint len) STEP(H2, hash[1], hash[2], hash[3], hash[0], W[15] + 0x6ed9eba1, 15); } -inline void cmp_final(uint gid, +INLINE void cmp_final(uint gid, uint iter, uint *hash, __global uint *offset_table, @@ -164,7 +164,7 @@ inline void cmp_final(uint gid, } } -inline void cmp(uint gid, +INLINE void cmp(uint gid, uint iter, uint *hash, #if USE_LOCAL_BITMAPS diff --git a/run/opencl/md5_kernel.cl b/run/opencl/md5_kernel.cl index 6d89578da1..7fdf4b0f8d 100644 --- a/run/opencl/md5_kernel.cl +++ b/run/opencl/md5_kernel.cl @@ -62,7 +62,7 @@ /* This handles an input of 0xffffffffU correctly */ #define BITMAP_SHIFT ((BITMAP_MASK >> 5) + 1) -inline void md5_encrypt(uint *hash, uint *W, uint len) +INLINE void md5_encrypt(uint *hash, uint *W, uint len) { hash[0] = 0x67452301; hash[1] = 0xefcdab89; @@ -142,7 +142,7 @@ inline void md5_encrypt(uint *hash, uint *W, uint len) STEP(I, hash[1], hash[2], hash[3], hash[0], W[9], 0xeb86d391, 21); } -inline void cmp_final(uint gid, +INLINE void cmp_final(uint gid, uint iter, uint *hash, __global uint *offset_table, @@ -188,7 +188,7 @@ inline void cmp_final(uint gid, } } -inline void cmp(uint gid, +INLINE void cmp(uint gid, uint iter, uint *hash, #if USE_LOCAL_BITMAPS diff --git a/run/opencl/md5x50.h b/run/opencl/md5x50.h index 89c85e70d3..a2bdb9b9e1 100644 --- a/run/opencl/md5x50.h +++ b/run/opencl/md5x50.h @@ -13,7 +13,7 @@ #define INIT_C 0x98badcfe #define INIT_D 0x10325476 -inline void md5x50_40(uint* msg) +INLINE void md5x50_40(uint* msg) { uint a, b, c, d; int i; @@ -130,7 +130,7 @@ inline void md5x50_40(uint* msg) msg[3] = 0; } -inline void md5x50_128(uint* msg) +INLINE void md5x50_128(uint* msg) { uint a, b, c, d; int i; diff --git a/run/opencl/mscash_kernel.cl b/run/opencl/mscash_kernel.cl index 13a314d949..31a41f65cc 100644 --- a/run/opencl/mscash_kernel.cl +++ b/run/opencl/mscash_kernel.cl @@ -20,7 +20,7 @@ #define SQRT_2 0x5a827999 #define SQRT_3 0x6ed9eba1 -inline void md4_crypt_a(uint *hash, uint *nt_buffer) +INLINE void md4_crypt_a(uint *hash, uint *nt_buffer) { unsigned int a = INIT_A; unsigned int b = INIT_B; @@ -146,7 +146,7 @@ inline void md4_crypt_a(uint *hash, uint *nt_buffer) hash[3] = d + INIT_D; } -inline void md4_crypt_b(uint *hash, constant uint *salt) +INLINE void md4_crypt_b(uint *hash, constant uint *salt) { unsigned int a = INIT_A; unsigned int b = INIT_B; @@ -281,7 +281,7 @@ inline void md4_crypt_b(uint *hash, constant uint *salt) #if UTF_8 -inline void prepare_key(__global uint *key, uint length, +INLINE void prepare_key(__global uint *key, uint length, MAYBE_VOLATILE uint *nt_buffer) { const __global UTF8 *source = (const __global uchar*)key; @@ -351,7 +351,7 @@ inline void prepare_key(__global uint *key, uint length, #else -inline void prepare_key(__global uint *key, uint length, uint *nt_buffer) +INLINE void prepare_key(__global uint *key, uint length, uint *nt_buffer) { uint i, nt_index, keychars; @@ -369,7 +369,7 @@ inline void prepare_key(__global uint *key, uint length, uint *nt_buffer) #endif /* UTF_8 */ -inline void cmp_final(uint gid, +INLINE void cmp_final(uint gid, uint iter, uint *hash, __global uint *offset_table, @@ -416,7 +416,7 @@ inline void cmp_final(uint gid, } } -inline void cmp(uint gid, +INLINE void cmp(uint gid, uint iter, uint *hash, __global uint *bitmaps, diff --git a/run/opencl/nt_kernel.cl b/run/opencl/nt_kernel.cl index 3be75fbeea..e64c656dd4 100644 --- a/run/opencl/nt_kernel.cl +++ b/run/opencl/nt_kernel.cl @@ -59,7 +59,7 @@ /* This handles an input of 0xffffffffU correctly */ #define BITMAP_SHIFT ((BITMAP_MASK >> 5) + 1) -inline int nt_crypt(uint *hash, uint *nt_buffer, uint md4_size, BITMAPS_TYPE uint *bitmaps) +INLINE int nt_crypt(uint *hash, uint *nt_buffer, uint md4_size, BITMAPS_TYPE uint *bitmaps) { MD4_G_VARS @@ -282,7 +282,7 @@ inline int nt_crypt(uint *hash, uint *nt_buffer, uint md4_size, BITMAPS_TYPE uin #if UTF_8 -inline uint prepare_key(__global uint *key, uint length, +INLINE uint prepare_key(__global uint *key, uint length, MAYBE_VOLATILE uint *nt_buffer) { const __global UTF8 *source = (const __global UTF8*)key; @@ -351,7 +351,7 @@ inline uint prepare_key(__global uint *key, uint length, #else -inline uint prepare_key(__global uint *key, uint length, uint *nt_buffer) +INLINE uint prepare_key(__global uint *key, uint length, uint *nt_buffer) { uint i, nt_index, keychars; @@ -369,7 +369,7 @@ inline uint prepare_key(__global uint *key, uint length, uint *nt_buffer) #endif /* UTF_8 */ -inline void cmp_final(uint gid, +INLINE void cmp_final(uint gid, uint iter, uint *hash, __global uint *offset_table, @@ -399,7 +399,7 @@ inline void cmp_final(uint gid, } } -inline void cmp(uint gid, +INLINE void cmp(uint gid, uint iter, uint *hash, BITMAPS_TYPE uint *bitmaps, diff --git a/run/opencl/ntlmv2_kernel.cl b/run/opencl/ntlmv2_kernel.cl index 2186acbbac..39544a9185 100644 --- a/run/opencl/ntlmv2_kernel.cl +++ b/run/opencl/ntlmv2_kernel.cl @@ -26,7 +26,7 @@ #ifdef UTF_8 -inline +INLINE void prepare_key(const __global uint *key, uint length, MAYBE_VOLATILE uint *nt_buffer) { @@ -97,7 +97,7 @@ void prepare_key(const __global uint *key, uint length, #else -inline +INLINE void prepare_key(const __global uint *key, uint length, uint *nt_buffer) { uint i, nt_index, keychars; @@ -116,7 +116,7 @@ void prepare_key(const __global uint *key, uint length, uint *nt_buffer) #endif /* encodings */ -inline +INLINE void ntlmv2_final(uint *nthash, MAYBE_CONSTANT uint *challenge, uint *output) { uint block[16]; @@ -209,7 +209,7 @@ void ntlmv2_final(uint *nthash, MAYBE_CONSTANT uint *challenge, uint *output) md5_block(uint, block, output); /* md5_update(hash, 16), md5_final() */ } -inline +INLINE void cmp_final(uint gid, uint iter, uint *hash, @@ -257,7 +257,7 @@ void cmp_final(uint gid, } } -inline +INLINE void cmp(uint gid, uint iter, uint *hash, diff --git a/run/opencl/office_kernel.cl b/run/opencl/office_kernel.cl index a796e705a3..8339698612 100644 --- a/run/opencl/office_kernel.cl +++ b/run/opencl/office_kernel.cl @@ -237,7 +237,7 @@ void Final2007(__global ms_office_state *state, out[gid].cracked = result; } -inline void Decrypt(__constant ms_office_salt *salt, +INLINE void Decrypt(__constant ms_office_salt *salt, const uchar *verifierInputKey, __constant uchar *encryptedVerifier, uchar *decryptedVerifier, diff --git a/run/opencl/oldoffice_kernel.cl b/run/opencl/oldoffice_kernel.cl index 0f35d86195..c072a54829 100644 --- a/run/opencl/oldoffice_kernel.cl +++ b/run/opencl/oldoffice_kernel.cl @@ -40,7 +40,7 @@ typedef struct { #ifdef UTF_8 -inline +INLINE void oldoffice_utf16(__global const uchar *source, __global const uint *index, nt_buffer_t *nt_buffer) @@ -116,7 +116,7 @@ void oldoffice_utf16(__global const uchar *source, #else -inline +INLINE void oldoffice_utf16(__global const uchar *password, __global const uint *index, nt_buffer_t *nt_buffer) @@ -148,7 +148,7 @@ void oldoffice_utf16(__global const uchar *password, #define MAYBE_VOLATILE #endif -inline +INLINE void oldoffice_md5(const nt_buffer_t *nt_buffer, __global salt_t *cs, __global uint *result, @@ -361,7 +361,7 @@ void oldoffice_md5(const nt_buffer_t *nt_buffer, } } -inline +INLINE void oldoffice_sha1(const nt_buffer_t *nt_buffer, __global salt_t *cs, __global uint *result, diff --git a/run/opencl/opencl_aes.h b/run/opencl/opencl_aes.h index e2b62c2e59..94260a3d94 100644 --- a/run/opencl/opencl_aes.h +++ b/run/opencl/opencl_aes.h @@ -62,7 +62,7 @@ */ #ifndef AES_ecb_encrypt -inline void +INLINE void AES_ecb_encrypt(AES_SRC_TYPE void *_in, AES_DST_TYPE void *_out, uint len, AES_KEY *akey) { @@ -93,7 +93,7 @@ AES_ecb_encrypt(AES_SRC_TYPE void *_in, AES_DST_TYPE void *_out, uint len, #endif } -inline void +INLINE void AES_ecb_encrypt_pp(const void *_in, void *_out, uint len, AES_KEY *akey) { const uchar *in = _in; @@ -110,7 +110,7 @@ AES_ecb_encrypt_pp(const void *_in, void *_out, uint len, AES_KEY *akey) #endif /* AES_ecb_encrypt */ #ifndef AES_ecb_decrypt -inline void +INLINE void AES_ecb_decrypt(AES_SRC_TYPE void *_in, AES_DST_TYPE void *_out, uint len, AES_KEY *akey) { @@ -141,7 +141,7 @@ AES_ecb_decrypt(AES_SRC_TYPE void *_in, AES_DST_TYPE void *_out, uint len, #endif } -inline void +INLINE void AES_ecb_decrypt_pp(const void *_in, void *_out, uint len, AES_KEY *akey) { const uchar *in = _in; @@ -157,7 +157,7 @@ AES_ecb_decrypt_pp(const void *_in, void *_out, uint len, AES_KEY *akey) } #endif /* AES_ecb_decrypt */ -inline void +INLINE void AES_cbc_encrypt(AES_SRC_TYPE void *_in, AES_DST_TYPE void *_out, uint len, AES_KEY *akey, void *_iv) { @@ -189,7 +189,7 @@ AES_cbc_encrypt(AES_SRC_TYPE void *_in, AES_DST_TYPE void *_out, * This function decrypts two blocks at a time, to utilize * that our bitsliced AES can do them in parallel. */ -inline void +INLINE void AES_cbc_decrypt(AES_SRC_TYPE void *_in, AES_DST_TYPE void *_out, uint len, AES_KEY *akey, void *iv) { @@ -220,7 +220,7 @@ AES_cbc_decrypt(AES_SRC_TYPE void *_in, AES_DST_TYPE void *_out, } } -inline void +INLINE void AES_cts_encrypt(AES_CTS_SRC_TYPE void *_in, AES_CTS_DST_TYPE void *_out, uint len, AES_KEY *akey, void *_iv) { @@ -252,7 +252,7 @@ AES_cts_encrypt(AES_CTS_SRC_TYPE void *_in, AES_CTS_DST_TYPE void *_out, memcpy_macro(iv, out - AES_BLOCK_SIZE, AES_BLOCK_SIZE); } -inline void +INLINE void AES_cts_decrypt(AES_CTS_SRC_TYPE void *_in, AES_CTS_DST_TYPE void *_out, uint len, AES_KEY *akey, void *_iv) { @@ -290,7 +290,7 @@ AES_cts_decrypt(AES_CTS_SRC_TYPE void *_in, AES_CTS_DST_TYPE void *_out, memcpy_macro(iv, tmp, AES_BLOCK_SIZE); } -inline void AES_cfb_decrypt(AES_SRC_TYPE void *_in, +INLINE void AES_cfb_decrypt(AES_SRC_TYPE void *_in, AES_DST_TYPE void *_out, uint len, AES_KEY *akey, void *_iv) { @@ -313,7 +313,7 @@ inline void AES_cfb_decrypt(AES_SRC_TYPE void *_in, } } -inline void AES_256_XTS_first_sector(AES_SRC_TYPE uint *in, AES_DST_TYPE uint *out, +INLINE void AES_256_XTS_first_sector(AES_SRC_TYPE uint *in, AES_DST_TYPE uint *out, AES_KEY_TYPE uchar *double_key, __local aes_local_t *lt) { uint tweak[4] = { 0 }; @@ -336,7 +336,7 @@ inline void AES_256_XTS_first_sector(AES_SRC_TYPE uint *in, AES_DST_TYPE uint *o out[i] = buf[i] ^ tweak[i]; } -inline void AES_256_XTS_DiskCryptor(AES_SRC_TYPE uchar *data, AES_DST_TYPE uchar *output, +INLINE void AES_256_XTS_DiskCryptor(AES_SRC_TYPE uchar *data, AES_DST_TYPE uchar *output, AES_KEY_TYPE uchar *double_key, int len, __local aes_local_t *lt) { uchar buf[16]; @@ -383,7 +383,7 @@ inline void AES_256_XTS_DiskCryptor(AES_SRC_TYPE uchar *data, AES_DST_TYPE uchar #define N_WORDS (AES_BLOCK_SIZE / sizeof(unsigned long)) -inline void +INLINE void AES_ige_decrypt(AES_SRC_TYPE void *_in, AES_DST_TYPE void *_out, uint length, AES_KEY *akey, uchar *_iv) { diff --git a/run/opencl/opencl_aes_bitslice.h b/run/opencl/opencl_aes_bitslice.h index c698c71f15..1bbc6af997 100644 --- a/run/opencl/opencl_aes_bitslice.h +++ b/run/opencl/opencl_aes_bitslice.h @@ -37,7 +37,7 @@ typedef struct aes_ctx { __local aes_local_t *lt; /* Not currently used yet but here for interoperability */ } AES_CTX; -inline void +INLINE void enc32le(void *dst, uint32_t x) { uchar *buf = dst; @@ -48,7 +48,7 @@ enc32le(void *dst, uint32_t x) buf[3] = (uchar)(x >> 24); } -inline void +INLINE void enc32le_dst(AES_DST_TYPE void *dst, uint32_t x) { AES_DST_TYPE uchar *buf = dst; @@ -59,7 +59,7 @@ enc32le_dst(AES_DST_TYPE void *dst, uint32_t x) buf[3] = (uchar)(x >> 24); } -inline uint32_t +INLINE uint32_t dec32le(const void *src) { const uchar *buf = (const uchar*)src; @@ -70,7 +70,7 @@ dec32le(const void *src) | ((uint32_t)buf[3] << 24); } -inline uint32_t +INLINE uint32_t dec32le_src(AES_SRC_TYPE void *src) { AES_SRC_TYPE uchar *buf = src; @@ -81,7 +81,7 @@ dec32le_src(AES_SRC_TYPE void *src) | ((uint32_t)buf[3] << 24); } -inline uint32_t +INLINE uint32_t dec32le_key(AES_KEY_TYPE void *src) { AES_KEY_TYPE uchar *buf = src; @@ -130,7 +130,7 @@ dec32le_key(AES_KEY_TYPE void *src) * parallel. Bits 0 to 7 of each S-box input (bit 0 is least significant) * are spread over the words 0 to 7, at the same rank. */ -inline void +INLINE void aes_ct_bitslice_Sbox(uint32_t *q) { /* @@ -314,7 +314,7 @@ aes_ct_bitslice_Sbox(uint32_t *q) * * This operation is an involution. */ -inline void +INLINE void aes_ct_ortho(uint32_t *q) { #define SWAPN(cl, ch, s, x, y) do { \ @@ -345,7 +345,7 @@ aes_ct_ortho(uint32_t *q) SWAP8(q[3], q[7]); } -inline uint32_t +INLINE uint32_t sub_word(uint32_t x) { uint32_t q[8]; @@ -365,7 +365,7 @@ sub_word(uint32_t x) * below. Subkeys are produced in little-endian convention (but not * bitsliced). Key length is expressed in bytes. */ -inline uint +INLINE uint aes_keysched_base(uint32_t *skey, AES_KEY_TYPE void *key, size_t key_len) { uint num_rounds; @@ -418,7 +418,7 @@ aes_keysched_base(uint32_t *skey, AES_KEY_TYPE void *key, size_t key_len) * on key size). The number of rounds is returned. If the key size is * invalid (not 16, 24 or 32), then 0 is returned. */ -inline uint +INLINE uint aes_ct_keysched(uint32_t *comp_skey, AES_KEY_TYPE void *key, size_t key_len) { uint32_t skey[60]; @@ -450,7 +450,7 @@ aes_ct_keysched(uint32_t *comp_skey, AES_KEY_TYPE void *key, size_t key_len) * a larger array suitable for aes_ct_bitslice_encrypt() and * aes_ct_bitslice_decrypt(). */ -inline void +INLINE void aes_ct_skey_expand(uint32_t *skey, uint num_rounds, const uint32_t *comp_skey) { @@ -468,7 +468,7 @@ aes_ct_skey_expand(uint32_t *skey, } } -inline void +INLINE void add_round_key(uint32_t *q, const uint32_t *sk) { q[0] ^= sk[0]; @@ -481,7 +481,7 @@ add_round_key(uint32_t *q, const uint32_t *sk) q[7] ^= sk[7]; } -inline void +INLINE void shift_rows(uint32_t *q) { int i; @@ -499,7 +499,7 @@ shift_rows(uint32_t *q) #define rotr16(x) rotate((uint32_t)x, 16U) -inline void +INLINE void mix_columns(uint32_t *q) { uint32_t q0, q1, q2, q3, q4, q5, q6, q7; @@ -537,7 +537,7 @@ mix_columns(uint32_t *q) * eight 32-bit words, two block encryptions are actually performed * in parallel. */ -inline void +INLINE void aes_ct_bitslice_encrypt(uint num_rounds, const uint32_t *skey, uint32_t *q) { @@ -558,7 +558,7 @@ aes_ct_bitslice_encrypt(uint num_rounds, /* * Like aes_ct_bitslice_Sbox(), but for the inverse S-box. */ -inline void +INLINE void aes_ct_bitslice_invSbox(uint32_t *q) { /* @@ -618,7 +618,7 @@ aes_ct_bitslice_invSbox(uint32_t *q) q[0] = q2 ^ q5 ^ q7; } -inline void +INLINE void inv_shift_rows(uint32_t *q) { int i; @@ -634,7 +634,7 @@ inv_shift_rows(uint32_t *q) } } -inline void +INLINE void inv_mix_columns(uint32_t *q) { uint32_t q0, q1, q2, q3, q4, q5, q6, q7; @@ -672,7 +672,7 @@ inv_mix_columns(uint32_t *q) * eight 32-bit words, two block decryptions are actually performed * in parallel. */ -inline void +INLINE void aes_ct_bitslice_decrypt(uint num_rounds, const uint32_t *skey, uint32_t *q) { @@ -690,7 +690,7 @@ aes_ct_bitslice_decrypt(uint num_rounds, add_round_key(q, skey); } -inline int +INLINE int AES_Setkey(AES_CTX *ctx, AES_KEY_TYPE void *key, int len) { ctx->num_rounds = aes_ct_keysched(ctx->sk, (AES_KEY_TYPE char*)key, len); @@ -700,7 +700,7 @@ AES_Setkey(AES_CTX *ctx, AES_KEY_TYPE void *key, int len) return 0; } -inline void +INLINE void AES_Encrypt_ECB_pp(AES_CTX *ctx, const void *_src, void *_dst, size_t num_blocks) { @@ -746,7 +746,7 @@ AES_Encrypt_ECB_pp(AES_CTX *ctx, const void *_src, } } -inline void +INLINE void AES_Encrypt_ECB(AES_CTX *ctx, AES_SRC_TYPE void *_src, AES_DST_TYPE void *_dst, size_t num_blocks) { @@ -792,7 +792,7 @@ AES_Encrypt_ECB(AES_CTX *ctx, AES_SRC_TYPE void *_src, } } -inline void +INLINE void AES_Decrypt_ECB_pp(AES_CTX *ctx, const void *_src, void *_dst, size_t num_blocks) { @@ -838,7 +838,7 @@ AES_Decrypt_ECB_pp(AES_CTX *ctx, const void *_src, } } -inline void +INLINE void AES_Decrypt_ECB(AES_CTX *ctx, AES_SRC_TYPE void *_src, AES_DST_TYPE void *_dst, size_t num_blocks) { @@ -884,7 +884,7 @@ AES_Decrypt_ECB(AES_CTX *ctx, AES_SRC_TYPE void *_src, } } -inline int +INLINE int AES_KeySetup_Encrypt(uint32_t *skey, AES_KEY_TYPE uint8_t *key, int len) { uint r, u; @@ -910,7 +910,7 @@ AES_KeySetup_Encrypt(uint32_t *skey, AES_KEY_TYPE uint8_t *key, int len) * Reduce value x modulo polynomial x^8+x^4+x^3+x+1. This works as * long as x fits on 12 bits at most. */ -inline uint32_t +INLINE uint32_t redgf256(uint32_t x) { uint32_t h; @@ -922,7 +922,7 @@ redgf256(uint32_t x) /* * Multiplication by 0x09 in GF(256). */ -inline uint32_t +INLINE uint32_t mul9(uint32_t x) { return redgf256(x ^ (x << 3)); @@ -931,7 +931,7 @@ mul9(uint32_t x) /* * Multiplication by 0x0B in GF(256). */ -inline uint32_t +INLINE uint32_t mulb(uint32_t x) { return redgf256(x ^ (x << 1) ^ (x << 3)); @@ -940,7 +940,7 @@ mulb(uint32_t x) /* * Multiplication by 0x0D in GF(256). */ -inline uint32_t +INLINE uint32_t muld(uint32_t x) { return redgf256(x ^ (x << 2) ^ (x << 3)); @@ -949,13 +949,13 @@ muld(uint32_t x) /* * Multiplication by 0x0E in GF(256). */ -inline uint32_t +INLINE uint32_t mule(uint32_t x) { return redgf256((x << 1) ^ (x << 2) ^ (x << 3)); } -inline int +INLINE int AES_KeySetup_Decrypt(uint32_t *skey, AES_KEY_TYPE uint8_t *key, int len) { uint i, r, u; diff --git a/run/opencl/opencl_aes_plain.h b/run/opencl/opencl_aes_plain.h index 4fac64fbec..a1a15d0948 100644 --- a/run/opencl/opencl_aes_plain.h +++ b/run/opencl/opencl_aes_plain.h @@ -80,7 +80,7 @@ typedef struct aes_key_st { /** * Copy tables to local memory */ -inline void aes_table_init(__local aes_local_t *lt) +INLINE void aes_table_init(__local aes_local_t *lt) { for (uint i = THREAD; i < 256; i += LWS) { lt->Te0[i] = Te0[i]; @@ -116,7 +116,7 @@ inline void aes_table_init(__local aes_local_t *lt) /** * Expand the cipher key into the encryption key schedule. */ -inline void AES_set_encrypt_key(AES_KEY_TYPE void *_userKey, +INLINE void AES_set_encrypt_key(AES_KEY_TYPE void *_userKey, const int bits, AES_KEY *key) { AES_KEY_TYPE uchar *userKey = _userKey; @@ -217,7 +217,7 @@ inline void AES_set_encrypt_key(AES_KEY_TYPE void *_userKey, /** * Expand the cipher key into the decryption key schedule. */ -inline void AES_set_decrypt_key(AES_KEY_TYPE void *_userKey, +INLINE void AES_set_decrypt_key(AES_KEY_TYPE void *_userKey, const int bits, AES_KEY *key) { AES_KEY_TYPE uchar *userKey = _userKey; @@ -267,7 +267,7 @@ inline void AES_set_decrypt_key(AES_KEY_TYPE void *_userKey, /* * Encrypt a single block. */ -inline void AES_encrypt(const uchar *in, uchar *out, const AES_KEY *key) +INLINE void AES_encrypt(const uchar *in, uchar *out, const AES_KEY *key) { const u32 *rk; u32 s0, s1, s2, s3, t0, t1, t2, t3; @@ -453,7 +453,7 @@ inline void AES_encrypt(const uchar *in, uchar *out, const AES_KEY *key) /* * Decrypt a single block. */ -inline void AES_decrypt(const uchar *in, uchar *out, const AES_KEY *key) +INLINE void AES_decrypt(const uchar *in, uchar *out, const AES_KEY *key) { const u32 *rk; u32 s0, s1, s2, s3, t0, t1, t2, t3; diff --git a/run/opencl/opencl_asn1.h b/run/opencl/opencl_asn1.h index 811f2dbdbb..4b57dfc6de 100644 --- a/run/opencl/opencl_asn1.h +++ b/run/opencl/opencl_asn1.h @@ -58,7 +58,7 @@ struct asn1_oid { size_t len; }; -inline +INLINE int asn1_get_next(const uint8_t *buf, size_t len, struct asn1_hdr *hdr) { const uint8_t *pos, *end; diff --git a/run/opencl/opencl_cast.h b/run/opencl/opencl_cast.h index a29c8e5c38..a5739c798e 100644 --- a/run/opencl/opencl_cast.h +++ b/run/opencl/opencl_cast.h @@ -586,7 +586,7 @@ __constant uint S[8][256] = { #define _CAST_F2(l, r, i, j) _CAST_f2(l, r, K[i], K[i+j]) #define _CAST_F3(l, r, i, j) _CAST_f3(l, r, K[i], K[i+j]) -inline void Cast5Encrypt(const uchar *inBlock, uchar *outBlock, CAST_KEY *key) +INLINE void Cast5Encrypt(const uchar *inBlock, uchar *outBlock, CAST_KEY *key) { uint l; GET_UINT32BE(l, inBlock, 0); uint r; GET_UINT32BE(r, inBlock, 4); @@ -616,7 +616,7 @@ inline void Cast5Encrypt(const uchar *inBlock, uchar *outBlock, CAST_KEY *key) PUT_UINT32BE(l, outBlock, 4); } -inline void Cast5Decrypt(const uchar *inBlock, uchar *outBlock, CAST_KEY *key) +INLINE void Cast5Decrypt(const uchar *inBlock, uchar *outBlock, CAST_KEY *key) { uint l; GET_UINT32BE(l, inBlock, 0); uint r; GET_UINT32BE(r, inBlock, 4); @@ -647,7 +647,7 @@ inline void Cast5Decrypt(const uchar *inBlock, uchar *outBlock, CAST_KEY *key) t = l = r = 0; } -inline void Cast5SetKey(CAST_KEY *key, uint keylength, const uchar *userKey) +INLINE void Cast5SetKey(CAST_KEY *key, uint keylength, const uchar *userKey) { uint i; uint *K = key->K; diff --git a/run/opencl/opencl_chacha.h b/run/opencl/opencl_chacha.h index b54132e832..04a7dd2718 100644 --- a/run/opencl/opencl_chacha.h +++ b/run/opencl/opencl_chacha.h @@ -77,7 +77,7 @@ typedef struct chacha_ctx_s { __constant char sigma[16] = "expand 32-byte k"; __constant char tau[16] = "expand 16-byte k"; -inline +INLINE void chacha_keysetup(chacha_ctx *x, CHACHA_KEY_TYPE uchar *k, uint kbits) { __constant char *constants; @@ -102,7 +102,7 @@ void chacha_keysetup(chacha_ctx *x, CHACHA_KEY_TYPE uchar *k, uint kbits) x->input[3] = U8TO32_LITTLE(constants + 12); } -inline +INLINE void chacha_ivsetup(chacha_ctx *x, CHACHA_IV_TYPE uchar *iv, const uchar *counter, uint length) { if (length == 0 || length == 8) { @@ -118,7 +118,7 @@ void chacha_ivsetup(chacha_ctx *x, CHACHA_IV_TYPE uchar *iv, const uchar *counte } } -inline +INLINE void chacha_encrypt_bytes(chacha_ctx *x, CHACHA_SRC_TYPE uchar *m, CHACHA_DST_TYPE uchar *c, uint bytes) { uint x0, x1, x2, x3, x4, x5, x6, x7, x8, x9, x10, x11, x12, x13, x14, x15; @@ -253,7 +253,7 @@ void chacha_encrypt_bytes(chacha_ctx *x, CHACHA_SRC_TYPE uchar *m, CHACHA_DST_TY } } -inline +INLINE void chacha_decrypt_bytes(chacha_ctx *x, CHACHA_SRC_TYPE uchar *c, CHACHA_DST_TYPE uchar *m, uint bytes) { chacha_encrypt_bytes(x, c, m, bytes); diff --git a/run/opencl/opencl_cmac.h b/run/opencl/opencl_cmac.h index 743eaf572c..01dd6d1fb6 100644 --- a/run/opencl/opencl_cmac.h +++ b/run/opencl/opencl_cmac.h @@ -46,7 +46,7 @@ typedef struct _AES_CMAC_CTX { (r)[i] ^= (v)[i]; \ } while (0) -inline void +INLINE void AES_CMAC_Init(AES_CMAC_CTX *ctx, __local aes_local_t *lt) { uint i; @@ -57,13 +57,13 @@ AES_CMAC_Init(AES_CMAC_CTX *ctx, __local aes_local_t *lt) ctx->aesctx.lt = lt; } -inline void +INLINE void AES_CMAC_SetKey(AES_CMAC_CTX *ctx, const uint8_t *key) { AES_set_encrypt_key(key, 128, &ctx->aesctx); } -inline void +INLINE void AES_CMAC_Update(AES_CMAC_CTX *ctx, MAYBE_CONSTANT uint8_t *data, uint len) { uint i; @@ -95,7 +95,7 @@ AES_CMAC_Update(AES_CMAC_CTX *ctx, MAYBE_CONSTANT uint8_t *data, uint len) ctx->M_n = len; } -inline void +INLINE void AES_CMAC_Final(uint8_t *digest, AES_CMAC_CTX *ctx) { uint8_t K[16] = { 0 }; diff --git a/run/opencl/opencl_crc32.h b/run/opencl/opencl_crc32.h index a3162b6aae..4e51ed73cd 100644 --- a/run/opencl/opencl_crc32.h +++ b/run/opencl/opencl_crc32.h @@ -47,12 +47,12 @@ __constant CRC32_t CRC32_table[256] = { #define POLY 0xEDB88320 #define ALL1 0xFFFFFFFF -inline void CRC32_Init(CRC32_t *value) +INLINE void CRC32_Init(CRC32_t *value) { *value = ALL1; } -inline void CRC32_Update(CRC32_t *value, void *data, uint count) +INLINE void CRC32_Update(CRC32_t *value, void *data, uint count) { uchar *ptr = (uchar*)data; CRC32_t result = *value; @@ -63,7 +63,7 @@ inline void CRC32_Update(CRC32_t *value, void *data, uint count) *value = result; } -inline void CRC32_Final(uchar *out, CRC32_t value) +INLINE void CRC32_Final(uchar *out, CRC32_t value) { value = ~value; out[0] = value; diff --git a/run/opencl/opencl_des.h b/run/opencl/opencl_des.h index acbcac5ad3..61f9cb122b 100644 --- a/run/opencl/opencl_des.h +++ b/run/opencl/opencl_des.h @@ -296,7 +296,7 @@ __constant uchar odd_parity_table[128] = { 1, 2, 4, 7, 8, 227, 229, 230, 233, 234, 236, 239, 241, 242, 244, 247, 248, 251, 253, 254 }; -inline void des_key_set_parity(uchar key[DES_KEY_SIZE]) +INLINE void des_key_set_parity(uchar key[DES_KEY_SIZE]) { int i; @@ -307,7 +307,7 @@ inline void des_key_set_parity(uchar key[DES_KEY_SIZE]) /* * Check the given key's parity, returns 1 on failure, 0 on SUCCESS */ -inline int des_key_check_key_parity(const uchar key[DES_KEY_SIZE]) +INLINE int des_key_check_key_parity(const uchar key[DES_KEY_SIZE]) { int i; @@ -319,7 +319,7 @@ inline int des_key_check_key_parity(const uchar key[DES_KEY_SIZE]) } #endif -inline void des_setkey(uint32_t SK[32], const uchar key[DES_KEY_SIZE]) +INLINE void des_setkey(uint32_t SK[32], const uchar key[DES_KEY_SIZE]) { int i; uint32_t X, Y, T; @@ -391,7 +391,7 @@ inline void des_setkey(uint32_t SK[32], const uchar key[DES_KEY_SIZE]) /* * DES key schedule (56-bit, encryption) */ -inline void des_setkey_enc(des_context *ctx, const uchar key[DES_KEY_SIZE]) +INLINE void des_setkey_enc(des_context *ctx, const uchar key[DES_KEY_SIZE]) { des_setkey(ctx->sk, key); } @@ -399,7 +399,7 @@ inline void des_setkey_enc(des_context *ctx, const uchar key[DES_KEY_SIZE]) /* * DES key schedule (56-bit, decryption) */ -inline void des_setkey_dec(des_context *ctx, const uchar key[DES_KEY_SIZE]) +INLINE void des_setkey_dec(des_context *ctx, const uchar key[DES_KEY_SIZE]) { int i; @@ -411,7 +411,7 @@ inline void des_setkey_dec(des_context *ctx, const uchar key[DES_KEY_SIZE]) } } -inline void des3_set2key(uint32_t esk[96], uint32_t dsk[96], +INLINE void des3_set2key(uint32_t esk[96], uint32_t dsk[96], const uchar key[DES_KEY_SIZE * 2]) { int i; @@ -437,7 +437,7 @@ inline void des3_set2key(uint32_t esk[96], uint32_t dsk[96], /* * Triple-DES key schedule (112-bit, encryption) */ -inline void des3_set2key_enc(des3_context *ctx, +INLINE void des3_set2key_enc(des3_context *ctx, const uchar key[DES_KEY_SIZE * 2]) { uint32_t sk[96]; @@ -448,7 +448,7 @@ inline void des3_set2key_enc(des3_context *ctx, /* * Triple-DES key schedule (112-bit, decryption) */ -inline void des3_set2key_dec(des3_context *ctx, +INLINE void des3_set2key_dec(des3_context *ctx, const uchar key[DES_KEY_SIZE * 2]) { uint32_t sk[96]; @@ -456,7 +456,7 @@ inline void des3_set2key_dec(des3_context *ctx, des3_set2key(sk, ctx->sk, key); } -inline void des3_set3key(uint32_t esk[96], +INLINE void des3_set3key(uint32_t esk[96], uint32_t dsk[96], const uchar key[24]) { @@ -481,7 +481,7 @@ inline void des3_set3key(uint32_t esk[96], /* * Triple-DES key schedule (168-bit, encryption) */ -inline void des3_set3key_enc(des3_context *ctx, +INLINE void des3_set3key_enc(des3_context *ctx, const uchar key[DES_KEY_SIZE * 3]) { uint32_t sk[96]; @@ -492,7 +492,7 @@ inline void des3_set3key_enc(des3_context *ctx, /* * Triple-DES key schedule (168-bit, decryption) */ -inline void des3_set3key_dec(des3_context *ctx, +INLINE void des3_set3key_dec(des3_context *ctx, const uchar key[DES_KEY_SIZE * 3]) { uint32_t sk[96]; @@ -503,7 +503,7 @@ inline void des3_set3key_dec(des3_context *ctx, /* * DES-ECB block encryption/decryption */ -inline void des_crypt_ecb(des_context *ctx, const uchar *input, +INLINE void des_crypt_ecb(des_context *ctx, const uchar *input, uchar *output) { int i; @@ -530,7 +530,7 @@ inline void des_crypt_ecb(des_context *ctx, const uchar *input, /* * DES-CBC buffer encryption/decryption */ -inline void des_crypt_cbc(des_context *ctx, int mode, size_t length, +INLINE void des_crypt_cbc(des_context *ctx, int mode, size_t length, uchar *iv, const uchar *input, uchar *output) { @@ -570,7 +570,7 @@ inline void des_crypt_cbc(des_context *ctx, int mode, size_t length, /* * 3DES-ECB block encryption/decryption */ -inline void des3_crypt_ecb(des3_context *ctx, const uchar *input, +INLINE void des3_crypt_ecb(des3_context *ctx, const uchar *input, uchar *output) { int i; @@ -639,7 +639,7 @@ inline void des3_crypt_ecb(des3_context *ctx, const uchar *input, } #if 0 /* the above macro replaces this function */ -inline void des3_crypt_cbc(des3_context *ctx, int mode, size_t length, +INLINE void des3_crypt_cbc(des3_context *ctx, int mode, size_t length, uchar *iv, const uchar *input, uchar *output) { diff --git a/run/opencl/opencl_gost94.h b/run/opencl/opencl_gost94.h index 3965d72f4f..b12dd1b4e7 100644 --- a/run/opencl/opencl_gost94.h +++ b/run/opencl/opencl_gost94.h @@ -76,7 +76,7 @@ typedef struct { * * @param ctx context to initialize */ -inline void gost94_init(gost94_ctx *ctx) +INLINE void gost94_init(gost94_ctx *ctx) { memset_p(ctx, 0, sizeof(gost94_ctx)); } @@ -124,7 +124,7 @@ inline void gost94_init(gost94_ctx *ctx) * @param hash intermediate message hash * @param block the message block to process */ -inline void rhash_gost94_block_compress(gost94_ctx *ctx, const uint* block, MAYBE_LOCAL const rhash_gost94_sbox *sbox) +INLINE void rhash_gost94_block_compress(gost94_ctx *ctx, const uint* block, MAYBE_LOCAL const rhash_gost94_sbox *sbox) { uint i; uint key[8], u[8], v[8], w[8], s[8]; @@ -266,7 +266,7 @@ inline void rhash_gost94_block_compress(gost94_ctx *ctx, const uint* block, MAYB * @param ctx algorithm context * @param block the 256-bit message block to process */ -inline void rhash_gost94_compute_sum_and_hash(gost94_ctx * ctx, const uint* block, MAYBE_LOCAL const rhash_gost94_sbox *sbox) +INLINE void rhash_gost94_compute_sum_and_hash(gost94_ctx * ctx, const uint* block, MAYBE_LOCAL const rhash_gost94_sbox *sbox) { #if !__ENDIAN_LITTLE__ uint block_le[8]; /* tmp buffer for little endian number */ @@ -300,7 +300,7 @@ inline void rhash_gost94_compute_sum_and_hash(gost94_ctx * ctx, const uint* bloc * @param msg message chunk * @param size length of the message chunk */ -__attribute__((noinline)) void gost94_update(gost94_ctx *ctx, const uchar* msg, uint size, MAYBE_LOCAL const rhash_gost94_sbox *sbox) +NOINLINE void gost94_update(gost94_ctx *ctx, const uchar* msg, uint size, MAYBE_LOCAL const rhash_gost94_sbox *sbox) { uint index = ctx->length & 31; ctx->length += size; @@ -339,7 +339,7 @@ __attribute__((noinline)) void gost94_update(gost94_ctx *ctx, const uchar* msg, } #if !__ENDIAN_LITTLE__ -inline void rhash_u32_swap_copy(void* to, const void* from, uint length) { +INLINE void rhash_u32_swap_copy(void* to, const void* from, uint length) { uint i; uint *pO, *pI; pO = (uint *)to; @@ -361,7 +361,7 @@ inline void rhash_u32_swap_copy(void* to, const void* from, uint length) { * @param ctx the algorithm context containing current hashing state * @param result calculated hash in binary form */ -__attribute__((noinline)) void gost94_final(gost94_ctx *ctx, uchar *result, MAYBE_LOCAL const rhash_gost94_sbox *sbox) +NOINLINE void gost94_final(gost94_ctx *ctx, uchar *result, MAYBE_LOCAL const rhash_gost94_sbox *sbox) { uint index = ctx->length & 31; uint* msg32 = (uint*)ctx->message; @@ -553,7 +553,7 @@ __constant uchar sbox[8][16] = { * it at run-time can save a little space in the executable file * in trade of consuming some time at program start. */ -inline void gost94_init_table(MAYBE_LOCAL rhash_gost94_sbox *cur_sbox) +INLINE void gost94_init_table(MAYBE_LOCAL rhash_gost94_sbox *cur_sbox) { uint i; #if GOST94_FLAT_INIT diff --git a/run/opencl/opencl_hmac_md5.h b/run/opencl/opencl_hmac_md5.h index 89f46e10a7..e2fcc083a5 100644 --- a/run/opencl/opencl_hmac_md5.h +++ b/run/opencl/opencl_hmac_md5.h @@ -24,7 +24,7 @@ #define HMAC_OUT_TYPE #endif -inline void hmac_md5(HMAC_KEY_TYPE void *_key, uint key_len, +INLINE void hmac_md5(HMAC_KEY_TYPE void *_key, uint key_len, HMAC_MSG_TYPE void *_data, uint data_len, HMAC_OUT_TYPE void *_digest, uint digest_len) { diff --git a/run/opencl/opencl_hmac_sha1.h b/run/opencl/opencl_hmac_sha1.h index aecb9967c3..2c27898c19 100644 --- a/run/opencl/opencl_hmac_sha1.h +++ b/run/opencl/opencl_hmac_sha1.h @@ -24,7 +24,7 @@ #define HMAC_OUT_TYPE #endif -inline void hmac_sha1(HMAC_KEY_TYPE void *_key, uint key_len, +INLINE void hmac_sha1(HMAC_KEY_TYPE void *_key, uint key_len, HMAC_MSG_TYPE void *_data, uint data_len, HMAC_OUT_TYPE void *_digest, uint digest_len) { diff --git a/run/opencl/opencl_hmac_sha256.h b/run/opencl/opencl_hmac_sha256.h index 89a42d3fae..b7d7999f15 100644 --- a/run/opencl/opencl_hmac_sha256.h +++ b/run/opencl/opencl_hmac_sha256.h @@ -24,7 +24,7 @@ #define HMAC_OUT_TYPE #endif -inline void hmac_sha256(HMAC_KEY_TYPE void *_key, uint key_len, +INLINE void hmac_sha256(HMAC_KEY_TYPE void *_key, uint key_len, HMAC_MSG_TYPE void *_data, uint data_len, HMAC_OUT_TYPE void *_digest, uint digest_len) { diff --git a/run/opencl/opencl_hmac_sha512.h b/run/opencl/opencl_hmac_sha512.h index 37b69b1b8c..3ba3be8021 100644 --- a/run/opencl/opencl_hmac_sha512.h +++ b/run/opencl/opencl_hmac_sha512.h @@ -24,7 +24,7 @@ #define HMAC_OUT_TYPE #endif -inline void hmac_sha512(HMAC_KEY_TYPE void *_key, uint key_len, +INLINE void hmac_sha512(HMAC_KEY_TYPE void *_key, uint key_len, HMAC_MSG_TYPE void *_data, uint data_len, HMAC_OUT_TYPE void *_digest, uint digest_len) { diff --git a/run/opencl/opencl_md5_ctx.h b/run/opencl/opencl_md5_ctx.h index 10be4789b5..94eb839b8b 100644 --- a/run/opencl/opencl_md5_ctx.h +++ b/run/opencl/opencl_md5_ctx.h @@ -19,7 +19,7 @@ typedef struct { uchar buffer[64]; /* data block being processed */ } MD5_CTX; -inline void _md5_process(MD5_CTX *ctx, const uchar data[64]) +INLINE void _md5_process(MD5_CTX *ctx, const uchar data[64]) { uint W[16], A, B, C, D; @@ -78,7 +78,7 @@ inline void _md5_process(MD5_CTX *ctx, const uchar data[64]) /* * MD5 context setup */ -inline void MD5_Init(MD5_CTX *ctx) +INLINE void MD5_Init(MD5_CTX *ctx) { ctx->total = 0; @@ -91,7 +91,7 @@ inline void MD5_Init(MD5_CTX *ctx) /* * MD5 process buffer */ -inline void MD5_Update(MD5_CTX *ctx, const uchar *input, uint ilen) +INLINE void MD5_Update(MD5_CTX *ctx, const uchar *input, uint ilen) { uint fill; uint left; @@ -129,7 +129,7 @@ inline void MD5_Update(MD5_CTX *ctx, const uchar *input, uint ilen) /* * MD5 final digest */ -inline void MD5_Final(uchar output[20], MD5_CTX *ctx) +INLINE void MD5_Final(uchar output[20], MD5_CTX *ctx) { uint last, padn; ulong bits; diff --git a/run/opencl/opencl_misc.h b/run/opencl/opencl_misc.h index 37fb8d0c75..c991e6c9df 100644 --- a/run/opencl/opencl_misc.h +++ b/run/opencl/opencl_misc.h @@ -18,7 +18,7 @@ #include "opencl_device_info.h" -/* Note: long is *always* 64-bit in OpenCL */ +/* long is always 64-bit in OpenCL while long long is reserved for 128 bits */ typedef uchar uint8_t; typedef char int8_t; typedef ushort uint16_t; @@ -52,21 +52,18 @@ typedef uint32_t host_size_t; #endif /* - * Some runtimes/drivers breaks on using inline, others breaks on lack of it, - * yet others require use of static as well. - * - * Only usable in device code + * Most runtimes will inline nearly everything without request. */ #if _OPENCL_COMPILER +#define NOINLINE __attribute__((noinline)) + #if __MESA__ -#define inline // empty! +#define INLINE #elif __POCL__ -// Do nothing (POCL complains if we redefine) -#elif gpu_amd(DEVICE_INFO) // We really target ROCM here -#define inline static inline +#define INLINE inline #else -// Do nothing +#define INLINE static inline #endif #endif /* _OPENCL_COMPILER */ @@ -101,7 +98,7 @@ typedef struct dyna_salt_t { #if SCALAR && 0 /* Used for testing */ #define HAVE_LUT3 1 -inline uint lut3(uint x, uint y, uint z, uchar m) +INLINE uint lut3(uint x, uint y, uint z, uchar m) { uint i; uint r = 0; @@ -132,7 +129,7 @@ inline uint lut3(uint x, uint y, uint z, uchar m) #if SCALAR && SM_MAJOR >= 5 && (DEV_VER_MAJOR > 352 || (DEV_VER_MAJOR == 352 && DEV_VER_MINOR >= 21)) #define HAVE_LUT3 1 -inline uint lut3(uint a, uint b, uint c, uint imm) +INLINE uint lut3(uint a, uint b, uint c, uint imm) { uint r; asm("lop3.b32 %0, %1, %2, %3, %4;" @@ -143,7 +140,7 @@ inline uint lut3(uint a, uint b, uint c, uint imm) #if 0 /* This does no good */ #define HAVE_LUT3_64 1 -inline ulong lut3_64(ulong a, ulong b, ulong c, uint imm) +INLINE ulong lut3_64(ulong a, ulong b, ulong c, uint imm) { ulong t, r; @@ -163,7 +160,7 @@ inline ulong lut3_64(ulong a, ulong b, ulong c, uint imm) #pragma OPENCL EXTENSION cl_amd_media_ops : enable #define BITALIGN(hi, lo, s) amd_bitalign((hi), (lo), (s)) #elif SCALAR && SM_MAJOR > 3 || (SM_MAJOR == 3 && SM_MINOR >= 2) -inline uint funnel_shift_right(uint hi, uint lo, uint s) +INLINE uint funnel_shift_right(uint hi, uint lo, uint s) { uint r; asm("shf.r.wrap.b32 %0, %1, %2, %3;" @@ -172,7 +169,7 @@ inline uint funnel_shift_right(uint hi, uint lo, uint s) return r; } -inline uint funnel_shift_right_imm(uint hi, uint lo, uint s) +INLINE uint funnel_shift_right_imm(uint hi, uint lo, uint s) { uint r; asm("shf.r.wrap.b32 %0, %1, %2, %3;" @@ -205,13 +202,13 @@ inline uint funnel_shift_right_imm(uint hi, uint lo, uint s) #define block_swap32(W, len) for (uint i = 0; i < len; i++) W[i] = SWAP32(W[i]) #define block_swap64(W, len) for (uint i = 0; i < len; i++) W[i] = SWAP64(W[i]) -inline ushort SWAP16(ushort x) +INLINE ushort SWAP16(ushort x) { return ((x << 8) + (x >> 8)); } #if USE_BITSELECT -inline uint SWAP32(uint x) +INLINE uint SWAP32(uint x) { return bitselect(rotate(x, 24U), rotate(x, 8U), 0x00FF00FFU); } @@ -223,7 +220,7 @@ inline uint SWAP32(uint x) rotate(n, 40UL), 0x00FF000000FF0000UL), \ 0xFFFF0000FFFF0000UL) #else -inline uint SWAP32(uint x) +INLINE uint SWAP32(uint x) { x = rotate(x, 16U); return ((x & 0x00FF00FF) << 8) + ((x >> 8) & 0x00FF00FF); @@ -241,7 +238,7 @@ inline uint SWAP32(uint x) #define VSWAP32 SWAP32 #else /* Vector-capable swap32() */ -inline MAYBE_VECTOR_UINT VSWAP32(MAYBE_VECTOR_UINT x) +INLINE MAYBE_VECTOR_UINT VSWAP32(MAYBE_VECTOR_UINT x) { x = rotate(x, 16U); return ((x & 0x00FF00FF) << 8) + ((x >> 8) & 0x00FF00FF); @@ -379,7 +376,7 @@ inline MAYBE_VECTOR_UINT VSWAP32(MAYBE_VECTOR_UINT x) #define XORCHAR_BE(buf, index, val) ((uchar*)(buf))[(index) ^ 3] ^= (val) #endif -inline int check_pkcs_pad(const uchar *data, int len, int blocksize) +INLINE int check_pkcs_pad(const uchar *data, int len, int blocksize) { int pad_len, padding, real_len; @@ -424,7 +421,7 @@ inline int check_pkcs_pad(const uchar *data, int len, int blocksize) */ /* src and dst are private mem */ -inline void memcpy_pp(void* restrict dst, const void* restrict src, uint count) +INLINE void memcpy_pp(void* restrict dst, const void* restrict src, uint count) { const char *s = src; char *d = dst; @@ -434,7 +431,7 @@ inline void memcpy_pp(void* restrict dst, const void* restrict src, uint count) } /* src is private mem, dst is global mem */ -inline void memcpy_pg(__global void* restrict dst, const void* restrict src, uint count) +INLINE void memcpy_pg(__global void* restrict dst, const void* restrict src, uint count) { const char *s = src; __global char *d = dst; @@ -444,7 +441,7 @@ inline void memcpy_pg(__global void* restrict dst, const void* restrict src, uin } /* src is global mem, dst is private mem */ -inline void memcpy_gp(void* restrict dst, __global const void* restrict src, uint count) +INLINE void memcpy_gp(void* restrict dst, __global const void* restrict src, uint count) { __global const char *s = src; char *d = dst; @@ -454,7 +451,7 @@ inline void memcpy_gp(void* restrict dst, __global const void* restrict src, uin } /* src is constant mem, dst is private mem */ -inline void memcpy_cp(void* restrict dst, __constant void* restrict src, uint count) +INLINE void memcpy_cp(void* restrict dst, __constant void* restrict src, uint count) { __constant char *s = src; char *d = dst; @@ -464,7 +461,7 @@ inline void memcpy_cp(void* restrict dst, __constant void* restrict src, uint co } /* src is MAYBE_CONSTANT mem, dst is private mem */ -inline void memcpy_mcp(void* restrict dst, MAYBE_CONSTANT void* restrict src, uint count) +INLINE void memcpy_mcp(void* restrict dst, MAYBE_CONSTANT void* restrict src, uint count) { MAYBE_CONSTANT char *s = src; char *d = dst; @@ -474,7 +471,7 @@ inline void memcpy_mcp(void* restrict dst, MAYBE_CONSTANT void* restrict src, ui } /* dst is private mem */ -inline void memset_p(void *p, uint val, uint count) +INLINE void memset_p(void *p, uint val, uint count) { char *d = p; @@ -483,7 +480,7 @@ inline void memset_p(void *p, uint val, uint count) } /* dst is global mem */ -inline void memset_g(__global void *p, uint val, uint count) +INLINE void memset_g(__global void *p, uint val, uint count) { __global char *d = p; @@ -492,7 +489,7 @@ inline void memset_g(__global void *p, uint val, uint count) } /* s1 and s2 are private mem */ -inline int memcmp_pp(const void *s1, const void *s2, uint size) +INLINE int memcmp_pp(const void *s1, const void *s2, uint size) { const uchar *a = s1; const uchar *b = s2; @@ -505,7 +502,7 @@ inline int memcmp_pp(const void *s1, const void *s2, uint size) } /* s1 is private mem, s2 is global mem */ -inline int memcmp_pg(const void *s1, __global const void *s2, uint size) +INLINE int memcmp_pg(const void *s1, __global const void *s2, uint size) { const uchar *a = s1; __global const uchar *b = s2; @@ -518,7 +515,7 @@ inline int memcmp_pg(const void *s1, __global const void *s2, uint size) } /* s1 is private mem, s2 is constant mem */ -inline int memcmp_pc(const void *s1, __constant const void *s2, uint size) +INLINE int memcmp_pc(const void *s1, __constant const void *s2, uint size) { const uchar *a = s1; __constant const uchar *b = s2; @@ -531,7 +528,7 @@ inline int memcmp_pc(const void *s1, __constant const void *s2, uint size) } /* s1 is global mem, s2 is constant mem */ -inline int memcmp_gc(__global const void *s1, __constant void *s2, uint size) +INLINE int memcmp_gc(__global const void *s1, __constant void *s2, uint size) { __global const uchar *a = s1; __constant uchar *b = s2; @@ -544,7 +541,7 @@ inline int memcmp_gc(__global const void *s1, __constant void *s2, uint size) } /* s1 is private mem, s2 is MAYBE_CONSTANT mem */ -inline int memcmp_pmc(const void *s1, MAYBE_CONSTANT void *s2, uint size) +INLINE int memcmp_pmc(const void *s1, MAYBE_CONSTANT void *s2, uint size) { const uchar *a = s1; MAYBE_CONSTANT uchar *b = s2; @@ -557,7 +554,7 @@ inline int memcmp_pmc(const void *s1, MAYBE_CONSTANT void *s2, uint size) } /* haystack is private mem, needle is constant mem */ -inline int memmem_pc(const void *haystack, size_t haystack_len, +INLINE int memmem_pc(const void *haystack, size_t haystack_len, __constant const void *needle, size_t needle_len) { const char *haystack_ = haystack; diff --git a/run/opencl/opencl_pkcs12.h b/run/opencl/opencl_pkcs12.h index e0251f7e64..03e901a5f9 100644 --- a/run/opencl/opencl_pkcs12.h +++ b/run/opencl/opencl_pkcs12.h @@ -29,7 +29,7 @@ #ifdef UTF_8 -inline uint enc2utf16be(const UTF8 *pwd, uint length, UTF16 *unipwd) +INLINE uint enc2utf16be(const UTF8 *pwd, uint length, UTF16 *unipwd) { const UTF8 *source = pwd; const UTF8 *sourceEnd = &source[length]; @@ -91,7 +91,7 @@ inline uint enc2utf16be(const UTF8 *pwd, uint length, UTF16 *unipwd) #else -inline uint enc2utf16be(const UTF8 *pwd, uint length, UTF16 *unipwd) +INLINE uint enc2utf16be(const UTF8 *pwd, uint length, UTF16 *unipwd) { uint l = length; @@ -106,7 +106,7 @@ inline uint enc2utf16be(const UTF8 *pwd, uint length, UTF16 *unipwd) #endif /* encodings */ -inline void pkcs12_fill_buffer(uint *data, uint data_len, +INLINE void pkcs12_fill_buffer(uint *data, uint data_len, const uint *filler, uint fill_len) { if ((fill_len & 0x03) == 0) { @@ -131,7 +131,7 @@ inline void pkcs12_fill_buffer(uint *data, uint data_len, } /* SHA-1 */ -inline void pkcs12_pbe_derive_key(uint iterations, int id, +INLINE void pkcs12_pbe_derive_key(uint iterations, int id, const uint *pwd, uint pwdlen, const uint *salt, uint saltlen, uint *key, uint keylen) @@ -251,7 +251,7 @@ inline void pkcs12_pbe_derive_key(uint iterations, int id, } /* SHA-256 */ -inline void pkcs12_pbe_derive_key_sha256(uint iterations, int id, +INLINE void pkcs12_pbe_derive_key_sha256(uint iterations, int id, const uint *pwd, uint pwdlen, const uint *salt, uint saltlen, uint *key, uint keylen) @@ -381,7 +381,7 @@ inline void pkcs12_pbe_derive_key_sha256(uint iterations, int id, /* SHA-512 */ -inline void pkcs12_pbe_derive_key_sha512(uint iterations, int id, +INLINE void pkcs12_pbe_derive_key_sha512(uint iterations, int id, const uint *pwd, uint pwdlen, const uint *salt, uint saltlen, uint *key, uint keylen) diff --git a/run/opencl/opencl_rc4.h b/run/opencl/opencl_rc4.h index 861b9a988c..5919e8aafc 100644 --- a/run/opencl/opencl_rc4.h +++ b/run/opencl/opencl_rc4.h @@ -98,7 +98,7 @@ typedef struct { * Set IV. Clever, compact 32-bit implementation nicked from hashcat, replacing * the (also 32-bit) constant array we had. No difference in speed though. */ -inline void rc4_init( +INLINE void rc4_init( #ifdef RC4_USE_LOCAL __local #endif @@ -114,7 +114,7 @@ inline void rc4_init( /* * Arbitrary length key */ -inline void rc4_set_key( +INLINE void rc4_set_key( #ifdef RC4_USE_LOCAL __local #endif @@ -138,7 +138,7 @@ inline void rc4_set_key( /* * Unrolled fixed keylen of 5 (40-bit). */ -inline void rc4_40_set_key( +INLINE void rc4_40_set_key( #ifdef RC4_USE_LOCAL __local #endif @@ -168,7 +168,7 @@ inline void rc4_40_set_key( /* * Unrolled fixed keylen of 16 (128-bit). */ -inline void rc4_128_set_key( +INLINE void rc4_128_set_key( #ifdef RC4_USE_LOCAL __local #endif @@ -210,7 +210,7 @@ inline void rc4_128_set_key( /* * Len is given in bytes but must be multiple of 4. */ -inline void rc4( +INLINE void rc4( #ifdef RC4_USE_LOCAL __local #endif diff --git a/run/opencl/opencl_ripemd.h b/run/opencl/opencl_ripemd.h index e184698f3e..35cd16c6a2 100644 --- a/run/opencl/opencl_ripemd.h +++ b/run/opencl/opencl_ripemd.h @@ -61,7 +61,7 @@ RR(a ## 2, b ## 2, c ## 2, d ## 2, e ## 2, f, s, r, K2 ## k) /* Input is raw Merkle Damgard */ -inline void ripemd160(uint *W, uint *ctx) +INLINE void ripemd160(uint *W, uint *ctx) { uint A1, B1, C1, D1, E1; uint A2, B2, C2, D2, E2; @@ -252,7 +252,7 @@ inline void ripemd160(uint *W, uint *ctx) } /* Input is last output; length is 160 bits */ -inline void ripemd160_160Z(uint *W, uint *ctx) +INLINE void ripemd160_160Z(uint *W, uint *ctx) { uint A1, B1, C1, D1, E1; uint A2, B2, C2, D2, E2; diff --git a/run/opencl/opencl_rotate.h b/run/opencl/opencl_rotate.h index e26f0e821c..09fb2c246d 100644 --- a/run/opencl/opencl_rotate.h +++ b/run/opencl/opencl_rotate.h @@ -23,7 +23,7 @@ * Leaving it here as a curious reference. - magnum */ #if 0 && gpu_nvidia(DEVICE_INFO) && SM_MAJOR >= 2 -inline uint byte_perm(uint a, uint b, uint imm) +INLINE uint byte_perm(uint a, uint b, uint imm) { uint r; asm("prmt.b32 %0, %1, %2, %3;" @@ -32,7 +32,7 @@ inline uint byte_perm(uint a, uint b, uint imm) return r; } -inline uint ror32(uint x, uint n) +INLINE uint ror32(uint x, uint n) { switch (n) { case 8: diff --git a/run/opencl/opencl_sha1_ctx.h b/run/opencl/opencl_sha1_ctx.h index 1358550684..173b2bd891 100644 --- a/run/opencl/opencl_sha1_ctx.h +++ b/run/opencl/opencl_sha1_ctx.h @@ -50,7 +50,7 @@ typedef struct { uchar buffer[64]; /* data block being processed */ } SHA_CTX; -inline void SHA1_Init(SHA_CTX *ctx) +INLINE void SHA1_Init(SHA_CTX *ctx) { ctx->total = 0; @@ -61,7 +61,7 @@ inline void SHA1_Init(SHA_CTX *ctx) ctx->state[4] = INIT_E; } -inline void _sha1_process(SHA_CTX *ctx, const uchar data[64]) +INLINE void _sha1_process(SHA_CTX *ctx, const uchar data[64]) { #if __OS_X__ && gpu_amd(DEVICE_INFO) volatile @@ -125,7 +125,7 @@ inline void _sha1_process(SHA_CTX *ctx, const uchar data[64]) /* * SHA-1 process buffer */ -inline void SHA1_Update(SHA_CTX *ctx, const uchar *input, uint ilen) +INLINE void SHA1_Update(SHA_CTX *ctx, const uchar *input, uint ilen) { uint fill; uint left; @@ -163,7 +163,7 @@ inline void SHA1_Update(SHA_CTX *ctx, const uchar *input, uint ilen) /* * SHA-1 final digest */ -inline void SHA1_Final(uchar output[20], SHA_CTX *ctx) +INLINE void SHA1_Final(uchar output[20], SHA_CTX *ctx) { uint last, padn; ulong bits; diff --git a/run/opencl/opencl_sha2.h b/run/opencl/opencl_sha2.h index 7413561f49..315c7fb406 100644 --- a/run/opencl/opencl_sha2.h +++ b/run/opencl/opencl_sha2.h @@ -687,7 +687,7 @@ __constant ulong K[] = { #else /* Raw'n'lean single-block SHA-512, no context[tm] */ -inline void sha512_single_s(ulong *W, ulong *output) +INLINE void sha512_single_s(ulong *W, ulong *output) { ulong A, B, C, D, E, F, G, H, t; @@ -748,7 +748,7 @@ inline void sha512_single_s(ulong *W, ulong *output) } /* Raw'n'lean single-block SHA-512, no context[tm] */ -inline void sha512_single(MAYBE_VECTOR_ULONG *W, MAYBE_VECTOR_ULONG *output) +INLINE void sha512_single(MAYBE_VECTOR_ULONG *W, MAYBE_VECTOR_ULONG *output) { MAYBE_VECTOR_ULONG A, B, C, D, E, F, G, H, t; @@ -773,7 +773,7 @@ inline void sha512_single(MAYBE_VECTOR_ULONG *W, MAYBE_VECTOR_ULONG *output) output[7] = H + SHA512_INIT_H; } -inline void sha512_single_zeros(MAYBE_VECTOR_ULONG *W, +INLINE void sha512_single_zeros(MAYBE_VECTOR_ULONG *W, MAYBE_VECTOR_ULONG *output) { MAYBE_VECTOR_ULONG A, B, C, D, E, F, G, H, t; diff --git a/run/opencl/opencl_sha2_ctx.h b/run/opencl/opencl_sha2_ctx.h index ec0e35aaa9..292c919254 100644 --- a/run/opencl/opencl_sha2_ctx.h +++ b/run/opencl/opencl_sha2_ctx.h @@ -26,7 +26,7 @@ typedef struct { uchar buffer[64]; /* data block being processed */ } SHA256_CTX; -inline +INLINE void SHA256_Init(SHA256_CTX *ctx) { uint i; @@ -36,7 +36,7 @@ void SHA256_Init(SHA256_CTX *ctx) { ctx->state[i] = h[i]; } -inline +INLINE void _sha256_process(SHA256_CTX *ctx, const uchar data[64]) { MAYBE_VOLATILE uint t, W[16], A, B, C, D, E, F, G, H; @@ -103,7 +103,7 @@ void _sha256_process(SHA256_CTX *ctx, const uchar data[64]) { /* * SHA-256 process buffer */ -inline +INLINE void SHA256_Update(SHA256_CTX *ctx, const uchar *input, uint ilen) { uint fill; uint left; @@ -141,7 +141,7 @@ void SHA256_Update(SHA256_CTX *ctx, const uchar *input, uint ilen) { /* * SHA-256 final digest */ -inline +INLINE void SHA256_Final(uchar output[32], SHA256_CTX *ctx) { uint last, padn; ulong bits; @@ -192,7 +192,7 @@ typedef struct { uchar buffer[128]; /* data block being processed */ } SHA512_CTX; -inline +INLINE void SHA384_Init(SHA512_CTX *ctx) { ctx->total = 0; ctx->state[0] = SHA384_INIT_A; @@ -205,7 +205,7 @@ void SHA384_Init(SHA512_CTX *ctx) { ctx->state[7] = SHA384_INIT_H; } -inline +INLINE void SHA512_Init(SHA512_CTX *ctx) { ctx->total = 0; ctx->state[0] = SHA512_INIT_A; @@ -218,7 +218,7 @@ void SHA512_Init(SHA512_CTX *ctx) { ctx->state[7] = SHA512_INIT_H; } -inline +INLINE void _sha512_process(SHA512_CTX *ctx, const uchar data[128]) { ulong t, W[16], A, B, C, D, E, F, G, H; @@ -285,7 +285,7 @@ void _sha512_process(SHA512_CTX *ctx, const uchar data[128]) { /* * SHA-512 process buffer */ -inline +INLINE void SHA512_Update(SHA512_CTX *ctx, const uchar *input, uint ilen) { uint fill; uint left; @@ -325,7 +325,7 @@ void SHA512_Update(SHA512_CTX *ctx, const uchar *input, uint ilen) { /* * SHA-384 final digest */ -inline +INLINE void SHA384_Final(uchar output[64], SHA512_CTX *ctx) { uint last, padn; ulong bits; @@ -366,7 +366,7 @@ void SHA384_Final(uchar output[64], SHA512_CTX *ctx) { /* * SHA-512 final digest */ -inline +INLINE void SHA512_Final(uchar output[64], SHA512_CTX *ctx) { uint last, padn; ulong bits; diff --git a/run/opencl/opencl_streebog.h b/run/opencl/opencl_streebog.h index 508d293e8e..306edc62f6 100644 --- a/run/opencl/opencl_streebog.h +++ b/run/opencl/opencl_streebog.h @@ -731,7 +731,7 @@ __constant const ulong Ax[8][256] = { } }; -inline void +INLINE void GOST34112012Init(GOST34112012Context *CTX, const uint digest_size) { CTX->buffer.VWORD = 0; @@ -752,7 +752,7 @@ GOST34112012Init(GOST34112012Context *CTX, const uint digest_size) CTX->bufsize = 0; } -inline void +INLINE void pad(GOST34112012Context *CTX) { if (CTX->bufsize > 63) @@ -764,7 +764,7 @@ pad(GOST34112012Context *CTX) } /* Let r = x + y modulo 2^512 */ -inline void +INLINE void add512(const uint512_u *x, const uint512_u *y, uint512_u *r) { uint CF; @@ -785,7 +785,7 @@ add512(const uint512_u *x, const uint512_u *y, uint512_u *r) } } -inline void +INLINE void g(uint512_u *h, const uint512_u *N, const uint512_u *m, __local localbuf *loc_buf) { uint512_u Ki, data; @@ -821,7 +821,7 @@ g(uint512_u *h, const uint512_u *N, const uint512_u *m, __local localbuf *loc_bu } // Special case of the above where N is all zeros -inline void +INLINE void g0(uint512_u *h, const uint512_u *m, __local localbuf *loc_buf) { uint512_u Ki, data; @@ -857,7 +857,7 @@ g0(uint512_u *h, const uint512_u *m, __local localbuf *loc_buf) XOR512(&data, m, h); } -inline void +INLINE void stage2(GOST34112012Context *CTX, const uint512_u *data, __local localbuf *loc_buf) { const uint512_u buffer512 = {{ 0x0000000000000200UL, 0, 0, 0, 0, 0, 0, 0 }}; @@ -872,7 +872,7 @@ stage2(GOST34112012Context *CTX, const uint512_u *data, __local localbuf *loc_bu add512(&(CTX->Sigma), &m, &(CTX->Sigma)); } -inline void +INLINE void stage2d(GOST34112012Context *CTX, const uchar *data, __local localbuf *loc_buf) { const uint512_u buffer512 = {{ 0x0000000000000200UL, 0, 0, 0, 0, 0, 0, 0 }}; @@ -888,7 +888,7 @@ stage2d(GOST34112012Context *CTX, const uchar *data, __local localbuf *loc_buf) add512(&(CTX->Sigma), &m, &(CTX->Sigma)); } -inline void +INLINE void stage3(GOST34112012Context *CTX, __local localbuf *loc_buf) { uint512_u buf = {{ 0 }}; @@ -906,7 +906,7 @@ stage3(GOST34112012Context *CTX, __local localbuf *loc_buf) g0(&(CTX->h), &(CTX->Sigma), loc_buf); } -__attribute__((noinline)) void +NOINLINE void GOST34112012Update(GOST34112012Context *CTX, const uchar *data, uint len, __local localbuf *loc_buf) { if (CTX->bufsize) { @@ -937,7 +937,7 @@ GOST34112012Update(GOST34112012Context *CTX, const uchar *data, uint len, __loca } } -__attribute__((noinline)) void +NOINLINE void GOST34112012Final(GOST34112012Context *CTX, #if STREEBOG512CRYPT uint512_u diff --git a/run/opencl/opencl_twofish.h b/run/opencl/opencl_twofish.h index 10703e6ab8..6a7ce9d1b5 100644 --- a/run/opencl/opencl_twofish.h +++ b/run/opencl/opencl_twofish.h @@ -381,7 +381,7 @@ __constant UInt32 mds_poly_divx_const[] = { 0, 0xb4 }; #define H24( y, L ) H23( q0[y]^L[26], L ) #define H34( y, L ) H33( q1[y]^L[27], L ) -inline +INLINE UInt32 two_h(int k, Byte L[], int kCycles) { switch (kCycles) { @@ -396,7 +396,7 @@ UInt32 two_h(int k, Byte L[], int kCycles) } } -inline +INLINE void fill_keyed_sboxes(Byte S[], int kCycles, Twofish_key *xkey) { int i; @@ -432,7 +432,7 @@ void fill_keyed_sboxes(Byte S[], int kCycles, Twofish_key *xkey) __constant uint rs_poly_const[] = { 0, 0x14d }; __constant uint rs_poly_div_const[] = { 0, 0xa6 }; -inline +INLINE void Twofish_prepare_key(Byte *key, int key_len, Twofish_key *xkey) { Byte K[32 + 32 + 4] = { 0 }; @@ -553,7 +553,7 @@ void Twofish_prepare_key(Byte *key, int key_len, Twofish_key *xkey) PUT32( A, dst ); PUT32( B, dst+ 4 ); \ PUT32( C, dst+8 ); PUT32( D, dst+12 ) -inline +INLINE void Twofish_encrypt(Twofish_key *xkey, Byte p[16], Byte c[16]) { UInt32 A, B, C, D, T0, T1; @@ -565,7 +565,7 @@ void Twofish_encrypt(Twofish_key *xkey, Byte p[16], Byte c[16]) PUT_OUTPUT(C, D, A, B, c, xkey, 4); } -inline +INLINE int Twofish_Encrypt(Twofish_key *m_key, Byte *pInput, Byte *pOutBuffer, int nInputOctets, Byte *m_pInitVector) { @@ -617,7 +617,7 @@ int Twofish_Encrypt(Twofish_key *m_key, Byte *pInput, Byte *pOutBuffer, return 16 * (numBlocks + 1); } -inline +INLINE void Twofish_decrypt(Twofish_key *xkey, Byte c[16], Byte p[16]) { UInt32 A, B, C, D, T0, T1; @@ -629,7 +629,7 @@ void Twofish_decrypt(Twofish_key *xkey, Byte c[16], Byte p[16]) PUT_OUTPUT(C, D, A, B, p, xkey, 0); } -inline +INLINE int Twofish_Decrypt(Twofish_key *m_key, Byte *pInput, Byte *pOutBuffer, int nInputOctets, Byte *m_pInitVector, int check_pad) { @@ -693,7 +693,7 @@ int Twofish_Decrypt(Twofish_key *m_key, Byte *pInput, Byte *pOutBuffer, } #if 0 -inline +INLINE int Twofish_Decrypt_cfb128(Twofish_key *m_key, Twofish_Byte *pInput, Twofish_Byte *pOutBuffer, int nInputOctets, Twofish_Byte *m_pInitVector) diff --git a/run/opencl/pbkdf1_hmac_sha1_kernel.cl b/run/opencl/pbkdf1_hmac_sha1_kernel.cl index e54fc0931c..a054006a6d 100644 --- a/run/opencl/pbkdf1_hmac_sha1_kernel.cl +++ b/run/opencl/pbkdf1_hmac_sha1_kernel.cl @@ -58,7 +58,7 @@ /* MAYBE_VECTOR_UINT need to be defined before this header */ #include "opencl_pbkdf1_hmac_sha1.h" -inline void hmac_sha1(__global MAYBE_VECTOR_UINT *state, +INLINE void hmac_sha1(__global MAYBE_VECTOR_UINT *state, __global MAYBE_VECTOR_UINT *ipad, __global MAYBE_VECTOR_UINT *opad, MAYBE_CONSTANT uchar *salt, uint saltlen) @@ -118,7 +118,7 @@ inline void hmac_sha1(__global MAYBE_VECTOR_UINT *state, state[i] = output[i]; } -inline void preproc(__global const MAYBE_VECTOR_UINT *key, +INLINE void preproc(__global const MAYBE_VECTOR_UINT *key, __global MAYBE_VECTOR_UINT *state, uint padding) { uint i; diff --git a/run/opencl/pbkdf2_hmac_md4_kernel.cl b/run/opencl/pbkdf2_hmac_md4_kernel.cl index d64ecb81a1..3ae9e1f801 100644 --- a/run/opencl/pbkdf2_hmac_md4_kernel.cl +++ b/run/opencl/pbkdf2_hmac_md4_kernel.cl @@ -58,7 +58,7 @@ /* MAYBE_VECTOR_UINT need to be defined before this header */ #include "opencl_pbkdf2_hmac_md4.h" -inline void hmac_md4(__global MAYBE_VECTOR_UINT *state, +INLINE void hmac_md4(__global MAYBE_VECTOR_UINT *state, __global MAYBE_VECTOR_UINT *ipad, __global MAYBE_VECTOR_UINT *opad, MAYBE_CONSTANT uchar *salt, uint saltlen, uchar add) @@ -133,7 +133,7 @@ inline void hmac_md4(__global MAYBE_VECTOR_UINT *state, state[i] = output[i]; } -inline void preproc(__global const MAYBE_VECTOR_UINT *key, +INLINE void preproc(__global const MAYBE_VECTOR_UINT *key, __global MAYBE_VECTOR_UINT *state, uint padding) { uint i; diff --git a/run/opencl/pbkdf2_hmac_md5_kernel.cl b/run/opencl/pbkdf2_hmac_md5_kernel.cl index 486b8032fc..6f6d9e998a 100644 --- a/run/opencl/pbkdf2_hmac_md5_kernel.cl +++ b/run/opencl/pbkdf2_hmac_md5_kernel.cl @@ -58,7 +58,7 @@ /* MAYBE_VECTOR_UINT need to be defined before this header */ #include "opencl_pbkdf2_hmac_md5.h" -inline void hmac_md5(__global MAYBE_VECTOR_UINT *state, +INLINE void hmac_md5(__global MAYBE_VECTOR_UINT *state, __global MAYBE_VECTOR_UINT *ipad, __global MAYBE_VECTOR_UINT *opad, MAYBE_CONSTANT uchar *salt, uint saltlen, uchar add) @@ -133,7 +133,7 @@ inline void hmac_md5(__global MAYBE_VECTOR_UINT *state, state[i] = output[i]; } -inline void preproc(__global const MAYBE_VECTOR_UINT *key, +INLINE void preproc(__global const MAYBE_VECTOR_UINT *key, __global MAYBE_VECTOR_UINT *state, uint padding) { uint i; diff --git a/run/opencl/pbkdf2_hmac_sha1_kernel.cl b/run/opencl/pbkdf2_hmac_sha1_kernel.cl index 3bf703b7f0..7fb4bb3bc3 100644 --- a/run/opencl/pbkdf2_hmac_sha1_kernel.cl +++ b/run/opencl/pbkdf2_hmac_sha1_kernel.cl @@ -59,7 +59,7 @@ /* MAYBE_VECTOR_UINT need to be defined before this header */ #include "opencl_pbkdf2_hmac_sha1.h" -inline void _phsk_hmac_sha1(__global MAYBE_VECTOR_UINT *state, +INLINE void _phsk_hmac_sha1(__global MAYBE_VECTOR_UINT *state, __global MAYBE_VECTOR_UINT *ipad, __global MAYBE_VECTOR_UINT *opad, MAYBE_CONSTANT uchar *salt, uint saltlen, uchar add) @@ -129,7 +129,7 @@ inline void _phsk_hmac_sha1(__global MAYBE_VECTOR_UINT *state, state[i] = output[i]; } -inline void _phsk_preproc(__global const MAYBE_VECTOR_UINT *key, +INLINE void _phsk_preproc(__global const MAYBE_VECTOR_UINT *key, __global MAYBE_VECTOR_UINT *state, uint padding) { uint i; diff --git a/run/opencl/pbkdf2_hmac_sha1_unsplit_kernel.cl b/run/opencl/pbkdf2_hmac_sha1_unsplit_kernel.cl index 5666062c59..84a0e357f7 100644 --- a/run/opencl/pbkdf2_hmac_sha1_unsplit_kernel.cl +++ b/run/opencl/pbkdf2_hmac_sha1_unsplit_kernel.cl @@ -44,7 +44,7 @@ typedef struct { uchar salt[SALTLEN]; } pbkdf2_salt; -inline void preproc(__global const uchar *key, uint keylen, +INLINE void preproc(__global const uchar *key, uint keylen, uint *state, uint padding) { uint i; @@ -72,7 +72,7 @@ inline void preproc(__global const uchar *key, uint keylen, state[4] = E + INIT_E; } -inline void hmac_sha1(uint *output, +INLINE void hmac_sha1(uint *output, uint *ipad_state, uint *opad_state, __constant uchar *salt, int saltlen, uchar add) @@ -133,7 +133,7 @@ inline void hmac_sha1(uint *output, output[4] = E; } -inline void big_hmac_sha1(uint *input, uint inputlen, +INLINE void big_hmac_sha1(uint *input, uint inputlen, uint *ipad_state, uint *opad_state, uint *tmp_out, uint iterations) { @@ -193,7 +193,7 @@ inline void big_hmac_sha1(uint *input, uint inputlen, } } -inline void pbkdf2(__global const uchar *pass, uint passlen, +INLINE void pbkdf2(__global const uchar *pass, uint passlen, __constant uchar *salt, uint saltlen, uint iterations, __global uint *out, uint outlen, uint skip_bytes) { diff --git a/run/opencl/pbkdf2_hmac_sha256_kernel.cl b/run/opencl/pbkdf2_hmac_sha256_kernel.cl index 19d666580f..919a539095 100644 --- a/run/opencl/pbkdf2_hmac_sha256_kernel.cl +++ b/run/opencl/pbkdf2_hmac_sha256_kernel.cl @@ -25,7 +25,7 @@ #define SALT_TYPE MAYBE_CONSTANT #endif -inline void _phsk_preproc(__global const uchar *key, uint keylen, +INLINE void _phsk_preproc(__global const uchar *key, uint keylen, __global uint *state, uint padding) { uint j, t; @@ -58,7 +58,7 @@ inline void _phsk_preproc(__global const uchar *key, uint keylen, } -inline void _phsk_hmac_sha256(__global uint *output, __global uint *ipad_state, +INLINE void _phsk_hmac_sha256(__global uint *output, __global uint *ipad_state, __global uint *opad_state, SALT_TYPE uchar *salt, uint saltlen, uchar add) { diff --git a/run/opencl/pbkdf2_hmac_sha512_kernel.cl b/run/opencl/pbkdf2_hmac_sha512_kernel.cl index 6f6f8223aa..e6e235fd40 100644 --- a/run/opencl/pbkdf2_hmac_sha512_kernel.cl +++ b/run/opencl/pbkdf2_hmac_sha512_kernel.cl @@ -13,7 +13,7 @@ #include "opencl_sha2.h" #include "opencl_pbkdf2_hmac_sha512.h" -inline void _phs512_preproc(__global const ulong *key, uint keylen, +INLINE void _phs512_preproc(__global const ulong *key, uint keylen, ulong *state, ulong mask) { uint i, j; @@ -48,7 +48,7 @@ inline void _phs512_preproc(__global const ulong *key, uint keylen, state[7] = H + SHA512_INIT_H; } -inline void _phs512_hmac(ulong *output, ulong *ipad_state, ulong *opad_state, +INLINE void _phs512_hmac(ulong *output, ulong *ipad_state, ulong *opad_state, __constant ulong *salt, uint saltlen) { uint i, j; diff --git a/run/opencl/pbkdf2_kernel.cl b/run/opencl/pbkdf2_kernel.cl index 98f10d786c..f52e10b44e 100644 --- a/run/opencl/pbkdf2_kernel.cl +++ b/run/opencl/pbkdf2_kernel.cl @@ -351,7 +351,7 @@ typedef struct { P(C, D, E, A, B, R6); \ P(B, C, D, E, A, R7); -inline void SHA1(uint *A, uint *W) { +INLINE void SHA1(uint *A, uint *W) { #if HAVE_LUT3 #define F(x, y, z) lut3(x, y, z, 0xca) #elif USE_BITSELECT @@ -403,7 +403,7 @@ inline void SHA1(uint *A, uint *W) { #undef F } -inline void SHA1_digest(uint *A, uint *W) { +INLINE void SHA1_digest(uint *A, uint *W) { #if HAVE_LUT3 #define F(x, y, z) lut3(x, y, z, 0xca) #elif USE_BITSELECT @@ -455,7 +455,7 @@ inline void SHA1_digest(uint *A, uint *W) { #undef F } -inline void sha1_pad(uint *pad, uint *state) { +INLINE void sha1_pad(uint *pad, uint *state) { uint A[5], W[16] ; GET_WORD_32_BE(W[0], pad, 0) ; @@ -496,7 +496,7 @@ inline void sha1_pad(uint *pad, uint *state) { state[4] = A[4] ; } -inline void hmac_sha1(uint *istate, uint *ostate, uint *buf) +INLINE void hmac_sha1(uint *istate, uint *ostate, uint *buf) { uint A[5], W[16] ; diff --git a/run/opencl/pbkdf2_ripemd160_kernel.cl b/run/opencl/pbkdf2_ripemd160_kernel.cl index b7dc9adcc8..6c862de10f 100644 --- a/run/opencl/pbkdf2_ripemd160_kernel.cl +++ b/run/opencl/pbkdf2_ripemd160_kernel.cl @@ -32,7 +32,7 @@ typedef struct { uint bin[(512 - 64) / 4]; } tc_salt; -inline void preproc(__global const uchar *key, uint keylen, uint *state, +INLINE void preproc(__global const uchar *key, uint keylen, uint *state, uint padding) { uint i; @@ -53,7 +53,7 @@ inline void preproc(__global const uchar *key, uint keylen, uint *state, ripemd160(W, state); } -inline void hmac_ripemd160(uint *output, uint *ipad_state, uint *opad_state, +INLINE void hmac_ripemd160(uint *output, uint *ipad_state, uint *opad_state, __constant uint *salt, uchar add) { uint i; @@ -85,7 +85,7 @@ inline void hmac_ripemd160(uint *output, uint *ipad_state, uint *opad_state, ripemd160_160Z(W, output); } -inline void big_hmac_ripemd160(uint *input, uint inputlen, uint *ipad_state, +INLINE void big_hmac_ripemd160(uint *input, uint inputlen, uint *ipad_state, uint *opad_state, uint *tmp_out) { uint i; @@ -119,7 +119,7 @@ inline void big_hmac_ripemd160(uint *input, uint inputlen, uint *ipad_state, } } -inline void pbkdf2(__global const uchar *pass, uint passlen, +INLINE void pbkdf2(__global const uchar *pass, uint passlen, __constant uint *salt, uint *out) { uint ipad_state[5]; diff --git a/run/opencl/pdf_kernel.cl b/run/opencl/pdf_kernel.cl index eb765898e4..eb03aa040e 100644 --- a/run/opencl/pdf_kernel.cl +++ b/run/opencl/pdf_kernel.cl @@ -41,7 +41,7 @@ __constant uint padding[8] = { 0xb6002e2e, 0x803e68d0, 0xfea90c2f, 0x7a695364 }; -inline uint prepare234(__global const uchar *pwbuf, __global const uint *index, uint *password) +INLINE uint prepare234(__global const uchar *pwbuf, __global const uint *index, uint *password) { uint i; uint gid = get_global_id(0); @@ -391,7 +391,7 @@ void pdf_r34(__global const uchar *pwbuf, } } -inline uint prepare56(__global const uchar *pwbuf, __global const uint *index, uint *password) +INLINE uint prepare56(__global const uchar *pwbuf, __global const uint *index, uint *password) { uint i; uint gid = get_global_id(0); diff --git a/run/opencl/pfx_kernel.cl b/run/opencl/pfx_kernel.cl index 2881f76e35..e2d159c5a1 100644 --- a/run/opencl/pfx_kernel.cl +++ b/run/opencl/pfx_kernel.cl @@ -48,7 +48,7 @@ typedef struct { } data; } pfx_salt; -inline void pfx_crypt(__global const uint *password, uint32_t password_length, +INLINE void pfx_crypt(__global const uint *password, uint32_t password_length, __constant pfx_salt *salt, __global uint *out) { uint i; diff --git a/run/opencl/pgpdisk_kernel.cl b/run/opencl/pgpdisk_kernel.cl index 0477a9e850..611a81f127 100644 --- a/run/opencl/pgpdisk_kernel.cl +++ b/run/opencl/pgpdisk_kernel.cl @@ -34,7 +34,7 @@ typedef struct { uchar salt[16]; } pgpdisk_salt; -inline void pgpdisk_kdf(__global const uchar *ipassword, const uint plen, +INLINE void pgpdisk_kdf(__global const uchar *ipassword, const uint plen, __constant uchar *isalt, const uint saltlen, const uint iterations, uchar *okey, uint bytesNeeded) { diff --git a/run/opencl/pgpsda_kernel.cl b/run/opencl/pgpsda_kernel.cl index d02f5e0995..15044e7cd5 100644 --- a/run/opencl/pgpsda_kernel.cl +++ b/run/opencl/pgpsda_kernel.cl @@ -29,7 +29,7 @@ typedef struct { uchar salt[8]; } pgpsda_salt; -inline void pgpsda_kdf(__global const uchar *ipassword, const uint plen, +INLINE void pgpsda_kdf(__global const uchar *ipassword, const uint plen, __constant uchar *isalt, const uint iterations, uchar *key) { diff --git a/run/opencl/pgpwde_kernel.cl b/run/opencl/pgpwde_kernel.cl index 0b4dd09766..d8eaa7377c 100644 --- a/run/opencl/pgpwde_kernel.cl +++ b/run/opencl/pgpwde_kernel.cl @@ -34,7 +34,7 @@ typedef struct { uchar esk[128]; } pgpwde_salt; -inline void pgpwde_kdf(__global const uchar *ipassword, const uint plen, +INLINE void pgpwde_kdf(__global const uchar *ipassword, const uint plen, __constant uchar *isalt, uint cbytes, uint *okey) { const uint saltlen = 16; @@ -78,7 +78,7 @@ inline void pgpwde_kdf(__global const uchar *ipassword, const uint plen, } } -inline int PKCS1oaepMGF1Unpack(uchar *in, uint32_t inlen) +INLINE int PKCS1oaepMGF1Unpack(uchar *in, uint32_t inlen) { const uint32_t hashlen = SHA1_DIGEST_LENGTH; const uchar nullhash[20] = { 0xda, 0x39, 0xa3, 0xee, 0x5e, 0x6b, 0x4b, 0x0d, @@ -137,7 +137,7 @@ inline int PKCS1oaepMGF1Unpack(uchar *in, uint32_t inlen) return memcmp_pp(nullhash, msg + hashlen / 4, hashlen); } -inline int pgpwde_decrypt_and_verify(uchar *key, __constant uchar *esk, __local aes_local_t *lt) +INLINE int pgpwde_decrypt_and_verify(uchar *key, __constant uchar *esk, __local aes_local_t *lt) { AES_KEY aes_key; aes_key.lt = lt; uchar iv[16] = { 8, 0 }; diff --git a/run/opencl/phpass_kernel.cl b/run/opencl/phpass_kernel.cl index b54885e293..35cda0bd45 100644 --- a/run/opencl/phpass_kernel.cl +++ b/run/opencl/phpass_kernel.cl @@ -123,7 +123,7 @@ typedef struct { #define ACF3 0x98badcfe #define ACF4 0x10325476 -inline void md5(MAYBE_VECTOR_UINT len, +INLINE void md5(MAYBE_VECTOR_UINT len, MAYBE_VECTOR_UINT *internal_ret, MAYBE_VECTOR_UINT *x) { diff --git a/run/opencl/pwsafe_kernel.cl b/run/opencl/pwsafe_kernel.cl index 5126c05a49..c7a6f80d3f 100644 --- a/run/opencl/pwsafe_kernel.cl +++ b/run/opencl/pwsafe_kernel.cl @@ -81,7 +81,7 @@ typedef struct { uint8_t salt[32]; } pwsafe_salt; -inline void sha256_transform(uint32_t *w, uint32_t *state) +INLINE void sha256_transform(uint32_t *w, uint32_t *state) { uint32_t a = state[0]; uint32_t b = state[1]; diff --git a/run/opencl/rar_kernel.cl b/run/opencl/rar_kernel.cl index b9d1cc49de..d6f008ed99 100644 --- a/run/opencl/rar_kernel.cl +++ b/run/opencl/rar_kernel.cl @@ -54,7 +54,7 @@ typedef struct { /* * This version does several blocks at a time */ -inline void sha1_mblock(uint *W, uint *out, uint blocks) +INLINE void sha1_mblock(uint *W, uint *out, uint blocks) { uint i; uint ctx[5]; @@ -71,7 +71,7 @@ inline void sha1_mblock(uint *W, uint *out, uint blocks) out[i] = ctx[i]; } -inline void sha1_empty_final(uint *W, uint *ctx, const uint tot_len) +INLINE void sha1_empty_final(uint *W, uint *ctx, const uint tot_len) { uint len = ((tot_len & 63) >> 2) + 1; @@ -178,7 +178,7 @@ __kernel void RarHashLoop(const __global uint *unicode_pw, const __global uint * * * Returns 0 for early rejection, 1 if passed */ -inline int check_huffman(uchar *next) { +INLINE int check_huffman(uchar *next) { uint bits, hold, i; int left; uint ncount[4] = { 0 }; @@ -239,7 +239,7 @@ inline int check_huffman(uchar *next) { /* * Returns 0 for early rejection, 1 if passed */ -inline int check_rar(__global rar_file *cur_file, __global uint *_key, __global uint *_iv, __local aes_local_t *lt) +INLINE int check_rar(__global rar_file *cur_file, __global uint *_key, __global uint *_iv, __local aes_local_t *lt) { AES_KEY aes_ctx; aes_ctx.lt = lt; uchar iv[16]; diff --git a/run/opencl/salted_sha_kernel.cl b/run/opencl/salted_sha_kernel.cl index 0314c0f533..9d36fb597d 100644 --- a/run/opencl/salted_sha_kernel.cl +++ b/run/opencl/salted_sha_kernel.cl @@ -21,7 +21,7 @@ #define SL3CONV #endif -inline void cmp_final(uint gid, +INLINE void cmp_final(uint gid, uint iter, uint *hash, __global uint *offset_table, @@ -70,7 +70,7 @@ inline void cmp_final(uint gid, } } -inline void cmp(uint gid, +INLINE void cmp(uint gid, uint iter, uint *hash, #if USE_LOCAL_BITMAPS diff --git a/run/opencl/sap_pse_kernel.cl b/run/opencl/sap_pse_kernel.cl index e4f23362cf..92554421fd 100644 --- a/run/opencl/sap_pse_kernel.cl +++ b/run/opencl/sap_pse_kernel.cl @@ -36,7 +36,7 @@ typedef struct { uint cracked; } sappse_out; -inline int sappse_crypt(__global const uint *password, uint32_t password_length, +INLINE int sappse_crypt(__global const uint *password, uint32_t password_length, __constant sappse_salt *salt, __global sappse_out *out) { uint i; diff --git a/run/opencl/sha1_kernel.cl b/run/opencl/sha1_kernel.cl index 736bb21530..08bf8f39dd 100644 --- a/run/opencl/sha1_kernel.cl +++ b/run/opencl/sha1_kernel.cl @@ -15,7 +15,7 @@ /* This handles an input of 0xffffffffU correctly */ #define BITMAP_SHIFT ((BITMAP_MASK >> 5) + 1) -inline void cmp_final(uint gid, +INLINE void cmp_final(uint gid, uint iter, uint *hash, __global uint *offset_table, @@ -64,7 +64,7 @@ inline void cmp_final(uint gid, } } -inline void cmp(uint gid, +INLINE void cmp(uint gid, uint iter, uint *hash, #if USE_LOCAL_BITMAPS diff --git a/run/opencl/sha512_kernel.cl b/run/opencl/sha512_kernel.cl index e9c61e4a23..4ed26f93c1 100644 --- a/run/opencl/sha512_kernel.cl +++ b/run/opencl/sha512_kernel.cl @@ -18,7 +18,7 @@ typedef struct { char v[PLAINTEXT_LENGTH+1]; } sha512_key; -inline void sha512(__global const char *password, uint8_t pass_len, +INLINE void sha512(__global const char *password, uint8_t pass_len, __global uint64_t *hash, uint32_t offset) { sha512_ctx ctx; diff --git a/run/opencl/ssh_kernel.cl b/run/opencl/ssh_kernel.cl index ba424395ed..0beb7188ee 100644 --- a/run/opencl/ssh_kernel.cl +++ b/run/opencl/ssh_kernel.cl @@ -40,7 +40,7 @@ typedef struct { uint cracked; } ssh_out; -inline void generate_key_bytes(int nbytes, uchar *password, uint32_t len, uchar *salt, uchar *key) +INLINE void generate_key_bytes(int nbytes, uchar *password, uint32_t len, uchar *salt, uchar *key) { uchar digest[16]; int keyidx = 0; @@ -70,7 +70,7 @@ inline void generate_key_bytes(int nbytes, uchar *password, uint32_t len, uchar } } -inline int check_padding_and_structure_EC(uchar *out, int length) +INLINE int check_padding_and_structure_EC(uchar *out, int length) { struct asn1_hdr hdr; const uint8_t *pos, *end; @@ -121,7 +121,7 @@ inline int check_padding_and_structure_EC(uchar *out, int length) return 1; } -inline int check_padding_and_structure(uchar *out, uint length, uint strict_mode, uint block_size) +INLINE int check_padding_and_structure(uchar *out, uint length, uint strict_mode, uint block_size) { struct asn1_hdr hdr; const uint8_t *pos, *end; @@ -194,7 +194,7 @@ inline int check_padding_and_structure(uchar *out, uint length, uint strict_mode return 1; } -inline void common_crypt_code(uchar *password, uint len, __constant ssh_salt *osalt, uchar *out, uint full_decrypt, __local aes_local_t *lt) +INLINE void common_crypt_code(uchar *password, uint len, __constant ssh_salt *osalt, uchar *out, uint full_decrypt, __local aes_local_t *lt) { AES_KEY akey; akey.lt = lt; uchar salt[16]; @@ -277,7 +277,7 @@ inline void common_crypt_code(uchar *password, uint len, __constant ssh_salt *os #define QUICK 0 #define FULL 1 -inline int ssh_decrypt(uchar *password, uint len, __constant ssh_salt *osalt, __global ssh_out *output, __local aes_local_t *lt) +INLINE int ssh_decrypt(uchar *password, uint len, __constant ssh_salt *osalt, __global ssh_out *output, __local aes_local_t *lt) { uchar out[CTLEN]; int block_size = osalt->cipher == 0 ? 8 : 16; diff --git a/run/opencl/timeroast_kernel.cl b/run/opencl/timeroast_kernel.cl index 8cd9f69fb2..f7003d8a3f 100644 --- a/run/opencl/timeroast_kernel.cl +++ b/run/opencl/timeroast_kernel.cl @@ -13,12 +13,12 @@ #include "opencl_unicode.h" #include "opencl_mask.h" -inline void md4_crypt(uint *hash, uint *nt_buffer) +INLINE void md4_crypt(uint *hash, uint *nt_buffer) { md4_single(uint, nt_buffer, hash); } -inline void md5_crypt(uint *hash, __constant uint *salt) +INLINE void md5_crypt(uint *hash, __constant uint *salt) { uint W[16]; @@ -43,7 +43,7 @@ inline void md5_crypt(uint *hash, __constant uint *salt) #if UTF_8 -inline void prepare_key(__global uint *key, uint length, +INLINE void prepare_key(__global uint *key, uint length, MAYBE_VOLATILE uint *nt_buffer) { const __global UTF8 *source = (const __global uchar*)key; @@ -113,7 +113,7 @@ inline void prepare_key(__global uint *key, uint length, #else -inline void prepare_key(__global uint *key, uint length, uint *nt_buffer) +INLINE void prepare_key(__global uint *key, uint length, uint *nt_buffer) { uint i, nt_index, keychars; @@ -131,7 +131,7 @@ inline void prepare_key(__global uint *key, uint length, uint *nt_buffer) #endif /* UTF_8 */ -inline void cmp_final(uint gid, +INLINE void cmp_final(uint gid, uint iter, uint *hash, __global uint *offset_table, @@ -178,7 +178,7 @@ inline void cmp_final(uint gid, } } -inline void cmp(uint gid, +INLINE void cmp(uint gid, uint iter, uint *hash, __global uint *bitmaps, diff --git a/run/opencl/wpapsk_kernel.cl b/run/opencl/wpapsk_kernel.cl index b5522a1e6a..b4049ad5b6 100644 --- a/run/opencl/wpapsk_kernel.cl +++ b/run/opencl/wpapsk_kernel.cl @@ -53,7 +53,7 @@ void wpapmk_init(__global const uint *inbuffer, #else -inline void hmac_sha1(__global MAYBE_VECTOR_UINT *state, +INLINE void hmac_sha1(__global MAYBE_VECTOR_UINT *state, __global MAYBE_VECTOR_UINT *ipad, __global MAYBE_VECTOR_UINT *opad, MAYBE_CONSTANT uchar *salt, uint saltlen, uchar add) @@ -87,7 +87,7 @@ inline void hmac_sha1(__global MAYBE_VECTOR_UINT *state, state[i] = output[i]; } -inline void preproc(__global const MAYBE_VECTOR_UINT *key, +INLINE void preproc(__global const MAYBE_VECTOR_UINT *key, __global MAYBE_VECTOR_UINT *state, uint padding) { uint i; @@ -193,7 +193,7 @@ void wpapsk_pass2(MAYBE_CONSTANT wpapsk_salt *salt, //__constant uint text[6] = { 0x72696150, 0x65736977, 0x79656b20, 0x70786520, 0x69736e61, 0x00006e6f }; __constant uint text[6] = { 0x50616972, 0x77697365, 0x206b6579, 0x20657870, 0x616e7369, 0x6f6e0000 }; -inline void prf_512(const MAYBE_VECTOR_UINT *key, +INLINE void prf_512(const MAYBE_VECTOR_UINT *key, MAYBE_CONSTANT uint *data, MAYBE_VECTOR_UINT *ret) { @@ -529,14 +529,14 @@ void wpapsk_final_pmkid(__global wpapsk_state *state, #define SHA256_MAC_LEN 32 -inline void +INLINE void WPA_PUT_LE16(uchar *a, uint val) { a[1] = (val >> 8) & 0xff; a[0] = val & 0xff; } -inline void +INLINE void sha256_vector(uint num_elem, const uchar *addr[], const uint *len, uchar *mac) { SHA256_CTX ctx; @@ -550,7 +550,7 @@ sha256_vector(uint num_elem, const uchar *addr[], const uint *len, uchar *mac) SHA256_Final(mac, &ctx); } -inline void +INLINE void hmac_sha256_vector(const uchar *key, uint key_len, uint num_elem, const uchar *addr[], const uint *len, uchar *mac) { @@ -596,7 +596,7 @@ hmac_sha256_vector(const uchar *key, uint key_len, uint num_elem, sha256_vector(2, _addr, _len, mac); } -inline void +INLINE void sha256_prf_bits(const uchar *key, uint key_len, MAYBE_CONSTANT uchar *data, uint data_len, uchar *buf, uint buf_len_bits) { diff --git a/run/opencl/xsha512_kernel.cl b/run/opencl/xsha512_kernel.cl index 02c2a88629..16c480d9f8 100644 --- a/run/opencl/xsha512_kernel.cl +++ b/run/opencl/xsha512_kernel.cl @@ -21,7 +21,7 @@ typedef struct { char v[PLAINTEXT_LENGTH+1]; } xsha512_key; -inline void xsha512(__global const char *password, uint8_t pass_len, +INLINE void xsha512(__global const char *password, uint8_t pass_len, __global uint64_t *hash, uint32_t offset, __constant uint32_t *salt) { xsha512_ctx ctx; diff --git a/run/opencl/zed_kernel.cl b/run/opencl/zed_kernel.cl index dc82ce6186..6e1ac42cfd 100644 --- a/run/opencl/zed_kernel.cl +++ b/run/opencl/zed_kernel.cl @@ -35,7 +35,7 @@ typedef struct { uint32_t salt[salt_len / 4]; } zed_salt; -inline void zed_crypt(__global const uint *password, uint32_t password_length, +INLINE void zed_crypt(__global const uint *password, uint32_t password_length, __constant zed_salt *salt, __global uint *out) { uint i; diff --git a/run/opencl/zip_kernel.cl b/run/opencl/zip_kernel.cl index d181eb2f3e..1e57c94e66 100644 --- a/run/opencl/zip_kernel.cl +++ b/run/opencl/zip_kernel.cl @@ -36,7 +36,7 @@ typedef struct { #define hmac_sha1 u_hmac_sha1 #define big_hmac_sha1 u_big_hmac_sha1 -inline void preproc(const uchar *key, uint keylen, uint *state, uint padding) +INLINE void preproc(const uchar *key, uint keylen, uint *state, uint padding) { uint i; uint W[16]; @@ -63,7 +63,7 @@ inline void preproc(const uchar *key, uint keylen, uint *state, uint padding) state[4] = E + INIT_E; } -inline void hmac_sha1(uint *output, uint *ipad_state, uint *opad_state, __constant uchar *salt, int saltlen, uchar add) +INLINE void hmac_sha1(uint *output, uint *ipad_state, uint *opad_state, __constant uchar *salt, int saltlen, uchar add) { int i; uint W[16]; @@ -121,7 +121,7 @@ inline void hmac_sha1(uint *output, uint *ipad_state, uint *opad_state, __consta output[4] = E; } -inline void big_hmac_sha1(uint *input, uint inputlen, uint *ipad_state, uint *opad_state, uint *tmp_out, uint iter) +INLINE void big_hmac_sha1(uint *input, uint inputlen, uint *ipad_state, uint *opad_state, uint *tmp_out, uint iter) { uint i; uint W[16]; @@ -179,7 +179,7 @@ inline void big_hmac_sha1(uint *input, uint inputlen, uint *ipad_state, uint *op } } -inline void pbkdf2_hmac_sha1(const uchar *pass, const uint passlen, +INLINE void pbkdf2_hmac_sha1(const uchar *pass, const uint passlen, __constant uchar *salt, const uint saltlen, const uint iterations, uchar *out, const uint outlen, uint skip_bytes) { @@ -218,7 +218,7 @@ inline void pbkdf2_hmac_sha1(const uchar *pass, const uint passlen, #undef hmac_sha1 #undef big_hmac_sha1 -inline uint prepare(__global const uchar *pwbuf, __global const uint *buf_idx, uint index, uchar *password) +INLINE uint prepare(__global const uchar *pwbuf, __global const uint *buf_idx, uint index, uchar *password) { uint i; uint base = buf_idx[index];