nv/amp_a1_v4.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define VECT_SIZE4
   7
   8 #include "include/constants.h"
   9 #include "types_nv.c"
  10
  11 __device__ static void switch_buffer_by_offset (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], const u32 offset)
  12 {
  13   #if __CUDA_ARCH__ >= 200
  14
  15   const int offset_minus_4 = 4 - (offset % 4);
  16
  17   int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff;
  18
  19   switch (offset / 4)
  20   {
  21     case 0:
  22       w3[1] = __byte_perm (w3[0], w3[1], selector);
  23       w3[0] = __byte_perm (w2[3], w3[0], selector);
  24       w2[3] = __byte_perm (w2[2], w2[3], selector);
  25       w2[2] = __byte_perm (w2[1], w2[2], selector);
  26       w2[1] = __byte_perm (w2[0], w2[1], selector);
  27       w2[0] = __byte_perm (w1[3], w2[0], selector);
  28       w1[3] = __byte_perm (w1[2], w1[3], selector);
  29       w1[2] = __byte_perm (w1[1], w1[2], selector);
  30       w1[1] = __byte_perm (w1[0], w1[1], selector);
  31       w1[0] = __byte_perm (w0[3], w1[0], selector);
  32       w0[3] = __byte_perm (w0[2], w0[3], selector);
  33       w0[2] = __byte_perm (w0[1], w0[2], selector);
  34       w0[1] = __byte_perm (w0[0], w0[1], selector);
  35       w0[0] = __byte_perm (    0, w0[0], selector);
  36
  37       break;
  38
  39     case 1:
  40       w3[1] = __byte_perm (w2[3], w3[0], selector);
  41       w3[0] = __byte_perm (w2[2], w2[3], selector);
  42       w2[3] = __byte_perm (w2[1], w2[2], selector);
  43       w2[2] = __byte_perm (w2[0], w2[1], selector);
  44       w2[1] = __byte_perm (w1[3], w2[0], selector);
  45       w2[0] = __byte_perm (w1[2], w1[3], selector);
  46       w1[3] = __byte_perm (w1[1], w1[2], selector);
  47       w1[2] = __byte_perm (w1[0], w1[1], selector);
  48       w1[1] = __byte_perm (w0[3], w1[0], selector);
  49       w1[0] = __byte_perm (w0[2], w0[3], selector);
  50       w0[3] = __byte_perm (w0[1], w0[2], selector);
  51       w0[2] = __byte_perm (w0[0], w0[1], selector);
  52       w0[1] = __byte_perm (    0, w0[0], selector);
  53       w0[0] = 0;
  54
  55       break;
  56
  57     case 2:
  58       w3[1] = __byte_perm (w2[2], w2[3], selector);
  59       w3[0] = __byte_perm (w2[1], w2[2], selector);
  60       w2[3] = __byte_perm (w2[0], w2[1], selector);
  61       w2[2] = __byte_perm (w1[3], w2[0], selector);
  62       w2[1] = __byte_perm (w1[2], w1[3], selector);
  63       w2[0] = __byte_perm (w1[1], w1[2], selector);
  64       w1[3] = __byte_perm (w1[0], w1[1], selector);
  65       w1[2] = __byte_perm (w0[3], w1[0], selector);
  66       w1[1] = __byte_perm (w0[2], w0[3], selector);
  67       w1[0] = __byte_perm (w0[1], w0[2], selector);
  68       w0[3] = __byte_perm (w0[0], w0[1], selector);
  69       w0[2] = __byte_perm (    0, w0[0], selector);
  70       w0[1] = 0;
  71       w0[0] = 0;
  72
  73       break;
  74
  75     case 3:
  76       w3[1] = __byte_perm (w2[1], w2[2], selector);
  77       w3[0] = __byte_perm (w2[0], w2[1], selector);
  78       w2[3] = __byte_perm (w1[3], w2[0], selector);
  79       w2[2] = __byte_perm (w1[2], w1[3], selector);
  80       w2[1] = __byte_perm (w1[1], w1[2], selector);
  81       w2[0] = __byte_perm (w1[0], w1[1], selector);
  82       w1[3] = __byte_perm (w0[3], w1[0], selector);
  83       w1[2] = __byte_perm (w0[2], w0[3], selector);
  84       w1[1] = __byte_perm (w0[1], w0[2], selector);
  85       w1[0] = __byte_perm (w0[0], w0[1], selector);
  86       w0[3] = __byte_perm (    0, w0[0], selector);
  87       w0[2] = 0;
  88       w0[1] = 0;
  89       w0[0] = 0;
  90
  91       break;
  92
  93     case 4:
  94       w3[1] = __byte_perm (w2[0], w2[1], selector);
  95       w3[0] = __byte_perm (w1[3], w2[0], selector);
  96       w2[3] = __byte_perm (w1[2], w1[3], selector);
  97       w2[2] = __byte_perm (w1[1], w1[2], selector);
  98       w2[1] = __byte_perm (w1[0], w1[1], selector);
  99       w2[0] = __byte_perm (w0[3], w1[0], selector);
 100       w1[3] = __byte_perm (w0[2], w0[3], selector);
 101       w1[2] = __byte_perm (w0[1], w0[2], selector);
 102       w1[1] = __byte_perm (w0[0], w0[1], selector);
 103       w1[0] = __byte_perm (    0, w0[0], selector);
 104       w0[3] = 0;
 105       w0[2] = 0;
 106       w0[1] = 0;
 107       w0[0] = 0;
 108
 109       break;
 110
 111     case 5:
 112       w3[1] = __byte_perm (w1[3], w2[0], selector);
 113       w3[0] = __byte_perm (w1[2], w1[3], selector);
 114       w2[3] = __byte_perm (w1[1], w1[2], selector);
 115       w2[2] = __byte_perm (w1[0], w1[1], selector);
 116       w2[1] = __byte_perm (w0[3], w1[0], selector);
 117       w2[0] = __byte_perm (w0[2], w0[3], selector);
 118       w1[3] = __byte_perm (w0[1], w0[2], selector);
 119       w1[2] = __byte_perm (w0[0], w0[1], selector);
 120       w1[1] = __byte_perm (    0, w0[0], selector);
 121       w1[0] = 0;
 122       w0[3] = 0;
 123       w0[2] = 0;
 124       w0[1] = 0;
 125       w0[0] = 0;
 126
 127       break;
 128
 129     case 6:
 130       w3[1] = __byte_perm (w1[2], w1[3], selector);
 131       w3[0] = __byte_perm (w1[1], w1[2], selector);
 132       w2[3] = __byte_perm (w1[0], w1[1], selector);
 133       w2[2] = __byte_perm (w0[3], w1[0], selector);
 134       w2[1] = __byte_perm (w0[2], w0[3], selector);
 135       w2[0] = __byte_perm (w0[1], w0[2], selector);
 136       w1[3] = __byte_perm (w0[0], w0[1], selector);
 137       w1[2] = __byte_perm (    0, w0[0], selector);
 138       w1[1] = 0;
 139       w1[0] = 0;
 140       w0[3] = 0;
 141       w0[2] = 0;
 142       w0[1] = 0;
 143       w0[0] = 0;
 144
 145       break;
 146
 147     case 7:
 148       w3[1] = __byte_perm (w1[1], w1[2], selector);
 149       w3[0] = __byte_perm (w1[0], w1[1], selector);
 150       w2[3] = __byte_perm (w0[3], w1[0], selector);
 151       w2[2] = __byte_perm (w0[2], w0[3], selector);
 152       w2[1] = __byte_perm (w0[1], w0[2], selector);
 153       w2[0] = __byte_perm (w0[0], w0[1], selector);
 154       w1[3] = __byte_perm (    0, w0[0], selector);
 155       w1[2] = 0;
 156       w1[1] = 0;
 157       w1[0] = 0;
 158       w0[3] = 0;
 159       w0[2] = 0;
 160       w0[1] = 0;
 161       w0[0] = 0;
 162
 163       break;
 164
 165     case 8:
 166       w3[1] = __byte_perm (w1[0], w1[1], selector);
 167       w3[0] = __byte_perm (w0[3], w1[0], selector);
 168       w2[3] = __byte_perm (w0[2], w0[3], selector);
 169       w2[2] = __byte_perm (w0[1], w0[2], selector);
 170       w2[1] = __byte_perm (w0[0], w0[1], selector);
 171       w2[0] = __byte_perm (    0, w0[0], selector);
 172       w1[3] = 0;
 173       w1[2] = 0;
 174       w1[1] = 0;
 175       w1[0] = 0;
 176       w0[3] = 0;
 177       w0[2] = 0;
 178       w0[1] = 0;
 179       w0[0] = 0;
 180
 181       break;
 182
 183     case 9:
 184       w3[1] = __byte_perm (w0[3], w1[0], selector);
 185       w3[0] = __byte_perm (w0[2], w0[3], selector);
 186       w2[3] = __byte_perm (w0[1], w0[2], selector);
 187       w2[2] = __byte_perm (w0[0], w0[1], selector);
 188       w2[1] = __byte_perm (    0, w0[0], selector);
 189       w2[0] = 0;
 190       w1[3] = 0;
 191       w1[2] = 0;
 192       w1[1] = 0;
 193       w1[0] = 0;
 194       w0[3] = 0;
 195       w0[2] = 0;
 196       w0[1] = 0;
 197       w0[0] = 0;
 198
 199       break;
 200
 201     case 10:
 202       w3[1] = __byte_perm (w0[2], w0[3], selector);
 203       w3[0] = __byte_perm (w0[1], w0[2], selector);
 204       w2[3] = __byte_perm (w0[0], w0[1], selector);
 205       w2[2] = __byte_perm (    0, w0[0], selector);
 206       w2[1] = 0;
 207       w2[0] = 0;
 208       w1[3] = 0;
 209       w1[2] = 0;
 210       w1[1] = 0;
 211       w1[0] = 0;
 212       w0[3] = 0;
 213       w0[2] = 0;
 214       w0[1] = 0;
 215       w0[0] = 0;
 216
 217       break;
 218
 219     case 11:
 220       w3[1] = __byte_perm (w0[1], w0[2], selector);
 221       w3[0] = __byte_perm (w0[0], w0[1], selector);
 222       w2[3] = __byte_perm (    0, w0[0], selector);
 223       w2[2] = 0;
 224       w2[1] = 0;
 225       w2[0] = 0;
 226       w1[3] = 0;
 227       w1[2] = 0;
 228       w1[1] = 0;
 229       w1[0] = 0;
 230       w0[3] = 0;
 231       w0[2] = 0;
 232       w0[1] = 0;
 233       w0[0] = 0;
 234
 235       break;
 236
 237     case 12:
 238       w3[1] = __byte_perm (w0[0], w0[1], selector);
 239       w3[0] = __byte_perm (    0, w0[0], selector);
 240       w2[3] = 0;
 241       w2[2] = 0;
 242       w2[1] = 0;
 243       w2[0] = 0;
 244       w1[3] = 0;
 245       w1[2] = 0;
 246       w1[1] = 0;
 247       w1[0] = 0;
 248       w0[3] = 0;
 249       w0[2] = 0;
 250       w0[1] = 0;
 251       w0[0] = 0;
 252
 253       break;
 254
 255     case 13:
 256       w3[1] = __byte_perm (    0, w0[0], selector);
 257       w3[0] = 0;
 258       w2[3] = 0;
 259       w2[2] = 0;
 260       w2[1] = 0;
 261       w2[0] = 0;
 262       w1[3] = 0;
 263       w1[2] = 0;
 264       w1[1] = 0;
 265       w1[0] = 0;
 266       w0[3] = 0;
 267       w0[2] = 0;
 268       w0[1] = 0;
 269       w0[0] = 0;
 270
 271       break;
 272   }
 273
 274   #else
 275
 276   u32x tmp0[4];
 277   u32x tmp1[4];
 278   u32x tmp2[1];
 279
 280   switch (offset % 4)
 281   {
 282                 case 0:
 283       tmp0[0] = w0[0];
 284       tmp0[1] = w0[1];
 285       tmp0[2] = w0[2];
 286       tmp0[3] = w0[3];
 287       tmp1[0] = w1[0];
 288       tmp1[1] = w1[1];
 289       tmp1[2] = w1[2];
 290       tmp1[3] = w1[3];
 291       tmp2[0] = 0;
 292                   break;
 293
 294                 case 1:
 295                   tmp0[0] =               w0[0] <<  8;
 296                   tmp0[1] = w0[0] >> 24 | w0[1] <<  8;
 297                   tmp0[2] = w0[1] >> 24 | w0[2] <<  8;
 298                   tmp0[3] = w0[2] >> 24 | w0[3] <<  8;
 299                   tmp1[0] = w0[3] >> 24 | w1[0] <<  8;
 300                   tmp1[1] = w1[0] >> 24 | w1[1] <<  8;
 301                   tmp1[2] = w1[1] >> 24 | w1[2] <<  8;
 302                   tmp1[3] = w1[2] >> 24 | w1[3] <<  8;
 303                   tmp2[0] = w1[3] >> 24;
 304                   break;
 305
 306                 case 2:
 307                   tmp0[0] =               w0[0] << 16;
 308                   tmp0[1] = w0[0] >> 16 | w0[1] << 16;
 309                   tmp0[2] = w0[1] >> 16 | w0[2] << 16;
 310                   tmp0[3] = w0[2] >> 16 | w0[3] << 16;
 311                   tmp1[0] = w0[3] >> 16 | w1[0] << 16;
 312                   tmp1[1] = w1[0] >> 16 | w1[1] << 16;
 313                   tmp1[2] = w1[1] >> 16 | w1[2] << 16;
 314                   tmp1[3] = w1[2] >> 16 | w1[3] << 16;
 315                   tmp2[0] = w1[3] >> 16;
 316                   break;
 317
 318                 case 3:
 319                   tmp0[0] =               w0[0] << 24;
 320                   tmp0[1] = w0[0] >>  8 | w0[1] << 24;
 321                   tmp0[2] = w0[1] >>  8 | w0[2] << 24;
 322                   tmp0[3] = w0[2] >>  8 | w0[3] << 24;
 323                   tmp1[0] = w0[3] >>  8 | w1[0] << 24;
 324                   tmp1[1] = w1[0] >>  8 | w1[1] << 24;
 325                   tmp1[2] = w1[1] >>  8 | w1[2] << 24;
 326                   tmp1[3] = w1[2] >>  8 | w1[3] << 24;
 327                   tmp2[0] = w1[3] >>  8;
 328                   break;
 329   }
 330
 331   switch (offset / 4)
 332   {
 333     case 0:
 334       w0[0] = tmp0[0];
 335       w0[1] = tmp0[1];
 336       w0[2] = tmp0[2];
 337       w0[3] = tmp0[3];
 338       w1[0] = tmp1[0];
 339       w1[1] = tmp1[1];
 340       w1[2] = tmp1[2];
 341       w1[3] = tmp1[3];
 342       w2[0] = tmp2[0];
 343       break;
 344
 345     case 1:
 346       w0[0] = 0;
 347       w0[1] = tmp0[0];
 348       w0[2] = tmp0[1];
 349       w0[3] = tmp0[2];
 350       w1[0] = tmp0[3];
 351       w1[1] = tmp1[0];
 352       w1[2] = tmp1[1];
 353       w1[3] = tmp1[2];
 354       w2[0] = tmp1[3];
 355       w2[1] = tmp2[0];
 356       break;
 357
 358     case 2:
 359       w0[0] = 0;
 360       w0[1] = 0;
 361       w0[2] = tmp0[0];
 362       w0[3] = tmp0[1];
 363       w1[0] = tmp0[2];
 364       w1[1] = tmp0[3];
 365       w1[2] = tmp1[0];
 366       w1[3] = tmp1[1];
 367       w2[0] = tmp1[2];
 368       w2[1] = tmp1[3];
 369       w2[2] = tmp2[0];
 370       break;
 371
 372     case 3:
 373       w0[0] = 0;
 374       w0[1] = 0;
 375       w0[2] = 0;
 376       w0[3] = tmp0[0];
 377       w1[0] = tmp0[1];
 378       w1[1] = tmp0[2];
 379       w1[2] = tmp0[3];
 380       w1[3] = tmp1[0];
 381       w2[0] = tmp1[1];
 382       w2[1] = tmp1[2];
 383       w2[2] = tmp1[3];
 384       w2[3] = tmp2[0];
 385       break;
 386
 387     case 4:
 388       w0[0] = 0;
 389       w0[1] = 0;
 390       w0[2] = 0;
 391       w0[3] = 0;
 392       w1[0] = tmp0[0];
 393       w1[1] = tmp0[1];
 394       w1[2] = tmp0[2];
 395       w1[3] = tmp0[3];
 396       w2[0] = tmp1[0];
 397       w2[1] = tmp1[1];
 398       w2[2] = tmp1[2];
 399       w2[3] = tmp1[3];
 400       w3[0] = tmp2[0];
 401       break;
 402
 403     case 5:
 404       w0[0] = 0;
 405       w0[1] = 0;
 406       w0[2] = 0;
 407       w0[3] = 0;
 408       w1[0] = 0;
 409       w1[1] = tmp0[0];
 410       w1[2] = tmp0[1];
 411       w1[3] = tmp0[2];
 412       w2[0] = tmp0[3];
 413       w2[1] = tmp1[0];
 414       w2[2] = tmp1[1];
 415       w2[3] = tmp1[2];
 416       w3[0] = tmp1[3];
 417       w3[1] = tmp2[0];
 418       break;
 419
 420     case 6:
 421       w0[0] = 0;
 422       w0[1] = 0;
 423       w0[2] = 0;
 424       w0[3] = 0;
 425       w1[0] = 0;
 426       w1[1] = 0;
 427       w1[2] = tmp0[0];
 428       w1[3] = tmp0[1];
 429       w2[0] = tmp0[2];
 430       w2[1] = tmp0[3];
 431       w2[2] = tmp1[0];
 432       w2[3] = tmp1[1];
 433       w3[0] = tmp1[2];
 434       w3[1] = tmp1[3];
 435       w3[2] = tmp2[0];
 436       break;
 437
 438     case 7:
 439       w0[0] = 0;
 440       w0[1] = 0;
 441       w0[2] = 0;
 442       w0[3] = 0;
 443       w1[0] = 0;
 444       w1[1] = 0;
 445       w1[2] = 0;
 446       w1[3] = tmp0[0];
 447       w2[0] = tmp0[1];
 448       w2[1] = tmp0[2];
 449       w2[2] = tmp0[3];
 450       w2[3] = tmp1[0];
 451       w3[0] = tmp1[1];
 452       w3[1] = tmp1[2];
 453       w3[2] = tmp1[3];
 454       w3[3] = tmp2[0];
 455       break;
 456
 457     case 8:
 458       w0[0] = 0;
 459       w0[1] = 0;
 460       w0[2] = 0;
 461       w0[3] = 0;
 462       w1[0] = 0;
 463       w1[1] = 0;
 464       w1[2] = 0;
 465       w1[3] = 0;
 466       w2[0] = tmp0[0];
 467       w2[1] = tmp0[1];
 468       w2[2] = tmp0[2];
 469       w2[3] = tmp0[3];
 470       w3[0] = tmp1[0];
 471       w3[1] = tmp1[1];
 472       w3[2] = tmp1[2];
 473       w3[3] = tmp1[3];
 474       break;
 475
 476     case 9:
 477       w0[0] = 0;
 478       w0[1] = 0;
 479       w0[2] = 0;
 480       w0[3] = 0;
 481       w1[0] = 0;
 482       w1[1] = 0;
 483       w1[2] = 0;
 484       w1[3] = 0;
 485       w2[0] = 0;
 486       w2[1] = tmp0[0];
 487       w2[2] = tmp0[1];
 488       w2[3] = tmp0[2];
 489       w3[0] = tmp0[3];
 490       w3[1] = tmp1[0];
 491       w3[2] = tmp1[1];
 492       w3[3] = tmp1[2];
 493       break;
 494
 495     case 10:
 496       w0[0] = 0;
 497       w0[1] = 0;
 498       w0[2] = 0;
 499       w0[3] = 0;
 500       w1[0] = 0;
 501       w1[1] = 0;
 502       w1[2] = 0;
 503       w1[3] = 0;
 504       w2[0] = 0;
 505       w2[1] = 0;
 506       w2[2] = tmp0[0];
 507       w2[3] = tmp0[1];
 508       w3[0] = tmp0[2];
 509       w3[1] = tmp0[3];
 510       w3[2] = tmp1[0];
 511       w3[3] = tmp1[1];
 512       break;
 513
 514     case 11:
 515       w0[0] = 0;
 516       w0[1] = 0;
 517       w0[2] = 0;
 518       w0[3] = 0;
 519       w1[0] = 0;
 520       w1[1] = 0;
 521       w1[2] = 0;
 522       w1[3] = 0;
 523       w2[0] = 0;
 524       w2[1] = 0;
 525       w2[2] = 0;
 526       w2[3] = tmp0[0];
 527       w3[0] = tmp0[1];
 528       w3[1] = tmp0[2];
 529       w3[2] = tmp0[3];
 530       w3[3] = tmp1[0];
 531       break;
 532
 533     case 12:
 534       w0[0] = 0;
 535       w0[1] = 0;
 536       w0[2] = 0;
 537       w0[3] = 0;
 538       w1[0] = 0;
 539       w1[1] = 0;
 540       w1[2] = 0;
 541       w1[3] = 0;
 542       w2[0] = 0;
 543       w2[1] = 0;
 544       w2[2] = 0;
 545       w2[3] = 0;
 546       w3[0] = tmp0[0];
 547       w3[1] = tmp0[1];
 548       w3[2] = tmp0[2];
 549       w3[3] = tmp0[3];
 550       break;
 551
 552     case 13:
 553       w0[0] = 0;
 554       w0[1] = 0;
 555       w0[2] = 0;
 556       w0[3] = 0;
 557       w1[0] = 0;
 558       w1[1] = 0;
 559       w1[2] = 0;
 560       w1[3] = 0;
 561       w2[0] = 0;
 562       w2[1] = 0;
 563       w2[2] = 0;
 564       w2[3] = 0;
 565       w3[0] = 0;
 566       w3[1] = tmp0[0];
 567       w3[2] = tmp0[1];
 568       w3[3] = tmp0[2];
 569       break;
 570
 571   }
 572
 573   #endif
 574 }
 575
 576 __device__ __constant__ comb_t c_combs[1024];
 577
 578 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)
 579 {
 580   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 581
 582   if (gid >= gid_max) return;
 583
 584   const u32 pw_l_len = pws[gid].pw_len;
 585
 586   u32x wordl0[4];
 587
 588   wordl0[0] = pws[gid].i[ 0];
 589   wordl0[1] = pws[gid].i[ 1];
 590   wordl0[2] = pws[gid].i[ 2];
 591   wordl0[3] = pws[gid].i[ 3];
 592
 593   u32x wordl1[4];
 594
 595   wordl1[0] = pws[gid].i[ 4];
 596   wordl1[1] = pws[gid].i[ 5];
 597   wordl1[2] = pws[gid].i[ 6];
 598   wordl1[3] = pws[gid].i[ 7];
 599
 600   u32x wordl2[4];
 601
 602   wordl2[0] = 0;
 603   wordl2[1] = 0;
 604   wordl2[2] = 0;
 605   wordl2[3] = 0;
 606
 607   u32x wordl3[4];
 608
 609   wordl3[0] = 0;
 610   wordl3[1] = 0;
 611   wordl3[2] = 0;
 612   wordl3[3] = 0;
 613
 614   const u32 pw_r_len = c_combs[0].pw_len;
 615
 616   u32x wordr0[4];
 617
 618   wordr0[0] = c_combs[0].i[0];
 619   wordr0[1] = c_combs[0].i[1];
 620   wordr0[2] = c_combs[0].i[2];
 621   wordr0[3] = c_combs[0].i[3];
 622
 623   u32x wordr1[4];
 624
 625   wordr1[0] = c_combs[0].i[4];
 626   wordr1[1] = c_combs[0].i[5];
 627   wordr1[2] = c_combs[0].i[6];
 628   wordr1[3] = c_combs[0].i[7];
 629
 630   u32x wordr2[4];
 631
 632   wordr2[0] = 0;
 633   wordr2[1] = 0;
 634   wordr2[2] = 0;
 635   wordr2[3] = 0;
 636
 637   u32x wordr3[4];
 638
 639   wordr3[0] = 0;
 640   wordr3[1] = 0;
 641   wordr3[2] = 0;
 642   wordr3[3] = 0;
 643
 644   if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
 645   {
 646     switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
 647   }
 648
 649   if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
 650   {
 651     switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, pw_r_len);
 652   }
 653
 654   u32x w0[4];
 655
 656   w0[0] = wordl0[0] | wordr0[0];
 657   w0[1] = wordl0[1] | wordr0[1];
 658   w0[2] = wordl0[2] | wordr0[2];
 659   w0[3] = wordl0[3] | wordr0[3];
 660
 661   u32x w1[4];
 662
 663   w1[0] = wordl1[0] | wordr1[0];
 664   w1[1] = wordl1[1] | wordr1[1];
 665   w1[2] = wordl1[2] | wordr1[2];
 666   w1[3] = wordl1[3] | wordr1[3];
 667
 668   u32x w2[4];
 669
 670   w2[0] = wordl2[0] | wordr2[0];
 671   w2[1] = wordl2[1] | wordr2[1];
 672   w2[2] = wordl2[2] | wordr2[2];
 673   w2[3] = wordl2[3] | wordr2[3];
 674
 675   u32x w3[4];
 676
 677   w3[0] = wordl3[0] | wordr3[0];
 678   w3[1] = wordl3[1] | wordr3[1];
 679   w3[2] = wordl3[2] | wordr3[2];
 680   w3[3] = wordl3[3] | wordr3[3];
 681
 682   const u32 pw_len = pw_l_len + pw_r_len;
 683
 684   pws_amp[gid].i[ 0] = w0[0];
 685   pws_amp[gid].i[ 1] = w0[1];
 686   pws_amp[gid].i[ 2] = w0[2];
 687   pws_amp[gid].i[ 3] = w0[3];
 688   pws_amp[gid].i[ 4] = w1[0];
 689   pws_amp[gid].i[ 5] = w1[1];
 690   pws_amp[gid].i[ 6] = w1[2];
 691   pws_amp[gid].i[ 7] = w1[3];
 692   pws_amp[gid].i[ 8] = w2[0];
 693   pws_amp[gid].i[ 9] = w2[1];
 694   pws_amp[gid].i[10] = w2[2];
 695   pws_amp[gid].i[11] = w2[3];
 696   pws_amp[gid].i[12] = w3[0];
 697   pws_amp[gid].i[13] = w3[1];
 698   pws_amp[gid].i[14] = w3[2];
 699   pws_amp[gid].i[15] = w3[3];
 700
 701   pws_amp[gid].pw_len = pw_len;
 702 }