nv/m00200_a0.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _MYSQL323_
   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 #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 #ifdef  VECT_SIZE4
  41 #define VECT_COMPARE_S "check_single_vect4_comp4.c"
  42 #define VECT_COMPARE_M "check_multi_vect4_comp4.c"
  43 #endif
  44
  45 __device__ __constant__ gpu_rule_t c_rules[1024];
  46
  47 extern "C" __global__ void __launch_bounds__ (256, 1) m00200_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)
  48 {
  49   /**
  50    * modifier
  51    */
  52
  53   const u32 lid = threadIdx.x;
  54
  55   /**
  56    * base
  57    */
  58
  59   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
  60
  61   if (gid >= gid_max) return;
  62
  63   u32x pw_buf0[4];
  64
  65   pw_buf0[0] = pws[gid].i[ 0];
  66   pw_buf0[1] = pws[gid].i[ 1];
  67   pw_buf0[2] = pws[gid].i[ 2];
  68   pw_buf0[3] = pws[gid].i[ 3];
  69
  70   u32x pw_buf1[4];
  71
  72   pw_buf1[0] = pws[gid].i[ 4];
  73   pw_buf1[1] = pws[gid].i[ 5];
  74   pw_buf1[2] = pws[gid].i[ 6];
  75   pw_buf1[3] = pws[gid].i[ 7];
  76
  77   const u32 pw_len = pws[gid].pw_len;
  78
  79   /**
  80    * loop
  81    */
  82
  83   for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
  84   {
  85     u32x w0[4];
  86
  87     w0[0] = pw_buf0[0];
  88     w0[1] = pw_buf0[1];
  89     w0[2] = pw_buf0[2];
  90     w0[3] = pw_buf0[3];
  91
  92     u32x w1[4];
  93
  94     w1[0] = pw_buf1[0];
  95     w1[1] = pw_buf1[1];
  96     w1[2] = pw_buf1[2];
  97     w1[3] = pw_buf1[3];
  98
  99     u32x w2[4];
 100
 101     w2[0] = 0;
 102     w2[1] = 0;
 103     w2[2] = 0;
 104     w2[3] = 0;
 105
 106     u32x w3[4];
 107
 108     w3[0] = 0;
 109     w3[1] = 0;
 110     w3[2] = 0;
 111     w3[3] = 0;
 112
 113     const u32 out_len = apply_rules (c_rules[il_pos].cmds, w0, w1, pw_len);
 114
 115     u32x w_t[16];
 116
 117     w_t[ 0] = w0[0];
 118     w_t[ 1] = w0[1];
 119     w_t[ 2] = w0[2];
 120     w_t[ 3] = w0[3];
 121     w_t[ 4] = w1[0];
 122     w_t[ 5] = w1[1];
 123     w_t[ 6] = w1[2];
 124     w_t[ 7] = w1[3];
 125     w_t[ 8] = 0;
 126     w_t[ 9] = 0;
 127     w_t[10] = 0;
 128     w_t[11] = 0;
 129     w_t[12] = 0;
 130     w_t[13] = 0;
 131     w_t[14] = 0;
 132     w_t[15] = 0;
 133
 134     u32x a = MYSQL323_A;
 135     u32x b = MYSQL323_B;
 136
 137     u32x add = 7;
 138
 139     #define ROUND(v)                              \
 140     {                                             \
 141       a ^= (((a & 0x3f) + add) * (v)) + (a << 8); \
 142       b += (b << 8) ^ a;                          \
 143       add += v;                                   \
 144     }
 145
 146     int i;
 147     int j;
 148
 149     for (i = 0, j = 0; i <= (int) out_len - 4; i += 4, j += 1)
 150     {
 151       const u32x wj = w_t[j];
 152
 153       ROUND ((wj >>  0) & 0xff);
 154       ROUND ((wj >>  8) & 0xff);
 155       ROUND ((wj >> 16) & 0xff);
 156       ROUND ((wj >> 24) & 0xff);
 157     }
 158
 159     const u32x wj = w_t[j];
 160
 161     const u32 left = out_len - i;
 162
 163     if (left == 3)
 164     {
 165       ROUND ((wj >>  0) & 0xff);
 166       ROUND ((wj >>  8) & 0xff);
 167       ROUND ((wj >> 16) & 0xff);
 168     }
 169     else if (left == 2)
 170     {
 171       ROUND ((wj >>  0) & 0xff);
 172       ROUND ((wj >>  8) & 0xff);
 173     }
 174     else if (left == 1)
 175     {
 176       ROUND ((wj >>  0) & 0xff);
 177     }
 178
 179     a &= 0x7fffffff;
 180     b &= 0x7fffffff;
 181
 182     const u32x r0 = a;
 183     const u32x r1 = b;
 184     const u32x r2 = 0;
 185     const u32x r3 = 0;
 186
 187     #include VECT_COMPARE_M
 188   }
 189 }
 190
 191 extern "C" __global__ void __launch_bounds__ (256, 1) m00200_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)
 192 {
 193 }
 194
 195 extern "C" __global__ void __launch_bounds__ (256, 1) m00200_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)
 196 {
 197 }
 198
 199 extern "C" __global__ void __launch_bounds__ (256, 1) m00200_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)
 200 {
 201   /**
 202    * modifier
 203    */
 204
 205   const u32 lid = threadIdx.x;
 206
 207   /**
 208    * base
 209    */
 210
 211   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 212
 213   if (gid >= gid_max) return;
 214
 215   u32x pw_buf0[4];
 216
 217   pw_buf0[0] = pws[gid].i[ 0];
 218   pw_buf0[1] = pws[gid].i[ 1];
 219   pw_buf0[2] = pws[gid].i[ 2];
 220   pw_buf0[3] = pws[gid].i[ 3];
 221
 222   u32x pw_buf1[4];
 223
 224   pw_buf1[0] = pws[gid].i[ 4];
 225   pw_buf1[1] = pws[gid].i[ 5];
 226   pw_buf1[2] = pws[gid].i[ 6];
 227   pw_buf1[3] = pws[gid].i[ 7];
 228
 229   const u32 pw_len = pws[gid].pw_len;
 230
 231   /**
 232    * digest
 233    */
 234
 235   const u32 search[4] =
 236   {
 237     digests_buf[digests_offset].digest_buf[DGST_R0],
 238     digests_buf[digests_offset].digest_buf[DGST_R1],
 239     digests_buf[digests_offset].digest_buf[DGST_R2],
 240     digests_buf[digests_offset].digest_buf[DGST_R3]
 241   };
 242
 243   /**
 244    * loop
 245    */
 246
 247   for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
 248   {
 249     u32x w0[4];
 250
 251     w0[0] = pw_buf0[0];
 252     w0[1] = pw_buf0[1];
 253     w0[2] = pw_buf0[2];
 254     w0[3] = pw_buf0[3];
 255
 256     u32x w1[4];
 257
 258     w1[0] = pw_buf1[0];
 259     w1[1] = pw_buf1[1];
 260     w1[2] = pw_buf1[2];
 261     w1[3] = pw_buf1[3];
 262
 263     u32x w2[4];
 264
 265     w2[0] = 0;
 266     w2[1] = 0;
 267     w2[2] = 0;
 268     w2[3] = 0;
 269
 270     u32x w3[4];
 271
 272     w3[0] = 0;
 273     w3[1] = 0;
 274     w3[2] = 0;
 275     w3[3] = 0;
 276
 277     const u32 out_len = apply_rules (c_rules[il_pos].cmds, w0, w1, pw_len);
 278
 279     u32x w_t[16];
 280
 281     w_t[ 0] = w0[0];
 282     w_t[ 1] = w0[1];
 283     w_t[ 2] = w0[2];
 284     w_t[ 3] = w0[3];
 285     w_t[ 4] = w1[0];
 286     w_t[ 5] = w1[1];
 287     w_t[ 6] = w1[2];
 288     w_t[ 7] = w1[3];
 289     w_t[ 8] = 0;
 290     w_t[ 9] = 0;
 291     w_t[10] = 0;
 292     w_t[11] = 0;
 293     w_t[12] = 0;
 294     w_t[13] = 0;
 295     w_t[14] = 0;
 296     w_t[15] = 0;
 297
 298     u32x a = MYSQL323_A;
 299     u32x b = MYSQL323_B;
 300
 301     u32x add = 7;
 302
 303     #define ROUND(v)                              \
 304     {                                             \
 305       a ^= (((a & 0x3f) + add) * (v)) + (a << 8); \
 306       b += (b << 8) ^ a;                          \
 307       add += v;                                   \
 308     }
 309
 310     int i;
 311     int j;
 312
 313     for (i = 0, j = 0; i <= (int) out_len - 4; i += 4, j += 1)
 314     {
 315       const u32x wj = w_t[j];
 316
 317       ROUND ((wj >>  0) & 0xff);
 318       ROUND ((wj >>  8) & 0xff);
 319       ROUND ((wj >> 16) & 0xff);
 320       ROUND ((wj >> 24) & 0xff);
 321     }
 322
 323     const u32x wj = w_t[j];
 324
 325     const u32 left = out_len - i;
 326
 327     if (left == 3)
 328     {
 329       ROUND ((wj >>  0) & 0xff);
 330       ROUND ((wj >>  8) & 0xff);
 331       ROUND ((wj >> 16) & 0xff);
 332     }
 333     else if (left == 2)
 334     {
 335       ROUND ((wj >>  0) & 0xff);
 336       ROUND ((wj >>  8) & 0xff);
 337     }
 338     else if (left == 1)
 339     {
 340       ROUND ((wj >>  0) & 0xff);
 341     }
 342
 343     a &= 0x7fffffff;
 344     b &= 0x7fffffff;
 345
 346     const u32x r0 = a;
 347     const u32x r1 = b;
 348     const u32x r2 = 0;
 349     const u32x r3 = 0;
 350
 351     #include VECT_COMPARE_S
 352   }
 353 }
 354
 355 extern "C" __global__ void __launch_bounds__ (256, 1) m00200_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)
 356 {
 357 }
 358
 359 extern "C" __global__ void __launch_bounds__ (256, 1) m00200_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)
 360 {
 361 }