--- /dev/null
+//============================================================================\r
+// Name : rcuda_md5.inc\r
+// Author : Jan Kyska\r
+// Version : 1.00\r
+// Description : MD5 hash kernel for Generator of FreeRainbowTables\r
+//============================================================================ \r
+\r
+namespace RC_MD5 {\r
+\r
+__device__ __constant__ unsigned int h[4] = { 0x67452301, 0xEFCDAB89, 0x98BADCFE, 0x10325476 };\r
+__device__ __constant__ unsigned char r[64] = { \\r
+ 7, 12, 17, 22, 7, 12, 17, 22, 7, 12, 17, 22, 7, 12, 17, 22, \\r
+ 5, 9, 14, 20, 5, 9, 14, 20, 5, 9, 14, 20, 5, 9, 14, 20, \\r
+ 4, 11, 16, 23, 4, 11, 16, 23, 4, 11, 16, 23, 4, 11, 16, 23, \\r
+ 6, 10, 15, 21, 6, 10, 15, 21, 6, 10, 15, 21, 6, 10, 15, 21 };\r
+__device__ __constant__ unsigned char g[64] = { \\r
+ 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, \\r
+ 1, 6, 11, 0, 5, 10, 15, 4, 9, 14, 3, 8, 13, 2, 7, 12, \\r
+ 5, 8, 11, 14, 1, 4, 7, 10, 13, 0, 3, 6, 9, 12, 15, 2, \\r
+ 0, 7, 14, 5, 12, 3, 10, 1, 8, 15, 6, 13, 4, 11, 2, 9 };\r
+__device__ __constant__ unsigned int ac[64] = { \\r
+ 0xd76aa478, 0xe8c7b756, 0x242070db, 0xc1bdceee, \\r
+ 0xf57c0faf, 0x4787c62a, 0xa8304613, 0xfd469501, \\r
+ 0x698098d8, 0x8b44f7af, 0xffff5bb1, 0x895cd7be, \\r
+ 0x6b901122, 0xfd987193, 0xa679438e, 0x49b40821, \\r
+ \\r
+ 0xf61e2562, 0xc040b340, 0x265e5a51, 0xe9b6c7aa, \\r
+ 0xd62f105d, 0x02441453, 0xd8a1e681, 0xe7d3fbc8, \\r
+ 0x21e1cde6, 0xc33707d6, 0xf4d50d87, 0x455a14ed, \\r
+ 0xa9e3e905, 0xfcefa3f8, 0x676f02d9, 0x8d2a4c8a, \\r
+ \\r
+ 0xfffa3942, 0x8771f681, 0x6d9d6122, 0xfde5380c, \\r
+ 0xa4beea44, 0x4bdecfa9, 0xf6bb4b60, 0xbebfbc70, \\r
+ 0x289b7ec6, 0xeaa127fa, 0xd4ef3085, 0x04881d05, \\r
+ 0xd9d4d039, 0xe6db99e5, 0x1fa27cf8, 0xc4ac5665, \\r
+ \\r
+ 0xf4292244, 0x432aff97, 0xab9423a7, 0xfc93a039, \\r
+ 0x655b59c3, 0x8f0ccc92, 0xffeff47d, 0x85845dd1, \\r
+ 0x6fa87e4f, 0xfe2ce6e0, 0xa3014314, 0x4e0811a1, \\r
+ 0xf7537e82, 0xbd3af235, 0x2ad7d2bb, 0xeb86d391 };\r
+\r
+__device__ unsigned int FF(unsigned int a, unsigned int b, unsigned int c, unsigned int d, int i, const unsigned int* data) {\r
+ unsigned int ret;\r
+ ret = a + ((b&c)|((~b)&d)) + ac[i] + data[SHIDX(g[i])];\r
+ ret = (ret<<r[i])|(ret>>(32-r[i]));\r
+ ret += b;\r
+ return ret;\r
+}\r
+\r
+__device__ unsigned int GG(unsigned int a, unsigned int b, unsigned int c, unsigned int d, int i, const unsigned int* data) {\r
+ unsigned int ret;\r
+ ret = a + ((b&d)|(c&(~d))) + ac[i] + data[SHIDX(g[i])];\r
+ ret = (ret<<r[i])|(ret>>(32-r[i]));\r
+ ret += b;\r
+ return ret;\r
+}\r
+\r
+__device__ unsigned int HH(unsigned int a, unsigned int b, unsigned int c, unsigned int d, int i, const unsigned int* data) {\r
+ unsigned int ret;\r
+ ret = a + (b^c^d) + ac[i] + data[SHIDX(g[i])];\r
+ ret = (ret<<r[i])|(ret>>(32-r[i]));\r
+ ret += b;\r
+ return ret;\r
+}\r
+\r
+__device__ unsigned int II(unsigned int a, unsigned int b, unsigned int c, unsigned int d, int i, const unsigned int* data) {\r
+ unsigned int ret;\r
+ ret = a + (c^(b|(~d))) + ac[i] + data[SHIDX(g[i])];\r
+ ret = (ret<<r[i])|(ret>>(32-r[i]));\r
+ ret += b;\r
+ return ret;\r
+}\r
+\r
+__device__ void MD5(unsigned int* dataHash) {\r
+ unsigned int a = h[0], b = h[1], c = h[2], d = h[3], x;\r
+ int ii;\r
+\r
+ // Round 1\r
+ for(ii = 0; ii < 16; ii++) {\r
+ x = b;\r
+ b = FF(a, b, c, d, ii, dataHash);\r
+ a = d; d = c; c = x;\r
+ }\r
+\r
+ // Round 2\r
+ for(; ii < 32; ii++) {\r
+ x = b;\r
+ b = GG(a, b, c, d, ii, dataHash);\r
+ a = d; d = c; c = x;\r
+ }\r
+ \r
+ // Round 3\r
+ for(; ii < 48; ii++) {\r
+ x = b;\r
+ b = HH(a, b, c, d, ii, dataHash);\r
+ a = d; d = c; c = x;\r
+ }\r
+ \r
+ // Round 4\r
+ for(; ii < 64; ii++) {\r
+ x = b;\r
+ b = II(a, b, c, d, ii, dataHash);\r
+ a = d; d = c; c = x;\r
+ }\r
+\r
+ dataHash[SHIDX(0)] = a + h[0];\r
+ dataHash[SHIDX(1)] = b + h[1];\r
+ dataHash[SHIDX(2)] = c + h[2];\r
+ dataHash[SHIDX(3)] = d + h[3];\r
+}\r
+\r
+}\r
+\r
+__global__ void RTGenMD5Kernel(unsigned int chainStart, unsigned int chainStop) {\r
+ uint3 dimItem;\r
+ uint64 uiDiv64, uiVal64, uiMul64;\r
+ unsigned int uiVal, uiDiv;\r
+ unsigned int size, jj;\r
+ unsigned int plain;\r
+\r
+ RTGEN_PROLOGUE;\r
+ \r
+ // transform to the plain text\r
+ plain = 0x80;\r
+ jj = (PLAIN_MAX_SIZE>>2)+1;\r
+\r
+ for(ii = 0; idx64 > 0xfffffff0ull && ii < PLAIN_MAX_SIZE; ii++) {\r
+ uiVal64 = idx64 + cplStart[ii];\r
+ uiVal64--;\r
+ dimItem = cplDimVec[ii];\r
+\r
+ uiMul64 = (uint64)dimItem.y<<32;\r
+ idx64 = __umul64hi(uiVal64, uiMul64);\r
+ uiDiv64 = uiVal64 - idx64*(uint64)dimItem.x;\r
+ uiVal = __umulhi((unsigned int)uiDiv64, dimItem.y);\r
+ uiDiv = (unsigned int)uiDiv64 - uiVal * dimItem.x;\r
+ idx64 += uiVal;\r
+ if(uiDiv >= dimItem.x) {\r
+ uiDiv -= dimItem.x;\r
+ idx64++;\r
+ }\r
+ plain = (plain<<8) | cplChrSet[dimItem.z + uiDiv];\r
+ if((ii&3) == 2) {\r
+ hData[SHIDX(jj--)] = plain;\r
+ plain = 0;\r
+ }\r
+ }\r
+\r
+ for(idx = (unsigned int)idx64; idx != 0 && ii < PLAIN_MAX_SIZE; ii++) {\r
+ uiVal = idx + cplStart[ii];\r
+ uiVal--;\r
+ dimItem = cplDimVec[ii];\r
+\r
+ idx = __umulhi(uiVal, dimItem.y);\r
+ uiDiv = uiVal - idx*dimItem.x;\r
+ if(uiDiv >= dimItem.x) {\r
+ uiDiv -= dimItem.x;\r
+ idx++;\r
+ }\r
+ plain = (plain<<8) | cplChrSet[dimItem.z + uiDiv];\r
+ if((ii&3) == 2) {\r
+ hData[SHIDX(jj--)] = plain;\r
+ plain = 0;\r
+ }\r
+ }\r
+\r
+ // prepare for MD5\r
+ size = ii;\r
+ ii = ((((3-(ii&3))<<3)-1)&0x1f)+1;\r
+ plain = plain<<ii;\r
+ for(jj++, idx = 0; jj <= (PLAIN_MAX_SIZE>>2)+1; plain = hData[SHIDX(jj++)], idx++)\r
+ hData[SHIDX(idx)] = (plain>>ii)|(hData[SHIDX(jj)]<<(32-ii));\r
+ hData[SHIDX(idx)] = plain>>ii;\r
+ for(idx++; idx < 14; idx++)\r
+ hData[SHIDX(idx)] = 0;\r
+ hData[SHIDX(idx++)] = size<<3;\r
+ hData[SHIDX(idx)] = 0;\r
+\r
+ // hash\r
+ RC_MD5::MD5(hData);\r
+\r
+ RTGEN_EPILOGUE\r
+}\r
projects / freerainbowtables / blobdiff