nv/m05000_a0.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _KECCAK_
   7
   8 #include "include/constants.h"
   9 #include "include/kernel_vendor.h"
  10
  11 #ifdef  VLIW1
  12 #define VECT_SIZE1
  13 #endif
  14
  15 #ifdef  VLIW2
  16 #define VECT_SIZE1
  17 #endif
  18
  19 #define DGST_R0 2
  20 #define DGST_R1 3
  21 #define DGST_R2 4
  22 #define DGST_R3 5
  23
  24 #include "include/kernel_functions.c"
  25 #include "types_nv.c"
  26 #include "common_nv.c"
  27 #include "include/rp_gpu.h"
  28 #include "rp_nv.c"
  29
  30 #ifdef  VECT_SIZE1
  31 #define VECT_COMPARE_S "check_single_vect1_comp4.c"
  32 #define VECT_COMPARE_M "check_multi_vect1_comp4.c"
  33 #endif
  34
  35 #ifdef  VECT_SIZE2
  36 #define VECT_COMPARE_S "check_single_vect2_comp4.c"
  37 #define VECT_COMPARE_M "check_multi_vect2_comp4.c"
  38 #endif
  39
  40 /**
  41  * constants
  42  */
  43
  44 #ifndef KECCAK_ROUNDS
  45 #define KECCAK_ROUNDS 24
  46 #endif
  47
  48 #define Theta1(s) (st[0 + s] ^ st[5 + s] ^ st[10 + s] ^ st[15 + s] ^ st[20 + s])
  49
  50 #define Theta2(s)               \
  51 {                               \
  52   st[ 0 + s] ^= t;              \
  53   st[ 5 + s] ^= t;              \
  54   st[10 + s] ^= t;              \
  55   st[15 + s] ^= t;              \
  56   st[20 + s] ^= t;              \
  57 }
  58
  59 #define Rho_Pi(s)               \
  60 {                               \
  61   u32 j = keccakf_piln[s];     \
  62   u32 k = keccakf_rotc[s];     \
  63   bc0 = st[j];                  \
  64   st[j] = rotl64 (t, k);        \
  65   t = bc0;                      \
  66 }
  67
  68 #define Chi(s)                  \
  69 {                               \
  70   bc0 = st[0 + s];              \
  71   bc1 = st[1 + s];              \
  72   bc2 = st[2 + s];              \
  73   bc3 = st[3 + s];              \
  74   bc4 = st[4 + s];              \
  75   st[0 + s] ^= ~bc1 & bc2;      \
  76   st[1 + s] ^= ~bc2 & bc3;      \
  77   st[2 + s] ^= ~bc3 & bc4;      \
  78   st[3 + s] ^= ~bc4 & bc0;      \
  79   st[4 + s] ^= ~bc0 & bc1;      \
  80 }
  81
  82 __device__ __constant__ gpu_rule_t c_rules[1024];
  83
  84 extern "C" __global__ void __launch_bounds__ (256, 1) m05000_m04 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
  85 {
  86   /**
  87    * modifier
  88    */
  89
  90   const u32 lid = threadIdx.x;
  91
  92   /**
  93    * base
  94    */
  95
  96   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
  97
  98   if (gid >= gid_max) return;
  99
 100   /**
 101    * constants
 102    */
 103
 104   const u64 keccakf_rndc[24] =
 105   {
 106     0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
 107     0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
 108     0x8000000080008081, 0x8000000000008009, 0x000000000000008a,
 109     0x0000000000000088, 0x0000000080008009, 0x000000008000000a,
 110     0x000000008000808b, 0x800000000000008b, 0x8000000000008089,
 111     0x8000000000008003, 0x8000000000008002, 0x8000000000000080,
 112     0x000000000000800a, 0x800000008000000a, 0x8000000080008081,
 113     0x8000000000008080, 0x0000000080000001, 0x8000000080008008
 114   };
 115
 116   const u32 keccakf_rotc[24] =
 117   {
 118      1,  3,  6, 10, 15, 21, 28, 36, 45, 55,  2, 14,
 119     27, 41, 56,  8, 25, 43, 62, 18, 39, 61, 20, 44
 120   };
 121
 122   const u32 keccakf_piln[24] =
 123   {
 124     10,  7, 11, 17, 18,  3,  5, 16,  8, 21, 24,  4,
 125     15, 23, 19, 13, 12,  2, 20, 14, 22,  9,  6,  1
 126   };
 127
 128   u32x pw_buf0[4];
 129
 130   pw_buf0[0] = pws[gid].i[ 0];
 131   pw_buf0[1] = pws[gid].i[ 1];
 132   pw_buf0[2] = pws[gid].i[ 2];
 133   pw_buf0[3] = pws[gid].i[ 3];
 134
 135   u32x pw_buf1[4];
 136
 137   pw_buf1[0] = pws[gid].i[ 4];
 138   pw_buf1[1] = pws[gid].i[ 5];
 139   pw_buf1[2] = pws[gid].i[ 6];
 140   pw_buf1[3] = pws[gid].i[ 7];
 141
 142   const u32 pw_len = pws[gid].pw_len;
 143
 144   /**
 145    * 0x80 keccak, very special
 146    */
 147
 148   const u32 mdlen = salt_bufs[salt_pos].keccak_mdlen;
 149
 150   const u32 rsiz = 200 - (2 * mdlen);
 151
 152   const u32 add80w = (rsiz - 1) / 8;
 153
 154   /**
 155    * loop
 156    */
 157
 158   for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
 159   {
 160     u32x w0[4];
 161
 162     w0[0] = pw_buf0[0];
 163     w0[1] = pw_buf0[1];
 164     w0[2] = pw_buf0[2];
 165     w0[3] = pw_buf0[3];
 166
 167     u32x w1[4];
 168
 169     w1[0] = pw_buf1[0];
 170     w1[1] = pw_buf1[1];
 171     w1[2] = pw_buf1[2];
 172     w1[3] = pw_buf1[3];
 173
 174     u32x w2[4];
 175
 176     w2[0] = 0;
 177     w2[1] = 0;
 178     w2[2] = 0;
 179     w2[3] = 0;
 180
 181     u32x w3[4];
 182
 183     w3[0] = 0;
 184     w3[1] = 0;
 185     w3[2] = 0;
 186     w3[3] = 0;
 187
 188     const u32 out_len = apply_rules (c_rules[il_pos].cmds, w0, w1, pw_len);
 189
 190     append_0x01_2 (w0, w1, out_len);
 191
 192     u64x st[25];
 193
 194     st[ 0] = (u64x) (w0[0]) | (u64x) (w0[1]) << 32;
 195     st[ 1] = (u64x) (w0[2]) | (u64x) (w0[3]) << 32;
 196     st[ 2] = (u64x) (w1[0]) | (u64x) (w1[1]) << 32;
 197     st[ 3] = (u64x) (w1[2]) | (u64x) (w1[3]) << 32;
 198     st[ 4] = 0;
 199     st[ 5] = 0;
 200     st[ 6] = 0;
 201     st[ 7] = 0;
 202     st[ 8] = 0;
 203     st[ 9] = 0;
 204     st[10] = 0;
 205     st[11] = 0;
 206     st[12] = 0;
 207     st[13] = 0;
 208     st[14] = 0;
 209     st[15] = 0;
 210     st[16] = 0;
 211     st[17] = 0;
 212     st[18] = 0;
 213     st[19] = 0;
 214     st[20] = 0;
 215     st[21] = 0;
 216     st[22] = 0;
 217     st[23] = 0;
 218     st[24] = 0;
 219
 220     st[add80w] |= 0x8000000000000000;
 221
 222     int round;
 223
 224     for (round = 0; round < KECCAK_ROUNDS; round++)
 225     {
 226       // Theta
 227
 228       u64x bc0 = Theta1 (0);
 229       u64x bc1 = Theta1 (1);
 230       u64x bc2 = Theta1 (2);
 231       u64x bc3 = Theta1 (3);
 232       u64x bc4 = Theta1 (4);
 233
 234       u64x t;
 235
 236       t = bc4 ^ rotl64 (bc1, 1); Theta2 (0);
 237       t = bc0 ^ rotl64 (bc2, 1); Theta2 (1);
 238       t = bc1 ^ rotl64 (bc3, 1); Theta2 (2);
 239       t = bc2 ^ rotl64 (bc4, 1); Theta2 (3);
 240       t = bc3 ^ rotl64 (bc0, 1); Theta2 (4);
 241
 242       // Rho Pi
 243
 244       t = st[1];
 245
 246       Rho_Pi (0);
 247       Rho_Pi (1);
 248       Rho_Pi (2);
 249       Rho_Pi (3);
 250       Rho_Pi (4);
 251       Rho_Pi (5);
 252       Rho_Pi (6);
 253       Rho_Pi (7);
 254       Rho_Pi (8);
 255       Rho_Pi (9);
 256       Rho_Pi (10);
 257       Rho_Pi (11);
 258       Rho_Pi (12);
 259       Rho_Pi (13);
 260       Rho_Pi (14);
 261       Rho_Pi (15);
 262       Rho_Pi (16);
 263       Rho_Pi (17);
 264       Rho_Pi (18);
 265       Rho_Pi (19);
 266       Rho_Pi (20);
 267       Rho_Pi (21);
 268       Rho_Pi (22);
 269       Rho_Pi (23);
 270
 271       //  Chi
 272
 273       Chi (0);
 274       Chi (5);
 275       Chi (10);
 276       Chi (15);
 277       Chi (20);
 278
 279       //  Iota
 280
 281       st[0] ^= keccakf_rndc[round];
 282     }
 283
 284     const u32x r0 = l32_from_64 (st[1]);
 285     const u32x r1 = h32_from_64 (st[1]);
 286     const u32x r2 = l32_from_64 (st[2]);
 287     const u32x r3 = h32_from_64 (st[2]);
 288
 289     #include VECT_COMPARE_M
 290   }
 291 }
 292
 293 extern "C" __global__ void __launch_bounds__ (256, 1) m05000_m08 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 294 {
 295 }
 296
 297 extern "C" __global__ void __launch_bounds__ (256, 1) m05000_m16 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 298 {
 299 }
 300
 301 extern "C" __global__ void __launch_bounds__ (256, 1) m05000_s04 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 302 {
 303   /**
 304    * modifier
 305    */
 306
 307   const u32 lid = threadIdx.x;
 308
 309   /**
 310    * base
 311    */
 312
 313   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 314
 315   if (gid >= gid_max) return;
 316
 317   /**
 318    * digest
 319    */
 320
 321   const u32 search[4] =
 322   {
 323     digests_buf[digests_offset].digest_buf[DGST_R0],
 324     digests_buf[digests_offset].digest_buf[DGST_R1],
 325     digests_buf[digests_offset].digest_buf[DGST_R2],
 326     digests_buf[digests_offset].digest_buf[DGST_R3]
 327   };
 328
 329   /**
 330    * constants
 331    */
 332
 333   const u64 keccakf_rndc[24] =
 334   {
 335     0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
 336     0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
 337     0x8000000080008081, 0x8000000000008009, 0x000000000000008a,
 338     0x0000000000000088, 0x0000000080008009, 0x000000008000000a,
 339     0x000000008000808b, 0x800000000000008b, 0x8000000000008089,
 340     0x8000000000008003, 0x8000000000008002, 0x8000000000000080,
 341     0x000000000000800a, 0x800000008000000a, 0x8000000080008081,
 342     0x8000000000008080, 0x0000000080000001, 0x8000000080008008
 343   };
 344
 345   const u32 keccakf_rotc[24] =
 346   {
 347      1,  3,  6, 10, 15, 21, 28, 36, 45, 55,  2, 14,
 348     27, 41, 56,  8, 25, 43, 62, 18, 39, 61, 20, 44
 349   };
 350
 351   const u32 keccakf_piln[24] =
 352   {
 353     10,  7, 11, 17, 18,  3,  5, 16,  8, 21, 24,  4,
 354     15, 23, 19, 13, 12,  2, 20, 14, 22,  9,  6,  1
 355   };
 356
 357   u32x pw_buf0[4];
 358
 359   pw_buf0[0] = pws[gid].i[ 0];
 360   pw_buf0[1] = pws[gid].i[ 1];
 361   pw_buf0[2] = pws[gid].i[ 2];
 362   pw_buf0[3] = pws[gid].i[ 3];
 363
 364   u32x pw_buf1[4];
 365
 366   pw_buf1[0] = pws[gid].i[ 4];
 367   pw_buf1[1] = pws[gid].i[ 5];
 368   pw_buf1[2] = pws[gid].i[ 6];
 369   pw_buf1[3] = pws[gid].i[ 7];
 370
 371   const u32 pw_len = pws[gid].pw_len;
 372
 373   /**
 374    * 0x80 keccak, very special
 375    */
 376
 377   const u32 mdlen = salt_bufs[salt_pos].keccak_mdlen;
 378
 379   const u32 rsiz = 200 - (2 * mdlen);
 380
 381   const u32 add80w = (rsiz - 1) / 8;
 382
 383   /**
 384    * loop
 385    */
 386
 387   for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
 388   {
 389     u32x w0[4];
 390
 391     w0[0] = pw_buf0[0];
 392     w0[1] = pw_buf0[1];
 393     w0[2] = pw_buf0[2];
 394     w0[3] = pw_buf0[3];
 395
 396     u32x w1[4];
 397
 398     w1[0] = pw_buf1[0];
 399     w1[1] = pw_buf1[1];
 400     w1[2] = pw_buf1[2];
 401     w1[3] = pw_buf1[3];
 402
 403     u32x w2[4];
 404
 405     w2[0] = 0;
 406     w2[1] = 0;
 407     w2[2] = 0;
 408     w2[3] = 0;
 409
 410     u32x w3[4];
 411
 412     w3[0] = 0;
 413     w3[1] = 0;
 414     w3[2] = 0;
 415     w3[3] = 0;
 416
 417     const u32 out_len = apply_rules (c_rules[il_pos].cmds, w0, w1, pw_len);
 418
 419     append_0x01_2 (w0, w1, out_len);
 420
 421     u64x st[25];
 422
 423     st[ 0] = (u64x) (w0[0]) | (u64x) (w0[1]) << 32;
 424     st[ 1] = (u64x) (w0[2]) | (u64x) (w0[3]) << 32;
 425     st[ 2] = (u64x) (w1[0]) | (u64x) (w1[1]) << 32;
 426     st[ 3] = (u64x) (w1[2]) | (u64x) (w1[3]) << 32;
 427     st[ 4] = 0;
 428     st[ 5] = 0;
 429     st[ 6] = 0;
 430     st[ 7] = 0;
 431     st[ 8] = 0;
 432     st[ 9] = 0;
 433     st[10] = 0;
 434     st[11] = 0;
 435     st[12] = 0;
 436     st[13] = 0;
 437     st[14] = 0;
 438     st[15] = 0;
 439     st[16] = 0;
 440     st[17] = 0;
 441     st[18] = 0;
 442     st[19] = 0;
 443     st[20] = 0;
 444     st[21] = 0;
 445     st[22] = 0;
 446     st[23] = 0;
 447     st[24] = 0;
 448
 449     st[add80w] |= 0x8000000000000000;
 450
 451     int round;
 452
 453     for (round = 0; round < KECCAK_ROUNDS; round++)
 454     {
 455       // Theta
 456
 457       u64x bc0 = Theta1 (0);
 458       u64x bc1 = Theta1 (1);
 459       u64x bc2 = Theta1 (2);
 460       u64x bc3 = Theta1 (3);
 461       u64x bc4 = Theta1 (4);
 462
 463       u64x t;
 464
 465       t = bc4 ^ rotl64 (bc1, 1); Theta2 (0);
 466       t = bc0 ^ rotl64 (bc2, 1); Theta2 (1);
 467       t = bc1 ^ rotl64 (bc3, 1); Theta2 (2);
 468       t = bc2 ^ rotl64 (bc4, 1); Theta2 (3);
 469       t = bc3 ^ rotl64 (bc0, 1); Theta2 (4);
 470
 471       // Rho Pi
 472
 473       t = st[1];
 474
 475       Rho_Pi (0);
 476       Rho_Pi (1);
 477       Rho_Pi (2);
 478       Rho_Pi (3);
 479       Rho_Pi (4);
 480       Rho_Pi (5);
 481       Rho_Pi (6);
 482       Rho_Pi (7);
 483       Rho_Pi (8);
 484       Rho_Pi (9);
 485       Rho_Pi (10);
 486       Rho_Pi (11);
 487       Rho_Pi (12);
 488       Rho_Pi (13);
 489       Rho_Pi (14);
 490       Rho_Pi (15);
 491       Rho_Pi (16);
 492       Rho_Pi (17);
 493       Rho_Pi (18);
 494       Rho_Pi (19);
 495       Rho_Pi (20);
 496       Rho_Pi (21);
 497       Rho_Pi (22);
 498       Rho_Pi (23);
 499
 500       //  Chi
 501
 502       Chi (0);
 503       Chi (5);
 504       Chi (10);
 505       Chi (15);
 506       Chi (20);
 507
 508       //  Iota
 509
 510       st[0] ^= keccakf_rndc[round];
 511     }
 512
 513     const u32x r0 = l32_from_64 (st[1]);
 514     const u32x r1 = h32_from_64 (st[1]);
 515     const u32x r2 = l32_from_64 (st[2]);
 516     const u32x r3 = h32_from_64 (st[2]);
 517
 518     #include VECT_COMPARE_S
 519   }
 520 }
 521
 522 extern "C" __global__ void __launch_bounds__ (256, 1) m05000_s08 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 523 {
 524 }
 525
 526 extern "C" __global__ void __launch_bounds__ (256, 1) m05000_s16 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 527 {
 528 }