nv/amp_a3_v2.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define VECT_SIZE2
   7
   8 #include "include/constants.h"
   9 #include "types_nv.c"
  10
  11 __device__ __constant__ bf_t c_bfs[1024];
  12
  13 extern "C" __global__ void __launch_bounds__ (256, 1) amp (pw_t *pws, pw_t *pws_amp, gpu_rule_t *rules_buf, comb_t *combs_buf, bf_t *bfs_buf, const u32 combs_mode, const u32 gid_max)
  14 {
  15   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
  16
  17   if (gid >= gid_max) return;
  18
  19   const u32 pw_len = pws[gid].pw_len;
  20
  21   u32x w0[4];
  22   u32x w1[4];
  23   u32x w2[4];
  24   u32x w3[4];
  25
  26   w0[0] = pws[gid].i[ 0];
  27   w0[1] = pws[gid].i[ 1];
  28   w0[2] = pws[gid].i[ 2];
  29   w0[3] = pws[gid].i[ 3];
  30   w1[0] = pws[gid].i[ 4];
  31   w1[1] = pws[gid].i[ 5];
  32   w1[2] = pws[gid].i[ 6];
  33   w1[3] = pws[gid].i[ 7];
  34   w2[0] = pws[gid].i[ 8];
  35   w2[1] = pws[gid].i[ 9];
  36   w2[2] = pws[gid].i[10];
  37   w2[3] = pws[gid].i[11];
  38   w3[0] = pws[gid].i[12];
  39   w3[1] = pws[gid].i[13];
  40   w3[2] = pws[gid].i[14];
  41   w3[3] = pws[gid].i[15];
  42
  43   const u32 w0r = c_bfs[0].i;
  44
  45   pws_amp[gid].i[ 0] = w0[0] | w0r;
  46   pws_amp[gid].i[ 1] = w0[1];
  47   pws_amp[gid].i[ 2] = w0[2];
  48   pws_amp[gid].i[ 3] = w0[3];
  49   pws_amp[gid].i[ 4] = w1[0];
  50   pws_amp[gid].i[ 5] = w1[1];
  51   pws_amp[gid].i[ 6] = w1[2];
  52   pws_amp[gid].i[ 7] = w1[3];
  53   pws_amp[gid].i[ 8] = w2[0];
  54   pws_amp[gid].i[ 9] = w2[1];
  55   pws_amp[gid].i[10] = w2[2];
  56   pws_amp[gid].i[11] = w2[3];
  57   pws_amp[gid].i[12] = w3[0];
  58   pws_amp[gid].i[13] = w3[1];
  59   pws_amp[gid].i[14] = w3[2];
  60   pws_amp[gid].i[15] = w3[3];
  61
  62   pws_amp[gid].pw_len = pw_len;
  63 }