Skip to content

Commit

Permalink
Bug fixes/tweaks for many OpenCL formats.
Browse files Browse the repository at this point in the history 
Support largest bitmaps (0xffffffff) without failing kernel build.
Fix a bug where some workgroup sizes would fail to properly copy bitmap
to local memory.
While at it, for DES/LM OpenCL formats fix a similar copy to local to
avoid modulo arithmetic.

Related to openwall#5246
  • Loading branch information
@magnumripper
magnumripper committed Apr 12, 2023
1 parent 49f9b33 commit 9db402a
Show file tree
Hide file tree
Showing 14 changed files with 170 additions and 171 deletions.
6 changes: 3 additions & 3 deletions run/opencl/DES_bs_kernel.cl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -150,15 +150,15 @@ __kernel void DES_bs_25_b(constant uint *key_map
__local DES_bs_vector s_des_bs_key[56 * WORK_GROUP_SIZE];
int s_key_offset = lid * 56;
for (i = 0; i < 56; i++)
s_des_bs_key[lid * 56 + i] = des_bs_key[section + i * gws];
s_des_bs_key[s_key_offset + i] = des_bs_key[section + i * gws];
#endif

#if USE_LOCAL_MEM
__local ushort s_key_map[768];
int lws = get_local_size(0);

for (i = 0; i < 768; i += lws)
s_key_map[(lid + i) % 768] = key_map[(lid + i) % 768];
for (i = lid; i < 768; i += lws)
s_key_map[i] = key_map[i];
#endif

#if USE_LOCAL_MEM || WORK_GROUP_SIZE > 0
Expand Down
6 changes: 3 additions & 3 deletions run/opencl/krb5tgs_kernel.cl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -189,10 +189,10 @@ __kernel void krb5tgs_init(__global const uchar *password,
#if USE_LOCAL_BITMAPS
uint lid = get_local_id(0);
uint lws = get_local_size(0);
uint __local s_bitmaps[(BITMAP_SIZE_BITS >> 5) * SELECT_CMP_STEPS];
__local uint s_bitmaps[BITMAP_SHIFT * SELECT_CMP_STEPS];

for (uint i = 0; i < (((BITMAP_SIZE_BITS >> 5) * SELECT_CMP_STEPS) / lws); i++)
s_bitmaps[i*lws + lid] = bitmaps[i*lws + lid];
for (i = lid; i < BITMAP_SHIFT * SELECT_CMP_STEPS; i+= lws)
s_bitmaps[i] = bitmaps[i];

barrier(CLK_LOCAL_MEM_FENCE);
#endif
Expand Down
4 changes: 2 additions & 2 deletions run/opencl/lm_kernel_b.cl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -181,8 +181,8 @@ __kernel void lm_bs_b(__global opencl_lm_transfer *lm_raw_keys,
#if USE_LOCAL_MEM
__local ushort s_key_idx[768];
unsigned int lws= get_local_size(0);
for (i = 0; i < 768; i += lws)
s_key_idx[(lid + i) % 768] = lm_key_idx[(lid + i) % 768];
for (i = lid; i < 768; i += lws)
s_key_idx[i] = lm_key_idx[i];
#endif
#if USE_LOCAL_MEM || WORK_GROUP_SIZE
barrier(CLK_LOCAL_MEM_FENCE);
Expand Down
65 changes: 31 additions & 34 deletions run/opencl/md4_kernel.cl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -53,11 +53,8 @@
(a) += f((b), (c), (d)) + (x); \
(a) = rotate((a), (uint)(s))

#if BITMAP_SIZE_BITS_LESS_ONE < 0xffffffff
#define BITMAP_SIZE_BITS (BITMAP_SIZE_BITS_LESS_ONE + 1)
#else
/*undefined, cause error.*/
#endif
/* This handles an input of 0xffffffffU correctly */
#define BITMAP_SHIFT ((BITMAP_MASK >> 5) + 1)

inline void md4_encrypt(uint *hash, uint *W, uint len)
{
Expand Down Expand Up @@ -189,38 +186,38 @@ inline void cmp(uint gid,
hash[3] += 0x10325476;

#if SELECT_CMP_STEPS > 4
bitmap_index = hash[0] & (BITMAP_SIZE_BITS - 1);
bitmap_index = hash[0] & BITMAP_MASK;
tmp &= (bitmaps[bitmap_index >> 5] >> (bitmap_index & 31)) & 1U;
bitmap_index = (hash[0] >> 16) & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 5) + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[1] & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 4) + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = (hash[1] >> 16) & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 5) * 3 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[2] & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 3) + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = (hash[2] >> 16) & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 5) * 5 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[3] & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 5) * 6 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = (hash[3] >> 16) & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 5) * 7 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = (hash[0] >> 16) & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[1] & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT * 2 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = (hash[1] >> 16) & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT * 3 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[2] & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT * 4 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = (hash[2] >> 16) & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT * 5 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[3] & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT * 6 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = (hash[3] >> 16) & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT * 7 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
#elif SELECT_CMP_STEPS > 2
bitmap_index = hash[3] & (BITMAP_SIZE_BITS - 1);
bitmap_index = hash[3] & BITMAP_MASK;
tmp &= (bitmaps[bitmap_index >> 5] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[2] & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 5) + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[1] & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 4) + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[0] & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 5) * 3 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[2] & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[1] & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT * 2 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[0] & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT * 3 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
#elif SELECT_CMP_STEPS > 1
bitmap_index = hash[3] & (BITMAP_SIZE_BITS - 1);
bitmap_index = hash[3] & BITMAP_MASK;
tmp &= (bitmaps[bitmap_index >> 5] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[2] & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 5) + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[2] & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
#else
bitmap_index = hash[3] & BITMAP_SIZE_BITS_LESS_ONE;
bitmap_index = hash[3] & BITMAP_MASK;
tmp &= (bitmaps[bitmap_index >> 5] >> (bitmap_index & 31)) & 1U;
#endif

Expand Down Expand Up @@ -295,10 +292,10 @@ __kernel void md4(__global uint *keys,
#if USE_LOCAL_BITMAPS
uint lid = get_local_id(0);
uint lws = get_local_size(0);
uint __local s_bitmaps[(BITMAP_SIZE_BITS >> 5) * SELECT_CMP_STEPS];
__local uint s_bitmaps[BITMAP_SHIFT * SELECT_CMP_STEPS];

for (i = 0; i < (((BITMAP_SIZE_BITS >> 5) * SELECT_CMP_STEPS) / lws); i++)
s_bitmaps[i*lws + lid] = bitmaps[i*lws + lid];
for (i = lid; i < BITMAP_SHIFT * SELECT_CMP_STEPS; i+= lws)
s_bitmaps[i] = bitmaps[i];

barrier(CLK_LOCAL_MEM_FENCE);
#endif
Expand Down
65 changes: 31 additions & 34 deletions run/opencl/md5_kernel.cl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -59,11 +59,8 @@
(a) = rotate((a), (uint)(s)); \
(a) += (b)

#if BITMAP_SIZE_BITS_LESS_ONE < 0xffffffff
#define BITMAP_SIZE_BITS (BITMAP_SIZE_BITS_LESS_ONE + 1)
#else
/*undefined, cause error.*/
#endif
/* This handles an input of 0xffffffffU correctly */
#define BITMAP_SHIFT ((BITMAP_MASK >> 5) + 1)

inline void md5_encrypt(uint *hash, uint *W, uint len)
{
Expand Down Expand Up @@ -213,38 +210,38 @@ inline void cmp(uint gid,
hash[3] += 0x10325476;

#if SELECT_CMP_STEPS > 4
bitmap_index = hash[0] & (BITMAP_SIZE_BITS - 1);
bitmap_index = hash[0] & BITMAP_MASK;
tmp &= (bitmaps[bitmap_index >> 5] >> (bitmap_index & 31)) & 1U;
bitmap_index = (hash[0] >> 16) & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 5) + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[1] & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 4) + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = (hash[1] >> 16) & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 5) * 3 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[2] & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 3) + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = (hash[2] >> 16) & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 5) * 5 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[3] & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 5) * 6 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = (hash[3] >> 16) & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 5) * 7 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = (hash[0] >> 16) & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[1] & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT * 2 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = (hash[1] >> 16) & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT * 3 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[2] & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT * 4 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = (hash[2] >> 16) & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT * 5 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[3] & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT * 6 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = (hash[3] >> 16) & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT * 7 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
#elif SELECT_CMP_STEPS > 2
bitmap_index = hash[3] & (BITMAP_SIZE_BITS - 1);
bitmap_index = hash[3] & BITMAP_MASK;
tmp &= (bitmaps[bitmap_index >> 5] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[2] & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 5) + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[1] & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 4) + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[0] & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 5) * 3 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[2] & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[1] & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT * 2 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[0] & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT * 3 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
#elif SELECT_CMP_STEPS > 1
bitmap_index = hash[3] & (BITMAP_SIZE_BITS - 1);
bitmap_index = hash[3] & BITMAP_MASK;
tmp &= (bitmaps[bitmap_index >> 5] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[2] & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 5) + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[2] & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
#else
bitmap_index = hash[3] & BITMAP_SIZE_BITS_LESS_ONE;
bitmap_index = hash[3] & BITMAP_MASK;
tmp &= (bitmaps[bitmap_index >> 5] >> (bitmap_index & 31)) & 1U;
#endif

Expand Down Expand Up @@ -319,10 +316,10 @@ __kernel void md5(__global uint *keys,
#if USE_LOCAL_BITMAPS
uint lid = get_local_id(0);
uint lws = get_local_size(0);
uint __local s_bitmaps[(BITMAP_SIZE_BITS >> 5) * SELECT_CMP_STEPS];
__local uint s_bitmaps[BITMAP_SHIFT * SELECT_CMP_STEPS];

for (i = 0; i < (((BITMAP_SIZE_BITS >> 5) * SELECT_CMP_STEPS) / lws); i++)
s_bitmaps[i*lws + lid] = bitmaps[i*lws + lid];
for (i = lid; i < BITMAP_SHIFT * SELECT_CMP_STEPS; i+= lws)
s_bitmaps[i] = bitmaps[i];

barrier(CLK_LOCAL_MEM_FENCE);
#endif
Expand Down
15 changes: 6 additions & 9 deletions run/opencl/opencl_lm_kernel_params.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -68,11 +68,8 @@ typedef unsigned WORD vtype;
lm_clear_block_8(48); \
lm_clear_block_8(56);

#if BITMAP_SIZE_BITS_LESS_ONE < 0xffffffff
#define BITMAP_SIZE_BITS (BITMAP_SIZE_BITS_LESS_ONE + 1)
#else
/*undefined, cause error.*/
#endif
/* This handles an input of 0xffffffffU correctly */
#define BITMAP_SHIFT ((BITMAP_MASK >> 5) + 1)

#define GET_HASH_0(hash, x, k, bits) \
for (bit = bits; bit < k; bit++) \
Expand Down Expand Up @@ -134,14 +131,14 @@ inline void cmp( unsigned lm_vector *B,
value[1] = 0;
GET_HASH_0(value[0], i, REQ_BITMAP_BITS, 0);
GET_HASH_1(value[1], i, REQ_BITMAP_BITS, 0);
bitmap_index = value[1] & (BITMAP_SIZE_BITS - 1);
bitmap_index = value[1] & BITMAP_MASK;
bit = (bitmaps[bitmap_index >> 5] >> (bitmap_index & 31)) & 1U;
bitmap_index = value[0] & (BITMAP_SIZE_BITS - 1);
bit &= (bitmaps[(BITMAP_SIZE_BITS >> 5) + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = value[0] & BITMAP_MASK;
bit &= (bitmaps[BITMAP_SHIFT + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
#else
value[1] = 0;
GET_HASH_1(value[1], i, REQ_BITMAP_BITS, 0);
bitmap_index = value[1] & BITMAP_SIZE_BITS_LESS_ONE;
bitmap_index = value[1] & BITMAP_MASK;
bit = (bitmaps[bitmap_index >> 5] >> (bitmap_index & 31)) & 1U;
#endif
if (bit)
Expand Down
65 changes: 31 additions & 34 deletions run/opencl/salted_sha_kernel.cl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -12,11 +12,8 @@
#include "opencl_mask.h"
#include "opencl_sha1.h"

#if BITMAP_SIZE_BITS_LESS_ONE < 0xffffffff
#define BITMAP_SIZE_BITS (BITMAP_SIZE_BITS_LESS_ONE + 1)
#else
#error BITMAP_SIZE_BITS is not defined
#endif
/* This handles an input of 0xffffffffU correctly */
#define BITMAP_SHIFT ((BITMAP_MASK >> 5) + 1)

#if SL3
#define SL3CONV - '0'
Expand Down Expand Up @@ -96,38 +93,38 @@ inline void cmp(uint gid,
hash[4] = SWAP32(hash[4]);

#if SELECT_CMP_STEPS > 4
bitmap_index = hash[0] & (BITMAP_SIZE_BITS - 1);
bitmap_index = hash[0] & BITMAP_MASK;
tmp &= (bitmaps[bitmap_index >> 5] >> (bitmap_index & 31)) & 1U;
bitmap_index = (hash[0] >> 16) & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 5) + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[1] & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 4) + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = (hash[1] >> 16) & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 5) * 3 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[2] & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 3) + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = (hash[2] >> 16) & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 5) * 5 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[3] & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 5) * 6 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = (hash[3] >> 16) & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 5) * 7 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = (hash[0] >> 16) & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[1] & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT * 2 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = (hash[1] >> 16) & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT * 3 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[2] & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT * 4 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = (hash[2] >> 16) & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT * 5 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[3] & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT * 6 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = (hash[3] >> 16) & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT * 7 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
#elif SELECT_CMP_STEPS > 2
bitmap_index = hash[3] & (BITMAP_SIZE_BITS - 1);
bitmap_index = hash[3] & BITMAP_MASK;
tmp &= (bitmaps[bitmap_index >> 5] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[2] & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 5) + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[1] & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 4) + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[0] & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 5) * 3 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[2] & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[1] & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT * 2 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[0] & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT * 3 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
#elif SELECT_CMP_STEPS > 1
bitmap_index = hash[3] & (BITMAP_SIZE_BITS - 1);
bitmap_index = hash[3] & BITMAP_MASK;
tmp &= (bitmaps[bitmap_index >> 5] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[2] & (BITMAP_SIZE_BITS - 1);
tmp &= (bitmaps[(BITMAP_SIZE_BITS >> 5) + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
bitmap_index = hash[2] & BITMAP_MASK;
tmp &= (bitmaps[BITMAP_SHIFT + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
#else
bitmap_index = hash[3] & BITMAP_SIZE_BITS_LESS_ONE;
bitmap_index = hash[3] & BITMAP_MASK;
tmp &= (bitmaps[bitmap_index >> 5] >> (bitmap_index & 31)) & 1U;
#endif

Expand Down Expand Up @@ -208,10 +205,10 @@ void sha1(__global uint *keys,
#if USE_LOCAL_BITMAPS
uint lid = get_local_id(0);
uint lws = get_local_size(0);
uint __local s_bitmaps[(BITMAP_SIZE_BITS >> 5) * SELECT_CMP_STEPS];
__local uint s_bitmaps[BITMAP_SHIFT * SELECT_CMP_STEPS];

for (i = 0; i < (((BITMAP_SIZE_BITS >> 5) * SELECT_CMP_STEPS) / lws); i++)
s_bitmaps[i*lws + lid] = bitmaps[i*lws + lid];
for (i = lid; i < BITMAP_SHIFT * SELECT_CMP_STEPS; i+= lws)
s_bitmaps[i] = bitmaps[i];

barrier(CLK_LOCAL_MEM_FENCE);
#endif
Expand Down
Loading

0 comments on commit 9db402a

Please sign in to comment.