nv/m06000_a1.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _RIPEMD160_
   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 __device__ static void ripemd160_transform (const u32x w[16], u32x dgst[5])
  39 {
  40   u32x a1 = dgst[0];
  41   u32x b1 = dgst[1];
  42   u32x c1 = dgst[2];
  43   u32x d1 = dgst[3];
  44   u32x e1 = dgst[4];
  45
  46   RIPEMD160_STEP (RIPEMD160_F , a1, b1, c1, d1, e1, w[ 0], RIPEMD160C00, RIPEMD160S00);
  47   RIPEMD160_STEP (RIPEMD160_F , e1, a1, b1, c1, d1, w[ 1], RIPEMD160C00, RIPEMD160S01);
  48   RIPEMD160_STEP (RIPEMD160_F , d1, e1, a1, b1, c1, w[ 2], RIPEMD160C00, RIPEMD160S02);
  49   RIPEMD160_STEP (RIPEMD160_F , c1, d1, e1, a1, b1, w[ 3], RIPEMD160C00, RIPEMD160S03);
  50   RIPEMD160_STEP (RIPEMD160_F , b1, c1, d1, e1, a1, w[ 4], RIPEMD160C00, RIPEMD160S04);
  51   RIPEMD160_STEP (RIPEMD160_F , a1, b1, c1, d1, e1, w[ 5], RIPEMD160C00, RIPEMD160S05);
  52   RIPEMD160_STEP (RIPEMD160_F , e1, a1, b1, c1, d1, w[ 6], RIPEMD160C00, RIPEMD160S06);
  53   RIPEMD160_STEP (RIPEMD160_F , d1, e1, a1, b1, c1, w[ 7], RIPEMD160C00, RIPEMD160S07);
  54   RIPEMD160_STEP (RIPEMD160_F , c1, d1, e1, a1, b1, w[ 8], RIPEMD160C00, RIPEMD160S08);
  55   RIPEMD160_STEP (RIPEMD160_F , b1, c1, d1, e1, a1, w[ 9], RIPEMD160C00, RIPEMD160S09);
  56   RIPEMD160_STEP (RIPEMD160_F , a1, b1, c1, d1, e1, w[10], RIPEMD160C00, RIPEMD160S0A);
  57   RIPEMD160_STEP (RIPEMD160_F , e1, a1, b1, c1, d1, w[11], RIPEMD160C00, RIPEMD160S0B);
  58   RIPEMD160_STEP (RIPEMD160_F , d1, e1, a1, b1, c1, w[12], RIPEMD160C00, RIPEMD160S0C);
  59   RIPEMD160_STEP (RIPEMD160_F , c1, d1, e1, a1, b1, w[13], RIPEMD160C00, RIPEMD160S0D);
  60   RIPEMD160_STEP (RIPEMD160_F , b1, c1, d1, e1, a1, w[14], RIPEMD160C00, RIPEMD160S0E);
  61   RIPEMD160_STEP (RIPEMD160_F , a1, b1, c1, d1, e1, w[15], RIPEMD160C00, RIPEMD160S0F);
  62
  63   RIPEMD160_STEP (RIPEMD160_Go, e1, a1, b1, c1, d1, w[ 7], RIPEMD160C10, RIPEMD160S10);
  64   RIPEMD160_STEP (RIPEMD160_Go, d1, e1, a1, b1, c1, w[ 4], RIPEMD160C10, RIPEMD160S11);
  65   RIPEMD160_STEP (RIPEMD160_Go, c1, d1, e1, a1, b1, w[13], RIPEMD160C10, RIPEMD160S12);
  66   RIPEMD160_STEP (RIPEMD160_Go, b1, c1, d1, e1, a1, w[ 1], RIPEMD160C10, RIPEMD160S13);
  67   RIPEMD160_STEP (RIPEMD160_Go, a1, b1, c1, d1, e1, w[10], RIPEMD160C10, RIPEMD160S14);
  68   RIPEMD160_STEP (RIPEMD160_Go, e1, a1, b1, c1, d1, w[ 6], RIPEMD160C10, RIPEMD160S15);
  69   RIPEMD160_STEP (RIPEMD160_Go, d1, e1, a1, b1, c1, w[15], RIPEMD160C10, RIPEMD160S16);
  70   RIPEMD160_STEP (RIPEMD160_Go, c1, d1, e1, a1, b1, w[ 3], RIPEMD160C10, RIPEMD160S17);
  71   RIPEMD160_STEP (RIPEMD160_Go, b1, c1, d1, e1, a1, w[12], RIPEMD160C10, RIPEMD160S18);
  72   RIPEMD160_STEP (RIPEMD160_Go, a1, b1, c1, d1, e1, w[ 0], RIPEMD160C10, RIPEMD160S19);
  73   RIPEMD160_STEP (RIPEMD160_Go, e1, a1, b1, c1, d1, w[ 9], RIPEMD160C10, RIPEMD160S1A);
  74   RIPEMD160_STEP (RIPEMD160_Go, d1, e1, a1, b1, c1, w[ 5], RIPEMD160C10, RIPEMD160S1B);
  75   RIPEMD160_STEP (RIPEMD160_Go, c1, d1, e1, a1, b1, w[ 2], RIPEMD160C10, RIPEMD160S1C);
  76   RIPEMD160_STEP (RIPEMD160_Go, b1, c1, d1, e1, a1, w[14], RIPEMD160C10, RIPEMD160S1D);
  77   RIPEMD160_STEP (RIPEMD160_Go, a1, b1, c1, d1, e1, w[11], RIPEMD160C10, RIPEMD160S1E);
  78   RIPEMD160_STEP (RIPEMD160_Go, e1, a1, b1, c1, d1, w[ 8], RIPEMD160C10, RIPEMD160S1F);
  79
  80   RIPEMD160_STEP (RIPEMD160_H , d1, e1, a1, b1, c1, w[ 3], RIPEMD160C20, RIPEMD160S20);
  81   RIPEMD160_STEP (RIPEMD160_H , c1, d1, e1, a1, b1, w[10], RIPEMD160C20, RIPEMD160S21);
  82   RIPEMD160_STEP (RIPEMD160_H , b1, c1, d1, e1, a1, w[14], RIPEMD160C20, RIPEMD160S22);
  83   RIPEMD160_STEP (RIPEMD160_H , a1, b1, c1, d1, e1, w[ 4], RIPEMD160C20, RIPEMD160S23);
  84   RIPEMD160_STEP (RIPEMD160_H , e1, a1, b1, c1, d1, w[ 9], RIPEMD160C20, RIPEMD160S24);
  85   RIPEMD160_STEP (RIPEMD160_H , d1, e1, a1, b1, c1, w[15], RIPEMD160C20, RIPEMD160S25);
  86   RIPEMD160_STEP (RIPEMD160_H , c1, d1, e1, a1, b1, w[ 8], RIPEMD160C20, RIPEMD160S26);
  87   RIPEMD160_STEP (RIPEMD160_H , b1, c1, d1, e1, a1, w[ 1], RIPEMD160C20, RIPEMD160S27);
  88   RIPEMD160_STEP (RIPEMD160_H , a1, b1, c1, d1, e1, w[ 2], RIPEMD160C20, RIPEMD160S28);
  89   RIPEMD160_STEP (RIPEMD160_H , e1, a1, b1, c1, d1, w[ 7], RIPEMD160C20, RIPEMD160S29);
  90   RIPEMD160_STEP (RIPEMD160_H , d1, e1, a1, b1, c1, w[ 0], RIPEMD160C20, RIPEMD160S2A);
  91   RIPEMD160_STEP (RIPEMD160_H , c1, d1, e1, a1, b1, w[ 6], RIPEMD160C20, RIPEMD160S2B);
  92   RIPEMD160_STEP (RIPEMD160_H , b1, c1, d1, e1, a1, w[13], RIPEMD160C20, RIPEMD160S2C);
  93   RIPEMD160_STEP (RIPEMD160_H , a1, b1, c1, d1, e1, w[11], RIPEMD160C20, RIPEMD160S2D);
  94   RIPEMD160_STEP (RIPEMD160_H , e1, a1, b1, c1, d1, w[ 5], RIPEMD160C20, RIPEMD160S2E);
  95   RIPEMD160_STEP (RIPEMD160_H , d1, e1, a1, b1, c1, w[12], RIPEMD160C20, RIPEMD160S2F);
  96
  97   RIPEMD160_STEP (RIPEMD160_Io, c1, d1, e1, a1, b1, w[ 1], RIPEMD160C30, RIPEMD160S30);
  98   RIPEMD160_STEP (RIPEMD160_Io, b1, c1, d1, e1, a1, w[ 9], RIPEMD160C30, RIPEMD160S31);
  99   RIPEMD160_STEP (RIPEMD160_Io, a1, b1, c1, d1, e1, w[11], RIPEMD160C30, RIPEMD160S32);
 100   RIPEMD160_STEP (RIPEMD160_Io, e1, a1, b1, c1, d1, w[10], RIPEMD160C30, RIPEMD160S33);
 101   RIPEMD160_STEP (RIPEMD160_Io, d1, e1, a1, b1, c1, w[ 0], RIPEMD160C30, RIPEMD160S34);
 102   RIPEMD160_STEP (RIPEMD160_Io, c1, d1, e1, a1, b1, w[ 8], RIPEMD160C30, RIPEMD160S35);
 103   RIPEMD160_STEP (RIPEMD160_Io, b1, c1, d1, e1, a1, w[12], RIPEMD160C30, RIPEMD160S36);
 104   RIPEMD160_STEP (RIPEMD160_Io, a1, b1, c1, d1, e1, w[ 4], RIPEMD160C30, RIPEMD160S37);
 105   RIPEMD160_STEP (RIPEMD160_Io, e1, a1, b1, c1, d1, w[13], RIPEMD160C30, RIPEMD160S38);
 106   RIPEMD160_STEP (RIPEMD160_Io, d1, e1, a1, b1, c1, w[ 3], RIPEMD160C30, RIPEMD160S39);
 107   RIPEMD160_STEP (RIPEMD160_Io, c1, d1, e1, a1, b1, w[ 7], RIPEMD160C30, RIPEMD160S3A);
 108   RIPEMD160_STEP (RIPEMD160_Io, b1, c1, d1, e1, a1, w[15], RIPEMD160C30, RIPEMD160S3B);
 109   RIPEMD160_STEP (RIPEMD160_Io, a1, b1, c1, d1, e1, w[14], RIPEMD160C30, RIPEMD160S3C);
 110   RIPEMD160_STEP (RIPEMD160_Io, e1, a1, b1, c1, d1, w[ 5], RIPEMD160C30, RIPEMD160S3D);
 111   RIPEMD160_STEP (RIPEMD160_Io, d1, e1, a1, b1, c1, w[ 6], RIPEMD160C30, RIPEMD160S3E);
 112   RIPEMD160_STEP (RIPEMD160_Io, c1, d1, e1, a1, b1, w[ 2], RIPEMD160C30, RIPEMD160S3F);
 113
 114   RIPEMD160_STEP (RIPEMD160_J , b1, c1, d1, e1, a1, w[ 4], RIPEMD160C40, RIPEMD160S40);
 115   RIPEMD160_STEP (RIPEMD160_J , a1, b1, c1, d1, e1, w[ 0], RIPEMD160C40, RIPEMD160S41);
 116   RIPEMD160_STEP (RIPEMD160_J , e1, a1, b1, c1, d1, w[ 5], RIPEMD160C40, RIPEMD160S42);
 117   RIPEMD160_STEP (RIPEMD160_J , d1, e1, a1, b1, c1, w[ 9], RIPEMD160C40, RIPEMD160S43);
 118   RIPEMD160_STEP (RIPEMD160_J , c1, d1, e1, a1, b1, w[ 7], RIPEMD160C40, RIPEMD160S44);
 119   RIPEMD160_STEP (RIPEMD160_J , b1, c1, d1, e1, a1, w[12], RIPEMD160C40, RIPEMD160S45);
 120   RIPEMD160_STEP (RIPEMD160_J , a1, b1, c1, d1, e1, w[ 2], RIPEMD160C40, RIPEMD160S46);
 121   RIPEMD160_STEP (RIPEMD160_J , e1, a1, b1, c1, d1, w[10], RIPEMD160C40, RIPEMD160S47);
 122   RIPEMD160_STEP (RIPEMD160_J , d1, e1, a1, b1, c1, w[14], RIPEMD160C40, RIPEMD160S48);
 123   RIPEMD160_STEP (RIPEMD160_J , c1, d1, e1, a1, b1, w[ 1], RIPEMD160C40, RIPEMD160S49);
 124   RIPEMD160_STEP (RIPEMD160_J , b1, c1, d1, e1, a1, w[ 3], RIPEMD160C40, RIPEMD160S4A);
 125   RIPEMD160_STEP (RIPEMD160_J , a1, b1, c1, d1, e1, w[ 8], RIPEMD160C40, RIPEMD160S4B);
 126   RIPEMD160_STEP (RIPEMD160_J , e1, a1, b1, c1, d1, w[11], RIPEMD160C40, RIPEMD160S4C);
 127   RIPEMD160_STEP (RIPEMD160_J , d1, e1, a1, b1, c1, w[ 6], RIPEMD160C40, RIPEMD160S4D);
 128   RIPEMD160_STEP (RIPEMD160_J , c1, d1, e1, a1, b1, w[15], RIPEMD160C40, RIPEMD160S4E);
 129   RIPEMD160_STEP (RIPEMD160_J , b1, c1, d1, e1, a1, w[13], RIPEMD160C40, RIPEMD160S4F);
 130
 131   u32x a2 = dgst[0];
 132   u32x b2 = dgst[1];
 133   u32x c2 = dgst[2];
 134   u32x d2 = dgst[3];
 135   u32x e2 = dgst[4];
 136
 137   //RIPEMD160_STEP_WORKAROUND_BUG (RIPEMD160_J , a2, b2, c2, d2, e2, w[ 5], RIPEMD160C50, RIPEMD160S50);
 138   RIPEMD160_STEP (RIPEMD160_J , a2, b2, c2, d2, e2, w[ 5], RIPEMD160C50, RIPEMD160S50);
 139   RIPEMD160_STEP (RIPEMD160_J , e2, a2, b2, c2, d2, w[14], RIPEMD160C50, RIPEMD160S51);
 140   RIPEMD160_STEP (RIPEMD160_J , d2, e2, a2, b2, c2, w[ 7], RIPEMD160C50, RIPEMD160S52);
 141   RIPEMD160_STEP (RIPEMD160_J , c2, d2, e2, a2, b2, w[ 0], RIPEMD160C50, RIPEMD160S53);
 142   RIPEMD160_STEP (RIPEMD160_J , b2, c2, d2, e2, a2, w[ 9], RIPEMD160C50, RIPEMD160S54);
 143   RIPEMD160_STEP (RIPEMD160_J , a2, b2, c2, d2, e2, w[ 2], RIPEMD160C50, RIPEMD160S55);
 144   RIPEMD160_STEP (RIPEMD160_J , e2, a2, b2, c2, d2, w[11], RIPEMD160C50, RIPEMD160S56);
 145   RIPEMD160_STEP (RIPEMD160_J , d2, e2, a2, b2, c2, w[ 4], RIPEMD160C50, RIPEMD160S57);
 146   RIPEMD160_STEP (RIPEMD160_J , c2, d2, e2, a2, b2, w[13], RIPEMD160C50, RIPEMD160S58);
 147   RIPEMD160_STEP (RIPEMD160_J , b2, c2, d2, e2, a2, w[ 6], RIPEMD160C50, RIPEMD160S59);
 148   RIPEMD160_STEP (RIPEMD160_J , a2, b2, c2, d2, e2, w[15], RIPEMD160C50, RIPEMD160S5A);
 149   RIPEMD160_STEP (RIPEMD160_J , e2, a2, b2, c2, d2, w[ 8], RIPEMD160C50, RIPEMD160S5B);
 150   RIPEMD160_STEP (RIPEMD160_J , d2, e2, a2, b2, c2, w[ 1], RIPEMD160C50, RIPEMD160S5C);
 151   RIPEMD160_STEP (RIPEMD160_J , c2, d2, e2, a2, b2, w[10], RIPEMD160C50, RIPEMD160S5D);
 152   RIPEMD160_STEP (RIPEMD160_J , b2, c2, d2, e2, a2, w[ 3], RIPEMD160C50, RIPEMD160S5E);
 153   RIPEMD160_STEP (RIPEMD160_J , a2, b2, c2, d2, e2, w[12], RIPEMD160C50, RIPEMD160S5F);
 154
 155   RIPEMD160_STEP (RIPEMD160_Io, e2, a2, b2, c2, d2, w[ 6], RIPEMD160C60, RIPEMD160S60);
 156   RIPEMD160_STEP (RIPEMD160_Io, d2, e2, a2, b2, c2, w[11], RIPEMD160C60, RIPEMD160S61);
 157   RIPEMD160_STEP (RIPEMD160_Io, c2, d2, e2, a2, b2, w[ 3], RIPEMD160C60, RIPEMD160S62);
 158   RIPEMD160_STEP (RIPEMD160_Io, b2, c2, d2, e2, a2, w[ 7], RIPEMD160C60, RIPEMD160S63);
 159   RIPEMD160_STEP (RIPEMD160_Io, a2, b2, c2, d2, e2, w[ 0], RIPEMD160C60, RIPEMD160S64);
 160   RIPEMD160_STEP (RIPEMD160_Io, e2, a2, b2, c2, d2, w[13], RIPEMD160C60, RIPEMD160S65);
 161   RIPEMD160_STEP (RIPEMD160_Io, d2, e2, a2, b2, c2, w[ 5], RIPEMD160C60, RIPEMD160S66);
 162   RIPEMD160_STEP (RIPEMD160_Io, c2, d2, e2, a2, b2, w[10], RIPEMD160C60, RIPEMD160S67);
 163   RIPEMD160_STEP (RIPEMD160_Io, b2, c2, d2, e2, a2, w[14], RIPEMD160C60, RIPEMD160S68);
 164   RIPEMD160_STEP (RIPEMD160_Io, a2, b2, c2, d2, e2, w[15], RIPEMD160C60, RIPEMD160S69);
 165   RIPEMD160_STEP (RIPEMD160_Io, e2, a2, b2, c2, d2, w[ 8], RIPEMD160C60, RIPEMD160S6A);
 166   RIPEMD160_STEP (RIPEMD160_Io, d2, e2, a2, b2, c2, w[12], RIPEMD160C60, RIPEMD160S6B);
 167   RIPEMD160_STEP (RIPEMD160_Io, c2, d2, e2, a2, b2, w[ 4], RIPEMD160C60, RIPEMD160S6C);
 168   RIPEMD160_STEP (RIPEMD160_Io, b2, c2, d2, e2, a2, w[ 9], RIPEMD160C60, RIPEMD160S6D);
 169   RIPEMD160_STEP (RIPEMD160_Io, a2, b2, c2, d2, e2, w[ 1], RIPEMD160C60, RIPEMD160S6E);
 170   RIPEMD160_STEP (RIPEMD160_Io, e2, a2, b2, c2, d2, w[ 2], RIPEMD160C60, RIPEMD160S6F);
 171
 172   RIPEMD160_STEP (RIPEMD160_H , d2, e2, a2, b2, c2, w[15], RIPEMD160C70, RIPEMD160S70);
 173   RIPEMD160_STEP (RIPEMD160_H , c2, d2, e2, a2, b2, w[ 5], RIPEMD160C70, RIPEMD160S71);
 174   RIPEMD160_STEP (RIPEMD160_H , b2, c2, d2, e2, a2, w[ 1], RIPEMD160C70, RIPEMD160S72);
 175   RIPEMD160_STEP (RIPEMD160_H , a2, b2, c2, d2, e2, w[ 3], RIPEMD160C70, RIPEMD160S73);
 176   RIPEMD160_STEP (RIPEMD160_H , e2, a2, b2, c2, d2, w[ 7], RIPEMD160C70, RIPEMD160S74);
 177   RIPEMD160_STEP (RIPEMD160_H , d2, e2, a2, b2, c2, w[14], RIPEMD160C70, RIPEMD160S75);
 178   RIPEMD160_STEP (RIPEMD160_H , c2, d2, e2, a2, b2, w[ 6], RIPEMD160C70, RIPEMD160S76);
 179   RIPEMD160_STEP (RIPEMD160_H , b2, c2, d2, e2, a2, w[ 9], RIPEMD160C70, RIPEMD160S77);
 180   RIPEMD160_STEP (RIPEMD160_H , a2, b2, c2, d2, e2, w[11], RIPEMD160C70, RIPEMD160S78);
 181   RIPEMD160_STEP (RIPEMD160_H , e2, a2, b2, c2, d2, w[ 8], RIPEMD160C70, RIPEMD160S79);
 182   RIPEMD160_STEP (RIPEMD160_H , d2, e2, a2, b2, c2, w[12], RIPEMD160C70, RIPEMD160S7A);
 183   RIPEMD160_STEP (RIPEMD160_H , c2, d2, e2, a2, b2, w[ 2], RIPEMD160C70, RIPEMD160S7B);
 184   RIPEMD160_STEP (RIPEMD160_H , b2, c2, d2, e2, a2, w[10], RIPEMD160C70, RIPEMD160S7C);
 185   RIPEMD160_STEP (RIPEMD160_H , a2, b2, c2, d2, e2, w[ 0], RIPEMD160C70, RIPEMD160S7D);
 186   RIPEMD160_STEP (RIPEMD160_H , e2, a2, b2, c2, d2, w[ 4], RIPEMD160C70, RIPEMD160S7E);
 187   RIPEMD160_STEP (RIPEMD160_H , d2, e2, a2, b2, c2, w[13], RIPEMD160C70, RIPEMD160S7F);
 188
 189   RIPEMD160_STEP (RIPEMD160_Go, c2, d2, e2, a2, b2, w[ 8], RIPEMD160C80, RIPEMD160S80);
 190   RIPEMD160_STEP (RIPEMD160_Go, b2, c2, d2, e2, a2, w[ 6], RIPEMD160C80, RIPEMD160S81);
 191   RIPEMD160_STEP (RIPEMD160_Go, a2, b2, c2, d2, e2, w[ 4], RIPEMD160C80, RIPEMD160S82);
 192   RIPEMD160_STEP (RIPEMD160_Go, e2, a2, b2, c2, d2, w[ 1], RIPEMD160C80, RIPEMD160S83);
 193   RIPEMD160_STEP (RIPEMD160_Go, d2, e2, a2, b2, c2, w[ 3], RIPEMD160C80, RIPEMD160S84);
 194   RIPEMD160_STEP (RIPEMD160_Go, c2, d2, e2, a2, b2, w[11], RIPEMD160C80, RIPEMD160S85);
 195   RIPEMD160_STEP (RIPEMD160_Go, b2, c2, d2, e2, a2, w[15], RIPEMD160C80, RIPEMD160S86);
 196   RIPEMD160_STEP (RIPEMD160_Go, a2, b2, c2, d2, e2, w[ 0], RIPEMD160C80, RIPEMD160S87);
 197   RIPEMD160_STEP (RIPEMD160_Go, e2, a2, b2, c2, d2, w[ 5], RIPEMD160C80, RIPEMD160S88);
 198   RIPEMD160_STEP (RIPEMD160_Go, d2, e2, a2, b2, c2, w[12], RIPEMD160C80, RIPEMD160S89);
 199   RIPEMD160_STEP (RIPEMD160_Go, c2, d2, e2, a2, b2, w[ 2], RIPEMD160C80, RIPEMD160S8A);
 200   RIPEMD160_STEP (RIPEMD160_Go, b2, c2, d2, e2, a2, w[13], RIPEMD160C80, RIPEMD160S8B);
 201   RIPEMD160_STEP (RIPEMD160_Go, a2, b2, c2, d2, e2, w[ 9], RIPEMD160C80, RIPEMD160S8C);
 202   RIPEMD160_STEP (RIPEMD160_Go, e2, a2, b2, c2, d2, w[ 7], RIPEMD160C80, RIPEMD160S8D);
 203   RIPEMD160_STEP (RIPEMD160_Go, d2, e2, a2, b2, c2, w[10], RIPEMD160C80, RIPEMD160S8E);
 204   RIPEMD160_STEP (RIPEMD160_Go, c2, d2, e2, a2, b2, w[14], RIPEMD160C80, RIPEMD160S8F);
 205
 206   RIPEMD160_STEP (RIPEMD160_F , b2, c2, d2, e2, a2, w[12], RIPEMD160C90, RIPEMD160S90);
 207   RIPEMD160_STEP (RIPEMD160_F , a2, b2, c2, d2, e2, w[15], RIPEMD160C90, RIPEMD160S91);
 208   RIPEMD160_STEP (RIPEMD160_F , e2, a2, b2, c2, d2, w[10], RIPEMD160C90, RIPEMD160S92);
 209   RIPEMD160_STEP (RIPEMD160_F , d2, e2, a2, b2, c2, w[ 4], RIPEMD160C90, RIPEMD160S93);
 210   RIPEMD160_STEP (RIPEMD160_F , c2, d2, e2, a2, b2, w[ 1], RIPEMD160C90, RIPEMD160S94);
 211   RIPEMD160_STEP (RIPEMD160_F , b2, c2, d2, e2, a2, w[ 5], RIPEMD160C90, RIPEMD160S95);
 212   RIPEMD160_STEP (RIPEMD160_F , a2, b2, c2, d2, e2, w[ 8], RIPEMD160C90, RIPEMD160S96);
 213   RIPEMD160_STEP (RIPEMD160_F , e2, a2, b2, c2, d2, w[ 7], RIPEMD160C90, RIPEMD160S97);
 214   RIPEMD160_STEP (RIPEMD160_F , d2, e2, a2, b2, c2, w[ 6], RIPEMD160C90, RIPEMD160S98);
 215   RIPEMD160_STEP (RIPEMD160_F , c2, d2, e2, a2, b2, w[ 2], RIPEMD160C90, RIPEMD160S99);
 216   RIPEMD160_STEP (RIPEMD160_F , b2, c2, d2, e2, a2, w[13], RIPEMD160C90, RIPEMD160S9A);
 217   RIPEMD160_STEP (RIPEMD160_F , a2, b2, c2, d2, e2, w[14], RIPEMD160C90, RIPEMD160S9B);
 218   RIPEMD160_STEP (RIPEMD160_F , e2, a2, b2, c2, d2, w[ 0], RIPEMD160C90, RIPEMD160S9C);
 219   RIPEMD160_STEP (RIPEMD160_F , d2, e2, a2, b2, c2, w[ 3], RIPEMD160C90, RIPEMD160S9D);
 220   RIPEMD160_STEP (RIPEMD160_F , c2, d2, e2, a2, b2, w[ 9], RIPEMD160C90, RIPEMD160S9E);
 221   RIPEMD160_STEP (RIPEMD160_F , b2, c2, d2, e2, a2, w[11], RIPEMD160C90, RIPEMD160S9F);
 222
 223   const u32x a = dgst[1] + c1 + d2;
 224   const u32x b = dgst[2] + d1 + e2;
 225   const u32x c = dgst[3] + e1 + a2;
 226   const u32x d = dgst[4] + a1 + b2;
 227   const u32x e = dgst[0] + b1 + c2;
 228
 229   dgst[0] = a;
 230   dgst[1] = b;
 231   dgst[2] = c;
 232   dgst[3] = d;
 233   dgst[4] = e;
 234 }
 235
 236 __device__ __constant__ comb_t c_combs[1024];
 237
 238 extern "C" __global__ void __launch_bounds__ (256, 1) m06000_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)
 239 {
 240   /**
 241    * modifier
 242    */
 243
 244   const u32 lid = threadIdx.x;
 245
 246   /**
 247    * base
 248    */
 249
 250   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 251
 252   if (gid >= gid_max) return;
 253
 254   u32x wordl0[4];
 255
 256   wordl0[0] = pws[gid].i[ 0];
 257   wordl0[1] = pws[gid].i[ 1];
 258   wordl0[2] = pws[gid].i[ 2];
 259   wordl0[3] = pws[gid].i[ 3];
 260
 261   u32x wordl1[4];
 262
 263   wordl1[0] = pws[gid].i[ 4];
 264   wordl1[1] = pws[gid].i[ 5];
 265   wordl1[2] = pws[gid].i[ 6];
 266   wordl1[3] = pws[gid].i[ 7];
 267
 268   u32x wordl2[4];
 269
 270   wordl2[0] = 0;
 271   wordl2[1] = 0;
 272   wordl2[2] = 0;
 273   wordl2[3] = 0;
 274
 275   u32x wordl3[4];
 276
 277   wordl3[0] = 0;
 278   wordl3[1] = 0;
 279   wordl3[2] = 0;
 280   wordl3[3] = 0;
 281
 282   const u32 pw_l_len = pws[gid].pw_len;
 283
 284   if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
 285   {
 286     append_0x80_2 (wordl0, wordl1, pw_l_len);
 287
 288     switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, c_combs[0].pw_len);
 289   }
 290
 291   /**
 292    * loop
 293    */
 294
 295   for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
 296   {
 297     const u32 pw_r_len = c_combs[il_pos].pw_len;
 298
 299     const u32 pw_len = pw_l_len + pw_r_len;
 300
 301     u32 wordr0[4];
 302
 303     wordr0[0] = c_combs[il_pos].i[0];
 304     wordr0[1] = c_combs[il_pos].i[1];
 305     wordr0[2] = c_combs[il_pos].i[2];
 306     wordr0[3] = c_combs[il_pos].i[3];
 307
 308     u32 wordr1[4];
 309
 310     wordr1[0] = c_combs[il_pos].i[4];
 311     wordr1[1] = c_combs[il_pos].i[5];
 312     wordr1[2] = c_combs[il_pos].i[6];
 313     wordr1[3] = c_combs[il_pos].i[7];
 314
 315     u32 wordr2[4];
 316
 317     wordr2[0] = 0;
 318     wordr2[1] = 0;
 319     wordr2[2] = 0;
 320     wordr2[3] = 0;
 321
 322     u32 wordr3[4];
 323
 324     wordr3[0] = 0;
 325     wordr3[1] = 0;
 326     wordr3[2] = 0;
 327     wordr3[3] = 0;
 328
 329     if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
 330     {
 331       switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
 332     }
 333
 334     u32x w0[4];
 335
 336     w0[0] = wordl0[0] | wordr0[0];
 337     w0[1] = wordl0[1] | wordr0[1];
 338     w0[2] = wordl0[2] | wordr0[2];
 339     w0[3] = wordl0[3] | wordr0[3];
 340
 341     u32x w1[4];
 342
 343     w1[0] = wordl1[0] | wordr1[0];
 344     w1[1] = wordl1[1] | wordr1[1];
 345     w1[2] = wordl1[2] | wordr1[2];
 346     w1[3] = wordl1[3] | wordr1[3];
 347
 348     u32x w2[4];
 349
 350     w2[0] = wordl2[0] | wordr2[0];
 351     w2[1] = wordl2[1] | wordr2[1];
 352     w2[2] = wordl2[2] | wordr2[2];
 353     w2[3] = wordl2[3] | wordr2[3];
 354
 355     u32x w3[4];
 356
 357     w3[0] = wordl3[0] | wordr3[0];
 358     w3[1] = wordl3[1] | wordr3[1];
 359     w3[2] = pw_len * 8;
 360     w3[3] = 0;
 361
 362     u32x wl[16];
 363
 364     wl[ 0] = w0[0];
 365     wl[ 1] = w0[1];
 366     wl[ 2] = w0[2];
 367     wl[ 3] = w0[3];
 368     wl[ 4] = w1[0];
 369     wl[ 5] = w1[1];
 370     wl[ 6] = w1[2];
 371     wl[ 7] = w1[3];
 372     wl[ 8] = 0;
 373     wl[ 9] = 0;
 374     wl[10] = 0;
 375     wl[11] = 0;
 376     wl[12] = 0;
 377     wl[13] = 0;
 378     wl[14] = pw_len * 8;
 379     wl[15] = 0;
 380
 381     u32x dgst[5];
 382
 383     dgst[0] = RIPEMD160M_A;
 384     dgst[1] = RIPEMD160M_B;
 385     dgst[2] = RIPEMD160M_C;
 386     dgst[3] = RIPEMD160M_D;
 387     dgst[4] = RIPEMD160M_E;
 388
 389     ripemd160_transform (wl, dgst);
 390
 391     const u32x r0 = dgst[0];
 392     const u32x r1 = dgst[1];
 393     const u32x r2 = dgst[2];
 394     const u32x r3 = dgst[3];
 395
 396     #include VECT_COMPARE_M
 397   }
 398 }
 399
 400 extern "C" __global__ void __launch_bounds__ (256, 1) m06000_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)
 401 {
 402 }
 403
 404 extern "C" __global__ void __launch_bounds__ (256, 1) m06000_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)
 405 {
 406 }
 407
 408 extern "C" __global__ void __launch_bounds__ (256, 1) m06000_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)
 409 {
 410   /**
 411    * modifier
 412    */
 413
 414   const u32 lid = threadIdx.x;
 415
 416   /**
 417    * base
 418    */
 419
 420   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 421
 422   if (gid >= gid_max) return;
 423
 424   u32x wordl0[4];
 425
 426   wordl0[0] = pws[gid].i[ 0];
 427   wordl0[1] = pws[gid].i[ 1];
 428   wordl0[2] = pws[gid].i[ 2];
 429   wordl0[3] = pws[gid].i[ 3];
 430
 431   u32x wordl1[4];
 432
 433   wordl1[0] = pws[gid].i[ 4];
 434   wordl1[1] = pws[gid].i[ 5];
 435   wordl1[2] = pws[gid].i[ 6];
 436   wordl1[3] = pws[gid].i[ 7];
 437
 438   u32x wordl2[4];
 439
 440   wordl2[0] = 0;
 441   wordl2[1] = 0;
 442   wordl2[2] = 0;
 443   wordl2[3] = 0;
 444
 445   u32x wordl3[4];
 446
 447   wordl3[0] = 0;
 448   wordl3[1] = 0;
 449   wordl3[2] = 0;
 450   wordl3[3] = 0;
 451
 452   const u32 pw_l_len = pws[gid].pw_len;
 453
 454   if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
 455   {
 456     append_0x80_2 (wordl0, wordl1, pw_l_len);
 457
 458     switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, c_combs[0].pw_len);
 459   }
 460
 461   /**
 462    * digest
 463    */
 464
 465   const u32 search[4] =
 466   {
 467     digests_buf[digests_offset].digest_buf[DGST_R0],
 468     digests_buf[digests_offset].digest_buf[DGST_R1],
 469     digests_buf[digests_offset].digest_buf[DGST_R2],
 470     digests_buf[digests_offset].digest_buf[DGST_R3]
 471   };
 472
 473   /**
 474    * loop
 475    */
 476
 477   for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
 478   {
 479     const u32 pw_r_len = c_combs[il_pos].pw_len;
 480
 481     const u32 pw_len = pw_l_len + pw_r_len;
 482
 483     u32 wordr0[4];
 484
 485     wordr0[0] = c_combs[il_pos].i[0];
 486     wordr0[1] = c_combs[il_pos].i[1];
 487     wordr0[2] = c_combs[il_pos].i[2];
 488     wordr0[3] = c_combs[il_pos].i[3];
 489
 490     u32 wordr1[4];
 491
 492     wordr1[0] = c_combs[il_pos].i[4];
 493     wordr1[1] = c_combs[il_pos].i[5];
 494     wordr1[2] = c_combs[il_pos].i[6];
 495     wordr1[3] = c_combs[il_pos].i[7];
 496
 497     u32 wordr2[4];
 498
 499     wordr2[0] = 0;
 500     wordr2[1] = 0;
 501     wordr2[2] = 0;
 502     wordr2[3] = 0;
 503
 504     u32 wordr3[4];
 505
 506     wordr3[0] = 0;
 507     wordr3[1] = 0;
 508     wordr3[2] = 0;
 509     wordr3[3] = 0;
 510
 511     if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
 512     {
 513       switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
 514     }
 515
 516     u32x w0[4];
 517
 518     w0[0] = wordl0[0] | wordr0[0];
 519     w0[1] = wordl0[1] | wordr0[1];
 520     w0[2] = wordl0[2] | wordr0[2];
 521     w0[3] = wordl0[3] | wordr0[3];
 522
 523     u32x w1[4];
 524
 525     w1[0] = wordl1[0] | wordr1[0];
 526     w1[1] = wordl1[1] | wordr1[1];
 527     w1[2] = wordl1[2] | wordr1[2];
 528     w1[3] = wordl1[3] | wordr1[3];
 529
 530     u32x w2[4];
 531
 532     w2[0] = wordl2[0] | wordr2[0];
 533     w2[1] = wordl2[1] | wordr2[1];
 534     w2[2] = wordl2[2] | wordr2[2];
 535     w2[3] = wordl2[3] | wordr2[3];
 536
 537     u32x w3[4];
 538
 539     w3[0] = wordl3[0] | wordr3[0];
 540     w3[1] = wordl3[1] | wordr3[1];
 541     w3[2] = pw_len * 8;
 542     w3[3] = 0;
 543
 544     u32x wl[16];
 545
 546     wl[ 0] = w0[0];
 547     wl[ 1] = w0[1];
 548     wl[ 2] = w0[2];
 549     wl[ 3] = w0[3];
 550     wl[ 4] = w1[0];
 551     wl[ 5] = w1[1];
 552     wl[ 6] = w1[2];
 553     wl[ 7] = w1[3];
 554     wl[ 8] = 0;
 555     wl[ 9] = 0;
 556     wl[10] = 0;
 557     wl[11] = 0;
 558     wl[12] = 0;
 559     wl[13] = 0;
 560     wl[14] = pw_len * 8;
 561     wl[15] = 0;
 562
 563     u32x dgst[5];
 564
 565     dgst[0] = RIPEMD160M_A;
 566     dgst[1] = RIPEMD160M_B;
 567     dgst[2] = RIPEMD160M_C;
 568     dgst[3] = RIPEMD160M_D;
 569     dgst[4] = RIPEMD160M_E;
 570
 571     ripemd160_transform (wl, dgst);
 572
 573     const u32x r0 = dgst[0];
 574     const u32x r1 = dgst[1];
 575     const u32x r2 = dgst[2];
 576     const u32x r3 = dgst[3];
 577
 578     #include VECT_COMPARE_S
 579   }
 580 }
 581
 582 extern "C" __global__ void __launch_bounds__ (256, 1) m06000_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)
 583 {
 584 }
 585
 586 extern "C" __global__ void __launch_bounds__ (256, 1) m06000_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)
 587 {
 588 }