1 //============================================================================
\r
2 // Name : rcuda_md5.inc
\r
3 // Author : Jan Kyska
\r
5 // Description : MD5 hash kernel for Generator of FreeRainbowTables
\r
6 //============================================================================
\r
10 __device__ __constant__ unsigned int h[4] = { 0x67452301, 0xEFCDAB89, 0x98BADCFE, 0x10325476 };
\r
11 __device__ __constant__ unsigned char r[64] = { \
\r
12 7, 12, 17, 22, 7, 12, 17, 22, 7, 12, 17, 22, 7, 12, 17, 22, \
\r
13 5, 9, 14, 20, 5, 9, 14, 20, 5, 9, 14, 20, 5, 9, 14, 20, \
\r
14 4, 11, 16, 23, 4, 11, 16, 23, 4, 11, 16, 23, 4, 11, 16, 23, \
\r
15 6, 10, 15, 21, 6, 10, 15, 21, 6, 10, 15, 21, 6, 10, 15, 21 };
\r
16 __device__ __constant__ unsigned char g[64] = { \
\r
17 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, \
\r
18 1, 6, 11, 0, 5, 10, 15, 4, 9, 14, 3, 8, 13, 2, 7, 12, \
\r
19 5, 8, 11, 14, 1, 4, 7, 10, 13, 0, 3, 6, 9, 12, 15, 2, \
\r
20 0, 7, 14, 5, 12, 3, 10, 1, 8, 15, 6, 13, 4, 11, 2, 9 };
\r
21 __device__ __constant__ unsigned int ac[64] = { \
\r
22 0xd76aa478, 0xe8c7b756, 0x242070db, 0xc1bdceee, \
\r
23 0xf57c0faf, 0x4787c62a, 0xa8304613, 0xfd469501, \
\r
24 0x698098d8, 0x8b44f7af, 0xffff5bb1, 0x895cd7be, \
\r
25 0x6b901122, 0xfd987193, 0xa679438e, 0x49b40821, \
\r
27 0xf61e2562, 0xc040b340, 0x265e5a51, 0xe9b6c7aa, \
\r
28 0xd62f105d, 0x02441453, 0xd8a1e681, 0xe7d3fbc8, \
\r
29 0x21e1cde6, 0xc33707d6, 0xf4d50d87, 0x455a14ed, \
\r
30 0xa9e3e905, 0xfcefa3f8, 0x676f02d9, 0x8d2a4c8a, \
\r
32 0xfffa3942, 0x8771f681, 0x6d9d6122, 0xfde5380c, \
\r
33 0xa4beea44, 0x4bdecfa9, 0xf6bb4b60, 0xbebfbc70, \
\r
34 0x289b7ec6, 0xeaa127fa, 0xd4ef3085, 0x04881d05, \
\r
35 0xd9d4d039, 0xe6db99e5, 0x1fa27cf8, 0xc4ac5665, \
\r
37 0xf4292244, 0x432aff97, 0xab9423a7, 0xfc93a039, \
\r
38 0x655b59c3, 0x8f0ccc92, 0xffeff47d, 0x85845dd1, \
\r
39 0x6fa87e4f, 0xfe2ce6e0, 0xa3014314, 0x4e0811a1, \
\r
40 0xf7537e82, 0xbd3af235, 0x2ad7d2bb, 0xeb86d391 };
\r
42 __device__ unsigned int FF(unsigned int a, unsigned int b, unsigned int c, unsigned int d, int i, const unsigned int* data) {
\r
44 ret = a + ((b&c)|((~b)&d)) + ac[i] + data[SHIDX(g[i])];
\r
45 ret = (ret<<r[i])|(ret>>(32-r[i]));
\r
50 __device__ unsigned int GG(unsigned int a, unsigned int b, unsigned int c, unsigned int d, int i, const unsigned int* data) {
\r
52 ret = a + ((b&d)|(c&(~d))) + ac[i] + data[SHIDX(g[i])];
\r
53 ret = (ret<<r[i])|(ret>>(32-r[i]));
\r
58 __device__ unsigned int HH(unsigned int a, unsigned int b, unsigned int c, unsigned int d, int i, const unsigned int* data) {
\r
60 ret = a + (b^c^d) + ac[i] + data[SHIDX(g[i])];
\r
61 ret = (ret<<r[i])|(ret>>(32-r[i]));
\r
66 __device__ unsigned int II(unsigned int a, unsigned int b, unsigned int c, unsigned int d, int i, const unsigned int* data) {
\r
68 ret = a + (c^(b|(~d))) + ac[i] + data[SHIDX(g[i])];
\r
69 ret = (ret<<r[i])|(ret>>(32-r[i]));
\r
74 __device__ void MD5(unsigned int* dataHash) {
\r
75 unsigned int a = h[0], b = h[1], c = h[2], d = h[3], x;
\r
79 for(ii = 0; ii < 16; ii++) {
\r
81 b = FF(a, b, c, d, ii, dataHash);
\r
82 a = d; d = c; c = x;
\r
86 for(; ii < 32; ii++) {
\r
88 b = GG(a, b, c, d, ii, dataHash);
\r
89 a = d; d = c; c = x;
\r
93 for(; ii < 48; ii++) {
\r
95 b = HH(a, b, c, d, ii, dataHash);
\r
96 a = d; d = c; c = x;
\r
100 for(; ii < 64; ii++) {
\r
102 b = II(a, b, c, d, ii, dataHash);
\r
103 a = d; d = c; c = x;
\r
106 dataHash[SHIDX(0)] = a + h[0];
\r
107 dataHash[SHIDX(1)] = b + h[1];
\r
108 dataHash[SHIDX(2)] = c + h[2];
\r
109 dataHash[SHIDX(3)] = d + h[3];
\r
114 __global__ void RTGenMD5Kernel(unsigned int chainStart, unsigned int chainStop) {
\r
116 uint64 uiDiv64, uiVal64, uiMul64;
\r
117 unsigned int uiVal, uiDiv;
\r
118 unsigned int size, jj;
\r
119 unsigned int plain;
\r
123 // transform to the plain text
\r
125 jj = (PLAIN_MAX_SIZE>>2)+1;
\r
127 for(ii = 0; idx64 > 0xfffffff0ull && ii < PLAIN_MAX_SIZE; ii++) {
\r
128 uiVal64 = idx64 + cplStart[ii];
\r
130 dimItem = cplDimVec[ii];
\r
132 uiMul64 = (uint64)dimItem.y<<32;
\r
133 idx64 = __umul64hi(uiVal64, uiMul64);
\r
134 uiDiv64 = uiVal64 - idx64*(uint64)dimItem.x;
\r
135 uiVal = __umulhi((unsigned int)uiDiv64, dimItem.y);
\r
136 uiDiv = (unsigned int)uiDiv64 - uiVal * dimItem.x;
\r
138 if(uiDiv >= dimItem.x) {
\r
139 uiDiv -= dimItem.x;
\r
142 plain = (plain<<8) | cplChrSet[dimItem.z + uiDiv];
\r
144 hData[SHIDX(jj--)] = plain;
\r
149 for(idx = (unsigned int)idx64; idx != 0 && ii < PLAIN_MAX_SIZE; ii++) {
\r
150 uiVal = idx + cplStart[ii];
\r
152 dimItem = cplDimVec[ii];
\r
154 idx = __umulhi(uiVal, dimItem.y);
\r
155 uiDiv = uiVal - idx*dimItem.x;
\r
156 if(uiDiv >= dimItem.x) {
\r
157 uiDiv -= dimItem.x;
\r
160 plain = (plain<<8) | cplChrSet[dimItem.z + uiDiv];
\r
162 hData[SHIDX(jj--)] = plain;
\r
169 ii = ((((3-(ii&3))<<3)-1)&0x1f)+1;
\r
171 for(jj++, idx = 0; jj <= (PLAIN_MAX_SIZE>>2)+1; plain = hData[SHIDX(jj++)], idx++)
\r
172 hData[SHIDX(idx)] = (plain>>ii)|(hData[SHIDX(jj)]<<(32-ii));
\r
173 hData[SHIDX(idx)] = plain>>ii;
\r
174 for(idx++; idx < 14; idx++)
\r
175 hData[SHIDX(idx)] = 0;
\r
176 hData[SHIDX(idx++)] = size<<3;
\r
177 hData[SHIDX(idx)] = 0;
\r
180 RC_MD5::MD5(hData);
\r