--- /dev/null
+//============================================================================\r
+// Name : rcuda_sha1.inc\r
+// Author : Jan Kyska\r
+// Version : 1.00\r
+// Description : SHA1 hash kernel for Generator of FreeRainbowTables\r
+//============================================================================ \r
+\r
+namespace RC_SHA1 {\r
+\r
+__device__ __constant__ unsigned int h[5] = { 0x67452301, 0xEFCDAB89, 0x98BADCFE, 0x10325476, 0xC3D2E1F0 };\r
+\r
+__device__ unsigned int SwapEndian(unsigned int n) {\r
+ return (n<<24)|((n&0x0000ff00)<<8)|((n>>8)&0x0000ff00)|(n>>24);\r
+}\r
+\r
+__device__ unsigned int FF(unsigned int a, unsigned int b, unsigned int c, unsigned int d, unsigned int e, int i, unsigned int* data) {\r
+ return ((a<<5)|(a>>27)) + ((b&c)|((~b)&d)) + e + 0x5A827999 + data[SHIDX(i)];\r
+}\r
+\r
+__device__ unsigned int FF2(unsigned int a, unsigned int b, unsigned int c, unsigned int d, unsigned int e, int i, unsigned int* data) {\r
+ unsigned int dt;\r
+ dt = data[SHIDX(i&15)]^data[SHIDX((i-3)&15)]^data[SHIDX((i-8)&15)]^data[SHIDX((i-14)&15)];\r
+ data[SHIDX(i&15)] = dt = ((dt<<1)|(dt>>31));\r
+ return ((a<<5)|(a>>27)) + ((b&c)|((~b)&d)) + e + 0x5A827999 + dt;\r
+}\r
+\r
+__device__ unsigned int GG(unsigned int a, unsigned int b, unsigned int c, unsigned int d, unsigned int e, int i, unsigned int* data) {\r
+ unsigned int dt;\r
+ dt = data[SHIDX(i&15)]^data[SHIDX((i-3)&15)]^data[SHIDX((i-8)&15)]^data[SHIDX((i-14)&15)];\r
+ data[SHIDX(i&15)] = dt = ((dt<<1)|(dt>>31));\r
+ return ((a<<5)|(a>>27)) + (b^c^d) + e + 0x6ED9EBA1 + dt;\r
+}\r
+\r
+__device__ unsigned int HH(unsigned int a, unsigned int b, unsigned int c, unsigned int d, unsigned int e, int i, unsigned int* data) {\r
+ unsigned int dt;\r
+ dt = data[SHIDX(i&15)]^data[SHIDX((i-3)&15)]^data[SHIDX((i-8)&15)]^data[SHIDX((i-14)&15)];\r
+ data[SHIDX(i&15)] = dt = ((dt<<1)|(dt>>31));\r
+ return ((a<<5)|(a>>27)) + ((b&c)|(b&d)|(c&d)) + e + 0x8F1BBCDC + dt;\r
+}\r
+\r
+__device__ unsigned int II(unsigned int a, unsigned int b, unsigned int c, unsigned int d, unsigned int e, int i, unsigned int* data) {\r
+ unsigned int dt;\r
+ dt = data[SHIDX(i&15)]^data[SHIDX((i-3)&15)]^data[SHIDX((i-8)&15)]^data[SHIDX((i-14)&15)];\r
+ data[SHIDX(i&15)] = dt = ((dt<<1)|(dt>>31));\r
+ return ((a<<5)|(a>>27)) + (b^c^d) + e + 0xCA62C1D6 + dt;\r
+}\r
+\r
+__device__ void SHA1(unsigned int* dataHash) {\r
+ unsigned int a = h[0], b = h[1], c = h[2], d = h[3], e = h[4], x;\r
+ int ii;\r
+\r
+ // Round 1\r
+ for(ii = 0; ii < 16; ii++) {\r
+ x = FF(a, b, c, d, e, ii, dataHash);\r
+ e = d; d = c; c = ((b<<30)|(b>>2)); b = a; a = x;\r
+ }\r
+ for(; ii < 20; ii++) {\r
+ x = FF2(a, b, c, d, e, ii, dataHash);\r
+ e = d; d = c; c = ((b<<30)|(b>>2)); b = a; a = x;\r
+ }\r
+ \r
+ // Round 2\r
+ for(; ii < 40; ii++) {\r
+ x = GG(a, b, c, d, e, ii, dataHash);\r
+ e = d; d = c; c = ((b<<30)|(b>>2)); b = a; a = x;\r
+ }\r
+\r
+ // Round 3\r
+ for(; ii < 60; ii++) {\r
+ x = HH(a, b, c, d, e, ii, dataHash);\r
+ e = d; d = c; c = ((b<<30)|(b>>2)); b = a; a = x;\r
+ }\r
+ \r
+ // Round 4\r
+ for(; ii < 80; ii++) {\r
+ x = II(a, b, c, d, e, ii, dataHash);\r
+ e = d; d = c; c = ((b<<30)|(b>>2)); b = a; a = 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
+ dataHash[SHIDX(4)] = e + h[4];\r
+}\r
+\r
+}\r
+\r
+__global__ void RTGenSHA1Kernel(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 SHA1\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)] = RC_SHA1::SwapEndian((plain>>ii)|(hData[SHIDX(jj)]<<(32-ii)));\r
+ hData[SHIDX(idx)] = RC_SHA1::SwapEndian(plain>>ii);\r
+ for(idx++; idx < 14; idx++)\r
+ hData[SHIDX(idx)] = 0;\r
+ hData[SHIDX(idx++)] = 0;\r
+ hData[SHIDX(idx)] = size<<3;\r
+\r
+ // hash\r
+ RC_SHA1::SHA1(hData);\r
+ \r
+ hData[SHIDX(0)] = RC_SHA1::SwapEndian(hData[SHIDX(0)]);\r
+ hData[SHIDX(1)] = RC_SHA1::SwapEndian(hData[SHIDX(1)]);\r
+ RTGEN_EPILOGUE\r
+}\r
projects / freerainbowtables / blobdiff