API tools faq
paste
Guest User

Untitled

a guest
Jun 28th, 2010
409
0
Never
Add comment
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
text 19.03 KB | None | 0 0
raw download clone embed print report
  1. #define PRODUCT_NAME "CUDA SHA-1 Tripper 0.2.1"
  2.  
  3. #include <stdio.h>
  4. #include <stdlib.h>
  5. #include <string.h>
  6.  
  7. #include <windows.h>
  8.  
  9. #include <search.h>
  10.  
  11.  
  12. // default files
  13. #define TARGET_FILE "target.txt"
  14. #define TRIP_FILE "trip.txt"
  15.  
  16. char target_file[256], trip_file[256];
  17.  
  18. typedef unsigned long DWORD; // 32bit
  19. typedef unsigned short WORD; // 16bit
  20. typedef unsigned char BYTE; // 8bit
  21.  
  22. struct st_sha1trip{
  23. unsigned int count;
  24. BYTE key[10];
  25. BYTE trip[10];
  26. };
  27.  
  28. // 128 threads per block, each block is divided into 4 warps
  29. #define BLOCK_SIZE_X 64
  30. #define BLOCK_SIZE_Y 2
  31.  
  32. // max is StreamingMultiprocessor * 8
  33. // StreamingMultiprocessor consists of 8 StreamingProcessor
  34. unsigned int blocks;
  35.  
  36. // size of array R (BLOCK_SIZE_X * BLOCK_SIZE_Y * blocks * sizeof(st_sha1trip))
  37. unsigned long R_size;
  38.  
  39. // SHA-1 input message block
  40. __device__ __constant__ BYTE Md[64];
  41.  
  42. #define TARGETS_MAX 5000
  43.  
  44. char target[TARGETS_MAX][14];
  45. unsigned int target_num;
  46. __device__ __constant__ DWORD target_dw[TARGETS_MAX];
  47. __device__ __constant__ unsigned int target_dw_num;
  48.  
  49.  
  50. const char b64t[] = {
  51. 'A', 'B', 'C', 'D', 'E', 'F', 'G', 'H', 'I', 'J', 'K', 'L', 'M', 'N', 'O', 'P',
  52. 'Q', 'R', 'S', 'T', 'U', 'V', 'W', 'X', 'Y', 'Z', 'a', 'b', 'c', 'd', 'e', 'f',
  53. 'g', 'h', 'i', 'j', 'k', 'l', 'm', 'n', 'o', 'p', 'q', 'r', 's', 't', 'u', 'v',
  54. 'w', 'x', 'y', 'z', '0', '1', '2', '3', '4', '5', '6', '7', '8', '9', '.', '/'
  55. };
  56.  
  57.  
  58. st_sha1trip *R, *Rd;
  59. FILE *fp_trip;
  60.  
  61.  
  62. // constants and initial values defined in SHA-1
  63. #define K0 0x5A827999
  64. #define K1 0x6ED9EBA1
  65. #define K2 0x8F1BBCDC
  66. #define K3 0xCA62C1D6
  67.  
  68. #define H0 0x67452301
  69. #define H1 0xEFCDAB89
  70. #define H2 0x98BADCFE
  71. #define H3 0x10325476
  72. #define H4 0xC3D2E1F0
  73.  
  74. ////////////////////////////////////////////////////////////////////////////////
  75.  
  76. #define ROL32(_val32, _nBits) (((_val32)<<(_nBits))|((_val32)>>(32-(_nBits))))
  77.  
  78. #define _R0(v,w,x,y,z,i) { z += ((w&(x^y))^y) + i + K0 + ROL32(v,5); w=ROL32(w,30); }
  79. #define _R00(v,w,x,y,z) { z += ((w&(x^y))^y) + K0 + ROL32(v,5); w=ROL32(w,30); }
  80. #define _R2(v,w,x,y,z,i) { z += (w^x^y) + i + K1 + ROL32(v,5); w=ROL32(w,30); }
  81. #define _R4(v,w,x,y,z,i) { z += (((w|x)&y)|(w&x)) + i + K2 + ROL32(v,5); w=ROL32(w,30); }
  82. #define _R6(v,w,x,y,z,i) { z += (w^x^y) + i + K3 + ROL32(v,5); w=ROL32(w,30); }
  83.  
  84. __device__ __constant__ char b64t_d[] = {
  85. 'A', 'B', 'C', 'D', 'E', 'F', 'G', 'H', 'I', 'J', 'K', 'L', 'M', 'N', 'O', 'P',
  86. 'Q', 'R', 'S', 'T', 'U', 'V', 'W', 'X', 'Y', 'Z', 'a', 'b', 'c', 'd', 'e', 'f',
  87. 'g', 'h', 'i', 'j', 'k', 'l', 'm', 'n', 'o', 'p', 'q', 'r', 's', 't', 'u', 'v',
  88. 'w', 'x', 'y', 'z', '0', '1', '2', '3', '4', '5', '6', '7', '8', '9', '.', '/'
  89. };
  90.  
  91. // sha1trip_search
  92. // R - Result array (BLOCK_SIZE_X * BLOCK_SIZE_Y * blocks array of st_sha1trip)
  93. __global__ void sha1trip_search(st_sha1trip *R)
  94. {
  95. DWORD W0, W1, W2, W3, W4, W5, W6, W7, W8, W9, W10, W11, W12, W13, W14, W15;
  96. DWORD a, b, c, d, e;
  97.  
  98.  
  99. // 32 * 64 = 2048
  100. for (unsigned int Npass = 0; Npass < 2048; Npass++){
  101. __syncthreads();
  102.  
  103. a = H0;
  104. b = H1;
  105. c = H2;
  106. d = H3;
  107. e = H4;
  108.  
  109. // copy message block
  110. W0 = Md[0] << 24 | Md[1] << 16 | Md[2] << 8 | Md[3];
  111.  
  112. W1 = Md[4] << 24 | Md[5] << 16 | Md[6] << 8;
  113. W1 |= b64t_d[(Md[7] + (Npass >> 5)) & 63];
  114.  
  115. W2 = b64t_d[(Md[8] + (blockIdx.x >> 6)) & 63] << 24;
  116. W2 |= b64t_d[(Md[9] + blockIdx.x) & 63] << 16;
  117. W2 |= b64t_d[(Npass & 31) * BLOCK_SIZE_Y + threadIdx.y] << 8;
  118. W2 |= b64t_d[threadIdx.x];
  119.  
  120. W3 = 0x80000000; // padding
  121.  
  122. /*
  123. W4 = 0;
  124. W5 = 0;
  125. W6 = 0;
  126. W7 = 0;
  127. W8 = 0;
  128. W9 = 0;
  129. W10 = 0;
  130. W11 = 0;
  131. W12 = 0;
  132. W13 = 0;
  133. W14 = 0;
  134. */
  135.  
  136. W15 = 96; // bits of Message Block (12 bytes * 8 bits)
  137.  
  138. // round 0 to 15
  139. _R0(a, b, c, d, e, W0);
  140. _R0(e, a, b, c, d, W1);
  141. _R0(d, e, a, b, c, W2);
  142. _R0(c, d, e, a, b, W3);
  143. _R00(b, c, d, e, a); // W4 == 0
  144. _R00(a, b, c, d, e); // W5 == 0
  145. _R00(e, a, b, c, d); // W6 == 0
  146. _R00(d, e, a, b, c); // W7 == 0
  147. _R00(c, d, e, a, b); // W8 == 0
  148. _R00(b, c, d, e, a); // W9 == 0
  149. _R00(a, b, c, d, e); // W10 == 0
  150. _R00(e, a, b, c, d); // W11 == 0
  151. _R00(d, e, a, b, c); // W12 == 0
  152. _R00(c, d, e, a, b); // W13 == 0
  153. _R00(b, c, d, e, a); // W14 == 0
  154. _R0(a, b, c, d, e, W15);
  155.  
  156. // W[t] = ROL32(W[t-3] ^ W[t-8] ^ W[t-14] ^ W[t-16], 1);
  157. W0 = ROL32(W2 ^ W0, 1); // (t, W[t-3], W[t-8], W[t-14], W[t-16]) = (16, W13==0, W8==0, W2, W0)
  158. W1 = ROL32(W3 ^ W1, 1); // 17, W14==0, W9==0, W3, W1
  159. W2 = ROL32(W15 ^ W2, 1); // 18, W15, W10==0, W4==0, W2
  160. W3 = ROL32(W0 ^ W3, 1); // 19, W0, W11==0, W5==0, W3
  161. W4 = ROL32(W1, 1); // 20, W1, W12==0, W6==0, W4==0
  162. W5 = ROL32(W2, 1); // 21, W2, W13==0, W7==0, W5==0
  163. W6 = ROL32(W3, 1); // 22, W3, W14==0, W8==0, W6==0
  164. W7 = ROL32(W4 ^ W15, 1); // 23, W4, W15, W9==0, W7==0
  165. W8 = ROL32(W5 ^ W0, 1); // 24, W5, W0, W10==0, W8==0
  166. W9 = ROL32(W6 ^ W1, 1); // 25, W6, W1, W11==0, W9==0
  167. W10 = ROL32(W7 ^ W2, 1); // 26, W7, W2, W12==0, W10==0
  168. W11 = ROL32(W8 ^ W3, 1); // 27, W8, W3, W13==0, W11==0
  169. W12 = ROL32(W9 ^ W4, 1); // 28, W9, W4, W14==0, W12==0
  170. W13 = ROL32(W10 ^ W5 ^ W15, 1); // 29, W10, W5, W15, W13==0
  171. W14 = ROL32(W11 ^ W6 ^ W0, 1); // 30, W11, W6, W0, W14==0
  172. W15 = ROL32(W12 ^ W7 ^ W1 ^ W15, 1); // 31, W12, W7, W1, W15
  173.  
  174. // round 16 to 19
  175. _R0(e, a, b, c, d, W0);
  176. _R0(d, e, a, b, c, W1);
  177. _R0(c, d, e, a, b, W2);
  178. _R0(b, c, d, e, a, W3);
  179.  
  180. // round 20 to 31
  181. _R2(a, b, c, d, e, W4);
  182. _R2(e, a, b, c, d, W5);
  183. _R2(d, e, a, b, c, W6);
  184. _R2(c, d, e, a, b, W7);
  185. _R2(b, c, d, e, a, W8);
  186. _R2(a, b, c, d, e, W9);
  187. _R2(e, a, b, c, d, W10);
  188. _R2(d, e, a, b, c, W11);
  189. _R2(c, d, e, a, b, W12);
  190. _R2(b, c, d, e, a, W13);
  191. _R2(a, b, c, d, e, W14);
  192. _R2(e, a, b, c, d, W15);
  193.  
  194. // W[t] = ROL32(W[t-3] ^ W[t-8] ^ W[t-14] ^ W[t-16], 1);
  195. W0 = ROL32(W13 ^ W8 ^ W2 ^ W0, 1); // t==32
  196. W1 = ROL32(W14 ^ W9 ^ W3 ^ W1, 1);
  197. W2 = ROL32(W15 ^ W10 ^ W4 ^ W2, 1);
  198. W3 = ROL32(W0 ^ W11 ^ W5 ^ W3, 1);
  199. W4 = ROL32(W1 ^ W12 ^ W6 ^ W4, 1);
  200. W5 = ROL32(W2 ^ W13 ^ W7 ^ W5, 1);
  201. W6 = ROL32(W3 ^ W14 ^ W8 ^ W6, 1);
  202. W7 = ROL32(W4 ^ W15 ^ W9 ^ W7, 1);
  203. W8 = ROL32(W5 ^ W0 ^ W10 ^ W8, 1);
  204. W9 = ROL32(W6 ^ W1 ^ W11 ^ W9, 1);
  205. W10 = ROL32(W7 ^ W2 ^ W12 ^ W10, 1);
  206. W11 = ROL32(W8 ^ W3 ^ W13 ^ W11, 1);
  207. W12 = ROL32(W9 ^ W4 ^ W14 ^ W12, 1);
  208. W13 = ROL32(W10 ^ W5 ^ W15 ^ W13, 1);
  209. W14 = ROL32(W11 ^ W6 ^ W0 ^ W14, 1);
  210. W15 = ROL32(W12 ^ W7 ^ W1 ^ W15, 1); // t==47
  211.  
  212. // round 32 to 39
  213. _R2(d, e, a, b, c, W0);
  214. _R2(c, d, e, a, b, W1);
  215. _R2(b, c, d, e, a, W2);
  216. _R2(a, b, c, d, e, W3);
  217. _R2(e, a, b, c, d, W4);
  218. _R2(d, e, a, b, c, W5);
  219. _R2(c, d, e, a, b, W6);
  220. _R2(b, c, d, e, a, W7);
  221.  
  222. // round 40 to 47
  223. _R4(a, b, c, d, e, W8);
  224. _R4(e, a, b, c, d, W9);
  225. _R4(d, e, a, b, c, W10);
  226. _R4(c, d, e, a, b, W11);
  227. _R4(b, c, d, e, a, W12);
  228. _R4(a, b, c, d, e, W13);
  229. _R4(e, a, b, c, d, W14);
  230. _R4(d, e, a, b, c, W15);
  231.  
  232. // W[t] = ROL32(W[t-3] ^ W[t-8] ^ W[t-14] ^ W[t-16], 1);
  233. W0 = ROL32(W13 ^ W8 ^ W2 ^ W0, 1); // t==48
  234. W1 = ROL32(W14 ^ W9 ^ W3 ^ W1, 1);
  235. W2 = ROL32(W15 ^ W10 ^ W4 ^ W2, 1);
  236. W3 = ROL32(W0 ^ W11 ^ W5 ^ W3, 1);
  237. W4 = ROL32(W1 ^ W12 ^ W6 ^ W4, 1);
  238. W5 = ROL32(W2 ^ W13 ^ W7 ^ W5, 1);
  239. W6 = ROL32(W3 ^ W14 ^ W8 ^ W6, 1);
  240. W7 = ROL32(W4 ^ W15 ^ W9 ^ W7, 1);
  241. W8 = ROL32(W5 ^ W0 ^ W10 ^ W8, 1);
  242. W9 = ROL32(W6 ^ W1 ^ W11 ^ W9, 1);
  243. W10 = ROL32(W7 ^ W2 ^ W12 ^ W10, 1);
  244. W11 = ROL32(W8 ^ W3 ^ W13 ^ W11, 1);
  245. W12 = ROL32(W9 ^ W4 ^ W14 ^ W12, 1);
  246. W13 = ROL32(W10 ^ W5 ^ W15 ^ W13, 1);
  247. W14 = ROL32(W11 ^ W6 ^ W0 ^ W14, 1);
  248. W15 = ROL32(W12 ^ W7 ^ W1 ^ W15, 1); // t==63
  249.  
  250. // round 48 to 59
  251. _R4(c, d, e, a, b, W0);
  252. _R4(b, c, d, e, a, W1);
  253. _R4(a, b, c, d, e, W2);
  254. _R4(e, a, b, c, d, W3);
  255. _R4(d, e, a, b, c, W4);
  256. _R4(c, d, e, a, b, W5);
  257. _R4(b, c, d, e, a, W6);
  258. _R4(a, b, c, d, e, W7);
  259. _R4(e, a, b, c, d, W8);
  260. _R4(d, e, a, b, c, W9);
  261. _R4(c, d, e, a, b, W10);
  262. _R4(b, c, d, e, a, W11);
  263.  
  264. // round 60 to 63
  265. _R6(a, b, c, d, e, W12);
  266. _R6(e, a, b, c, d, W13);
  267. _R6(d, e, a, b, c, W14);
  268. _R6(c, d, e, a, b, W15);
  269.  
  270. // W[t] = ROL32(W[t-3] ^ W[t-8] ^ W[t-14] ^ W[t-16], 1);
  271. W0 = ROL32(W13 ^ W8 ^ W2 ^ W0, 1); // t==64
  272. W1 = ROL32(W14 ^ W9 ^ W3 ^ W1, 1);
  273. W2 = ROL32(W15 ^ W10 ^ W4 ^ W2, 1);
  274. W3 = ROL32(W0 ^ W11 ^ W5 ^ W3, 1);
  275. W4 = ROL32(W1 ^ W12 ^ W6 ^ W4, 1);
  276. W5 = ROL32(W2 ^ W13 ^ W7 ^ W5, 1);
  277. W6 = ROL32(W3 ^ W14 ^ W8 ^ W6, 1);
  278. W7 = ROL32(W4 ^ W15 ^ W9 ^ W7, 1);
  279. W8 = ROL32(W5 ^ W0 ^ W10 ^ W8, 1);
  280. W9 = ROL32(W6 ^ W1 ^ W11 ^ W9, 1);
  281. W10 = ROL32(W7 ^ W2 ^ W12 ^ W10, 1);
  282. W11 = ROL32(W8 ^ W3 ^ W13 ^ W11, 1);
  283. W12 = ROL32(W9 ^ W4 ^ W14 ^ W12, 1);
  284. W13 = ROL32(W10 ^ W5 ^ W15 ^ W13, 1);
  285. W14 = ROL32(W11 ^ W6 ^ W0 ^ W14, 1);
  286. W15 = ROL32(W12 ^ W7 ^ W1 ^ W15, 1); // t==79
  287.  
  288. // round 64 to 79
  289. _R6(b, c, d, e, a, W0);
  290. _R6(a, b, c, d, e, W1);
  291. _R6(e, a, b, c, d, W2);
  292. _R6(d, e, a, b, c, W3);
  293. _R6(c, d, e, a, b, W4);
  294. _R6(b, c, d, e, a, W5);
  295. _R6(a, b, c, d, e, W6);
  296. _R6(e, a, b, c, d, W7);
  297. _R6(d, e, a, b, c, W8);
  298. _R6(c, d, e, a, b, W9);
  299. _R6(b, c, d, e, a, W10);
  300. _R6(a, b, c, d, e, W11);
  301. _R6(e, a, b, c, d, W12);
  302. _R6(d, e, a, b, c, W13);
  303. _R6(c, d, e, a, b, W14);
  304. _R6(b, c, d, e, a, W15);
  305.  
  306. a += H0;
  307.  
  308. // need 30 bits for compare
  309. DWORD tmp_dw = a & 0xfffffffc;
  310.  
  311. for (unsigned int i = 0; i < target_dw_num; i++){
  312. if (target_dw[i] == tmp_dw){
  313. b += H1;
  314. c += H2;
  315.  
  316. st_sha1trip *out = &R[(blockIdx.x * BLOCK_SIZE_Y + threadIdx.y) * BLOCK_SIZE_X + threadIdx.x];
  317.  
  318. out->count = Npass + 1;
  319.  
  320. /*
  321. out->key[0] = Md[0];
  322. out->key[1] = Md[1];
  323. out->key[2] = Md[2];
  324. out->key[3] = Md[3];
  325. out->key[4] = Md[4];
  326. out->key[5] = Md[5];
  327. out->key[6] = Md[6];
  328. */
  329.  
  330. out->key[7] = b64t_d[(Md[7] + (Npass >> 5)) & 63];
  331. out->key[8] = b64t_d[(Md[8] + (blockIdx.x >> 6)) & 63];
  332. out->key[9] = b64t_d[(Md[9] + blockIdx.x) & 63];
  333. out->key[10] = b64t_d[(Npass & 31) * BLOCK_SIZE_Y + threadIdx.y];
  334. out->key[11] = b64t_d[threadIdx.x];
  335.  
  336. out->trip[0] = b64t_d[a >> 26];
  337. out->trip[1] = b64t_d[(a >> 20) & 63];
  338. out->trip[2] = b64t_d[(a >> 14) & 63];
  339. out->trip[3] = b64t_d[(a >> 8) & 63];
  340. out->trip[4] = b64t_d[(a >> 2) & 63];
  341. out->trip[5] = b64t_d[(a << 4 | b >> 28) & 63];
  342. out->trip[6] = b64t_d[(b >> 22) & 63];
  343. out->trip[7] = b64t_d[(b >> 16) & 63];
  344. out->trip[8] = b64t_d[(b >> 10) & 63];
  345. out->trip[9] = b64t_d[(b >> 4) & 63];
  346. out->trip[10] = b64t_d[(b << 2 | c >> 30) & 63];
  347. out->trip[11] = b64t_d[(c >> 24) & 63];
  348.  
  349. return;
  350. }
  351. }
  352. }
  353.  
  354. R[(blockIdx.x * BLOCK_SIZE_Y + threadIdx.y) * BLOCK_SIZE_X + threadIdx.x].count = 2048;
  355.  
  356. return;
  357. }
  358.  
  359.  
  360. ////////////////////////////////////////////////////////////////////////////////
  361.  
  362. // for qsort()
  363. int compare_dw(const void *_a, const void *_b)
  364. {
  365. DWORD a = *(DWORD *)_a;
  366. DWORD b = *(DWORD *)_b;
  367.  
  368. if (a < b){ return -1;}
  369. if (a > b){ return 1;}
  370.  
  371. return 0;
  372. }
  373.  
  374.  
  375. // for qsort()
  376. int compare_str(const void *a, const void *b)
  377. {
  378. return strcmp((char *)a, (char *)b);
  379. }
  380.  
  381.  
  382. size_t uniq(void *array, size_t num, size_t size, int (*comp)(const void *, const void *))
  383. {
  384. char *start = (char *)array;
  385. char *stop = (char *)array + (size * (num - 1));
  386.  
  387. char *dst = start;
  388. char *p = start;
  389.  
  390. while (p <= stop){
  391. char *q = p;
  392. while ((q < stop) && comp(q, q + size)){
  393. q += size;
  394. }
  395.  
  396. size_t elements = ((q - p) / size) + 1;
  397. memmove(dst, p, size * elements);
  398. dst += size * elements;
  399.  
  400. p = q + size;
  401. while ((p <= stop) && comp(q, p) == 0){
  402. p += size;
  403. }
  404. }
  405.  
  406. return (dst - start) / size;
  407. }
  408.  
  409.  
  410. int set_targets()
  411. {
  412. FILE *fp;
  413. char buf[256];
  414. DWORD target_dw_h[TARGETS_MAX];
  415. unsigned int target_dw_num_h;
  416.  
  417. memset(target, 0, sizeof(target));
  418.  
  419. if ((fp = fopen(target_file, "r")) == NULL){
  420. printf("Cannot open target file \"%s\"\n", target_file);
  421. exit(-1);
  422. }
  423.  
  424. for (target_num = 0; target_num < TARGETS_MAX; ){
  425. fgets(buf, sizeof(buf), fp);
  426. if (feof(fp)){ break;}
  427.  
  428. if (ferror(fp)){
  429. printf("Read error target file \"%s\"\n", target_file);
  430. fclose(fp);
  431. exit(-1);
  432. }
  433.  
  434. if (buf[0] == '#' || buf[0] == '\n'){ continue;}
  435.  
  436. for (int i = 0; i < 12; i++){
  437. if (buf[i] == '\n' || buf[i] == ' ' || buf[i] == '\t'){ break;}
  438.  
  439. target[target_num][i] = buf[i];
  440. }
  441.  
  442. if (strlen(target[target_num]) < 5){ continue;}
  443.  
  444. // Base64 decode
  445. for (int i = 0; i < 5; i++){
  446. if (buf[i] >= 'A' && buf[i] <= 'Z'){
  447. buf[i] = buf[i] - 'A';
  448. }
  449. else if (buf[i] >= 'a' && buf[i] <= 'z'){
  450. buf[i] = buf[i] - 'a' + 26;
  451. }
  452. else if (buf[i] >= '0' && buf[i] <= '9'){
  453. buf[i] = buf[i] - '0' + 52;
  454. }
  455. else if (buf[i] == '.'){
  456. buf[i] = 62;
  457. }
  458. else if (buf[i] == '/'){
  459. buf[i] = 63;
  460. }
  461. }
  462.  
  463. target_dw_h[target_num] = ((buf[0]<<2 | buf[1]>>4) & 0xff) << 24;
  464. target_dw_h[target_num] |= ((buf[1]<<4 | buf[2]>>2) & 0xff) << 16;
  465. target_dw_h[target_num] |= ((buf[2]<<6 | buf[3]) & 0xff) << 8;
  466. target_dw_h[target_num] |= (buf[4]<<2) & 0xff;
  467.  
  468. printf("target %2d: %s\t\t%08x\n", target_num, target[target_num], target_dw_h[target_num]);
  469.  
  470. target_num++;
  471. }
  472.  
  473. fclose(fp);
  474.  
  475. if (target_num == 0){
  476. printf("There is no target string\n");
  477. exit(-1);
  478. }
  479.  
  480. qsort(target, target_num, sizeof(target[0]), compare_str);
  481.  
  482. target_dw_num_h = target_num;
  483. qsort(target_dw_h, target_dw_num_h, sizeof(DWORD), compare_dw);
  484. target_dw_num_h = uniq(target_dw_h, target_dw_num_h, sizeof(DWORD), compare_dw);
  485.  
  486. printf("\n%d targets found, target_dw_num is %d\n\n", target_num, target_dw_num_h);
  487.  
  488. cudaMemcpyToSymbol(target_dw, target_dw_h, sizeof(DWORD) * target_dw_num_h);
  489. cudaMemcpyToSymbol("target_dw_num", &target_dw_num_h, sizeof(int));
  490.  
  491. return 0;
  492. }
  493.  
  494.  
  495. // pseudo-random number generator
  496. unsigned long xor128()
  497. {
  498. // static unsigned long x=123456789, y=362436069, z=521288629, w=88675123;
  499. static unsigned long x=123456789, y=362436069, z=521288629, w=GetTickCount();
  500. unsigned long t;
  501. t=(x^(x<<11)); x=y; y=z; z=w;
  502. return( w=(w^(w>>19))^(t^(t>>8)) );
  503. }
  504.  
  505.  
  506. // ConsoleCtrlHandler
  507. BOOL handler(DWORD dwCtrlType)
  508. {
  509. fclose(fp_trip);
  510.  
  511. cudaFree(Rd);
  512. free(R);
  513.  
  514. return FALSE;
  515. }
  516.  
  517.  
  518. void usage(char *argv0)
  519. {
  520. printf("Usage:\n");
  521. printf(" %s [-d device] [-x blocks_per_sm] [-o trip_file] [-f target_file]\n\n", argv0);
  522.  
  523. exit(-1);
  524. }
  525.  
  526.  
  527. int main(int argc, char **argv)
  528. {
  529. int deviceCount;
  530. BYTE M[64] = {0};
  531. char key[16] = {0}, trip[16] = {0};
  532. DWORD t1, t2;
  533. double delta_sec;
  534. unsigned long count;
  535. int device = 0, blocks_per_sm = 2;
  536. cudaDeviceProp deviceProp;
  537.  
  538.  
  539. printf("%s\n\n", PRODUCT_NAME);
  540.  
  541. cudaGetDeviceCount(&deviceCount);
  542. if (deviceCount == 0){
  543. printf("There is no device supporting CUDA\n\n");
  544. exit(-1);
  545. }
  546.  
  547. for (int i = 0; i < deviceCount; i++){
  548. cudaGetDeviceProperties(&deviceProp, i);
  549.  
  550. printf("Device %d: \"%s\"\n", i, deviceProp.name);
  551. printf(" Revision number: %d.%d\n",
  552. deviceProp.major, deviceProp.minor);
  553. printf(" Total amount of global memory: %u Mbytes\n",
  554. deviceProp.totalGlobalMem / (1024 * 1024));
  555. printf(" Number of multiprocessors: %d\n",
  556. deviceProp.multiProcessorCount);
  557. printf(" Number of cores: %d\n",
  558. deviceProp.multiProcessorCount * 8);
  559. printf(" Clock rate: %.2f GHz\n\n",
  560. deviceProp.clockRate * 1e-6f);
  561.  
  562. }
  563.  
  564.  
  565. strncpy(trip_file, TRIP_FILE, sizeof(trip_file));
  566. strncpy(target_file, TARGET_FILE, sizeof(target_file));
  567.  
  568. if (argc > 1){
  569. for (int i = 1; i < argc; i++){
  570. if (argv[i][0] == '-' && argv[i][1] == 'd'){
  571. if (++i == argc){ usage(argv[0]);}
  572. device = atoi(argv[i]);
  573. }
  574. else if (argv[i][0] == '-' && argv[i][1] == 'x'){
  575. if (++i == argc){ usage(argv[0]);}
  576. blocks_per_sm = atoi(argv[i]);
  577. }
  578. else if (argv[i][0] == '-' && argv[i][1] == 'o'){
  579. if (++i == argc){ usage(argv[0]);}
  580. strncpy(trip_file, argv[i], sizeof(trip_file));
  581. trip_file[sizeof(trip_file) - 1] = 0;
  582. }
  583. else if (argv[i][0] == '-' && argv[i][1] == 'f'){
  584. if (++i == argc){ usage(argv[0]);}
  585. strncpy(target_file, argv[i], sizeof(target_file));
  586. target_file[sizeof(target_file) - 1] = 0;
  587. }
  588. else {
  589. usage(argv[0]);
  590. }
  591. }
  592. }
  593.  
  594.  
  595. if (device >= deviceCount){
  596. printf("Only %d devices supporting CUDA\n\n", deviceCount);
  597. exit(-1);
  598. }
  599.  
  600. if (blocks_per_sm < 1 || blocks_per_sm > 8){
  601. printf("Blocks_per_SM must be 1 to 8\n\n");
  602. exit(-1);
  603. }
  604.  
  605. cudaGetDeviceProperties(&deviceProp, device);
  606. blocks = deviceProp.multiProcessorCount * blocks_per_sm;
  607.  
  608. cudaSetDevice(device);
  609. printf("Use device %d, grid is %d blocks\n\n", device, blocks);
  610.  
  611. R_size = BLOCK_SIZE_X * BLOCK_SIZE_Y * blocks * sizeof(st_sha1trip);
  612.  
  613. set_targets();
  614.  
  615. if ((fp_trip = fopen(trip_file, "a")) == NULL){
  616. printf("Cannot open trip file \"%s\"\n", trip_file);
  617. exit(-1);
  618. }
  619.  
  620. R = (st_sha1trip *) malloc(R_size);
  621. if (! R){
  622. fprintf(stderr, "malloc() error!\n");
  623. exit(-1);
  624. }
  625. memset(R, 0, R_size);
  626.  
  627. cudaMalloc((void **)&Rd, R_size);
  628.  
  629. M[12] = 0x80; // start of padding
  630. M[63] = 96; // bits of Message Block (12 bytes * 8 bits)
  631.  
  632. SetPriorityClass(GetCurrentProcess(), IDLE_PRIORITY_CLASS);
  633. SetConsoleCtrlHandler((PHANDLER_ROUTINE)handler, TRUE);
  634.  
  635. // main loop
  636. for ( ; ; ){
  637. t1 = GetTickCount();
  638.  
  639. count = 0;
  640.  
  641. for (int i = 0; i < 6; i++){
  642. key[i] = b64t[xor128() & 63];
  643. }
  644.  
  645. BYTE c7 = (BYTE)(xor128() & 63);
  646. key[7] = c7;
  647.  
  648. BYTE c8 = (BYTE)(xor128() & 63);
  649. key[8] = c8;
  650.  
  651. BYTE c9 = (BYTE)(xor128() & 63);
  652. key[9] = c9;
  653.  
  654. for (BYTE c6 = 0; c6 < 64; c6++){
  655. key[6] = b64t[c6];
  656.  
  657. memcpy(M, key, 10);
  658. cudaMemcpyToSymbol(Md, M, 64);
  659.  
  660. cudaMemcpy(Rd, R, R_size, cudaMemcpyHostToDevice);
  661.  
  662. dim3 dimBlock(BLOCK_SIZE_X, BLOCK_SIZE_Y);
  663. dim3 dimGrid(blocks);
  664.  
  665. sha1trip_search<<<dimGrid, dimBlock>>>(Rd);
  666.  
  667. cudaMemcpy(R, Rd, R_size, cudaMemcpyDeviceToHost);
  668.  
  669. for (unsigned int r_index = 0; r_index < BLOCK_SIZE_X * BLOCK_SIZE_Y * blocks; r_index++){
  670. st_sha1trip *out = &R[r_index];
  671.  
  672. count += out->count;
  673.  
  674. if (! out->trip[0]){
  675. continue;
  676. }
  677.  
  678. memcpy(trip, out->trip, 12);
  679.  
  680. for (int lower = 0, upper = target_num -1; lower <= upper; ){
  681. int middle = (lower + upper) >> 1;
  682.  
  683. if (strncmp(target[middle], trip, strlen(target[middle])) < 0){
  684. lower = middle + 1;
  685. }
  686. else if (strncmp(target[middle], trip, strlen(target[middle])) > 0){
  687. upper = middle - 1;
  688. }
  689. else {
  690. memcpy(key + 7, out->key + 7, 5);
  691.  
  692. fprintf(fp_trip, "%s\t#%s\n", trip, key);
  693. fflush(fp_trip);
  694. break;
  695. }
  696. }
  697.  
  698. }
  699.  
  700. memset(R, 0, R_size);
  701. }
  702.  
  703. t2 = GetTickCount();
  704.  
  705. if (t2 > t1){
  706. delta_sec = (t2 - t1) * 0.001;
  707. }
  708. else {
  709. delta_sec = 4294967.296 + t2 * 0.001 - t1 * 0.001;
  710. }
  711.  
  712. printf("%d kTrips in %.3f sec - %.3f MTrips/sec\n",
  713. count / 1000, delta_sec, (count / delta_sec) * 0.000001);
  714.  
  715. }
  716.  
  717. // never reach here
  718. fclose(fp_trip);
  719.  
  720. cudaFree(Rd);
  721. free(R);
  722.  
  723. return 0;
  724. }
Advertisement
Add Comment
Please, Sign In to add comment
Public Pastes
We use cookies for various purposes including analytics. By continuing to use Pastebin, you agree to our use of cookies as described in the Cookies Policy.  OK, I Understand
Not a member of Pastebin yet?
Sign Up, it unlocks many cool features!