Advertisement
Guest User

Untitled

a guest
Nov 1st, 2014
175
0
Never
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
text 15.19 KB | None | 0 0
  1. #define uint8 char
  2. #define int8 char
  3. #define uint16 ushort
  4. #define int16 short
  5. #define uint32 uint
  6. #define int32 int
  7. #define uint64 ulong
  8. #define int64 long
  9.  
  10. #define FASTSHA
  11.  
  12. GENERATED__CONSTANTS
  13.  
  14. // FNV hash: http://isthe.com/chongo/tech/comp/fnv/#FNV-source
  15. #define OFFSET_BASIS 2166136261u
  16. #define FNV_PRIME 16777619u
  17. #define fnv_hash_w3(w1,w2,w3) (uint)((((((OFFSET_BASIS ^ rotate5(w1)) * FNV_PRIME) ^ rotate5(w2)) * FNV_PRIME) ^ rotate5(w3)) * FNV_PRIME)
  18. #define fnv_hash_w5(w1,w2,w3,w4,w5) (uint)((((((((((OFFSET_BASIS ^ rotate5(w1)) * FNV_PRIME) ^ rotate5(w2)) * FNV_PRIME) ^ rotate5(w3)) * FNV_PRIME) ^ rotate5(w4)) * FNV_PRIME) ^ rotate5(w5)) * FNV_PRIME)
  19.  
  20. #ifdef FASTSHA
  21. inline uint32 andnot(uint32 a,uint32 b) { return a & ~b; }
  22. inline uint32 rotate1(uint32 a) { return (a << 1) | (a >> 31); }
  23. inline uint32 rotate5(uint32 a) { return (a << 5) | (a >> 27); }
  24. inline uint32 rotate30(uint32 a) { return (a << 30) | (a >> 2); }
  25.  
  26. // block size: 512b = 64B = 16W
  27. // in (80W long) is the 16W of work + scratch space (prepadded)
  28. // H (5W long) is the current hash state
  29. // this function taken from NearSHA, http://cr.yp.to/nearsha.html
  30. void sha1_block(uint32 *in, uint32 *H)
  31. {
  32. unsigned int a = H[0];
  33. unsigned int b = H[1];
  34. unsigned int c = H[2];
  35. unsigned int d = H[3];
  36. unsigned int e = H[4];
  37. unsigned int f;
  38. unsigned int x0,x1,x2,x3,x4,x5,x6,x7,x8,x9,x10,x11,x12,x13,x14,x15;
  39.  
  40. x0 = in[0];
  41. f = (c & b) | andnot(d,b);
  42. e = rotate5(a) + f + e + 0x5a827999 + x0;
  43. b = rotate30(b);
  44. x1 = in[1];
  45. f = (b & a) | andnot(c,a);
  46. d = rotate5(e) + f + d + 0x5a827999 + x1;
  47. a = rotate30(a);
  48. x2 = in[2];
  49. f = (a & e) | andnot(b,e);
  50. c = rotate5(d) + f + c + 0x5a827999 + x2;
  51. e = rotate30(e);
  52. x3 = in[3];
  53. f = (e & d) | andnot(a,d);
  54. b = rotate5(c) + f + b + 0x5a827999 + x3;
  55. d = rotate30(d);
  56. x4 = in[4];
  57. f = (d & c) | andnot(e,c);
  58. a = rotate5(b) + f + a + 0x5a827999 + x4;
  59. c = rotate30(c);
  60. x5 = in[5];
  61. f = (c & b) | andnot(d,b);
  62. e = rotate5(a) + f + e + 0x5a827999 + x5;
  63. b = rotate30(b);
  64. x6 = in[6];
  65. f = (b & a) | andnot(c,a);
  66. d = rotate5(e) + f + d + 0x5a827999 + x6;
  67. a = rotate30(a);
  68. x7 = in[7];
  69. f = (a & e) | andnot(b,e);
  70. c = rotate5(d) + f + c + 0x5a827999 + x7;
  71. e = rotate30(e);
  72. x8 = in[8];
  73. f = (e & d) | andnot(a,d);
  74. b = rotate5(c) + f + b + 0x5a827999 + x8;
  75. d = rotate30(d);
  76. x9 = in[9];
  77. f = (d & c) | andnot(e,c);
  78. a = rotate5(b) + f + a + 0x5a827999 + x9;
  79. c = rotate30(c);
  80. x10 = in[10];
  81. f = (c & b) | andnot(d,b);
  82. e = rotate5(a) + f + e + 0x5a827999 + x10;
  83. b = rotate30(b);
  84. x11 = in[11];
  85. f = (b & a) | andnot(c,a);
  86. d = rotate5(e) + f + d + 0x5a827999 + x11;
  87. a = rotate30(a);
  88. x12 = in[12];
  89. f = (a & e) | andnot(b,e);
  90. c = rotate5(d) + f + c + 0x5a827999 + x12;
  91. e = rotate30(e);
  92. x13 = in[13];
  93. f = (e & d) | andnot(a,d);
  94. b = rotate5(c) + f + b + 0x5a827999 + x13;
  95. d = rotate30(d);
  96. x14 = in[14];
  97. f = (d & c) | andnot(e,c);
  98. a = rotate5(b) + f + a + 0x5a827999 + x14;
  99. c = rotate30(c);
  100. x15 = in[15];
  101. f = (c & b) | andnot(d,b);
  102. e = rotate5(a) + f + e + 0x5a827999 + x15;
  103. b = rotate30(b);
  104. x0 = rotate1(x13 ^ x8 ^ x2 ^ x0);
  105. f = (b & a) | andnot(c,a);
  106. d = rotate5(e) + f + d + 0x5a827999 + x0;
  107. a = rotate30(a);
  108. x1 = rotate1(x14 ^ x9 ^ x3 ^ x1);
  109. f = (a & e) | andnot(b,e);
  110. c = rotate5(d) + f + c + 0x5a827999 + x1;
  111. e = rotate30(e);
  112. x2 = rotate1(x15 ^ x10 ^ x4 ^ x2);
  113. f = (e & d) | andnot(a,d);
  114. b = rotate5(c) + f + b + 0x5a827999 + x2;
  115. d = rotate30(d);
  116. x3 = rotate1(x0 ^ x11 ^ x5 ^ x3);
  117. f = (d & c) | andnot(e,c);
  118. a = rotate5(b) + f + a + 0x5a827999 + x3;
  119. c = rotate30(c);
  120. x4 = rotate1(x1 ^ x12 ^ x6 ^ x4);
  121. f = b ^ c ^ d;
  122. e = rotate5(a) + f + e + 0x6ed9eba1 + x4;
  123. b = rotate30(b);
  124. x5 = rotate1(x2 ^ x13 ^ x7 ^ x5);
  125. f = a ^ b ^ c;
  126. d = rotate5(e) + f + d + 0x6ed9eba1 + x5;
  127. a = rotate30(a);
  128. x6 = rotate1(x3 ^ x14 ^ x8 ^ x6);
  129. f = e ^ a ^ b;
  130. c = rotate5(d) + f + c + 0x6ed9eba1 + x6;
  131. e = rotate30(e);
  132. x7 = rotate1(x4 ^ x15 ^ x9 ^ x7);
  133. f = d ^ e ^ a;
  134. b = rotate5(c) + f + b + 0x6ed9eba1 + x7;
  135. d = rotate30(d);
  136. x8 = rotate1(x5 ^ x0 ^ x10 ^ x8);
  137. f = c ^ d ^ e;
  138. a = rotate5(b) + f + a + 0x6ed9eba1 + x8;
  139. c = rotate30(c);
  140. x9 = rotate1(x6 ^ x1 ^ x11 ^ x9);
  141. f = b ^ c ^ d;
  142. e = rotate5(a) + f + e + 0x6ed9eba1 + x9;
  143. b = rotate30(b);
  144. x10 = rotate1(x7 ^ x2 ^ x12 ^ x10);
  145. f = a ^ b ^ c;
  146. d = rotate5(e) + f + d + 0x6ed9eba1 + x10;
  147. a = rotate30(a);
  148. x11 = rotate1(x8 ^ x3 ^ x13 ^ x11);
  149. f = e ^ a ^ b;
  150. c = rotate5(d) + f + c + 0x6ed9eba1 + x11;
  151. e = rotate30(e);
  152. x12 = rotate1(x9 ^ x4 ^ x14 ^ x12);
  153. f = d ^ e ^ a;
  154. b = rotate5(c) + f + b + 0x6ed9eba1 + x12;
  155. d = rotate30(d);
  156. x13 = rotate1(x10 ^ x5 ^ x15 ^ x13);
  157. f = c ^ d ^ e;
  158. a = rotate5(b) + f + a + 0x6ed9eba1 + x13;
  159. c = rotate30(c);
  160. x14 = rotate1(x11 ^ x6 ^ x0 ^ x14);
  161. f = b ^ c ^ d;
  162. e = rotate5(a) + f + e + 0x6ed9eba1 + x14;
  163. b = rotate30(b);
  164. x15 = rotate1(x12 ^ x7 ^ x1 ^ x15);
  165. f = a ^ b ^ c;
  166. d = rotate5(e) + f + d + 0x6ed9eba1 + x15;
  167. a = rotate30(a);
  168. x0 = rotate1(x13 ^ x8 ^ x2 ^ x0);
  169. f = e ^ a ^ b;
  170. c = rotate5(d) + f + c + 0x6ed9eba1 + x0;
  171. e = rotate30(e);
  172. x1 = rotate1(x14 ^ x9 ^ x3 ^ x1);
  173. f = d ^ e ^ a;
  174. b = rotate5(c) + f + b + 0x6ed9eba1 + x1;
  175. d = rotate30(d);
  176. x2 = rotate1(x15 ^ x10 ^ x4 ^ x2);
  177. f = c ^ d ^ e;
  178. a = rotate5(b) + f + a + 0x6ed9eba1 + x2;
  179. c = rotate30(c);
  180. x3 = rotate1(x0 ^ x11 ^ x5 ^ x3);
  181. f = b ^ c ^ d;
  182. e = rotate5(a) + f + e + 0x6ed9eba1 + x3;
  183. b = rotate30(b);
  184. x4 = rotate1(x1 ^ x12 ^ x6 ^ x4);
  185. f = a ^ b ^ c;
  186. d = rotate5(e) + f + d + 0x6ed9eba1 + x4;
  187. a = rotate30(a);
  188. x5 = rotate1(x2 ^ x13 ^ x7 ^ x5);
  189. f = e ^ a ^ b;
  190. c = rotate5(d) + f + c + 0x6ed9eba1 + x5;
  191. e = rotate30(e);
  192. x6 = rotate1(x3 ^ x14 ^ x8 ^ x6);
  193. f = d ^ e ^ a;
  194. b = rotate5(c) + f + b + 0x6ed9eba1 + x6;
  195. d = rotate30(d);
  196. x7 = rotate1(x4 ^ x15 ^ x9 ^ x7);
  197. f = c ^ d ^ e;
  198. a = rotate5(b) + f + a + 0x6ed9eba1 + x7;
  199. c = rotate30(c);
  200. x8 = rotate1(x5 ^ x0 ^ x10 ^ x8);
  201. f = (b & c) | (b & d) | (c & d);
  202. e = rotate5(a) + f + e + 0x8f1bbcdc + x8;
  203. b = rotate30(b);
  204. x9 = rotate1(x6 ^ x1 ^ x11 ^ x9);
  205. f = (a & b) | (a & c) | (b & c);
  206. d = rotate5(e) + f + d + 0x8f1bbcdc + x9;
  207. a = rotate30(a);
  208. x10 = rotate1(x7 ^ x2 ^ x12 ^ x10);
  209. f = (e & a) | (e & b) | (a & b);
  210. c = rotate5(d) + f + c + 0x8f1bbcdc + x10;
  211. e = rotate30(e);
  212. x11 = rotate1(x8 ^ x3 ^ x13 ^ x11);
  213. f = (d & e) | (d & a) | (e & a);
  214. b = rotate5(c) + f + b + 0x8f1bbcdc + x11;
  215. d = rotate30(d);
  216. x12 = rotate1(x9 ^ x4 ^ x14 ^ x12);
  217. f = (c & d) | (c & e) | (d & e);
  218. a = rotate5(b) + f + a + 0x8f1bbcdc + x12;
  219. c = rotate30(c);
  220. x13 = rotate1(x10 ^ x5 ^ x15 ^ x13);
  221. f = (b & c) | (b & d) | (c & d);
  222. e = rotate5(a) + f + e + 0x8f1bbcdc + x13;
  223. b = rotate30(b);
  224. x14 = rotate1(x11 ^ x6 ^ x0 ^ x14);
  225. f = (a & b) | (a & c) | (b & c);
  226. d = rotate5(e) + f + d + 0x8f1bbcdc + x14;
  227. a = rotate30(a);
  228. x15 = rotate1(x12 ^ x7 ^ x1 ^ x15);
  229. f = (e & a) | (e & b) | (a & b);
  230. c = rotate5(d) + f + c + 0x8f1bbcdc + x15;
  231. e = rotate30(e);
  232. x0 = rotate1(x13 ^ x8 ^ x2 ^ x0);
  233. f = (d & e) | (d & a) | (e & a);
  234. b = rotate5(c) + f + b + 0x8f1bbcdc + x0;
  235. d = rotate30(d);
  236. x1 = rotate1(x14 ^ x9 ^ x3 ^ x1);
  237. f = (c & d) | (c & e) | (d & e);
  238. a = rotate5(b) + f + a + 0x8f1bbcdc + x1;
  239. c = rotate30(c);
  240. x2 = rotate1(x15 ^ x10 ^ x4 ^ x2);
  241. f = (b & c) | (b & d) | (c & d);
  242. e = rotate5(a) + f + e + 0x8f1bbcdc + x2;
  243. b = rotate30(b);
  244. x3 = rotate1(x0 ^ x11 ^ x5 ^ x3);
  245. f = (a & b) | (a & c) | (b & c);
  246. d = rotate5(e) + f + d + 0x8f1bbcdc + x3;
  247. a = rotate30(a);
  248. x4 = rotate1(x1 ^ x12 ^ x6 ^ x4);
  249. f = (e & a) | (e & b) | (a & b);
  250. c = rotate5(d) + f + c + 0x8f1bbcdc + x4;
  251. e = rotate30(e);
  252. x5 = rotate1(x2 ^ x13 ^ x7 ^ x5);
  253. f = (d & e) | (d & a) | (e & a);
  254. b = rotate5(c) + f + b + 0x8f1bbcdc + x5;
  255. d = rotate30(d);
  256. x6 = rotate1(x3 ^ x14 ^ x8 ^ x6);
  257. f = (c & d) | (c & e) | (d & e);
  258. a = rotate5(b) + f + a + 0x8f1bbcdc + x6;
  259. c = rotate30(c);
  260. x7 = rotate1(x4 ^ x15 ^ x9 ^ x7);
  261. f = (b & c) | (b & d) | (c & d);
  262. e = rotate5(a) + f + e + 0x8f1bbcdc + x7;
  263. b = rotate30(b);
  264. x8 = rotate1(x5 ^ x0 ^ x10 ^ x8);
  265. f = (a & b) | (a & c) | (b & c);
  266. d = rotate5(e) + f + d + 0x8f1bbcdc + x8;
  267. a = rotate30(a);
  268. x9 = rotate1(x6 ^ x1 ^ x11 ^ x9);
  269. f = (e & a) | (e & b) | (a & b);
  270. c = rotate5(d) + f + c + 0x8f1bbcdc + x9;
  271. e = rotate30(e);
  272. x10 = rotate1(x7 ^ x2 ^ x12 ^ x10);
  273. f = (d & e) | (d & a) | (e & a);
  274. b = rotate5(c) + f + b + 0x8f1bbcdc + x10;
  275. d = rotate30(d);
  276. x11 = rotate1(x8 ^ x3 ^ x13 ^ x11);
  277. f = (c & d) | (c & e) | (d & e);
  278. a = rotate5(b) + f + a + 0x8f1bbcdc + x11;
  279. c = rotate30(c);
  280. x12 = rotate1(x9 ^ x4 ^ x14 ^ x12);
  281. f = b ^ c ^ d;
  282. e = rotate5(a) + f + e + 0xca62c1d6 + x12;
  283. b = rotate30(b);
  284. x13 = rotate1(x10 ^ x5 ^ x15 ^ x13);
  285. f = a ^ b ^ c;
  286. d = rotate5(e) + f + d + 0xca62c1d6 + x13;
  287. a = rotate30(a);
  288. x14 = rotate1(x11 ^ x6 ^ x0 ^ x14);
  289. f = e ^ a ^ b;
  290. c = rotate5(d) + f + c + 0xca62c1d6 + x14;
  291. e = rotate30(e);
  292. x15 = rotate1(x12 ^ x7 ^ x1 ^ x15);
  293. f = d ^ e ^ a;
  294. b = rotate5(c) + f + b + 0xca62c1d6 + x15;
  295. d = rotate30(d);
  296. x0 = rotate1(x13 ^ x8 ^ x2 ^ x0);
  297. f = c ^ d ^ e;
  298. a = rotate5(b) + f + a + 0xca62c1d6 + x0;
  299. c = rotate30(c);
  300. x1 = rotate1(x14 ^ x9 ^ x3 ^ x1);
  301. f = b ^ c ^ d;
  302. e = rotate5(a) + f + e + 0xca62c1d6 + x1;
  303. b = rotate30(b);
  304. x2 = rotate1(x15 ^ x10 ^ x4 ^ x2);
  305. f = a ^ b ^ c;
  306. d = rotate5(e) + f + d + 0xca62c1d6 + x2;
  307. a = rotate30(a);
  308. x3 = rotate1(x0 ^ x11 ^ x5 ^ x3);
  309. f = e ^ a ^ b;
  310. c = rotate5(d) + f + c + 0xca62c1d6 + x3;
  311. e = rotate30(e);
  312. x4 = rotate1(x1 ^ x12 ^ x6 ^ x4);
  313. f = d ^ e ^ a;
  314. b = rotate5(c) + f + b + 0xca62c1d6 + x4;
  315. d = rotate30(d);
  316. x5 = rotate1(x2 ^ x13 ^ x7 ^ x5);
  317. f = c ^ d ^ e;
  318. a = rotate5(b) + f + a + 0xca62c1d6 + x5;
  319. c = rotate30(c);
  320. x6 = rotate1(x3 ^ x14 ^ x8 ^ x6);
  321. f = b ^ c ^ d;
  322. e = rotate5(a) + f + e + 0xca62c1d6 + x6;
  323. b = rotate30(b);
  324. x7 = rotate1(x4 ^ x15 ^ x9 ^ x7);
  325. f = a ^ b ^ c;
  326. d = rotate5(e) + f + d + 0xca62c1d6 + x7;
  327. a = rotate30(a);
  328. x8 = rotate1(x5 ^ x0 ^ x10 ^ x8);
  329. f = e ^ a ^ b;
  330. c = rotate5(d) + f + c + 0xca62c1d6 + x8;
  331. e = rotate30(e);
  332. x9 = rotate1(x6 ^ x1 ^ x11 ^ x9);
  333. f = d ^ e ^ a;
  334. b = rotate5(c) + f + b + 0xca62c1d6 + x9;
  335. d = rotate30(d);
  336. x10 = rotate1(x7 ^ x2 ^ x12 ^ x10);
  337. f = c ^ d ^ e;
  338. a = rotate5(b) + f + a + 0xca62c1d6 + x10;
  339. c = rotate30(c);
  340. x11 = rotate1(x8 ^ x3 ^ x13 ^ x11);
  341. f = b ^ c ^ d;
  342. e = rotate5(a) + f + e + 0xca62c1d6 + x11;
  343. b = rotate30(b);
  344. x12 = rotate1(x9 ^ x4 ^ x14 ^ x12);
  345. f = a ^ b ^ c;
  346. d = rotate5(e) + f + d + 0xca62c1d6 + x12;
  347. a = rotate30(a);
  348. x13 = rotate1(x10 ^ x5 ^ x15 ^ x13);
  349. f = e ^ a ^ b;
  350. c = rotate5(d) + f + c + 0xca62c1d6 + x13;
  351. e = rotate30(e);
  352. x14 = rotate1(x11 ^ x6 ^ x0 ^ x14);
  353. f = d ^ e ^ a;
  354. b = rotate5(c) + f + b + 0xca62c1d6 + x14;
  355. d = rotate30(d);
  356. x15 = rotate1(x12 ^ x7 ^ x1 ^ x15);
  357. f = c ^ d ^ e;
  358. a = rotate5(b) + f + a + 0xca62c1d6 + x15;
  359. c = rotate30(c);
  360.  
  361. a = a + H[0];
  362. b = b + H[1];
  363. c = c + H[2];
  364. d = d + H[3];
  365. e = e + H[4];
  366. H[0] = a;
  367. H[1] = b;
  368. H[2] = c;
  369. H[3] = d;
  370. H[4] = e;
  371. }
  372. #endif
  373.  
  374. #ifdef SAFESHA
  375. uint rotateLeft(uint32 x, int32 n)
  376. {
  377. return (x << n) | (x >> (32-n));
  378. }
  379.  
  380. // block size: 512b = 64B = 16W
  381. // W (80W long) is the 16W of work + scratch space (prepadded)
  382. // H (5W long) is the current hash state
  383. void sha1_block(uint32 *W, uint32 *H)
  384. {
  385. uint32 A,B,C,D,E,K0,K1,K2,K3,temp;
  386. int i;
  387.  
  388. K0 = 0x5A827999;
  389. K1 = 0x6ED9EBA1;
  390. K2 = 0x8F1BBCDC;
  391. K3 = 0xCA62C1D6;
  392.  
  393. A = H[0];
  394. B = H[1];
  395. C = H[2];
  396. D = H[3];
  397. E = H[4];
  398.  
  399. for(i = 16; i < 80; i++)
  400. {
  401. W[i] = rotateLeft(W[i-3] ^ W[i-8] ^ W[i-14] ^ W[i-16], 1);
  402. }
  403.  
  404. for(i = 0; i < 20; i++)
  405. {
  406. temp = rotateLeft(A,5) + ((B & C) | ((~ B) & D)) + E + W[i] + K0;
  407. E = D;
  408. D = C;
  409. C = rotateLeft(B, 30);
  410. B = A;
  411. A = temp;
  412. }
  413.  
  414. for(i = 20; i < 40; i++)
  415. {
  416. temp = rotateLeft(A, 5) + (B ^ C ^ D) + E + W[i] + K1;
  417. E = D;
  418. D = C;
  419. C = rotateLeft(B, 30);
  420. B = A;
  421. A = temp;
  422. }
  423.  
  424. for(i = 40; i < 60; i++)
  425. {
  426. temp = rotateLeft(A, 5) + ((B & C) | (B & D) | (C & D)) + E + W[i] + K2;
  427. E = D;
  428. D = C;
  429. C = rotateLeft(B, 30);
  430. B = A;
  431. A = temp;
  432. }
  433.  
  434. for(i = 60; i < 80; i++)
  435. {
  436. temp = rotateLeft(A, 5) + (B ^ C ^ D) + E + W[i] + K3;
  437. E = D;
  438. D = C;
  439. C = rotateLeft(B, 30);
  440. B = A;
  441. A = temp;
  442. }
  443.  
  444. H[0] = (H[0] + A);
  445. H[1] = (H[1] + B);
  446. H[2] = (H[2] + C);
  447. H[3] = (H[3] + D);
  448. H[4] = (H[4] + E);
  449. }
  450. #endif
  451.  
  452. // Must set the right define for W packing code
  453. // Only works with certain sized keys (1024 and 2048/4096 tested) and 4 byte exponents.
  454. // Base_exp must be >= 0x01000001 and global work size must be <= (0x7FFFFFFF-base_exp)/2
  455. __kernel void optimized(__constant uint32* LastWs, __constant uint32* Midstates, __global uint32* Results, uint32 BaseExp,
  456. uint8 LenStart, __constant int32* ExpIndexes, // Not used - for compat.
  457. __constant uint32* BitmaskArray, __constant uint16* HashTable, __constant uint32* DataArray)
  458. {
  459. uint64 exp;
  460. uint32 fnv,fnv10;
  461.  
  462. uint16 dataaddr;
  463.  
  464. int i;
  465.  
  466. uint32 W[16];
  467. uint32 H[5];
  468.  
  469. /*GENERATED__ARRAYS*/
  470.  
  471. exp = get_global_id(0) * 2 + BaseExp;
  472.  
  473. // Load Ws and Midstates into private variables
  474. for(i=0; i<16; i++) W[i] = LastWs[i];
  475. for(i=0; i<5; i++) H[i] = Midstates[i];
  476.  
  477. // Load the exponent into the W
  478. GENERATED__EXP_LOADING_CODE
  479.  
  480. // Take the last part of the hash
  481. sha1_block(W,H);
  482.  
  483. // Get and check the FNV hash for each bitmask
  484. // Uses code generated on the C# side
  485. GENERATED__CHECKING_CODE
  486. }
  487.  
  488. // Works with any exp index and starting length
  489. // Still requires that all of the exponent lie in the last SHA1 block.
  490. __kernel void normal(__constant uint32* LastWs, __constant uint32* Midstates, __global uint32* Results, uint32 BaseExp,
  491. uint8 LenStart, __constant int32* ExpIndexes,
  492. __constant uint32* BitmaskArray, __constant uint16* HashTable, __constant uint32* DataArray)
  493. {
  494. }
  495.  
  496. // Test the SHA hash code
  497. __kernel void shaTest(__global uint32* success)
  498. {
  499. int i;
  500. uint32 W[80];
  501. uint32 H[5];
  502.  
  503. // Zero out W
  504. for(i=0;i<80;i++) {
  505. W[i] = 0;
  506. }
  507.  
  508. // Init the SHA state
  509. H[0] = 0x67452301;
  510. H[1] = 0xEFCDAB89;
  511. H[2] = 0x98BADCFE;
  512. H[3] = 0x10325476;
  513. H[4] = 0xC3D2E1F0;
  514.  
  515. // Load our (pre-padded) test block: "Hello world!"
  516. W[0] = 0x48656c6cu; // Hell
  517. W[1] = 0x6f20776fu; // o wo
  518. W[2] = 0x726c6421u; // rld!
  519. W[3] = 0x80000000u; // (bit 1)
  520. W[15] = 0x00000060u; // m-length in bits (not including bit '1')
  521.  
  522. // Take the SHA
  523. sha1_block(W, H);
  524.  
  525. // Check for success
  526. *success = 0;
  527. if (H[0] == 0xd3486ae9 && H[1] == 0x136e7856 && H[2] == 0xbc422123 && H[3] == 0x85ea7970 && H[4] == 0x94475802) {
  528. *success = 1;
  529. }
  530.  
  531. success[0] = H[0];
  532. success[1] = H[1];
  533. success[2] = H[2];
  534. success[3] = H[3];
  535. success[4] = H[4];
  536. }
Advertisement
Add Comment
Please, Sign In to add comment
Advertisement