nv/markov_be_v4.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define CHARSIZ 256
   7
   8 #define VECT_SIZE4
   9
  10 #include "types_nv.c"
  11
  12 __device__ static void generate_pw (u32 pw_buf[16], cs_t *root_css_buf, cs_t *markov_css_buf, const u32 pw_l_len, const u32 pw_r_len, const u32 mask80, const u32 bits14, const u32 bits15, u64 val)
  13 {
  14   pw_buf[ 0] = 0;
  15   pw_buf[ 1] = 0;
  16   pw_buf[ 2] = 0;
  17   pw_buf[ 3] = 0;
  18   pw_buf[ 4] = 0;
  19   pw_buf[ 5] = 0;
  20   pw_buf[ 6] = 0;
  21   pw_buf[ 7] = 0;
  22   pw_buf[ 8] = 0;
  23   pw_buf[ 9] = 0;
  24   pw_buf[10] = 0;
  25   pw_buf[11] = 0;
  26   pw_buf[12] = 0;
  27   pw_buf[13] = 0;
  28   pw_buf[14] = 0;
  29   pw_buf[15] = 0;
  30
  31   cs_t *cs = &root_css_buf[pw_r_len];
  32
  33   u32 i;
  34   u32 j;
  35
  36   for (i = 0, j = pw_r_len; i < pw_l_len; i++, j++)
  37   {
  38     const u32 len = cs->cs_len;
  39
  40     const u64 next = val / len;
  41     const u64 pos  = val % len;
  42
  43     val = next;
  44
  45     const u32 key = cs->cs_buf[pos];
  46
  47     const u32 jd4 = j / 4;
  48     const u32 jm4 = j % 4;
  49
  50     pw_buf[jd4] |= key << ((3 - jm4) * 8);
  51
  52     cs = &markov_css_buf[(j * CHARSIZ) + key];
  53   }
  54
  55   const u32 jd4 = j / 4;
  56   const u32 jm4 = j % 4;
  57
  58   pw_buf[jd4] |= (0xff << ((3 - jm4) * 8)) & mask80;
  59
  60   if (bits14) pw_buf[14] = (pw_l_len + pw_r_len) * 8;
  61   if (bits15) pw_buf[15] = (pw_l_len + pw_r_len) * 8;
  62 }
  63
  64 extern "C" __global__ void __launch_bounds__ (256, 1) l_markov (pw_t *pws_buf_l, cs_t *root_css_buf, cs_t *markov_css_buf, const u64 off, const u32 pw_l_len, const u32 pw_r_len, const u32 mask80, const u32 bits14, const u32 bits15, const u32 gid_max)
  65 {
  66   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
  67
  68   if (gid >= gid_max) return;
  69
  70   const u32 gid4 = gid * 4;
  71
  72   u32 pw_buf0[16];
  73   u32 pw_buf1[16];
  74   u32 pw_buf2[16];
  75   u32 pw_buf3[16];
  76
  77   generate_pw (pw_buf0, root_css_buf, markov_css_buf, pw_l_len, pw_r_len, mask80, bits14, bits15, off + gid4 + 0);
  78   generate_pw (pw_buf1, root_css_buf, markov_css_buf, pw_l_len, pw_r_len, mask80, bits14, bits15, off + gid4 + 1);
  79   generate_pw (pw_buf2, root_css_buf, markov_css_buf, pw_l_len, pw_r_len, mask80, bits14, bits15, off + gid4 + 2);
  80   generate_pw (pw_buf3, root_css_buf, markov_css_buf, pw_l_len, pw_r_len, mask80, bits14, bits15, off + gid4 + 3);
  81
  82   #pragma unroll 16
  83   for (int i = 0; i < 16; i++)
  84   {
  85     pws_buf_l[gid].i[i].x = pw_buf0[i];
  86     pws_buf_l[gid].i[i].y = pw_buf1[i];
  87     pws_buf_l[gid].i[i].z = pw_buf2[i];
  88     pws_buf_l[gid].i[i].w = pw_buf3[i];
  89   }
  90
  91   pws_buf_l[gid].pw_len = pw_l_len + pw_r_len;
  92 }
  93
  94 extern "C" __global__ void __launch_bounds__ (256, 1) r_markov (bf_t *pws_buf_r, cs_t *root_css_buf, cs_t *markov_css_buf, const u64 off, const u32 pw_r_len, const u32 mask80, const u32 bits14, const u32 bits15, const u32 gid_max)
  95 {
  96   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
  97
  98   if (gid >= gid_max) return;
  99
 100   const u32 gid4 = gid * 4;
 101
 102   u32 pw_buf[16];
 103
 104   generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_r_len, 0, 0, 0, 0, off + gid4 + 0);
 105
 106   pws_buf_r[gid4 + 0].i = pw_buf[0];
 107
 108   generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_r_len, 0, 0, 0, 0, off + gid4 + 1);
 109
 110   pws_buf_r[gid4 + 1].i = pw_buf[0];
 111
 112   generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_r_len, 0, 0, 0, 0, off + gid4 + 2);
 113
 114   pws_buf_r[gid4 + 2].i = pw_buf[0];
 115
 116   generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_r_len, 0, 0, 0, 0, off + gid4 + 3);
 117
 118   pws_buf_r[gid4 + 3].i = pw_buf[0];
 119 }
 120
 121 extern "C" __global__ void __launch_bounds__ (256, 1) C_markov (comb_t *pws_buf, cs_t *root_css_buf, cs_t *markov_css_buf, const u64 off, const u32 pw_len, const u32 mask80, const u32 bits14, const u32 bits15, const u32 gid_max)
 122 {
 123   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 124
 125   if (gid >= gid_max) return;
 126
 127   const u32 gid4 = gid * 4;
 128
 129   u32 pw_buf[16];
 130
 131   generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_len, 0, mask80, bits14, bits15, off + gid4 + 0);
 132
 133   pws_buf[gid4 + 0].i[ 0] = pw_buf[ 0];
 134   pws_buf[gid4 + 0].i[ 1] = pw_buf[ 1];
 135   pws_buf[gid4 + 0].i[ 2] = pw_buf[ 2];
 136   pws_buf[gid4 + 0].i[ 3] = pw_buf[ 3];
 137   pws_buf[gid4 + 0].i[ 4] = pw_buf[ 4];
 138   pws_buf[gid4 + 0].i[ 5] = pw_buf[ 5];
 139   pws_buf[gid4 + 0].i[ 6] = pw_buf[ 6];
 140   pws_buf[gid4 + 0].i[ 7] = pw_buf[ 7];
 141
 142   pws_buf[gid4 + 0].pw_len = pw_len;
 143
 144   generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_len, 0, mask80, bits14, bits15, off + gid4 + 1);
 145
 146   pws_buf[gid4 + 1].i[ 0] = pw_buf[ 0];
 147   pws_buf[gid4 + 1].i[ 1] = pw_buf[ 1];
 148   pws_buf[gid4 + 1].i[ 2] = pw_buf[ 2];
 149   pws_buf[gid4 + 1].i[ 3] = pw_buf[ 3];
 150   pws_buf[gid4 + 1].i[ 4] = pw_buf[ 4];
 151   pws_buf[gid4 + 1].i[ 5] = pw_buf[ 5];
 152   pws_buf[gid4 + 1].i[ 6] = pw_buf[ 6];
 153   pws_buf[gid4 + 1].i[ 7] = pw_buf[ 7];
 154
 155   pws_buf[gid4 + 1].pw_len = pw_len;
 156
 157   generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_len, 0, mask80, bits14, bits15, off + gid4 + 2);
 158
 159   pws_buf[gid4 + 2].i[ 0] = pw_buf[ 0];
 160   pws_buf[gid4 + 2].i[ 1] = pw_buf[ 1];
 161   pws_buf[gid4 + 2].i[ 2] = pw_buf[ 2];
 162   pws_buf[gid4 + 2].i[ 3] = pw_buf[ 3];
 163   pws_buf[gid4 + 2].i[ 4] = pw_buf[ 4];
 164   pws_buf[gid4 + 2].i[ 5] = pw_buf[ 5];
 165   pws_buf[gid4 + 2].i[ 6] = pw_buf[ 6];
 166   pws_buf[gid4 + 2].i[ 7] = pw_buf[ 7];
 167
 168   pws_buf[gid4 + 2].pw_len = pw_len;
 169
 170   generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_len, 0, mask80, bits14, bits15, off + gid4 + 3);
 171
 172   pws_buf[gid4 + 3].i[ 0] = pw_buf[ 0];
 173   pws_buf[gid4 + 3].i[ 1] = pw_buf[ 1];
 174   pws_buf[gid4 + 3].i[ 2] = pw_buf[ 2];
 175   pws_buf[gid4 + 3].i[ 3] = pw_buf[ 3];
 176   pws_buf[gid4 + 3].i[ 4] = pw_buf[ 4];
 177   pws_buf[gid4 + 3].i[ 5] = pw_buf[ 5];
 178   pws_buf[gid4 + 3].i[ 6] = pw_buf[ 6];
 179   pws_buf[gid4 + 3].i[ 7] = pw_buf[ 7];
 180
 181   pws_buf[gid4 + 3].pw_len = pw_len;
 182 }