1 //============================================================================
\r
2 // Name : rcuda_md4.inc
\r
3 // Author : Jan Kyska
\r
5 // Description : MD4 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[48] = { \
\r
12 3, 7, 11, 19, 3, 7, 11, 19, 3, 7, 11, 19, 3, 7, 11, 19, \
\r
13 3, 5, 9, 13, 3, 5, 9, 13, 3, 5, 9, 13, 3, 5, 9, 13, \
\r
14 3, 9, 11, 15, 3, 9, 11, 15, 3, 9, 11, 15, 3, 9, 11, 15 };
\r
15 __device__ __constant__ unsigned char g[48] = { \
\r
16 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, \
\r
17 0, 4, 8, 12, 1, 5, 9, 13, 2, 6, 10, 14, 3, 7, 11, 15, \
\r
18 0, 8, 4, 12, 2, 10, 6, 14, 1, 9, 5, 13, 3, 11, 7, 15 };
\r
20 __device__ unsigned int FF(unsigned int a, unsigned int b, unsigned int c, unsigned int d, int i, const unsigned int* data) {
\r
22 ret = a + ((b&c)|((~b)&d)) + data[SHIDX(g[i])];
\r
23 ret = (ret<<r[i])|(ret>>(32-r[i]));
\r
27 __device__ unsigned int GG(unsigned int a, unsigned int b, unsigned int c, unsigned int d, int i, const unsigned int* data) {
\r
29 ret = a + ((b&c)|(b&d)|(c&d)) + data[SHIDX(g[i])] + 0x5a827999u;
\r
30 ret = (ret<<r[i])|(ret>>(32-r[i]));
\r
34 __device__ unsigned int HH(unsigned int a, unsigned int b, unsigned int c, unsigned int d, int i, const unsigned int* data) {
\r
36 ret = a + (b^c^d) + data[SHIDX(g[i])] + 0x6ed9eba1u;
\r
37 ret = (ret<<r[i])|(ret>>(32-r[i]));
\r
41 __device__ void MD4(unsigned int* dataHash) {
\r
42 unsigned int a = h[0], b = h[1], c = h[2], d = h[3], x;
\r
46 for(ii = 0; ii < 16; ii++) {
\r
48 b = FF(a, b, c, d, ii, dataHash);
\r
49 a = d; d = c; c = x;
\r
53 for(; ii < 32; ii++) {
\r
55 b = GG(a, b, c, d, ii, dataHash);
\r
56 a = d; d = c; c = x;
\r
60 for(; ii < 48; ii++) {
\r
62 b = HH(a, b, c, d, ii, dataHash);
\r
63 a = d; d = c; c = x;
\r
66 dataHash[SHIDX(0)] = a + h[0];
\r
67 dataHash[SHIDX(1)] = b + h[1];
\r
68 dataHash[SHIDX(2)] = c + h[2];
\r
69 dataHash[SHIDX(3)] = d + h[3];
\r
74 __global__ void RTGenMD4Kernel(unsigned int chainStart, unsigned int chainStop) {
\r
76 uint64 uiDiv64, uiVal64, uiMul64;
\r
77 unsigned int uiVal, uiDiv;
\r
78 unsigned int size, jj;
\r
83 // transform to the plain text
\r
85 jj = (PLAIN_MAX_SIZE>>2)+1;
\r
87 for(ii = 0; idx64 > 0xfffffff0ull && ii < PLAIN_MAX_SIZE; ii++) {
\r
88 uiVal64 = idx64 + cplStart[ii];
\r
90 dimItem = cplDimVec[ii];
\r
92 uiMul64 = (uint64)dimItem.y<<32;
\r
93 idx64 = __umul64hi(uiVal64, uiMul64);
\r
94 uiDiv64 = uiVal64 - idx64*(uint64)dimItem.x;
\r
95 uiVal = __umulhi((unsigned int)uiDiv64, dimItem.y);
\r
96 uiDiv = (unsigned int)uiDiv64 - uiVal * dimItem.x;
\r
98 if(uiDiv >= dimItem.x) {
\r
102 plain = (plain<<8) | cplChrSet[dimItem.z + uiDiv];
\r
104 hData[SHIDX(jj--)] = plain;
\r
109 for(idx = (unsigned int)idx64; idx != 0 && ii < PLAIN_MAX_SIZE; ii++) {
\r
110 uiVal = idx + cplStart[ii];
\r
112 dimItem = cplDimVec[ii];
\r
114 idx = __umulhi(uiVal, dimItem.y);
\r
115 uiDiv = uiVal - idx*dimItem.x;
\r
116 if(uiDiv >= dimItem.x) {
\r
117 uiDiv -= dimItem.x;
\r
120 plain = (plain<<8) | cplChrSet[dimItem.z + uiDiv];
\r
122 hData[SHIDX(jj--)] = plain;
\r
129 ii = ((((3-(ii&3))<<3)-1)&0x1f)+1;
\r
131 for(jj++, idx = 0; jj <= (PLAIN_MAX_SIZE>>2)+1; plain = hData[SHIDX(jj++)], idx++)
\r
132 hData[SHIDX(idx)] = (plain>>ii)|(hData[SHIDX(jj)]<<(32-ii));
\r
133 hData[SHIDX(idx)] = plain>>ii;
\r
134 for(idx++; idx < 14; idx++)
\r
135 hData[SHIDX(idx)] = 0;
\r
136 hData[SHIDX(idx++)] = size<<3;
\r
137 hData[SHIDX(idx)] = 0;
\r
140 RC_MD4::MD4(hData);
\r