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