int charsetLen, uint32_t numberOfPasswords, struct start_positions *DEVICE_Start_Positions, \ unsigned int count, unsigned char * DEVICE_Hashes, unsigned char *DEVICE_HashTable) { \ const int pass_length = length; \ uint32_t b0,b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15; \ uint32_t a,b,c,d; \ uint32_t thread_index = blockIdx.x*blockDim.x + threadIdx.x; \ uint32_t *DEVICE_Hashes_32 = (uint32_t *)DEVICE_Hashes; \ unsigned char p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15; \ UINT4 password_count = 0; \ __shared__ __align__(16) unsigned char sharedCharset[MAX_CHARSET_LENGTH * MAX_PASSWORD_LEN]; \ __shared__ __align__(16) unsigned char sharedBitmap[8192]; \ __shared__ unsigned char sharedLengths[16]; \ copyCharsetAndBitmap(sharedCharset, sharedBitmap, sharedLengths, charsetLen, pass_length); \ loadStartPositions(pass_length, thread_index, DEVICE_Start_Positions, \ p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15); \ while (password_count < count) { \ initMD(pass_length, sharedCharset, \ p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, \ b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15); \ CUDA_MD4(b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15, a, b, c, d); \ checkHashMulti(pass_length, sharedBitmap, DEVICE_HashTable, numberOfPasswords, \ DEVICE_Hashes_32, success, OutputPassword, \ p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, \ a, b, c, d, b0, b1, b2, b3, b4, b5); \ password_count++; \ incrementCounters##length##Multi(); \ } \ }