API tools faq
paste
Advertisement
Guest User

Faster Hamsi

a guest
Nov 24th, 2014
161
0
Never
Add comment
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
C++ 33.53 KB | None | 0 0
raw download clone embed print report
  1. #include <stdio.h>
  2. #include <stdint.h>
  3. #include <memory.h>
  4.  
  5. extern cudaError_t MyStreamSynchronize(cudaStream_t stream, int situation, int thr_id);
  6.  
  7. typedef unsigned char BitSequence;
  8.  
  9.  
  10. #include "cuda_helper.h"
  11.  
  12. #define SPH_C32(x)    (x)
  13. #define SPH_T32(x)    (x)
  14.  
  15. static __device__ __forceinline__ uint32_t ROTL32( const uint32_t in, const int offset )
  16. {
  17.     uint32_t ret;
  18. #if __CUDA_ARCH__ >= 320
  19.     asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(ret) : "r"(in), "r"(in), "r"(offset));
  20. #else
  21.     ret  = (in << offset) | (in >> (32-offset));
  22. #endif
  23.     return ret;
  24. }
  25.  
  26. static __constant__ uint32_t d_alpha_n[32];
  27. static __constant__ uint32_t d_alpha_f[32];
  28. static __constant__ uint32_t d_T512[64][16];
  29.  
  30.  
  31.  
  32. static const uint32_t alpha_n[] = {
  33.     SPH_C32(0xff00f0f0), SPH_C32(0xccccaaaa), SPH_C32(0xf0f0cccc),
  34.     SPH_C32(0xff00aaaa), SPH_C32(0xccccaaaa), SPH_C32(0xf0f0ff00),
  35.     SPH_C32(0xaaaacccc), SPH_C32(0xf0f0ff00), SPH_C32(0xf0f0cccc),
  36.     SPH_C32(0xaaaaff00), SPH_C32(0xccccff00), SPH_C32(0xaaaaf0f0),
  37.     SPH_C32(0xaaaaf0f0), SPH_C32(0xff00cccc), SPH_C32(0xccccf0f0),
  38.     SPH_C32(0xff00aaaa), SPH_C32(0xccccaaaa), SPH_C32(0xff00f0f0),
  39.     SPH_C32(0xff00aaaa), SPH_C32(0xf0f0cccc), SPH_C32(0xf0f0ff00),
  40.     SPH_C32(0xccccaaaa), SPH_C32(0xf0f0ff00), SPH_C32(0xaaaacccc),
  41.     SPH_C32(0xaaaaff00), SPH_C32(0xf0f0cccc), SPH_C32(0xaaaaf0f0),
  42.     SPH_C32(0xccccff00), SPH_C32(0xff00cccc), SPH_C32(0xaaaaf0f0),
  43.     SPH_C32(0xff00aaaa), SPH_C32(0xccccf0f0)
  44. };
  45.  
  46. static const uint32_t alpha_f[] = {
  47.     SPH_C32(0xcaf9639c), SPH_C32(0x0ff0f9c0), SPH_C32(0x639c0ff0),
  48.     SPH_C32(0xcaf9f9c0), SPH_C32(0x0ff0f9c0), SPH_C32(0x639ccaf9),
  49.     SPH_C32(0xf9c00ff0), SPH_C32(0x639ccaf9), SPH_C32(0x639c0ff0),
  50.     SPH_C32(0xf9c0caf9), SPH_C32(0x0ff0caf9), SPH_C32(0xf9c0639c),
  51.     SPH_C32(0xf9c0639c), SPH_C32(0xcaf90ff0), SPH_C32(0x0ff0639c),
  52.     SPH_C32(0xcaf9f9c0), SPH_C32(0x0ff0f9c0), SPH_C32(0xcaf9639c),
  53.     SPH_C32(0xcaf9f9c0), SPH_C32(0x639c0ff0), SPH_C32(0x639ccaf9),
  54.     SPH_C32(0x0ff0f9c0), SPH_C32(0x639ccaf9), SPH_C32(0xf9c00ff0),
  55.     SPH_C32(0xf9c0caf9), SPH_C32(0x639c0ff0), SPH_C32(0xf9c0639c),
  56.     SPH_C32(0x0ff0caf9), SPH_C32(0xcaf90ff0), SPH_C32(0xf9c0639c),
  57.     SPH_C32(0xcaf9f9c0), SPH_C32(0x0ff0639c)
  58. };
  59.  
  60. #define hamsi_s00   m0
  61. #define hamsi_s01   m1
  62. #define hamsi_s02   c0
  63. #define hamsi_s03   c1
  64. #define hamsi_s04   m2
  65. #define hamsi_s05   m3
  66. #define hamsi_s06   c2
  67. #define hamsi_s07   c3
  68. #define hamsi_s08   c4
  69. #define hamsi_s09   c5
  70. #define hamsi_s0A   m4
  71. #define hamsi_s0B   m5
  72. #define hamsi_s0C   c6
  73. #define hamsi_s0D   c7
  74. #define hamsi_s0E   m6
  75. #define hamsi_s0F   m7
  76. #define hamsi_s10   m8
  77. #define hamsi_s11   m9
  78. #define hamsi_s12   c8
  79. #define hamsi_s13   c9
  80. #define hamsi_s14   mA
  81. #define hamsi_s15   mB
  82. #define hamsi_s16   cA
  83. #define hamsi_s17   cB
  84. #define hamsi_s18   cC
  85. #define hamsi_s19   cD
  86. #define hamsi_s1A   mC
  87. #define hamsi_s1B   mD
  88. #define hamsi_s1C   cE
  89. #define hamsi_s1D   cF
  90. #define hamsi_s1E   mE
  91. #define hamsi_s1F   mF
  92.  
  93. #define SBOX(a, b, c, d) { \
  94.         uint32_t t; \
  95.         t = (a); \
  96.         (a) &= (c); \
  97.         (a) ^= (d); \
  98.         (c) ^= (b); \
  99.         (c) ^= (a); \
  100.         (d) |= t; \
  101.         (d) ^= (b); \
  102.         t ^= (c); \
  103.         (b) = (d); \
  104.         (d) |= t; \
  105.         (d) ^= (a); \
  106.         (a) &= (b); \
  107.         t ^= (a); \
  108.         (b) ^= (d); \
  109.         (b) ^= t; \
  110.         (a) = (c); \
  111.         (c) = (b); \
  112.         (b) = (d); \
  113.         (d) = SPH_T32(~t); \
  114.     }
  115.  
  116. #define HAMSI_L(a, b, c, d) { \
  117.         (a) = ROTL32(a, 13); \
  118.         (c) = ROTL32(c, 3); \
  119.         (b) ^= (a) ^ (c); \
  120.         (d) ^= (c) ^ ((a) << 3); \
  121.         (b) = ROTL32(b, 1); \
  122.         (d) = ROTL32(d, 7); \
  123.         (a) ^= (b) ^ (d); \
  124.         (c) ^= (d) ^ ((b) << 7); \
  125.         (a) = ROTL32(a, 5); \
  126.         (c) = ROTL32(c, 22); \
  127.     }
  128.  
  129. #define ROUND_BIG(rc, alpha) { \
  130.         hamsi_s00 ^= alpha[0x00]; \
  131.         hamsi_s08 ^= alpha[0x08]; \
  132.         hamsi_s10 ^= alpha[0x10]; \
  133.         hamsi_s18 ^= alpha[0x18]; \
  134.         hamsi_s01 ^= alpha[0x01] ^ (uint32_t)(rc); \
  135.         hamsi_s09 ^= alpha[0x09]; \
  136.         hamsi_s11 ^= alpha[0x11]; \
  137.         hamsi_s19 ^= alpha[0x19]; \
  138.         hamsi_s02 ^= alpha[0x02]; \
  139.         hamsi_s0A ^= alpha[0x0A]; \
  140.         hamsi_s12 ^= alpha[0x12]; \
  141.         hamsi_s1A ^= alpha[0x1A]; \
  142.         hamsi_s03 ^= alpha[0x03]; \
  143.         hamsi_s0B ^= alpha[0x0B]; \
  144.         hamsi_s13 ^= alpha[0x13]; \
  145.         hamsi_s1B ^= alpha[0x1B]; \
  146.         hamsi_s04 ^= alpha[0x04]; \
  147.         hamsi_s0C ^= alpha[0x0C]; \
  148.         hamsi_s14 ^= alpha[0x14]; \
  149.         hamsi_s1C ^= alpha[0x1C]; \
  150.         hamsi_s05 ^= alpha[0x05]; \
  151.         hamsi_s0D ^= alpha[0x0D]; \
  152.         hamsi_s15 ^= alpha[0x15]; \
  153.         hamsi_s1D ^= alpha[0x1D]; \
  154.         hamsi_s06 ^= alpha[0x06]; \
  155.         hamsi_s0E ^= alpha[0x0E]; \
  156.         hamsi_s16 ^= alpha[0x16]; \
  157.         hamsi_s1E ^= alpha[0x1E]; \
  158.         hamsi_s07 ^= alpha[0x07]; \
  159.         hamsi_s0F ^= alpha[0x0F]; \
  160.         hamsi_s17 ^= alpha[0x17]; \
  161.         hamsi_s1F ^= alpha[0x1F]; \
  162.         SBOX(hamsi_s00, hamsi_s08, hamsi_s10, hamsi_s18); \
  163.         SBOX(hamsi_s01, hamsi_s09, hamsi_s11, hamsi_s19); \
  164.         SBOX(hamsi_s02, hamsi_s0A, hamsi_s12, hamsi_s1A); \
  165.         SBOX(hamsi_s03, hamsi_s0B, hamsi_s13, hamsi_s1B); \
  166.         SBOX(hamsi_s04, hamsi_s0C, hamsi_s14, hamsi_s1C); \
  167.         SBOX(hamsi_s05, hamsi_s0D, hamsi_s15, hamsi_s1D); \
  168.         SBOX(hamsi_s06, hamsi_s0E, hamsi_s16, hamsi_s1E); \
  169.         SBOX(hamsi_s07, hamsi_s0F, hamsi_s17, hamsi_s1F); \
  170.         HAMSI_L(hamsi_s00, hamsi_s09, hamsi_s12, hamsi_s1B); \
  171.         HAMSI_L(hamsi_s01, hamsi_s0A, hamsi_s13, hamsi_s1C); \
  172.         HAMSI_L(hamsi_s02, hamsi_s0B, hamsi_s14, hamsi_s1D); \
  173.         HAMSI_L(hamsi_s03, hamsi_s0C, hamsi_s15, hamsi_s1E); \
  174.         HAMSI_L(hamsi_s04, hamsi_s0D, hamsi_s16, hamsi_s1F); \
  175.         HAMSI_L(hamsi_s05, hamsi_s0E, hamsi_s17, hamsi_s18); \
  176.         HAMSI_L(hamsi_s06, hamsi_s0F, hamsi_s10, hamsi_s19); \
  177.         HAMSI_L(hamsi_s07, hamsi_s08, hamsi_s11, hamsi_s1A); \
  178.         HAMSI_L(hamsi_s00, hamsi_s02, hamsi_s05, hamsi_s07); \
  179.         HAMSI_L(hamsi_s10, hamsi_s13, hamsi_s15, hamsi_s16); \
  180.         HAMSI_L(hamsi_s09, hamsi_s0B, hamsi_s0C, hamsi_s0E); \
  181.         HAMSI_L(hamsi_s19, hamsi_s1A, hamsi_s1C, hamsi_s1F); \
  182.     }
  183.  
  184.  
  185. #define P_BIG  { \
  186.         for( int r = 0; r < 6; r++ ) \
  187.             ROUND_BIG(r, d_alpha_n); \
  188.     }
  189.  
  190. #define PF_BIG  { \
  191.         for( int r = 0; r < 12; r++ ) \
  192.             ROUND_BIG(r, d_alpha_f); \
  193.     }
  194.  
  195. #define T_BIG   { \
  196.         /* order is important */ \
  197.         cF = (h[0xF] ^= hamsi_s17); \
  198.         cE = (h[0xE] ^= hamsi_s16); \
  199.         cD = (h[0xD] ^= hamsi_s15); \
  200.         cC = (h[0xC] ^= hamsi_s14); \
  201.         cB = (h[0xB] ^= hamsi_s13); \
  202.         cA = (h[0xA] ^= hamsi_s12); \
  203.         c9 = (h[0x9] ^= hamsi_s11); \
  204.         c8 = (h[0x8] ^= hamsi_s10); \
  205.         c7 = (h[0x7] ^= hamsi_s07); \
  206.         c6 = (h[0x6] ^= hamsi_s06); \
  207.         c5 = (h[0x5] ^= hamsi_s05); \
  208.         c4 = (h[0x4] ^= hamsi_s04); \
  209.         c3 = (h[0x3] ^= hamsi_s03); \
  210.         c2 = (h[0x2] ^= hamsi_s02); \
  211.         c1 = (h[0x1] ^= hamsi_s01); \
  212.         c0 = (h[0x0] ^= hamsi_s00); \
  213.     }#pragma unroll 2
  214.             for (int u = 0; u < 8; u ++) {
  215.                 db = h1[i+u];
  216. #pragma unroll 2
  217.                 for (int v = 0; v < 8; v ++, db >>= 1) {
  218.                     dm = -(uint32_t)(db & 1);
  219.                     m0 ^= dm & *(tp+ 0); m1 ^= dm & *(tp+ 1);
  220.                     m2 ^= dm & *(tp+ 2); m3 ^= dm & *(tp+ 3);
  221.                     m4 ^= dm & *(tp+ 4); m5 ^= dm & *(tp+ 5);
  222.                     m6 ^= dm & *(tp+ 6); m7 ^= dm & *(tp+ 7);
  223.                     m8 ^= dm & *(tp+ 8); m9 ^= dm & *(tp+ 9);
  224.                     mA ^= dm & *(tp+10); mB ^= dm & *(tp+11);
  225.                     mC ^= dm & *(tp+12); mD ^= dm & *(tp+13);
  226.                     mE ^= dm & *(tp+14); mF ^= dm & *(tp+15);
  227.                     tp += 16;
  228.                 }
  229.             }
  230.  
  231.             for( int r = 0; r < 6; r += 2 ) {
  232.                 ROUND_BIG(r, d_alpha_n);
  233.                 ROUND_BIG(r+1, d_alpha_n);
  234.             }
  235.             T_BIG;
  236.         }
  237.  
  238.  
  239.         tp = &d_T512[0][0] + 112;
  240.  
  241.         m0 = *(tp+ 0); m1 = *(tp+ 1);
  242.         m2 = *(tp+ 2); m3 = *(tp+ 3);
  243.         m4 = *(tp+ 4); m5 = *(tp+ 5);
  244.         m6 = *(tp+ 6); m7 = *(tp+ 7);
  245.         m8 = *(tp+ 8); m9 = *(tp+ 9);
  246.         mA = *(tp+10); mB = *(tp+11);
  247.         mC = *(tp+12); mD = *(tp+13);
  248.         mE = *(tp+14); mF = *(tp+15);
  249.  
  250.         for( int r = 0; r < 6; r += 2 ) {
  251.             ROUND_BIG(r, d_alpha_n);
  252.             ROUND_BIG(r+1, d_alpha_n);
  253.         }
  254.         T_BIG;
  255.  
  256.         tp = &d_T512[0][0] + 784;
  257.  
  258.         m0 = *(tp+ 0); m1 = *(tp+ 1);
  259.         m2 = *(tp+ 2); m3 = *(tp+ 3);
  260.         m4 = *(tp+ 4); m5 = *(tp+ 5);
  261.         m6 = *(tp+ 6); m7 = *(tp+ 7);
  262.         m8 = *(tp+ 8); m9 = *(tp+ 9);
  263.         mA = *(tp+10); mB = *(tp+11);
  264.         mC = *(tp+12); mD = *(tp+13);
  265.         mE = *(tp+14); mF = *(tp+15);
  266.  
  267.         for( int r = 0; r < 12; r += 2 ) {
  268.             ROUND_BIG(r, d_alpha_f);
  269.             ROUND_BIG(r+1, d_alpha_f);
  270.         }
  271.         T_BIG;
  272.  
  273. #pragma unroll 16
  274.         for (int i = 0; i < 16; i++)
  275.             Hash[i] = cuda_swab32(h[i]);
  276.     }
  277. }
  278.  
  279.  
  280. __host__ void x13_hamsi512_cpu_init(int thr_id, int threads)
  281. {
  282.     cudaMemcpyToSymbol( d_alpha_n, alpha_n, sizeof(uint32_t)*32, 0, cudaMemcpyHostToDevice);
  283.     cudaMemcpyToSymbol( d_alpha_f, alpha_f, sizeof(uint32_t)*32, 0, cudaMemcpyHostToDevice);
  284.     cudaMemcpyToSymbol( d_T512, T512, sizeof(uint32_t)*64*16, 0, cudaMemcpyHostToDevice);
  285. }
  286.  
  287. __host__ void x13_hamsi512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order)
  288. {
  289.     const int threadsperblock = 256;
  290.  
  291.     dim3 grid((threads + threadsperblock-1)/threadsperblock);
  292.     dim3 block(threadsperblock);
  293.  
  294.     x13_hamsi512_gpu_hash_64<<<grid, block>>>(threads, startNounce, (uint64_t*)d_hash, d_nonceVector);
  295.     MyStreamSynchronize(NULL, order, thr_id);
  296. }
Advertisement
Add Comment
Please, Sign In to add comment
Advertisement
Public Pastes
Advertisement
We use cookies for various purposes including analytics. By continuing to use Pastebin, you agree to our use of cookies as described in the Cookies Policy.  OK, I Understand
Not a member of Pastebin yet?
Sign Up, it unlocks many cool features!