diff --git a/src/opencl/md5_kernel.cl b/src/opencl/md5_kernel.cl index c3e68e1..a3940fc 100644 --- a/src/opencl/md5_kernel.cl +++ b/src/opencl/md5_kernel.cl @@ -35,7 +35,7 @@ (a) = rotate(a, (uint)s); \ (a) += (b); -#define GET(i) (key[(i)]) +#define GET(i) (((i) < 8) ? key[(i)] : 0) __constant char alpha_set[] = { 'a', 'b', 'c', 'd', 'e', 'f', 'g', 'h', 'i', 'j', 'k', 'l', 'm', 'n', 'o', 'p', 'q', 'r', 's', 't', 'u', 'v', 'w', 'x', 'y', 'z', 'A', 'B', 'C', 'D', 'E', 'F', 'G', 'H', 'I', 'J', 'K', 'L', 'M', 'N', 'O', 'P', 'Q', 'R', 'S', 'T', 'U', 'V', 'W', 'X', 'Y', 'Z' @@ -47,8 +47,8 @@ __constant char alpha_set[] = { #define MD5_PASSWORD_HASH_SIZE_1 0x10000 // 64K #define MD5_PASSWORD_HASH_SIZE_2 0x1000000 // 16M -static uint get_hash_low(uint hash) { return hash & 0xFFFF; } -static uint get_hash_high(uint hash) { return hash & 0xFFFFFF; } +#define MASK_HASH_SIZE_1 0xFFFF +#define MASK_HASH_SIZE_2 0xFFFFFF // kernel void create_bitmaps(global const uint *data_info, global const uint* loaded_hash, global uint* bitmaps, global int* hashtable, global int* loaded_next_hash, global int* semaphor) // { @@ -81,19 +81,15 @@ static uint get_hash_high(uint hash) { return hash & 0xFFFFFF; } /* some constants used below magically appear after make */ //#define KEY_LENGTH (MD5_PLAINTEXT_LENGTH + 1) -inline int bitmaps_val(uint hash, bool use_local, global uint* bitmaps[4], local uint* local_bitmaps, int index, uint bitmaps_num) +inline int bitmaps_val(uint hash, bool use_local, global uint* bitmap, local uint* local_bitmaps, int index, uint bitmaps_num) { - int val = 0; - if (use_local) - hash = get_hash_low(hash); - else - hash = get_hash_high(hash); + int val; if (use_local) - val = local_bitmaps[index * bitmaps_num + hash / (sizeof(*bitmaps[0]) *8)] & (1U << (hash % (sizeof(*bitmaps[0]) *8))); + val = local_bitmaps[index * bitmaps_num + hash / (sizeof(bitmap[0]) *8)] & (1U << (hash % (sizeof(bitmap[0]) *8))); else - val = bitmaps[index][hash / (sizeof(*bitmaps[0]) *8)] & - (1U << (hash % (sizeof(*bitmaps[0]) *8))); + val = bitmap[hash / (sizeof(bitmap[0]) *8)] & + (1U << (hash % (sizeof(bitmap[0]) *8))); return val; } @@ -114,7 +110,6 @@ __kernel void md5(__global uint *data_info, __global const uint * keys, __global uint loaded_count = data_info[2]; int base = id * (KEY_LENGTH / 4); uint hash_num = data_info[3]; - global uint* bitmaps[4] = {bitmaps0, bitmaps1, bitmaps2, bitmaps3}; #if 0 __private uint p_loaded_hash[2]; @@ -128,23 +123,25 @@ __kernel void md5(__global uint *data_info, __global const uint * keys, __global p_loaded_hash[i] = loaded_hash[i]; } else #endif + uint hash_mask = MASK_HASH_SIZE_2; if (hash_num == MD5_PASSWORD_HASH_SIZE_1) { + hash_mask = MASK_HASH_SIZE_1; bitmaps_num = (hash_num+sizeof(int)*8-1)/(sizeof(int)*8); for (int i = 0; i < bitmaps_num; i+=lws) { uint index = i+lid; if (index < bitmaps_num) { - local_bitmaps[index] = bitmaps[0][index]; - local_bitmaps[index+bitmaps_num] = bitmaps[1][index]; - local_bitmaps[index+bitmaps_num*2] = bitmaps[2][index]; - local_bitmaps[index+bitmaps_num*3] = bitmaps[3][index]; + local_bitmaps[index] = bitmaps0[index]; + local_bitmaps[index+bitmaps_num] = bitmaps1[index]; + local_bitmaps[index+bitmaps_num*2] = bitmaps2[index]; + local_bitmaps[index+bitmaps_num*3] = bitmaps3[index]; } } use_local = true; barrier(CLK_LOCAL_MEM_FENCE); } - - uint key[16] = { 0 }; + + uint key[17] = { 0 }; uint i; for (i = 0; i != (KEY_LENGTH / 4) && keys[base + i]; i++) @@ -154,31 +151,28 @@ __kernel void md5(__global uint *data_info, __global const uint * keys, __global char *p = (char *) key; for (i = 0; i != 64 && p[i]; i++); - int origin_i = i; - int loop_num = loaded_count == 0 ? 0 : ALPHA_SET_SIZE; - // -1 for add none character - for (int alpha_i = -1; alpha_i < loop_num; ++alpha_i) { - for (int alpha_j = -1; alpha_j < loop_num; ++alpha_j) { + uint origin_i = i; + uint length = i + 2; + int loop_num = ALPHA_SET_SIZE; - i = origin_i; - // Generate key - if (alpha_i != -1 && alpha_j != -1) { - PUTCHAR(key, i, alpha_set[alpha_i]); - ++i; - PUTCHAR(key, i, alpha_set[alpha_j]); - ++i; - } else if ((alpha_i != -1 && alpha_j == -1) || (alpha_i == -1 && alpha_j != -1)) - continue; - uint length = i; - - //p[i] = 0x80; - //p[56] = i << 3; - //p[57] = i >> 5; + if (loaded_count == 0) { + loop_num = 1; + length = i; + origin_i = 64; + } + + PUTCHAR(key, length, 0x80); +// PUTCHAR(key, 56, length << 3); +// PUTCHAR(key, 57, length >> 5); - PUTCHAR(key, i, 0x80); - PUTCHAR(key, 56, i << 3); - PUTCHAR(key, 57, i >> 5); + length <<= 3; + for (int alpha_i = 0; alpha_i < loop_num; ++alpha_i) { + for (int alpha_j = 0; alpha_j < loop_num; ++alpha_j) { + // Generate key + PUTCHAR(key, origin_i, alpha_set[alpha_i]); + PUTCHAR(key, origin_i + 1, alpha_set[alpha_j]); + uint a, b, c, d; a = 0x67452301; b = 0xefcdab89; @@ -204,7 +198,7 @@ __kernel void md5(__global uint *data_info, __global const uint * keys, __global STEP(F, b, c, d, a, GET(11), 0x895cd7be, 22); STEP(F, a, b, c, d, GET(12), 0x6b901122, 7); STEP(F, d, a, b, c, GET(13), 0xfd987193, 12); - STEP(F, c, d, a, b, GET(14), 0xa679438e, 17); + STEP(F, c, d, a, b, length, 0xa679438e, 17); STEP(F, b, c, d, a, 0, 0x49b40821, 22); /* Round 2 */ @@ -217,7 +211,7 @@ __kernel void md5(__global uint *data_info, __global const uint * keys, __global STEP(G, c, d, a, b, 0, 0xd8a1e681, 14); STEP(G, b, c, d, a, GET(4), 0xe7d3fbc8, 20); STEP(G, a, b, c, d, GET(9), 0x21e1cde6, 5); - STEP(G, d, a, b, c, GET(14), 0xc33707d6, 9); + STEP(G, d, a, b, c, length, 0xc33707d6, 9); STEP(G, c, d, a, b, GET(3), 0xf4d50d87, 14); STEP(G, b, c, d, a, GET(8), 0x455a14ed, 20); STEP(G, a, b, c, d, GET(13), 0xa9e3e905, 5); @@ -229,7 +223,7 @@ __kernel void md5(__global uint *data_info, __global const uint * keys, __global STEP(H, a, b, c, d, GET(5), 0xfffa3942, 4); STEP(H, d, a, b, c, GET(8), 0x8771f681, 11); STEP(H, c, d, a, b, GET(11), 0x6d9d6122, 16); - STEP(H, b, c, d, a, GET(14), 0xfde5380c, 23); + STEP(H, b, c, d, a, length, 0xfde5380c, 23); STEP(H, a, b, c, d, GET(1), 0xa4beea44, 4); STEP(H, d, a, b, c, GET(4), 0x4bdecfa9, 11); STEP(H, c, d, a, b, GET(7), 0xf6bb4b60, 16); @@ -246,7 +240,7 @@ __kernel void md5(__global uint *data_info, __global const uint * keys, __global /* Round 4 */ STEP(I, a, b, c, d, GET(0), 0xf4292244, 6); STEP(I, d, a, b, c, GET(7), 0x432aff97, 10); - STEP(I, c, d, a, b, GET(14), 0xab9423a7, 15); + STEP(I, c, d, a, b, length, 0xab9423a7, 15); STEP(I, b, c, d, a, GET(5), 0xfc93a039, 21); STEP(I, a, b, c, d, GET(12), 0x655b59c3, 6); STEP(I, d, a, b, c, GET(3), 0x8f0ccc92, 10); @@ -263,18 +257,7 @@ __kernel void md5(__global uint *data_info, __global const uint * keys, __global // Compare the hashes and return the matched count if (loaded_count != 0) { - uint hash; - - if (hash_num == MD5_PASSWORD_HASH_SIZE_2) { - hash = get_hash_high(a); - } - else { - hash = get_hash_low(a); - } - - int val[4] = {0}; - val[0] = bitmaps_val(a, use_local, bitmaps, local_bitmaps, 0, bitmaps_num); - if (val[0]) { + if (bitmaps_val(a & hash_mask, use_local, bitmaps0, local_bitmaps, 0, bitmaps_num)) { bool have_b = 0; if (!have_b) { a -= 0x67452301; @@ -287,37 +270,31 @@ __kernel void md5(__global uint *data_info, __global const uint * keys, __global d += 0x10325476; have_b = 1; - val[1] = bitmaps_val(b, use_local, bitmaps, local_bitmaps, 1, bitmaps_num); - if (val[1]) { - val[2] = bitmaps_val(c, use_local, bitmaps, local_bitmaps, 2, bitmaps_num); - if (val[2]) { - val[3] = bitmaps_val(d, use_local, bitmaps, local_bitmaps, 3, bitmaps_num); - if (val[3]) { - if (use_local) - hash = get_hash_low(d); - else - hash = get_hash_high(d); - - int hash_index = hashtable[hash >> MD5_HASH_SHR]; + if (bitmaps_val(b & hash_mask, use_local, bitmaps1, local_bitmaps, 1, bitmaps_num)) + if (bitmaps_val(c & hash_mask, use_local, bitmaps2, local_bitmaps, 2, bitmaps_num)) + if (bitmaps_val(d & hash_mask, use_local, bitmaps3, local_bitmaps, 3, bitmaps_num)) { + int hash_index = hashtable[(d & hash_mask) >> MD5_HASH_SHR]; if ( hash_index != -1) do { if ( a == loaded_hash[hash_index] && b == loaded_hash[hash_index+loaded_count] && c == loaded_hash[hash_index+loaded_count*2] && d == loaded_hash[hash_index+loaded_count*3] ){ uint index = atom_inc(matched_count); uint m_base = index * KEY_LENGTH; - PUTCHAR(key, length, '\0'); + PUTCHAR(key, length >> 3, '\0'); char *q = (char*)key; - for (int j = 0; j <= length; ++j) + for (int j = 0; j <= length >> 3; ++j) PUTCHAR(matched_keys, m_base+j, q[j]); + PUTCHAR(key, length >> 3, 0x80); + hashes[index] = a; hashes[loaded_count + index] = b; hashes[2 * loaded_count + index] = c; hashes[3 * loaded_count + index] = d; } } while ((hash_index = loaded_next_hash[hash_index])!=-1); - }}} + } } } } diff --git a/src/opencl_rawmd5_fmt.c b/src/opencl_rawmd5_fmt.c index 33dba8b..92692ff 100644 --- a/src/opencl_rawmd5_fmt.c +++ b/src/opencl_rawmd5_fmt.c @@ -43,9 +43,9 @@ static cl_int *bitmaps[4], *hashtable, *loaded_next_hash; #define MD5_PASSWORD_HASH_SIZE_1 0x10000 // 64K #define MD5_PASSWORD_HASH_SIZE_2 0x1000000 // 16M -#define MD5_PASSWORD_HASH_THRESHOLD_0 0 -#define MD5_PASSWORD_HASH_THRESHOLD_1 204 -#define MD5_PASSWORD_HASH_THRESHOLD_2 6553 +//#define MD5_PASSWORD_HASH_THRESHOLD_0 0 +//#define MD5_PASSWORD_HASH_THRESHOLD_1 204 +#define MD5_PASSWORD_HASH_THRESHOLD_2 MD5_PASSWORD_HASH_SIZE_1 #define MD5_HASH_SHR 2 @@ -110,12 +110,12 @@ static void release_hash() static void create_clobj(int kpc){ pinned_saved_keys = clCreateBuffer(context[gpu_id], CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, - (PLAINTEXT_LENGTH + 1) * kpc, NULL, &ret_code); + (PLAINTEXT_LENGTH + 1) * kpc + 7, NULL, &ret_code); HANDLE_CLERROR(ret_code, "Error creating page-locked memory pinned_saved_keys"); saved_plain = (char *) clEnqueueMapBuffer(queue[gpu_id], pinned_saved_keys, CL_TRUE, CL_MAP_WRITE | CL_MAP_READ, 0, - (PLAINTEXT_LENGTH + 1) * kpc, 0, NULL, NULL, &ret_code); + (PLAINTEXT_LENGTH + 1) * kpc + 7, 0, NULL, NULL, &ret_code); HANDLE_CLERROR(ret_code, "Error mapping page-locked memory saved_plain"); // create hashes @@ -463,7 +463,7 @@ static void reset(struct db_main *db) int index = 0; int loaded_hash_size; size_t lws, gws; - size_t hash_num; + static size_t hash_num = 0; size_t bitmaps_num; size_t hashtable_num; int i, j; @@ -522,6 +522,7 @@ static void reset(struct db_main *db) HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], buffer_loaded_hash, CL_TRUE, 0, loaded_hash_size, loaded_hash, 0, NULL, NULL), "failed in clEnqueueWriteBuffer buffer_loaded_hash"); +printf("hash_num = %d\n", hash_num); // CPU bitmaps bitmaps_num = (hash_num+sizeof(cl_uint)*8-1)/(sizeof(cl_uint)*8); hashtable_num = (hash_num >> MD5_HASH_SHR); @@ -674,7 +675,7 @@ static int crypt_all(int *pcount, struct db_salt *salt) /* for (i = 0; i < matched_count; ++i) */ /* printf("mcount: %d mpass %s\n", matched_count, matched_keys+i*(PLAINTEXT_LENGTH+1)); */ - *pcount *= 53*53; + *pcount *= 52*52; } // read back partial hashes