GOSTSec
/
ccminer
mirror of https://github.com/GOSTSec/ccminer
1
0
Fork 0
Code Issues Releases Wiki Activity
Browse Source

myr-gr: remove unused allocated memory + pascal tweak 

+ cleanup...
2upstream
Tanguy Pruvot 8 years ago
parent
39d0e42b11
commit
3fe4dda4c1
2 changed files with 310 additions and 323 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Unified View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 616
      cuda_myriadgroestl.cu
  2. 17
      myriadgroestl.cpp

616
cuda_myriadgroestl.cu
Unescape View File

@ -5,6 +5,11 @@
#include "cuda_helper.h" #include "cuda_helper.h"
#ifdef __INTELLISENSE__
#define __CUDA_ARCH__ 500
#define __funnelshift_r(x,y,n) (x >> n)
#endif
#if __CUDA_ARCH__ >= 300 #if __CUDA_ARCH__ >= 300
// 64 Registers Variant for Compute 3.0 // 64 Registers Variant for Compute 3.0
#include "quark/groestl_functions_quad.h" #include "quark/groestl_functions_quad.h"
@ -21,238 +26,240 @@ __constant__ uint32_t myriadgroestl_gpu_msg[32];
// muss expandiert werden // muss expandiert werden
__constant__ uint32_t myr_sha256_gpu_constantTable[64]; __constant__ uint32_t myr_sha256_gpu_constantTable[64];
__constant__ uint32_t myr_sha256_gpu_constantTable2[64]; __constant__ uint32_t myr_sha256_gpu_constantTable2[64];
__constant__ uint32_t myr_sha256_gpu_hashTable[8];
const uint32_t myr_sha256_cpu_constantTable[] = {
uint32_t myr_sha256_cpu_hashTable[] = { 0x428a2f98, 0x71374491, 0xb5c0fbcf, 0xe9b5dba5, 0x3956c25b, 0x59f111f1, 0x923f82a4, 0xab1c5ed5,
0x6a09e667, 0xbb67ae85, 0x3c6ef372, 0xa54ff53a, 0x510e527f, 0x9b05688c, 0x1f83d9ab, 0x5be0cd19 }; 0xd807aa98, 0x12835b01, 0x243185be, 0x550c7dc3, 0x72be5d74, 0x80deb1fe, 0x9bdc06a7, 0xc19bf174,
uint32_t myr_sha256_cpu_constantTable[] = { 0xe49b69c1, 0xefbe4786, 0x0fc19dc6, 0x240ca1cc, 0x2de92c6f, 0x4a7484aa, 0x5cb0a9dc, 0x76f988da,
0x428a2f98, 0x71374491, 0xb5c0fbcf, 0xe9b5dba5, 0x3956c25b, 0x59f111f1, 0x923f82a4, 0xab1c5ed5, 0x983e5152, 0xa831c66d, 0xb00327c8, 0xbf597fc7, 0xc6e00bf3, 0xd5a79147, 0x06ca6351, 0x14292967,
0xd807aa98, 0x12835b01, 0x243185be, 0x550c7dc3, 0x72be5d74, 0x80deb1fe, 0x9bdc06a7, 0xc19bf174, 0x27b70a85, 0x2e1b2138, 0x4d2c6dfc, 0x53380d13, 0x650a7354, 0x766a0abb, 0x81c2c92e, 0x92722c85,
0xe49b69c1, 0xefbe4786, 0x0fc19dc6, 0x240ca1cc, 0x2de92c6f, 0x4a7484aa, 0x5cb0a9dc, 0x76f988da, 0xa2bfe8a1, 0xa81a664b, 0xc24b8b70, 0xc76c51a3, 0xd192e819, 0xd6990624, 0xf40e3585, 0x106aa070,
0x983e5152, 0xa831c66d, 0xb00327c8, 0xbf597fc7, 0xc6e00bf3, 0xd5a79147, 0x06ca6351, 0x14292967, 0x19a4c116, 0x1e376c08, 0x2748774c, 0x34b0bcb5, 0x391c0cb3, 0x4ed8aa4a, 0x5b9cca4f, 0x682e6ff3,
0x27b70a85, 0x2e1b2138, 0x4d2c6dfc, 0x53380d13, 0x650a7354, 0x766a0abb, 0x81c2c92e, 0x92722c85, 0x748f82ee, 0x78a5636f, 0x84c87814, 0x8cc70208, 0x90befffa, 0xa4506ceb, 0xbef9a3f7, 0xc67178f2,
0xa2bfe8a1, 0xa81a664b, 0xc24b8b70, 0xc76c51a3, 0xd192e819, 0xd6990624, 0xf40e3585, 0x106aa070,
0x19a4c116, 0x1e376c08, 0x2748774c, 0x34b0bcb5, 0x391c0cb3, 0x4ed8aa4a, 0x5b9cca4f, 0x682e6ff3,
0x748f82ee, 0x78a5636f, 0x84c87814, 0x8cc70208, 0x90befffa, 0xa4506ceb, 0xbef9a3f7, 0xc67178f2,
}; };
uint32_t myr_sha256_cpu_w2Table[] = { const uint32_t myr_sha256_cpu_w2Table[] = {
0x80000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x80000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000200, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000200,
0x80000000, 0x01400000, 0x00205000, 0x00005088, 0x22000800, 0x22550014, 0x05089742, 0xa0000020, 0x80000000, 0x01400000, 0x00205000, 0x00005088, 0x22000800, 0x22550014, 0x05089742, 0xa0000020,
0x5a880000, 0x005c9400, 0x0016d49d, 0xfa801f00, 0xd33225d0, 0x11675959, 0xf6e6bfda, 0xb30c1549, 0x5a880000, 0x005c9400, 0x0016d49d, 0xfa801f00, 0xd33225d0, 0x11675959, 0xf6e6bfda, 0xb30c1549,
0x08b2b050, 0x9d7c4c27, 0x0ce2a393, 0x88e6e1ea, 0xa52b4335, 0x67a16f49, 0xd732016f, 0x4eeb2e91, 0x08b2b050, 0x9d7c4c27, 0x0ce2a393, 0x88e6e1ea, 0xa52b4335, 0x67a16f49, 0xd732016f, 0x4eeb2e91,
0x5dbf55e5, 0x8eee2335, 0xe2bc5ec2, 0xa83f4394, 0x45ad78f7, 0x36f3d0cd, 0xd99c05e8, 0xb0511dc7, 0x5dbf55e5, 0x8eee2335, 0xe2bc5ec2, 0xa83f4394, 0x45ad78f7, 0x36f3d0cd, 0xd99c05e8, 0xb0511dc7,
0x69bc7ac4, 0xbd11375b, 0xe3ba71e5, 0x3b209ff2, 0x18feee17, 0xe25ad9e7, 0x13375046, 0x0515089d, 0x69bc7ac4, 0xbd11375b, 0xe3ba71e5, 0x3b209ff2, 0x18feee17, 0xe25ad9e7, 0x13375046, 0x0515089d,
0x4f0d0f04, 0x2627484e, 0x310128d2, 0xc668b434, 0x420841cc, 0x62d311b8, 0xe59ba771, 0x85a7a484 }; 0x4f0d0f04, 0x2627484e, 0x310128d2, 0xc668b434, 0x420841cc, 0x62d311b8, 0xe59ba771, 0x85a7a484
};
#define SWAB32(x) ( ((x & 0x000000FF) << 24) | ((x & 0x0000FF00) << 8) | ((x & 0x00FF0000) >> 8) | ((x & 0xFF000000) >> 24) ) #define SWAB32(x) cuda_swab32(x)
#if __CUDA_ARCH__ < 320 #if __CUDA_ARCH__ < 320
// Kepler (Compute 3.0) // Kepler (Compute 3.0)
#define ROTR32(x, n) (((x) >> (n)) | ((x) << (32 - (n)))) #define ROTR32(x, n) (((x) >> (n)) | ((x) << (32 - (n))))
#else #else
// Kepler (Compute 3.5) // Kepler (Compute 3.5)
#define ROTR32(x, n) __funnelshift_r( (x), (x), (n) ) #define ROTR32(x, n) __funnelshift_r( (x), (x), (n) )
#endif #endif
#define R(x, n) ((x) >> (n))
#define Ch(x, y, z) ((x & (y ^ z)) ^ z) #define R(x, n) ((x) >> (n))
#define Ch(x, y, z) ((x & (y ^ z)) ^ z)
#define Maj(x, y, z) ((x & (y | z)) | (y & z)) #define Maj(x, y, z) ((x & (y | z)) | (y & z))
#define S0(x) (ROTR32(x, 2) ^ ROTR32(x, 13) ^ ROTR32(x, 22)) #define S0(x) (ROTR32(x, 2) ^ ROTR32(x, 13) ^ ROTR32(x, 22))
#define S1(x) (ROTR32(x, 6) ^ ROTR32(x, 11) ^ ROTR32(x, 25)) #define S1(x) (ROTR32(x, 6) ^ ROTR32(x, 11) ^ ROTR32(x, 25))
#define s0(x) (ROTR32(x, 7) ^ ROTR32(x, 18) ^ R(x, 3)) #define s0(x) (ROTR32(x, 7) ^ ROTR32(x, 18) ^ R(x, 3))
#define s1(x) (ROTR32(x, 17) ^ ROTR32(x, 19) ^ R(x, 10)) #define s1(x) (ROTR32(x, 17) ^ ROTR32(x, 19) ^ R(x, 10))
__device__ void myriadgroestl_gpu_sha256(uint32_t *message) __device__ void myriadgroestl_gpu_sha256(uint32_t *message)
{ {
uint32_t W1[16]; uint32_t regs[8], hash[8];
uint32_t W2[16]; const uint32_t myr_sha256_gpu_hashTable[8] = {
0x6a09e667, 0xbb67ae85, 0x3c6ef372, 0xa54ff53a, 0x510e527f, 0x9b05688c, 0x1f83d9ab, 0x5be0cd19
// Initialisiere die register a bis h mit der Hash-Tabelle };
uint32_t regs[8];
uint32_t hash[8]; // pre
#pragma unroll 8
// pre for (int k=0; k < 8; k++)
#pragma unroll 8 {
for (int k=0; k < 8; k++) regs[k] = myr_sha256_gpu_hashTable[k];
{ hash[k] = regs[k];
regs[k] = myr_sha256_gpu_hashTable[k]; }
hash[k] = regs[k];
} uint32_t W1[16];
#pragma unroll 16
#pragma unroll 16 for(int k=0; k<16; k++)
for(int k=0;k<16;k++) W1[k] = SWAB32(message[k]);
W1[k] = SWAB32(message[k]);
// Progress W1
// Progress W1 #pragma unroll 16
#pragma unroll 16 for(int j=0; j<16; j++)
for(int j=0;j<16;j++) {
{ uint32_t T1, T2;
uint32_t T1, T2; T1 = regs[7] + S1(regs[4]) + Ch(regs[4], regs[5], regs[6]) + myr_sha256_gpu_constantTable[j] + W1[j];
T1 = regs[7] + S1(regs[4]) + Ch(regs[4], regs[5], regs[6]) + myr_sha256_gpu_constantTable[j] + W1[j]; T2 = S0(regs[0]) + Maj(regs[0], regs[1], regs[2]);
T2 = S0(regs[0]) + Maj(regs[0], regs[1], regs[2]);
#pragma unroll 7
#pragma unroll 7 for (int k=6; k >= 0; k--) regs[k+1] = regs[k];
for (int k=6; k >= 0; k--) regs[k+1] = regs[k]; regs[0] = T1 + T2;
regs[0] = T1 + T2; regs[4] += T1;
regs[4] += T1; }
}
// Progress W2...W3
// Progress W2...W3 uint32_t W2[16];
////// PART 1
#pragma unroll 2 ////// PART 1
for(int j=0;j<2;j++) #pragma unroll 2
W2[j] = s1(W1[14+j]) + W1[9+j] + s0(W1[1+j]) + W1[j]; for(int j=0; j<2; j++)
#pragma unroll 5 W2[j] = s1(W1[14+j]) + W1[9+j] + s0(W1[1+j]) + W1[j];
for(int j=2;j<7;j++)
W2[j] = s1(W2[j-2]) + W1[9+j] + s0(W1[1+j]) + W1[j]; #pragma unroll 5
for(int j=2;j<7;j++)
#pragma unroll 8 W2[j] = s1(W2[j-2]) + W1[9+j] + s0(W1[1+j]) + W1[j];
for(int j=7;j<15;j++)
W2[j] = s1(W2[j-2]) + W2[j-7] + s0(W1[1+j]) + W1[j]; #pragma unroll 8
for(int j=7; j<15; j++)
W2[15] = s1(W2[13]) + W2[8] + s0(W2[0]) + W1[15]; W2[j] = s1(W2[j-2]) + W2[j-7] + s0(W1[1+j]) + W1[j];
// Rundenfunktion W2[15] = s1(W2[13]) + W2[8] + s0(W2[0]) + W1[15];
#pragma unroll 16
for(int j=0;j<16;j++) // Round function
{ #pragma unroll 16
uint32_t T1, T2; for(int j=0; j<16; j++)
T1 = regs[7] + S1(regs[4]) + Ch(regs[4], regs[5], regs[6]) + myr_sha256_gpu_constantTable[j + 16] + W2[j]; {
T2 = S0(regs[0]) + Maj(regs[0], regs[1], regs[2]); uint32_t T1, T2;
T1 = regs[7] + S1(regs[4]) + Ch(regs[4], regs[5], regs[6]) + myr_sha256_gpu_constantTable[j + 16] + W2[j];
#pragma unroll 7 T2 = S0(regs[0]) + Maj(regs[0], regs[1], regs[2]);
for (int l=6; l >= 0; l--) regs[l+1] = regs[l];
regs[0] = T1 + T2; #pragma unroll 7
regs[4] += T1; for (int l=6; l >= 0; l--) regs[l+1] = regs[l];
} regs[0] = T1 + T2;
regs[4] += T1;
////// PART 2 }
#pragma unroll 2
for(int j=0;j<2;j++) ////// PART 2
W1[j] = s1(W2[14+j]) + W2[9+j] + s0(W2[1+j]) + W2[j]; #pragma unroll 2
#pragma unroll 5 for(int j=0; j<2; j++)
for(int j=2;j<7;j++) W1[j] = s1(W2[14+j]) + W2[9+j] + s0(W2[1+j]) + W2[j];
W1[j] = s1(W1[j-2]) + W2[9+j] + s0(W2[1+j]) + W2[j]; #pragma unroll 5
for(int j=2; j<7; j++)
#pragma unroll 8 W1[j] = s1(W1[j-2]) + W2[9+j] + s0(W2[1+j]) + W2[j];
for(int j=7;j<15;j++)
W1[j] = s1(W1[j-2]) + W1[j-7] + s0(W2[1+j]) + W2[j]; #pragma unroll 8
for(int j=7; j<15; j++)
W1[15] = s1(W1[13]) + W1[8] + s0(W1[0]) + W2[15]; W1[j] = s1(W1[j-2]) + W1[j-7] + s0(W2[1+j]) + W2[j];
// Rundenfunktion W1[15] = s1(W1[13]) + W1[8] + s0(W1[0]) + W2[15];
#pragma unroll 16
for(int j=0;j<16;j++) // Round function
{ #pragma unroll 16
uint32_t T1, T2; for(int j=0; j<16; j++)
T1 = regs[7] + S1(regs[4]) + Ch(regs[4], regs[5], regs[6]) + myr_sha256_gpu_constantTable[j + 32] + W1[j]; {
T2 = S0(regs[0]) + Maj(regs[0], regs[1], regs[2]); uint32_t T1, T2;
T1 = regs[7] + S1(regs[4]) + Ch(regs[4], regs[5], regs[6]) + myr_sha256_gpu_constantTable[j + 32] + W1[j];
#pragma unroll 7 T2 = S0(regs[0]) + Maj(regs[0], regs[1], regs[2]);
for (int l=6; l >= 0; l--) regs[l+1] = regs[l];
regs[0] = T1 + T2; #pragma unroll 7
regs[4] += T1; for (int l=6; l >= 0; l--) regs[l+1] = regs[l];
} regs[0] = T1 + T2;
regs[4] += T1;
////// PART 3 }
#pragma unroll 2
for(int j=0;j<2;j++) ////// PART 3
W2[j] = s1(W1[14+j]) + W1[9+j] + s0(W1[1+j]) + W1[j]; #pragma unroll 2
#pragma unroll 5 for(int j=0; j<2; j++)
for(int j=2;j<7;j++) W2[j] = s1(W1[14+j]) + W1[9+j] + s0(W1[1+j]) + W1[j];
W2[j] = s1(W2[j-2]) + W1[9+j] + s0(W1[1+j]) + W1[j]; #pragma unroll 5
for(int j=2; j<7; j++)
#pragma unroll 8 W2[j] = s1(W2[j-2]) + W1[9+j] + s0(W1[1+j]) + W1[j];
for(int j=7;j<15;j++)
W2[j] = s1(W2[j-2]) + W2[j-7] + s0(W1[1+j]) + W1[j]; #pragma unroll 8
for(int j=7; j<15; j++)
W2[15] = s1(W2[13]) + W2[8] + s0(W2[0]) + W1[15]; W2[j] = s1(W2[j-2]) + W2[j-7] + s0(W1[1+j]) + W1[j];
// Rundenfunktion W2[15] = s1(W2[13]) + W2[8] + s0(W2[0]) + W1[15];
#pragma unroll 16
for(int j=0;j<16;j++) // Round function
{ #pragma unroll 16
uint32_t T1, T2; for(int j=0; j<16; j++)
T1 = regs[7] + S1(regs[4]) + Ch(regs[4], regs[5], regs[6]) + myr_sha256_gpu_constantTable[j + 48] + W2[j]; {
T2 = S0(regs[0]) + Maj(regs[0], regs[1], regs[2]); uint32_t T1, T2;
T1 = regs[7] + S1(regs[4]) + Ch(regs[4], regs[5], regs[6]) + myr_sha256_gpu_constantTable[j + 48] + W2[j];
#pragma unroll 7 T2 = S0(regs[0]) + Maj(regs[0], regs[1], regs[2]);
for (int l=6; l >= 0; l--) regs[l+1] = regs[l];
regs[0] = T1 + T2; #pragma unroll 7
regs[4] += T1; for (int l=6; l >= 0; l--) regs[l+1] = regs[l];
} regs[0] = T1 + T2;
regs[4] += T1;
#pragma unroll 8 }
for(int k=0;k<8;k++)
hash[k] += regs[k]; #pragma unroll 8
for(int k=0; k<8; k++)
///// hash[k] += regs[k];
///// Zweite Runde (wegen Msg-Padding)
///// /////
#pragma unroll 8 ///// 2nd Round (wegen Msg-Padding)
for(int k=0;k<8;k++) /////
regs[k] = hash[k]; #pragma unroll 8
for(int k=0; k<8; k++)
// Progress W1 regs[k] = hash[k];
#pragma unroll 64
for(int j=0;j<64;j++) // Progress W1
{ #pragma unroll 64
uint32_t T1, T2; for(int j=0; j<64; j++)
T1 = regs[7] + S1(regs[4]) + Ch(regs[4], regs[5], regs[6]) + myr_sha256_gpu_constantTable2[j]; {
T2 = S0(regs[0]) + Maj(regs[0], regs[1], regs[2]); uint32_t T1, T2;
T1 = regs[7] + S1(regs[4]) + Ch(regs[4], regs[5], regs[6]) + myr_sha256_gpu_constantTable2[j];
#pragma unroll 7 T2 = S0(regs[0]) + Maj(regs[0], regs[1], regs[2]);
for (int k=6; k >= 0; k--) regs[k+1] = regs[k];
regs[0] = T1 + T2; #pragma unroll 7
regs[4] += T1; for (int k=6; k >= 0; k--) regs[k+1] = regs[k];
} regs[0] = T1 + T2;
regs[4] += T1;
#pragma unroll 8 }
for(int k=0;k<8;k++)
hash[k] += regs[k]; #pragma unroll 8
for(int k=0; k<8; k++)
//// FERTIG hash[k] += regs[k];
#pragma unroll 8 //// Close
for(int k=0;k<8;k++)
message[k] = SWAB32(hash[k]); #pragma unroll 8
for(int k=0; k<8; k++)
message[k] = SWAB32(hash[k]);
} }
__global__ void __launch_bounds__(256, 4) __global__ void __launch_bounds__(256, 4)
myriadgroestl_gpu_hash_quad(uint32_t threads, uint32_t startNounce, uint32_t *hashBuffer) myriadgroestl_gpu_hash_quad(uint32_t threads, uint32_t startNounce, uint32_t *hashBuffer)
{ {
#if __CUDA_ARCH__ >= 300 #if __CUDA_ARCH__ >= 300
// durch 4 dividieren, weil jeweils 4 Threads zusammen ein Hash berechnen // durch 4 dividieren, weil jeweils 4 Threads zusammen ein Hash berechnen
uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x) / 4; uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x) / 4;
if (thread < threads) if (thread < threads)
{ {
// GROESTL // GROESTL
uint32_t paddedInput[8]; uint32_t paddedInput[8];
#pragma unroll 8 #pragma unroll 8
for(int k=0;k<8;k++) paddedInput[k] = myriadgroestl_gpu_msg[4*k+threadIdx.x%4]; for(int k=0; k<8; k++)
paddedInput[k] = myriadgroestl_gpu_msg[4*k+threadIdx.x%4];
uint32_t nounce = startNounce + thread;
if ((threadIdx.x % 4) == 3) uint32_t nounce = startNounce + thread;
paddedInput[4] = SWAB32(nounce); // 4*4+3 = 19 if ((threadIdx.x % 4) == 3)
paddedInput[4] = SWAB32(nounce); // 4*4+3 = 19
uint32_t msgBitsliced[8];
to_bitslice_quad(paddedInput, msgBitsliced); uint32_t msgBitsliced[8];
to_bitslice_quad(paddedInput, msgBitsliced);
uint32_t state[8];
uint32_t state[8];
groestl512_progressMessage_quad(state, msgBitsliced);
groestl512_progressMessage_quad(state, msgBitsliced);
uint32_t out_state[16];
from_bitslice_quad(state, out_state); uint32_t out_state[16];
from_bitslice_quad(state, out_state);
if ((threadIdx.x & 0x03) == 0)
{ if ((threadIdx.x & 0x03) == 0)
uint32_t *outpHash = &hashBuffer[16 * thread]; {
#pragma unroll 16 uint32_t *outpHash = &hashBuffer[16 * thread];
for(int k=0;k<16;k++) outpHash[k] = out_state[k]; #pragma unroll 16
} for(int k=0; k<16; k++) outpHash[k] = out_state[k];
} }
}
#endif #endif
} }
@ -260,42 +267,42 @@ __global__ void
myriadgroestl_gpu_hash_quad2(uint32_t threads, uint32_t startNounce, uint32_t *resNounce, uint32_t *hashBuffer) myriadgroestl_gpu_hash_quad2(uint32_t threads, uint32_t startNounce, uint32_t *resNounce, uint32_t *hashBuffer)
{ {
#if __CUDA_ARCH__ >= 300 #if __CUDA_ARCH__ >= 300
uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x); uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x);
if (thread < threads) if (thread < threads)
{ {
uint32_t nounce = startNounce + thread; uint32_t nounce = startNounce + thread;
uint32_t out_state[16]; uint32_t out_state[16];
uint32_t *inpHash = &hashBuffer[16 * thread]; uint32_t *inpHash = &hashBuffer[16 * thread];
#pragma unroll 16
for (int i=0; i < 16; i++) #pragma unroll 16
out_state[i] = inpHash[i]; for (int i=0; i < 16; i++)
out_state[i] = inpHash[i];
myriadgroestl_gpu_sha256(out_state);
myriadgroestl_gpu_sha256(out_state);
int i, position = -1;
bool rc = true; int i, position = -1;
bool rc = true;
#pragma unroll 8
for (i = 7; i >= 0; i--) { #pragma unroll 8
if (out_state[i] > pTarget[i]) { for (i = 7; i >= 0; i--) {
if(position < i) { if (out_state[i] > pTarget[i]) {
position = i; if(position < i) {
rc = false; position = i;
} rc = false;
} }
if (out_state[i] < pTarget[i]) { }
if(position < i) { if (out_state[i] < pTarget[i]) {
position = i; if(position < i) {
rc = true; position = i;
} rc = true;
} }
} }
}
if(rc == true)
if(resNounce[0] > nounce) if(rc && resNounce[0] > nounce)
resNounce[0] = nounce; resNounce[0] = nounce;
} }
#endif #endif
} }
@ -303,93 +310,76 @@ __global__ void
__host__ __host__
void myriadgroestl_cpu_init(int thr_id, uint32_t threads) void myriadgroestl_cpu_init(int thr_id, uint32_t threads)
{ {
cudaMemcpyToSymbol( myr_sha256_gpu_hashTable, uint32_t temp[64];
myr_sha256_cpu_hashTable, for(int i=0; i<64; i++)
sizeof(uint32_t) * 8 ); temp[i] = myr_sha256_cpu_w2Table[i] + myr_sha256_cpu_constantTable[i];
cudaMemcpyToSymbol( myr_sha256_gpu_constantTable, cudaMemcpyToSymbol( myr_sha256_gpu_constantTable2,
myr_sha256_cpu_constantTable, temp,
sizeof(uint32_t) * 64 ); sizeof(uint32_t) * 64 );
// zweite CPU-Tabelle bauen und auf die GPU laden cudaMemcpyToSymbol( myr_sha256_gpu_constantTable,
uint32_t temp[64]; myr_sha256_cpu_constantTable,
for(int i=0;i<64;i++) sizeof(uint32_t) * 64 );
temp[i] = myr_sha256_cpu_w2Table[i] + myr_sha256_cpu_constantTable[i];
cudaMemcpyToSymbol( myr_sha256_gpu_constantTable2, cudaMalloc(&d_outputHashes[thr_id], (size_t) 64 * threads);
temp, cudaMalloc(&d_resultNonce[thr_id], sizeof(uint32_t));
sizeof(uint32_t) * 64 );
// Speicher für Gewinner-Nonce belegen
cudaMalloc(&d_resultNonce[thr_id], sizeof(uint32_t));
// Speicher für temporäreHashes
cudaMalloc(&d_outputHashes[thr_id], 16*sizeof(uint32_t)*threads);
} }
__host__ __host__
void myriadgroestl_cpu_free(int thr_id) void myriadgroestl_cpu_free(int thr_id)
{ {
cudaFree(d_resultNonce[thr_id]); cudaFree(d_outputHashes[thr_id]);
cudaFree(d_outputHashes[thr_id]); cudaFree(d_resultNonce[thr_id]);
} }
__host__ __host__
void myriadgroestl_cpu_setBlock(int thr_id, void *data, void *pTargetIn) void myriadgroestl_cpu_setBlock(int thr_id, void *data, void *pTargetIn)
{ {
// Nachricht expandieren und setzen // Nachricht expandieren und setzen
uint32_t msgBlock[32]; uint32_t msgBlock[32] = { 0 };
memcpy(&msgBlock[0], data, 80);
memset(msgBlock, 0, sizeof(uint32_t) * 32); // Erweitere die Nachricht auf den Nachrichtenblock (padding)
memcpy(&msgBlock[0], data, 80); // Unsere Nachricht hat 80 Byte
msgBlock[20] = 0x80;
msgBlock[31] = 0x01000000;
// Erweitere die Nachricht auf den Nachrichtenblock (padding) // groestl512 braucht hierfür keinen CPU-Code (die einzige Runde wird
// Unsere Nachricht hat 80 Byte // auf der GPU ausgeführt)
msgBlock[20] = 0x80;
msgBlock[31] = 0x01000000;
// groestl512 braucht hierfür keinen CPU-Code (die einzige Runde wird // Blockheader setzen (korrekte Nonce und Hefty Hash fehlen da drin noch)
// auf der GPU ausgeführt) cudaMemcpyToSymbol(myriadgroestl_gpu_msg, msgBlock, 128);
// Blockheader setzen (korrekte Nonce und Hefty Hash fehlen da drin noch) cudaMemset(d_resultNonce[thr_id], 0xFF, sizeof(uint32_t));
cudaMemcpyToSymbol( myriadgroestl_gpu_msg, cudaMemcpyToSymbol(pTarget, pTargetIn, 32);
msgBlock,
128);
cudaMemset(d_resultNonce[thr_id], 0xFF, sizeof(uint32_t));
cudaMemcpyToSymbol( pTarget,
pTargetIn,
sizeof(uint32_t) * 8 );
} }
__host__ __host__
void myriadgroestl_cpu_hash(int thr_id, uint32_t threads, uint32_t startNounce, void *outputHashes, uint32_t *nounce) void myriadgroestl_cpu_hash(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *resNounce)
{ {
uint32_t threadsperblock = 256; uint32_t threadsperblock = 256;
// Compute 3.0 benutzt die registeroptimierte Quad Variante mit Warp Shuffle
// mit den Quad Funktionen brauchen wir jetzt 4 threads pro Hash, daher Faktor 4 bei der Blockzahl
const int factor=4;
// Größe des dynamischen Shared Memory Bereichs // Compute 3.0 benutzt die registeroptimierte Quad Variante mit Warp Shuffle
size_t shared_size = 0; // mit den Quad Funktionen brauchen wir jetzt 4 threads pro Hash, daher Faktor 4 bei der Blockzahl
const int factor = 4;
cudaMemset(d_resultNonce[thr_id], 0xFF, sizeof(uint32_t)); cudaMemset(d_resultNonce[thr_id], 0xFF, sizeof(uint32_t));
// berechne wie viele Thread Blocks wir brauchen // berechne wie viele Thread Blocks wir brauchen
dim3 grid(factor*((threads + threadsperblock-1)/threadsperblock)); dim3 grid(factor*((threads + threadsperblock-1)/threadsperblock));
dim3 block(threadsperblock); dim3 block(threadsperblock);
if (device_sm[device_map[thr_id]] < 300) { if (device_sm[device_map[thr_id]] < 300) {
printf("Sorry, This algo is not supported by this GPU arch (SM 3.0 required)"); printf("Sorry, This algo is not supported by this GPU arch (SM 3.0 required)");
return; return;
} }
myriadgroestl_gpu_hash_quad<<<grid, block, shared_size>>>(threads, startNounce, d_outputHashes[thr_id]); myriadgroestl_gpu_hash_quad <<< grid, block >>> (threads, startNounce, d_outputHashes[thr_id]);
dim3 grid2((threads + threadsperblock-1)/threadsperblock); dim3 grid2((threads + threadsperblock-1)/threadsperblock);
myriadgroestl_gpu_hash_quad2<<<grid2, block, shared_size>>>(threads, startNounce, d_resultNonce[thr_id], d_outputHashes[thr_id]); myriadgroestl_gpu_hash_quad2 <<< grid2, block >>> (threads, startNounce, d_resultNonce[thr_id], d_outputHashes[thr_id]);
// Strategisches Sleep Kommando zur Senkung der CPU Last // Strategisches Sleep Kommando zur Senkung der CPU Last
MyStreamSynchronize(NULL, 0, thr_id); MyStreamSynchronize(NULL, 0, thr_id);
cudaMemcpy(nounce, d_resultNonce[thr_id], sizeof(uint32_t), cudaMemcpyDeviceToHost); cudaMemcpy(resNounce, d_resultNonce[thr_id], sizeof(uint32_t), cudaMemcpyDeviceToHost);
} }

17
myriadgroestl.cpp
Unescape View File

@ -10,7 +10,7 @@
void myriadgroestl_cpu_init(int thr_id, uint32_t threads); void myriadgroestl_cpu_init(int thr_id, uint32_t threads);
void myriadgroestl_cpu_free(int thr_id); void myriadgroestl_cpu_free(int thr_id);
void myriadgroestl_cpu_setBlock(int thr_id, void *data, void *pTargetIn); void myriadgroestl_cpu_setBlock(int thr_id, void *data, void *pTargetIn);
void myriadgroestl_cpu_hash(int thr_id, uint32_t threads, uint32_t startNounce, void *outputHashes, uint32_t *nounce); void myriadgroestl_cpu_hash(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *nounce);
void myriadhash(void *state, const void *input) void myriadhash(void *state, const void *input)
{ {
@ -37,18 +37,18 @@ int scanhash_myriad(int thr_id, struct work *work, uint32_t max_nonce, unsigned
uint32_t *pdata = work->data; uint32_t *pdata = work->data;
uint32_t *ptarget = work->target; uint32_t *ptarget = work->target;
uint32_t start_nonce = pdata[19]; uint32_t start_nonce = pdata[19];
uint32_t throughput = cuda_default_throughput(thr_id, 1U << 17); int dev_id = device_map[thr_id];
int intensity = (device_sm[dev_id] >= 600) ? 20 : 18;
uint32_t throughput = cuda_default_throughput(thr_id, 1U << intensity);
if (init[thr_id]) throughput = min(throughput, max_nonce - start_nonce); if (init[thr_id]) throughput = min(throughput, max_nonce - start_nonce);
uint32_t *outputHash = (uint32_t*)malloc(throughput * 64);
if (opt_benchmark) if (opt_benchmark)
ptarget[7] = 0x0000ff; ptarget[7] = 0x0000ff;
// init // init
if(!init[thr_id]) if(!init[thr_id])
{ {
cudaSetDevice(device_map[thr_id]); cudaSetDevice(dev_id);
if (opt_cudaschedule == -1 && gpu_threads == 1) { if (opt_cudaschedule == -1 && gpu_threads == 1) {
cudaDeviceReset(); cudaDeviceReset();
// reduce cpu usage // reduce cpu usage
@ -62,14 +62,13 @@ int scanhash_myriad(int thr_id, struct work *work, uint32_t max_nonce, unsigned
for (int k=0; k < 20; k++) for (int k=0; k < 20; k++)
be32enc(&endiandata[k], pdata[k]); be32enc(&endiandata[k], pdata[k]);
// Context mit dem Endian gedrehten Blockheader vorbereiten (Nonce wird später ersetzt)
myriadgroestl_cpu_setBlock(thr_id, endiandata, (void*)ptarget); myriadgroestl_cpu_setBlock(thr_id, endiandata, (void*)ptarget);
do { do {
// GPU // GPU
uint32_t foundNounce = UINT32_MAX; uint32_t foundNounce = UINT32_MAX;
myriadgroestl_cpu_hash(thr_id, throughput, pdata[19], outputHash, &foundNounce); myriadgroestl_cpu_hash(thr_id, throughput, pdata[19], &foundNounce);
*hashes_done = pdata[19] - start_nonce + throughput; *hashes_done = pdata[19] - start_nonce + throughput;
@ -81,9 +80,8 @@ int scanhash_myriad(int thr_id, struct work *work, uint32_t max_nonce, unsigned
if (vhash[7] <= ptarget[7] && fulltest(vhash, ptarget)) { if (vhash[7] <= ptarget[7] && fulltest(vhash, ptarget)) {
work_set_target_ratio(work, vhash); work_set_target_ratio(work, vhash);
pdata[19] = foundNounce; pdata[19] = foundNounce;
free(outputHash);
return 1; return 1;
} else { } else if (vhash[7] > ptarget[7]) {
gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNounce); gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNounce);
} }
} }
@ -98,7 +96,6 @@ int scanhash_myriad(int thr_id, struct work *work, uint32_t max_nonce, unsigned
*hashes_done = max_nonce - start_nonce; *hashes_done = max_nonce - start_nonce;
free(outputHash);
return 0; return 0;
} }

Write Preview
Loading…
Cancel
Save