diff --git a/x11/cuda_x11_simd512.cu b/x11/cuda_x11_simd512.cu index 5c742ce..940818e 100644 --- a/x11/cuda_x11_simd512.cu +++ b/x11/cuda_x11_simd512.cu @@ -44,79 +44,6 @@ const int h_FFT[256] = -93, -125, 125, -86, 86, 117, -117, -91, 91, 42, -42, 87, -87, -117, 117, 102, -102 }; -__constant__ int c_P4[32][4]; -static const int h_P4[32][4] = { -{ 2, 34, 18, 50 }, -{ 6, 38, 22, 54 }, -{ 0, 32, 16, 48 }, -{ 4, 36, 20, 52 }, -{ 14, 46, 30, 62 }, -{ 10, 42, 26, 58 }, -{ 12, 44, 28, 60 }, -{ 8, 40, 24, 56 }, -{ 15, 47, 31, 63 }, -{ 13, 45, 29, 61 }, -{ 3, 35, 19, 51 }, -{ 1, 33, 17, 49 }, -{ 9, 41, 25, 57 }, -{ 11, 43, 27, 59 }, -{ 5, 37, 21, 53 }, -{ 7, 39, 23, 55 }, -{ 8, 40, 24, 56 }, -{ 4, 36, 20, 52 }, -{ 14, 46, 30, 62 }, -{ 2, 34, 18, 50 }, -{ 6, 38, 22, 54 }, -{ 10, 42, 26, 58 }, -{ 0, 32, 16, 48 }, -{ 12, 44, 28, 60 }, -{ 70, 102, 86, 118 }, -{ 64, 96, 80, 112 }, -{ 72, 104, 88, 120 }, -{ 78, 110, 94, 126 }, -{ 76, 108, 92, 124 }, -{ 74, 106, 90, 122 }, -{ 66, 98, 82, 114 }, -{ 68, 100, 84, 116 } -}; - -__constant__ int c_Q4[32][4]; -static const int h_Q4[32][4] = { -{ 66, 98, 82, 114 }, -{ 70, 102, 86, 118 }, -{ 64, 96, 80, 112 }, -{ 68, 100, 84, 116 }, -{ 78, 110, 94, 126 }, -{ 74, 106, 90, 122 }, -{ 76, 108, 92, 124 }, -{ 72, 104, 88, 120 }, -{ 79, 111, 95, 127 }, -{ 77, 109, 93, 125 }, -{ 67, 99, 83, 115 }, -{ 65, 97, 81, 113 }, -{ 73, 105, 89, 121 }, -{ 75, 107, 91, 123 }, -{ 69, 101, 85, 117 }, -{ 71, 103, 87, 119 }, -{ 9, 41, 25, 57 }, -{ 5, 37, 21, 53 }, -{ 15, 47, 31, 63 }, -{ 3, 35, 19, 51 }, -{ 7, 39, 23, 55 }, -{ 11, 43, 27, 59 }, -{ 1, 33, 17, 49 }, -{ 13, 45, 29, 61 }, -{ 71, 103, 87, 119 }, -{ 65, 97, 81, 113 }, -{ 73, 105, 89, 121 }, -{ 79, 111, 95, 127 }, -{ 77, 109, 93, 125 }, -{ 75, 107, 91, 123 }, -{ 67, 99, 83, 115 }, -{ 69, 101, 85, 117 } -}; - - __constant__ int c_P8[32][8]; static const int h_P8[32][8] = { { 2, 66, 34, 98, 18, 82, 50, 114 }, @@ -189,49 +116,6 @@ static const int h_Q8[32][8] = { { 133, 197, 165, 229, 149, 213, 181, 245 }, }; - __constant__ int c_FFT64_8_8_Twiddle[64]; - static const int h_FFT64_8_8_Twiddle[64] = { - 1, 1, 1, 1, 1, 1, 1, 1, - 1, 2, 4, 8, 16, 32, 64, 128, - 1, 60, 2, 120, 4, -17, 8, -34, - 1, 120, 8, -68, 64, -30, -2, 17, - 1, 46, 60, -67, 2, 92, 120, 123, - 1, 92, -17, -22, 32, 117, -30, 67, - 1, -67, 120, -73, 8, -22, -68, -70, - 1, 123, -34, -70, 128, 67, 17, 35, - }; - - __constant__ int c_FFT128_2_64_Twiddle[64]; - static const int h_FFT128_2_64_Twiddle[64] = { - 1, -118, 46, -31, 60, 116, -67, -61, - 2, 21, 92, -62, 120, -25, 123, -122, - 4, 42, -73, -124, -17, -50, -11, 13, - 8, 84, 111, 9, -34, -100, -22, 26, - 16, -89, -35, 18, -68, 57, -44, 52, - 32, 79, -70, 36, 121, 114, -88, 104, - 64, -99, 117, 72, -15, -29, 81, -49, - 128, 59, -23, -113, -30, -58, -95, -98 - }; - -__constant__ int c_FFT128_16_8_Twiddle[128]; -static const int h_FFT128_16_8_Twiddle[128] = { -1, 1, 1, 1, 1, 1, 1, 1, -1, 2, 4, 8, 16, 32, 64, 128, -1, 60, 2, 120, 4, -17, 8, -34, -1, 120, 8, -68, 64, -30, -2, 17, -1, 46, 60, -67, 2, 92, 120, 123, -1, 92, -17, -22, 32, 117, -30, 67, -1, -67, 120, -73, 8, -22, -68, -70, -1, 123, -34, -70, 128, 67, 17, 35, -1, -118, 46, -31, 60, 116, -67, -61, -1, 21, -73, 9, -68, 114, 81, -98, -1, 116, 92, -122, -17, 84, -22, 18, -1, -25, 111, 52, -15, 118, -123, -9, -1, -31, -67, 21, 120, -122, -73, -50, -1, -62, -11, -89, 121, -49, -46, 25, -1, -61, 123, -50, -34, 18, -70, -99, -1, -122, -22, 114, -30, 62, -111, -79 }; - __constant__ int c_FFT128_8_16_Twiddle[128]; static const int h_FFT128_8_16_Twiddle[128] = { 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, @@ -309,7 +193,7 @@ __device__ __forceinline__ void STEP8_MAJ(const uint32_t *w, const int i, const } } -__device__ __forceinline__ void Round8(uint32_t A[128], const int y[128], int i, +__device__ __forceinline__ void Round8(uint32_t A[32], const int y[256], int i, int r, int s, int t, int u) { int code = i<2? 185: 233; uint32_t w[8][8]; @@ -464,32 +348,36 @@ X(j) = (u-v) << n; \ DO_REDUCE(15); BUTTERFLY( 0, 4, 0); - BUTTERFLY( 8, 12, 0); BUTTERFLY( 1, 5, 2); - BUTTERFLY( 9, 13, 2); BUTTERFLY( 2, 6, 4); - BUTTERFLY(10, 14, 4); BUTTERFLY( 3, 7, 6); + + BUTTERFLY( 8, 12, 0); + BUTTERFLY( 9, 13, 2); + BUTTERFLY(10, 14, 4); BUTTERFLY(11, 15, 6); DO_REDUCE(5); DO_REDUCE(7); - DO_REDUCE(13); + DO_REDUCE(13); DO_REDUCE(15); BUTTERFLY( 0, 2, 0); + BUTTERFLY( 1, 3, 4); BUTTERFLY( 4, 6, 0); + BUTTERFLY( 5, 7, 4); + BUTTERFLY( 8, 10, 0); BUTTERFLY(12, 14, 0); - BUTTERFLY( 1, 3, 4); - BUTTERFLY( 5, 7, 4); BUTTERFLY( 9, 11, 4); BUTTERFLY(13, 15, 4); + BUTTERFLY( 0, 1, 0); BUTTERFLY( 2, 3, 0); BUTTERFLY( 4, 5, 0); BUTTERFLY( 6, 7, 0); + BUTTERFLY( 8, 9, 0); BUTTERFLY(10, 11, 0); BUTTERFLY(12, 13, 0); @@ -518,65 +406,6 @@ X(j) = (u-v) << n; \ #undef BUTTERFLY } -__device__ __forceinline__ void FFT_64(int *y) { - - /* - * FFT_64 using w=46 as 64th root of unity - * decimation in frequency (DIF) radix-8 NTT. - * Output data is in revbin_permuted order. - */ - - int i; - - /* - * Begin with 8 parallels DIF FFT_8. - */ -#pragma unroll 8 - for (i=0; i<8; i++) { - FFT_8(y+i,8); - } - - /* - * Multiply by twiddle factors - */ -#pragma unroll 56 - for (i=8; i<64; i++) - if (i & 7) y[i] = REDUCE(y[i]*c_FFT64_8_8_Twiddle[i]); - - /* - * Finish with 8 paralles DIF FFT_8. - */ -#pragma unroll 8 - for (i=0; i<8; i++) { - FFT_8(y+8*i,1); - } - -} - -__device__ __forceinline__ void FFT_128_halfzero(int *y) { - - /* - * FFT_128 using w=139 as 128th root of unity. - * Decimation in frequency (DIF) NTT. - * Output data is in revbin_permuted order. - * In place. - */ - - const int tmp = y[63]; - int i; - -#pragma unroll 63 - for (i=0; i<63; i++) - y[64+i] = REDUCE(y[i] * c_FFT128_2_64_Twiddle[i]); - - /* handle X^127 */ - y[63] = REDUCE(tmp + 1); - y[127] = REDUCE((tmp - 1) * c_FFT128_2_64_Twiddle[63]); - - FFT_64(y); - FFT_64(y+64); -} - __device__ __forceinline__ void FFT_128_full(int *y) { int i; @@ -596,7 +425,7 @@ __device__ __forceinline__ void FFT_128_full(int *y) { } -__device__ __forceinline__ void FFT_256_halfzero(int *y) { +__device__ __forceinline__ void FFT_256_halfzero(int y[256]) { int i; @@ -621,7 +450,7 @@ __device__ __forceinline__ void FFT_256_halfzero(int *y) { } -__device__ __forceinline__ void SIMD_Compress(uint32_t A[128], const int *expanded, const uint32_t *M) { +__device__ __forceinline__ void SIMD_Compress(uint32_t A[32], const int *expanded, const uint32_t *M) { uint32_t IV[4][8]; int i; @@ -666,7 +495,7 @@ __device__ __forceinline__ void SIMD_Compress(uint32_t A[128], const int *expand __device__ __forceinline__ void SIMDHash(const uint32_t *data, uint32_t *hashval) { - uint32_t A[128]; + uint32_t A[32]; int i; uint32_t buffer[16]; @@ -699,7 +528,7 @@ __device__ __forceinline__ void SIMDHash(const uint32_t *data, uint32_t *hashval } /* Compression Function */ - SIMD_Compress(A, expanded, buffer); + SIMD_Compress(A, expanded, buffer); /* Padding Round with known input (hence the FFT can be precomputed) */ buffer[0] = 512; @@ -735,13 +564,8 @@ __host__ void x11_simd512_cpu_init(int thr_id, int threads) { cudaMemcpyToSymbol( c_IV_512, h_IV_512, sizeof(h_IV_512), 0, cudaMemcpyHostToDevice); cudaMemcpyToSymbol( c_FFT, h_FFT, sizeof(h_FFT), 0, cudaMemcpyHostToDevice); - cudaMemcpyToSymbol( c_P4, h_P4, sizeof(h_P4), 0, cudaMemcpyHostToDevice); - cudaMemcpyToSymbol( c_Q4, h_Q4, sizeof(h_Q4), 0, cudaMemcpyHostToDevice); cudaMemcpyToSymbol( c_P8, h_P8, sizeof(h_P8), 0, cudaMemcpyHostToDevice); cudaMemcpyToSymbol( c_Q8, h_Q8, sizeof(h_Q8), 0, cudaMemcpyHostToDevice); - cudaMemcpyToSymbol( c_FFT64_8_8_Twiddle, h_FFT64_8_8_Twiddle, sizeof(h_FFT64_8_8_Twiddle), 0, cudaMemcpyHostToDevice); - cudaMemcpyToSymbol( c_FFT128_2_64_Twiddle, h_FFT128_2_64_Twiddle, sizeof(h_FFT128_2_64_Twiddle), 0, cudaMemcpyHostToDevice); - cudaMemcpyToSymbol( c_FFT128_16_8_Twiddle, h_FFT128_16_8_Twiddle, sizeof(h_FFT128_16_8_Twiddle), 0, cudaMemcpyHostToDevice); cudaMemcpyToSymbol( c_FFT128_8_16_Twiddle, h_FFT128_8_16_Twiddle, sizeof(h_FFT128_8_16_Twiddle), 0, cudaMemcpyHostToDevice); cudaMemcpyToSymbol( c_FFT256_2_128_Twiddle, h_FFT256_2_128_Twiddle, sizeof(h_FFT256_2_128_Twiddle), 0, cudaMemcpyHostToDevice); } diff --git a/x11/x11.cu b/x11/x11.cu index adaa3a6..f5382ea 100644 --- a/x11/x11.cu +++ b/x11/x11.cu @@ -168,7 +168,7 @@ extern "C" int scanhash_x11(int thr_id, uint32_t *pdata, const uint32_t Htarg = ptarget[7]; - const int throughput = 256*256; // 100; + const int throughput = 256*256*16; static bool init[8] = {0,0,0,0,0,0,0,0}; if (!init[thr_id])