opencl_misc.h: Define INLINE and NOINLINE macros

Define NOINLINE as "__attribute__((noinline))" which has been seen working. Replace all uses of the latter with the macro. Drop the questionable "inline" macro and instead define INLINE as 'static inline', which should be the right thing, for anything put POCL and MESA which we've seen problems with in the past. Then use this macro only for inlines that were replaced by the old inline macro (for now). Closes #5618
openwall · Dec 25, 2024 · 433dcac · 433dcac
1 parent 8e52f7d
commit 433dcac
Show file tree

Hide file tree

Showing 72 changed files with 299 additions and 302 deletions.
diff --git 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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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,
-Original file line number
+Diff line change
@@ Expand Up / @@ -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;
@@ Expand Down @@