#define CREATE_SHA1_GEN_KERNEL(length) \ __global__ void MakeSHA1ChainLen##length(unsigned char *InitialPasswordArray, unsigned char *OutputHashArray, \ uint32_t PasswordSpaceOffset, uint32_t StartChainIndex, uint32_t StepsToRun, uint32_t charset_offset) { \ const int pass_length = length; \ uint32_t CurrentStep, PassCount, password_index; \ uint32_t b0,b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15; \ uint32_t a,b,c,d,e; \ uint32_t *InitialArray32; \ uint32_t *OutputArray32; \ InitialArray32 = (uint32_t *)InitialPasswordArray; \ OutputArray32 = (uint32_t *)OutputHashArray; \ __shared__ char charset[512]; \ copySingleCharsetToShared(charset, SHA1_Generate_Device_Charset_Constant); \ password_index = ((blockIdx.x*blockDim.x + threadIdx.x) + (PasswordSpaceOffset * SHA1_Generate_Device_Number_Of_Threads)); \ if (password_index >= SHA1_Generate_Device_Number_Of_Chains) { \ return; \ } \ clearB0toB15(b0,b1,b2,b3,b4,b5,b6,b7,b8,b9,b10,b11,b12,b13,b14,b15); \ LoadMD5RegistersFromGlobalMemory(b0,b1,b2,b3,b4,b5,b6,b7,b8,b9,b10,b11,b12,b13,b14,b15, \ InitialArray32, SHA1_Generate_Device_Number_Of_Chains, password_index, pass_length); \ for (PassCount = 0; PassCount < StepsToRun; PassCount++) { \ CurrentStep = PassCount + StartChainIndex; \ b15 = ((pass_length * 8) & 0xff) << 24 | (((pass_length * 8) >> 8) & 0xff) << 16; \ SetCharacterAtPosition(0x80, pass_length, b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15 ); \ SHA_TRANSFORM(a, b, c, d, e, b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15); \ a = reverse(a);b = reverse(b);c = reverse(c);d = reverse(d);e = reverse(e); \ clearB0toB15(b0,b1,b2,b3,b4,b5,b6,b7,b8,b9,b10,b11,b12,b13,b14,b15); \ reduceSingleCharsetNormal(b0, b1, b2, a, b, c, d, CurrentStep, charset, charset_offset, pass_length, SHA1_Generate_Device_Table_Index); \ charset_offset++; \ if (charset_offset >= SHA1_Generate_Device_Charset_Length) { \ charset_offset = 0; \ } \ } \ if (CurrentStep >= (SHA1_Generate_Device_Chain_Length - 1)) { \ OutputArray32[0 * SHA1_Generate_Device_Number_Of_Chains + password_index] = a; \ OutputArray32[1 * SHA1_Generate_Device_Number_Of_Chains + password_index] = b; \ OutputArray32[2 * SHA1_Generate_Device_Number_Of_Chains + password_index] = c; \ OutputArray32[3 * SHA1_Generate_Device_Number_Of_Chains + password_index] = d; \ OutputArray32[4 * SHA1_Generate_Device_Number_Of_Chains + password_index] = e; \ } \ else { \ SaveMD5RegistersIntoGlobalMemory(b0,b1,b2,b3,b4,b5,b6,b7,b8,b9,b10,b11,b12,b13,b14,b15, \ InitialArray32, SHA1_Generate_Device_Number_Of_Chains, password_index, pass_length); \ } \ }