1 //============================================================================
\r
2 // Name : rcuda_sha1.inc
\r
3 // Author : Jan Kyska
\r
5 // Description : SHA1 hash kernel for Generator of FreeRainbowTables
\r
6 //============================================================================
\r
10 __device__ __constant__ unsigned int h[5] = { 0x67452301, 0xEFCDAB89, 0x98BADCFE, 0x10325476, 0xC3D2E1F0 };
\r
12 __device__ unsigned int SwapEndian(unsigned int n) {
\r
13 return (n<<24)|((n&0x0000ff00)<<8)|((n>>8)&0x0000ff00)|(n>>24);
\r
16 __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
17 return ((a<<5)|(a>>27)) + ((b&c)|((~b)&d)) + e + 0x5A827999 + data[SHIDX(i)];
\r
20 __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
22 dt = data[SHIDX(i&15)]^data[SHIDX((i-3)&15)]^data[SHIDX((i-8)&15)]^data[SHIDX((i-14)&15)];
\r
23 data[SHIDX(i&15)] = dt = ((dt<<1)|(dt>>31));
\r
24 return ((a<<5)|(a>>27)) + ((b&c)|((~b)&d)) + e + 0x5A827999 + dt;
\r
27 __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
29 dt = data[SHIDX(i&15)]^data[SHIDX((i-3)&15)]^data[SHIDX((i-8)&15)]^data[SHIDX((i-14)&15)];
\r
30 data[SHIDX(i&15)] = dt = ((dt<<1)|(dt>>31));
\r
31 return ((a<<5)|(a>>27)) + (b^c^d) + e + 0x6ED9EBA1 + dt;
\r
34 __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
36 dt = data[SHIDX(i&15)]^data[SHIDX((i-3)&15)]^data[SHIDX((i-8)&15)]^data[SHIDX((i-14)&15)];
\r
37 data[SHIDX(i&15)] = dt = ((dt<<1)|(dt>>31));
\r
38 return ((a<<5)|(a>>27)) + ((b&c)|(b&d)|(c&d)) + e + 0x8F1BBCDC + dt;
\r
41 __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
43 dt = data[SHIDX(i&15)]^data[SHIDX((i-3)&15)]^data[SHIDX((i-8)&15)]^data[SHIDX((i-14)&15)];
\r
44 data[SHIDX(i&15)] = dt = ((dt<<1)|(dt>>31));
\r
45 return ((a<<5)|(a>>27)) + (b^c^d) + e + 0xCA62C1D6 + dt;
\r
48 __device__ void SHA1(unsigned int* dataHash) {
\r
49 unsigned int a = h[0], b = h[1], c = h[2], d = h[3], e = h[4], x;
\r
53 for(ii = 0; ii < 16; ii++) {
\r
54 x = FF(a, b, c, d, e, ii, dataHash);
\r
55 e = d; d = c; c = ((b<<30)|(b>>2)); b = a; a = x;
\r
57 for(; ii < 20; ii++) {
\r
58 x = FF2(a, b, c, d, e, ii, dataHash);
\r
59 e = d; d = c; c = ((b<<30)|(b>>2)); b = a; a = x;
\r
63 for(; ii < 40; ii++) {
\r
64 x = GG(a, b, c, d, e, ii, dataHash);
\r
65 e = d; d = c; c = ((b<<30)|(b>>2)); b = a; a = x;
\r
69 for(; ii < 60; ii++) {
\r
70 x = HH(a, b, c, d, e, ii, dataHash);
\r
71 e = d; d = c; c = ((b<<30)|(b>>2)); b = a; a = x;
\r
75 for(; ii < 80; ii++) {
\r
76 x = II(a, b, c, d, e, ii, dataHash);
\r
77 e = d; d = c; c = ((b<<30)|(b>>2)); b = a; a = x;
\r
80 dataHash[SHIDX(0)] = a + h[0];
\r
81 dataHash[SHIDX(1)] = b + h[1];
\r
82 dataHash[SHIDX(2)] = c + h[2];
\r
83 dataHash[SHIDX(3)] = d + h[3];
\r
84 dataHash[SHIDX(4)] = e + h[4];
\r
89 __global__ void RTGenSHA1Kernel(unsigned int chainStart, unsigned int chainStop) {
\r
91 uint64 uiDiv64, uiVal64, uiMul64;
\r
92 unsigned int uiVal, uiDiv;
\r
93 unsigned int size, jj;
\r
98 // transform to the plain text
\r
100 jj = (PLAIN_MAX_SIZE>>2)+1;
\r
102 for(ii = 0; idx64 > 0xfffffff0ull && ii < PLAIN_MAX_SIZE; ii++) {
\r
103 uiVal64 = idx64 + cplStart[ii];
\r
105 dimItem = cplDimVec[ii];
\r
107 uiMul64 = (uint64)dimItem.y<<32;
\r
108 idx64 = __umul64hi(uiVal64, uiMul64);
\r
109 uiDiv64 = uiVal64 - idx64*(uint64)dimItem.x;
\r
110 uiVal = __umulhi((unsigned int)uiDiv64, dimItem.y);
\r
111 uiDiv = (unsigned int)uiDiv64 - uiVal * dimItem.x;
\r
113 if(uiDiv >= dimItem.x) {
\r
114 uiDiv -= dimItem.x;
\r
117 plain = (plain<<8) | cplChrSet[dimItem.z + uiDiv];
\r
119 hData[SHIDX(jj--)] = plain;
\r
124 for(idx = (unsigned int)idx64; idx != 0 && ii < PLAIN_MAX_SIZE; ii++) {
\r
125 uiVal = idx + cplStart[ii];
\r
127 dimItem = cplDimVec[ii];
\r
129 idx = __umulhi(uiVal, dimItem.y);
\r
130 uiDiv = uiVal - idx*dimItem.x;
\r
131 if(uiDiv >= dimItem.x) {
\r
132 uiDiv -= dimItem.x;
\r
135 plain = (plain<<8) | cplChrSet[dimItem.z + uiDiv];
\r
137 hData[SHIDX(jj--)] = plain;
\r
142 // prepare for SHA1
\r
144 ii = ((((3-(ii&3))<<3)-1)&0x1f)+1;
\r
146 for(jj++, idx = 0; jj <= (PLAIN_MAX_SIZE>>2)+1; plain = hData[SHIDX(jj++)], idx++)
\r
147 hData[SHIDX(idx)] = RC_SHA1::SwapEndian((plain>>ii)|(hData[SHIDX(jj)]<<(32-ii)));
\r
148 hData[SHIDX(idx)] = RC_SHA1::SwapEndian(plain>>ii);
\r
149 for(idx++; idx < 14; idx++)
\r
150 hData[SHIDX(idx)] = 0;
\r
151 hData[SHIDX(idx++)] = 0;
\r
152 hData[SHIDX(idx)] = size<<3;
\r
155 RC_SHA1::SHA1(hData);
\r
157 hData[SHIDX(0)] = RC_SHA1::SwapEndian(hData[SHIDX(0)]);
\r
158 hData[SHIDX(1)] = RC_SHA1::SwapEndian(hData[SHIDX(1)]);
\r