nv/m10100_a0.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 #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 #define SIPROUND(v0,v1,v2,v3) \
  46   (v0) += (v1);               \
  47   (v1)  = rotl64 ((v1), 13);  \
  48   (v1) ^= (v0);               \
  49   (v0)  = rotl64 ((v0), 32);  \
  50   (v2) += (v3);               \
  51   (v3)  = rotl64 ((v3), 16);  \
  52   (v3) ^= (v2);               \
  53   (v0) += (v3);               \
  54   (v3)  = rotl64 ((v3), 21);  \
  55   (v3) ^= (v0);               \
  56   (v2) += (v1);               \
  57   (v1)  = rotl64 ((v1), 17);  \
  58   (v1) ^= (v2);               \
  59   (v2)  = rotl64 ((v2), 32);
  60
  61 __device__ __constant__ gpu_rule_t c_rules[1024];
  62
  63 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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
  64 {
  65   /**
  66    * modifier
  67    */
  68
  69   const u32 lid = threadIdx.x;
  70
  71   /**
  72    * base
  73    */
  74
  75   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
  76
  77   if (gid >= gid_max) return;
  78
  79   u32x pw_buf0[4];
  80
  81   pw_buf0[0] = pws[gid].i[ 0];
  82   pw_buf0[1] = pws[gid].i[ 1];
  83   pw_buf0[2] = pws[gid].i[ 2];
  84   pw_buf0[3] = pws[gid].i[ 3];
  85
  86   u32x pw_buf1[4];
  87
  88   pw_buf1[0] = pws[gid].i[ 4];
  89   pw_buf1[1] = pws[gid].i[ 5];
  90   pw_buf1[2] = pws[gid].i[ 6];
  91   pw_buf1[3] = pws[gid].i[ 7];
  92
  93   const u32 pw_len = pws[gid].pw_len;
  94
  95   /**
  96    * base
  97    */
  98
  99   u64 v0p = SIPHASHM_0;
 100   u64 v1p = SIPHASHM_1;
 101   u64 v2p = SIPHASHM_2;
 102   u64 v3p = SIPHASHM_3;
 103
 104   v0p ^= hl32_to_64 (salt_bufs[salt_pos].salt_buf[1], salt_bufs[salt_pos].salt_buf[0]);
 105   v1p ^= hl32_to_64 (salt_bufs[salt_pos].salt_buf[3], salt_bufs[salt_pos].salt_buf[2]);
 106   v2p ^= hl32_to_64 (salt_bufs[salt_pos].salt_buf[1], salt_bufs[salt_pos].salt_buf[0]);
 107   v3p ^= hl32_to_64 (salt_bufs[salt_pos].salt_buf[3], salt_bufs[salt_pos].salt_buf[2]);
 108
 109   /**
 110    * loop
 111    */
 112
 113   for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
 114   {
 115     u32x w[16];
 116
 117     w[ 0] = pw_buf0[0];
 118     w[ 1] = pw_buf0[1];
 119     w[ 2] = pw_buf0[2];
 120     w[ 3] = pw_buf0[3];
 121     w[ 4] = pw_buf1[0];
 122     w[ 5] = pw_buf1[1];
 123     w[ 6] = pw_buf1[2];
 124     w[ 7] = pw_buf1[3];
 125     w[ 8] = 0;
 126     w[ 9] = 0;
 127     w[10] = 0;
 128     w[11] = 0;
 129     w[12] = 0;
 130     w[13] = 0;
 131     w[14] = 0;
 132     w[15] = 0;
 133
 134     const u32x out_len = apply_rules (c_rules[il_pos].cmds, &w[0], &w[4], pw_len);
 135
 136     u64 *w_ptr = (u64 *) w;
 137
 138     w_ptr[out_len / 8] |= (u64) out_len << 56;
 139
 140     u64x v0 = v0p;
 141     u64x v1 = v1p;
 142     u64x v2 = v2p;
 143     u64x v3 = v3p;
 144
 145     int i;
 146     int j;
 147
 148     for (i = 0, j = 0; i <= pw_len; i += 8, j += 2)
 149     {
 150       u64x m = hl32_to_64 (w[j + 1], w[j + 0]);
 151
 152       v3 ^= m;
 153
 154       SIPROUND (v0, v1, v2, v3);
 155       SIPROUND (v0, v1, v2, v3);
 156
 157       v0 ^= m;
 158     }
 159
 160     v2 ^= 0xff;
 161
 162     SIPROUND (v0, v1, v2, v3);
 163     SIPROUND (v0, v1, v2, v3);
 164     SIPROUND (v0, v1, v2, v3);
 165     SIPROUND (v0, v1, v2, v3);
 166
 167     const u64x v = v0 ^ v1 ^ v2 ^ v3;
 168
 169     const u32x a = l32_from_64 (v);
 170     const u32x b = h32_from_64 (v);
 171
 172     const u32x r0 = a;
 173     const u32x r1 = b;
 174     const u32x r2 = 0;
 175     const u32x r3 = 0;
 176
 177     #include VECT_COMPARE_M
 178   }
 179 }
 180
 181 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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 182 {
 183 }
 184
 185 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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 186 {
 187 }
 188
 189 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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 190 {
 191   /**
 192    * modifier
 193    */
 194
 195   const u32 lid = threadIdx.x;
 196
 197   /**
 198    * base
 199    */
 200
 201   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 202
 203   if (gid >= gid_max) return;
 204
 205   u32x pw_buf0[4];
 206
 207   pw_buf0[0] = pws[gid].i[ 0];
 208   pw_buf0[1] = pws[gid].i[ 1];
 209   pw_buf0[2] = pws[gid].i[ 2];
 210   pw_buf0[3] = pws[gid].i[ 3];
 211
 212   u32x pw_buf1[4];
 213
 214   pw_buf1[0] = pws[gid].i[ 4];
 215   pw_buf1[1] = pws[gid].i[ 5];
 216   pw_buf1[2] = pws[gid].i[ 6];
 217   pw_buf1[3] = pws[gid].i[ 7];
 218
 219   const u32 pw_len = pws[gid].pw_len;
 220
 221   /**
 222    * base
 223    */
 224
 225   u64 v0p = SIPHASHM_0;
 226   u64 v1p = SIPHASHM_1;
 227   u64 v2p = SIPHASHM_2;
 228   u64 v3p = SIPHASHM_3;
 229
 230   v0p ^= hl32_to_64 (salt_bufs[salt_pos].salt_buf[1], salt_bufs[salt_pos].salt_buf[0]);
 231   v1p ^= hl32_to_64 (salt_bufs[salt_pos].salt_buf[3], salt_bufs[salt_pos].salt_buf[2]);
 232   v2p ^= hl32_to_64 (salt_bufs[salt_pos].salt_buf[1], salt_bufs[salt_pos].salt_buf[0]);
 233   v3p ^= hl32_to_64 (salt_bufs[salt_pos].salt_buf[3], salt_bufs[salt_pos].salt_buf[2]);
 234
 235   /**
 236    * digest
 237    */
 238
 239   const u32 search[4] =
 240   {
 241     digests_buf[digests_offset].digest_buf[DGST_R0],
 242     digests_buf[digests_offset].digest_buf[DGST_R1],
 243     digests_buf[digests_offset].digest_buf[DGST_R2],
 244     digests_buf[digests_offset].digest_buf[DGST_R3]
 245   };
 246
 247   /**
 248    * loop
 249    */
 250
 251   for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
 252   {
 253     u32x w[16];
 254
 255     w[ 0] = pw_buf0[0];
 256     w[ 1] = pw_buf0[1];
 257     w[ 2] = pw_buf0[2];
 258     w[ 3] = pw_buf0[3];
 259     w[ 4] = pw_buf1[0];
 260     w[ 5] = pw_buf1[1];
 261     w[ 6] = pw_buf1[2];
 262     w[ 7] = pw_buf1[3];
 263     w[ 8] = 0;
 264     w[ 9] = 0;
 265     w[10] = 0;
 266     w[11] = 0;
 267     w[12] = 0;
 268     w[13] = 0;
 269     w[14] = 0;
 270     w[15] = 0;
 271
 272     const u32x out_len = apply_rules (c_rules[il_pos].cmds, &w[0], &w[4], pw_len);
 273
 274     u64 *w_ptr = (u64 *) w;
 275
 276     w_ptr[out_len / 8] |= (u64) out_len << 56;
 277
 278     u64x v0 = v0p;
 279     u64x v1 = v1p;
 280     u64x v2 = v2p;
 281     u64x v3 = v3p;
 282
 283     int i;
 284     int j;
 285
 286     for (i = 0, j = 0; i <= pw_len; i += 8, j += 2)
 287     {
 288       u64x m = hl32_to_64 (w[j + 1], w[j + 0]);
 289
 290       v3 ^= m;
 291
 292       SIPROUND (v0, v1, v2, v3);
 293       SIPROUND (v0, v1, v2, v3);
 294
 295       v0 ^= m;
 296     }
 297
 298     v2 ^= 0xff;
 299
 300     SIPROUND (v0, v1, v2, v3);
 301     SIPROUND (v0, v1, v2, v3);
 302     SIPROUND (v0, v1, v2, v3);
 303     SIPROUND (v0, v1, v2, v3);
 304
 305     const u64x v = v0 ^ v1 ^ v2 ^ v3;
 306
 307     const u32x a = l32_from_64 (v);
 308     const u32x b = h32_from_64 (v);
 309
 310     const u32x r0 = a;
 311     const u32x r1 = b;
 312     const u32x r2 = 0;
 313     const u32x r3 = 0;
 314
 315     #include VECT_COMPARE_S
 316   }
 317 }
 318
 319 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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 320 {
 321 }
 322
 323 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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 324 {
 325 }