nv/m10100_a1.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _SIPHASH_
   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 0
  20 #define DGST_R1 1
  21 #define DGST_R2 2
  22 #define DGST_R3 3
  23
  24 #include "include/kernel_functions.c"
  25 #include "types_nv.c"
  26 #include "common_nv.c"
  27
  28 #ifdef  VECT_SIZE1
  29 #define VECT_COMPARE_S "check_single_vect1_comp4.c"
  30 #define VECT_COMPARE_M "check_multi_vect1_comp4.c"
  31 #endif
  32
  33 #ifdef  VECT_SIZE2
  34 #define VECT_COMPARE_S "check_single_vect2_comp4.c"
  35 #define VECT_COMPARE_M "check_multi_vect2_comp4.c"
  36 #endif
  37
  38 #ifdef  VECT_SIZE4
  39 #define VECT_COMPARE_S "check_single_vect4_comp4.c"
  40 #define VECT_COMPARE_M "check_multi_vect4_comp4.c"
  41 #endif
  42
  43 #define SIPROUND(v0,v1,v2,v3) \
  44   (v0) += (v1);               \
  45   (v1)  = rotl64 ((v1), 13);  \
  46   (v1) ^= (v0);               \
  47   (v0)  = rotl64 ((v0), 32);  \
  48   (v2) += (v3);               \
  49   (v3)  = rotl64 ((v3), 16);  \
  50   (v3) ^= (v2);               \
  51   (v0) += (v3);               \
  52   (v3)  = rotl64 ((v3), 21);  \
  53   (v3) ^= (v0);               \
  54   (v2) += (v1);               \
  55   (v1)  = rotl64 ((v1), 17);  \
  56   (v1) ^= (v2);               \
  57   (v2)  = rotl64 ((v2), 32);
  58
  59 __device__ __constant__ comb_t c_combs[1024];
  60
  61 extern "C" __global__ void __launch_bounds__ (256, 1) m10100_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
  62 {
  63   /**
  64    * modifier
  65    */
  66
  67   const u32 lid = threadIdx.x;
  68
  69   /**
  70    * base
  71    */
  72
  73   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
  74
  75   if (gid >= gid_max) return;
  76
  77   u32x wordl0[4];
  78
  79   wordl0[0] = pws[gid].i[ 0];
  80   wordl0[1] = pws[gid].i[ 1];
  81   wordl0[2] = pws[gid].i[ 2];
  82   wordl0[3] = pws[gid].i[ 3];
  83
  84   u32x wordl1[4];
  85
  86   wordl1[0] = pws[gid].i[ 4];
  87   wordl1[1] = pws[gid].i[ 5];
  88   wordl1[2] = pws[gid].i[ 6];
  89   wordl1[3] = pws[gid].i[ 7];
  90
  91   u32x wordl2[4];
  92
  93   wordl2[0] = 0;
  94   wordl2[1] = 0;
  95   wordl2[2] = 0;
  96   wordl2[3] = 0;
  97
  98   u32x wordl3[4];
  99
 100   wordl3[0] = 0;
 101   wordl3[1] = 0;
 102   wordl3[2] = 0;
 103   wordl3[3] = 0;
 104
 105   const u32 pw_l_len = pws[gid].pw_len;
 106
 107   if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
 108   {
 109     switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, c_combs[0].pw_len);
 110   }
 111
 112   /**
 113    * base
 114    */
 115
 116   u64 v0p = SIPHASHM_0;
 117   u64 v1p = SIPHASHM_1;
 118   u64 v2p = SIPHASHM_2;
 119   u64 v3p = SIPHASHM_3;
 120
 121   v0p ^= hl32_to_64 (salt_bufs[salt_pos].salt_buf[1], salt_bufs[salt_pos].salt_buf[0]);
 122   v1p ^= hl32_to_64 (salt_bufs[salt_pos].salt_buf[3], salt_bufs[salt_pos].salt_buf[2]);
 123   v2p ^= hl32_to_64 (salt_bufs[salt_pos].salt_buf[1], salt_bufs[salt_pos].salt_buf[0]);
 124   v3p ^= hl32_to_64 (salt_bufs[salt_pos].salt_buf[3], salt_bufs[salt_pos].salt_buf[2]);
 125
 126   /**
 127    * loop
 128    */
 129
 130   for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
 131   {
 132     const u32 pw_r_len = c_combs[il_pos].pw_len;
 133
 134     const u32 pw_len = pw_l_len + pw_r_len;
 135
 136     u32 wordr0[4];
 137
 138     wordr0[0] = c_combs[il_pos].i[0];
 139     wordr0[1] = c_combs[il_pos].i[1];
 140     wordr0[2] = c_combs[il_pos].i[2];
 141     wordr0[3] = c_combs[il_pos].i[3];
 142
 143     u32 wordr1[4];
 144
 145     wordr1[0] = c_combs[il_pos].i[4];
 146     wordr1[1] = c_combs[il_pos].i[5];
 147     wordr1[2] = c_combs[il_pos].i[6];
 148     wordr1[3] = c_combs[il_pos].i[7];
 149
 150     u32 wordr2[4];
 151
 152     wordr2[0] = 0;
 153     wordr2[1] = 0;
 154     wordr2[2] = 0;
 155     wordr2[3] = 0;
 156
 157     u32 wordr3[4];
 158
 159     wordr3[0] = 0;
 160     wordr3[1] = 0;
 161     wordr3[2] = 0;
 162     wordr3[3] = 0;
 163
 164     if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
 165     {
 166       switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
 167     }
 168
 169     u32x w[16];
 170
 171     w[ 0] = wordl0[0] | wordr0[0];
 172     w[ 1] = wordl0[1] | wordr0[1];
 173     w[ 2] = wordl0[2] | wordr0[2];
 174     w[ 3] = wordl0[3] | wordr0[3];
 175     w[ 4] = wordl1[0] | wordr1[0];
 176     w[ 5] = wordl1[1] | wordr1[1];
 177     w[ 6] = wordl1[2] | wordr1[2];
 178     w[ 7] = wordl1[3] | wordr1[3];
 179     w[ 8] = wordl2[0] | wordr2[0];
 180     w[ 9] = wordl2[1] | wordr2[1];
 181     w[10] = wordl2[2] | wordr2[2];
 182     w[11] = wordl2[3] | wordr2[3];
 183     w[12] = wordl3[0] | wordr3[0];
 184     w[13] = wordl3[1] | wordr3[1];
 185     w[14] = wordl3[2] | wordr3[2];
 186     w[15] = wordl3[3] | wordr3[3];
 187
 188     u64 *w_ptr = (u64 *) w;
 189
 190     w_ptr[pw_len / 8] |= (u64) pw_len << 56;
 191
 192     u64x v0 = v0p;
 193     u64x v1 = v1p;
 194     u64x v2 = v2p;
 195     u64x v3 = v3p;
 196
 197     int i;
 198     int j;
 199
 200     for (i = 0, j = 0; i <= pw_len; i += 8, j += 2)
 201     {
 202       u64x m = hl32_to_64 (w[j + 1], w[j + 0]);
 203
 204       v3 ^= m;
 205
 206       SIPROUND (v0, v1, v2, v3);
 207       SIPROUND (v0, v1, v2, v3);
 208
 209       v0 ^= m;
 210     }
 211
 212     v2 ^= 0xff;
 213
 214     SIPROUND (v0, v1, v2, v3);
 215     SIPROUND (v0, v1, v2, v3);
 216     SIPROUND (v0, v1, v2, v3);
 217     SIPROUND (v0, v1, v2, v3);
 218
 219     const u64x v = v0 ^ v1 ^ v2 ^ v3;
 220
 221     const u32x a = l32_from_64 (v);
 222     const u32x b = h32_from_64 (v);
 223
 224     const u32x r0 = a;
 225     const u32x r1 = b;
 226     const u32x r2 = 0;
 227     const u32x r3 = 0;
 228
 229     #include VECT_COMPARE_M
 230   }
 231 }
 232
 233 extern "C" __global__ void __launch_bounds__ (256, 1) m10100_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 234 {
 235 }
 236
 237 extern "C" __global__ void __launch_bounds__ (256, 1) m10100_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 238 {
 239 }
 240
 241 extern "C" __global__ void __launch_bounds__ (256, 1) m10100_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 242 {
 243   /**
 244    * modifier
 245    */
 246
 247   const u32 lid = threadIdx.x;
 248
 249   /**
 250    * base
 251    */
 252
 253   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 254
 255   if (gid >= gid_max) return;
 256
 257   u32x wordl0[4];
 258
 259   wordl0[0] = pws[gid].i[ 0];
 260   wordl0[1] = pws[gid].i[ 1];
 261   wordl0[2] = pws[gid].i[ 2];
 262   wordl0[3] = pws[gid].i[ 3];
 263
 264   u32x wordl1[4];
 265
 266   wordl1[0] = pws[gid].i[ 4];
 267   wordl1[1] = pws[gid].i[ 5];
 268   wordl1[2] = pws[gid].i[ 6];
 269   wordl1[3] = pws[gid].i[ 7];
 270
 271   u32x wordl2[4];
 272
 273   wordl2[0] = 0;
 274   wordl2[1] = 0;
 275   wordl2[2] = 0;
 276   wordl2[3] = 0;
 277
 278   u32x wordl3[4];
 279
 280   wordl3[0] = 0;
 281   wordl3[1] = 0;
 282   wordl3[2] = 0;
 283   wordl3[3] = 0;
 284
 285   const u32 pw_l_len = pws[gid].pw_len;
 286
 287   if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
 288   {
 289     switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, c_combs[0].pw_len);
 290   }
 291
 292   /**
 293    * base
 294    */
 295
 296   u64 v0p = SIPHASHM_0;
 297   u64 v1p = SIPHASHM_1;
 298   u64 v2p = SIPHASHM_2;
 299   u64 v3p = SIPHASHM_3;
 300
 301   v0p ^= hl32_to_64 (salt_bufs[salt_pos].salt_buf[1], salt_bufs[salt_pos].salt_buf[0]);
 302   v1p ^= hl32_to_64 (salt_bufs[salt_pos].salt_buf[3], salt_bufs[salt_pos].salt_buf[2]);
 303   v2p ^= hl32_to_64 (salt_bufs[salt_pos].salt_buf[1], salt_bufs[salt_pos].salt_buf[0]);
 304   v3p ^= hl32_to_64 (salt_bufs[salt_pos].salt_buf[3], salt_bufs[salt_pos].salt_buf[2]);
 305
 306   /**
 307    * digest
 308    */
 309
 310   const u32 search[4] =
 311   {
 312     digests_buf[digests_offset].digest_buf[DGST_R0],
 313     digests_buf[digests_offset].digest_buf[DGST_R1],
 314     digests_buf[digests_offset].digest_buf[DGST_R2],
 315     digests_buf[digests_offset].digest_buf[DGST_R3]
 316   };
 317
 318   /**
 319    * loop
 320    */
 321
 322   for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
 323   {
 324     const u32 pw_r_len = c_combs[il_pos].pw_len;
 325
 326     const u32 pw_len = pw_l_len + pw_r_len;
 327
 328     u32 wordr0[4];
 329
 330     wordr0[0] = c_combs[il_pos].i[0];
 331     wordr0[1] = c_combs[il_pos].i[1];
 332     wordr0[2] = c_combs[il_pos].i[2];
 333     wordr0[3] = c_combs[il_pos].i[3];
 334
 335     u32 wordr1[4];
 336
 337     wordr1[0] = c_combs[il_pos].i[4];
 338     wordr1[1] = c_combs[il_pos].i[5];
 339     wordr1[2] = c_combs[il_pos].i[6];
 340     wordr1[3] = c_combs[il_pos].i[7];
 341
 342     u32 wordr2[4];
 343
 344     wordr2[0] = 0;
 345     wordr2[1] = 0;
 346     wordr2[2] = 0;
 347     wordr2[3] = 0;
 348
 349     u32 wordr3[4];
 350
 351     wordr3[0] = 0;
 352     wordr3[1] = 0;
 353     wordr3[2] = 0;
 354     wordr3[3] = 0;
 355
 356     if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
 357     {
 358       switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
 359     }
 360
 361     u32x w[16];
 362
 363     w[ 0] = wordl0[0] | wordr0[0];
 364     w[ 1] = wordl0[1] | wordr0[1];
 365     w[ 2] = wordl0[2] | wordr0[2];
 366     w[ 3] = wordl0[3] | wordr0[3];
 367     w[ 4] = wordl1[0] | wordr1[0];
 368     w[ 5] = wordl1[1] | wordr1[1];
 369     w[ 6] = wordl1[2] | wordr1[2];
 370     w[ 7] = wordl1[3] | wordr1[3];
 371     w[ 8] = wordl2[0] | wordr2[0];
 372     w[ 9] = wordl2[1] | wordr2[1];
 373     w[10] = wordl2[2] | wordr2[2];
 374     w[11] = wordl2[3] | wordr2[3];
 375     w[12] = wordl3[0] | wordr3[0];
 376     w[13] = wordl3[1] | wordr3[1];
 377     w[14] = wordl3[2] | wordr3[2];
 378     w[15] = wordl3[3] | wordr3[3];
 379
 380     u64 *w_ptr = (u64 *) w;
 381
 382     w_ptr[pw_len / 8] |= (u64) pw_len << 56;
 383
 384     u64x v0 = v0p;
 385     u64x v1 = v1p;
 386     u64x v2 = v2p;
 387     u64x v3 = v3p;
 388
 389     int i;
 390     int j;
 391
 392     for (i = 0, j = 0; i <= pw_len; i += 8, j += 2)
 393     {
 394       u64x m = hl32_to_64 (w[j + 1], w[j + 0]);
 395
 396       v3 ^= m;
 397
 398       SIPROUND (v0, v1, v2, v3);
 399       SIPROUND (v0, v1, v2, v3);
 400
 401       v0 ^= m;
 402     }
 403
 404     v2 ^= 0xff;
 405
 406     SIPROUND (v0, v1, v2, v3);
 407     SIPROUND (v0, v1, v2, v3);
 408     SIPROUND (v0, v1, v2, v3);
 409     SIPROUND (v0, v1, v2, v3);
 410
 411     const u64x v = v0 ^ v1 ^ v2 ^ v3;
 412
 413     const u32x a = l32_from_64 (v);
 414     const u32x b = h32_from_64 (v);
 415
 416     const u32x r0 = a;
 417     const u32x r1 = b;
 418     const u32x r2 = 0;
 419     const u32x r3 = 0;
 420
 421     #include VECT_COMPARE_S
 422   }
 423 }
 424
 425 extern "C" __global__ void __launch_bounds__ (256, 1) m10100_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 426 {
 427 }
 428
 429 extern "C" __global__ void __launch_bounds__ (256, 1) m10100_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 430 {
 431 }