1// freerainbowtables is a project for generating, distributing, and using 2// perfect rainbow tables 3// 4// Copyright 2010, 2011 Jan Kyska 5// Copyright 2010 Martin Westergaard Jørgensen <martinwj2005@gmail.com> 6// Copyright 2010, 2011, 2012 James Nobis <quel@quelrod.net> 7// 8// This file is part of freerainbowtables. 9// 10// freerainbowtables is free software: you can redistribute it and/or modify 11// it under the terms of the GNU General Public License as published by 12// the Free Software Foundation, either version 2 of the License, or 13// (at your option) any later version. 14// 15// freerainbowtables is distributed in the hope that it will be useful, 16// but WITHOUT ANY WARRANTY; without even the implied warranty of 17// MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the 18// GNU General Public License for more details. 19// 20// You should have received a copy of the GNU General Public License 21// along with freerainbowtables. If not, see <http://www.gnu.org/licenses/>. 22 23namespace RC_NTLM { 24 25__device__ __constant__ unsigned int h[4] = { 0x67452301, 0xEFCDAB89, 0x98BADCFE, 0x10325476 }; 26__device__ __constant__ unsigned char r[48] = { \ 27 3, 7, 11, 19, 3, 7, 11, 19, 3, 7, 11, 19, 3, 7, 11, 19, \ 28 3, 5, 9, 13, 3, 5, 9, 13, 3, 5, 9, 13, 3, 5, 9, 13, \ 29 3, 9, 11, 15, 3, 9, 11, 15, 3, 9, 11, 15, 3, 9, 11, 15 }; 30__device__ __constant__ unsigned char g[48] = { \ 31 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, \ 32 0, 4, 8, 12, 1, 5, 9, 13, 2, 6, 10, 14, 3, 7, 11, 15, \ 33 0, 8, 4, 12, 2, 10, 6, 14, 1, 9, 5, 13, 3, 11, 7, 15 }; 34 35__device__ unsigned int FF(unsigned int a, unsigned int b, unsigned int c, unsigned int d, int i, const unsigned int* data) { 36 unsigned int ret; 37 ret = a + ((b&c)|((~b)&d)) + data[SHIDX(g[i])]; 38 ret = (ret<<r[i])|(ret>>(32-r[i])); 39 return ret; 40} 41 42__device__ unsigned int GG(unsigned int a, unsigned int b, unsigned int c, unsigned int d, int i, const unsigned int* data) { 43 unsigned int ret; 44 ret = a + ((b&c)|(b&d)|(c&d)) + data[SHIDX(g[i])] + 0x5a827999u; 45 ret = (ret<<r[i])|(ret>>(32-r[i])); 46 return ret; 47} 48 49__device__ unsigned int HH(unsigned int a, unsigned int b, unsigned int c, unsigned int d, int i, const unsigned int* data) { 50 unsigned int ret; 51 ret = a + (b^c^d) + data[SHIDX(g[i])] + 0x6ed9eba1u; 52 ret = (ret<<r[i])|(ret>>(32-r[i])); 53 return ret; 54} 55 56__device__ void MD4(unsigned int* dataHash) { 57 unsigned int a = h[0], b = h[1], c = h[2], d = h[3], x; 58 int ii; 59 60 // Round 1 61 for(ii = 0; ii < 16; ii++) { 62 x = b; 63 b = FF(a, b, c, d, ii, dataHash); 64 a = d; d = c; c = x; 65 } 66 67 // Round 2 68 for(; ii < 32; ii++) { 69 x = b; 70 b = GG(a, b, c, d, ii, dataHash); 71 a = d; d = c; c = x; 72 } 73 74 // Round 3 75 for(; ii < 48; ii++) { 76 x = b; 77 b = HH(a, b, c, d, ii, dataHash); 78 a = d; d = c; c = x; 79 } 80 81 dataHash[SHIDX(0)] = a + h[0]; 82 dataHash[SHIDX(1)] = b + h[1]; 83 dataHash[SHIDX(2)] = c + h[2]; 84 dataHash[SHIDX(3)] = d + h[3]; 85} 86 87} 88 89#define RT_NTLM_KERNEL(kerName, kerPrologue, kerEpilogue) \ 90__global__ void kerName(unsigned int chainStart, unsigned int chainStop) { \ 91 unsigned int size, jj, kk; \ 92 unsigned int plain; \ 93 \ 94 kerPrologue; \ 95 \ 96 /* transform to the plain text */ \ 97 plain = 0x80; \ 98 jj = (PLAIN_MAX_SIZE>>1)+1; \ 99 \ 100 kk = 0; \ 101 RTGEN_IDX2PLAIN_BEGIN; \ 102 RTGEN_I2P_LOOP64_BEGIN; \ 103 plain = (plain<<8); \ 104 if((kk++)&1) { \ 105 hData[SHIDX(jj--)] = plain; \ 106 plain = 0; \ 107 } \ 108 plain = (plain<<8) | cplChrSet[(dimItem.z&255u) + uiDiv]; \ 109 RTGEN_I2P_LOOP64_END; \ 110 \ 111 RTGEN_I2P_LOOP32_BEGIN; \ 112 plain = (plain<<8); \ 113 if((kk++)&1) { \ 114 hData[SHIDX(jj--)] = plain; \ 115 plain = 0; \ 116 } \ 117 plain = (plain<<8) | cplChrSet[(dimItem.z&255u) + uiDiv]; \ 118 RTGEN_I2P_LOOP32_END; \ 119 RTGEN_IDX2PLAIN_END; \ 120 \ 121 /* prepare for MD4 */ \ 122 size = (ii<<1); \ 123 ii = (((kk^1)&1)<<4)+8; \ 124 plain = plain<<ii; \ 125 for(jj++, idx = 0; jj <= (PLAIN_MAX_SIZE>>1)+1; plain = hData[SHIDX(jj++)], idx++) \ 126 hData[SHIDX(idx)] = (plain>>ii)|(hData[SHIDX(jj)]<<(32-ii)); \ 127 hData[SHIDX(idx)] = plain>>ii; \ 128 for(idx++; idx < 14; idx++) \ 129 hData[SHIDX(idx)] = 0; \ 130 hData[SHIDX(idx++)] = size<<3; \ 131 hData[SHIDX(idx)] = 0; \ 132 \ 133 /* hash */ \ 134 RC_NTLM::MD4(hData); \ 135 \ 136 kerEpilogue; \ 137} 138 139 140RT_NTLM_KERNEL(RTGenNTLMKernel, RTGEN_PROLOGUE, RTGEN_EPILOGUE) 141RT_NTLM_KERNEL(RTPreCalcNTLMKernel, RTPRECALC_PROLOGUE, RTPRECALC_EPILOGUE) 142RT_NTLM_KERNEL(RTCheckAlarmNTLMKernel, RTCHKALARM_PROLOGUE, RTCHKALARM_EPILOGUE) 143