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

lyra2: cleanup 

quickly tested with a SM 3.0 binary...
2upstream
Tanguy Pruvot 10 years ago
parent
da2e2528a7
commit
b3188669e2
4 changed files with 93 additions and 328 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. 4
      Makefile.am
  2. 13
      cuda_helper.h
  3. 402
      lyra2/cuda_lyra2.cu
  4. 2
      lyra2/lyra2RE.cu

4
Makefile.am
Unescape View File

@ -1,9 +1,10 @@
# allow to use Host cuda functions in C/C++ # allow to use Host cuda functions in C/C++
DEF_INCLUDES = @CUDA_INCLUDES@ DEF_INCLUDES = @CUDA_INCLUDES@
JANSSON_INCLUDES=
if WANT_JANSSON if WANT_JANSSON
JANSSON_INCLUDES= -I$(top_srcdir)/compat/jansson JANSSON_INCLUDES= -I$(top_srcdir)/compat/jansson
else
JANSSON_INCLUDES=
endif endif
EXTRA_DIST = autogen.sh README.txt LICENSE.txt \ EXTRA_DIST = autogen.sh README.txt LICENSE.txt \
@ -69,6 +70,7 @@ ccminer_LDADD = @LIBCURL@ @JANSSON_LIBS@ @PTHREAD_LIBS@ @WS2_LIBS@ @CUDA_LIBS
ccminer_CPPFLAGS = @LIBCURL_CPPFLAGS@ @OPENMP_CFLAGS@ $(CPPFLAGS) $(PTHREAD_FLAGS) -fno-strict-aliasing $(JANSSON_INCLUDES) $(DEF_INCLUDES) $(nvml_defs) -DSCRYPT_KECCAK512 -DSCRYPT_CHACHA -DSCRYPT_CHOOSE_COMPILETIME ccminer_CPPFLAGS = @LIBCURL_CPPFLAGS@ @OPENMP_CFLAGS@ $(CPPFLAGS) $(PTHREAD_FLAGS) -fno-strict-aliasing $(JANSSON_INCLUDES) $(DEF_INCLUDES) $(nvml_defs) -DSCRYPT_KECCAK512 -DSCRYPT_CHACHA -DSCRYPT_CHOOSE_COMPILETIME
nvcc_ARCH = -gencode=arch=compute_50,code=\"sm_50,compute_50\" nvcc_ARCH = -gencode=arch=compute_50,code=\"sm_50,compute_50\"
#nvcc_ARCH += -gencode=arch=compute_52,code=\"sm_52,compute_52\"
#nvcc_ARCH += -gencode=arch=compute_35,code=\"sm_35,compute_35\" #nvcc_ARCH += -gencode=arch=compute_35,code=\"sm_35,compute_35\"
#nvcc_ARCH += -gencode=arch=compute_30,code=\"sm_30,compute_30\" #nvcc_ARCH += -gencode=arch=compute_30,code=\"sm_30,compute_30\"

13
cuda_helper.h
Unescape View File

@ -428,7 +428,8 @@ static __device__ __forceinline__ uint2 operator* (uint2 a, uint2 b)
} }
// uint2 ROR/ROL methods // uint2 ROR/ROL methods
__device__ __inline__ uint2 ROR2(const uint2 a, const int offset) __device__ __forceinline__
uint2 ROR2(const uint2 a, const int offset)
{ {
uint2 result; uint2 result;
#if __CUDA_ARCH__ > 300 #if __CUDA_ARCH__ > 300
@ -457,8 +458,8 @@ __device__ __inline__ uint2 ROR2(const uint2 a, const int offset)
return result; return result;
} }
__device__ __forceinline__
__inline__ __device__ uint2 ROL2(const uint2 a, const int offset) uint2 ROL2(const uint2 a, const int offset)
{ {
uint2 result; uint2 result;
#if __CUDA_ARCH__ > 300 #if __CUDA_ARCH__ > 300
@ -479,4 +480,10 @@ __inline__ __device__ uint2 ROL2(const uint2 a, const int offset)
return result; return result;
} }
__device__ __forceinline__
uint2 SWAPUINT2(uint2 value)
{
return make_uint2(value.y, value.x);
}
#endif // #ifndef CUDA_HELPER_H #endif // #ifndef CUDA_HELPER_H

402
lyra2/cuda_lyra2.cu
Unescape View File

@ -14,35 +14,12 @@ static __constant__ uint2 blake2b_IV[8] = {
{ 0xfb41bd6b, 0x1f83d9ab }, { 0xfb41bd6b, 0x1f83d9ab },
{ 0x137e2179, 0x5be0cd19 } { 0x137e2179, 0x5be0cd19 }
}; };
// data: 0-4 outputhash 4-8 outputhash 8-16 basil
#define reduceDuplexRowSetup(rowIn, rowInOut, rowOut) { \
for (int i = 0; i < 8; i++) { \
for (int j = 0; j < 12; j++) \
state[j] ^= Matrix[12 * i + j][rowIn] + Matrix[12 * i + j][rowInOut]; \
round_lyra_v35(state); \
for (int j = 0; j < 12; j++) \
Matrix[j + 84 - 12 * i][rowOut] = Matrix[12 * i + j][rowIn] ^ state[j]; \
Matrix[0 + 12 * i][rowInOut] ^= state[11]; \
Matrix[1 + 12 * i][rowInOut] ^= state[0]; \
Matrix[2 + 12 * i][rowInOut] ^= state[1]; \
Matrix[3 + 12 * i][rowInOut] ^= state[2]; \
Matrix[4 + 12 * i][rowInOut] ^= state[3]; \
Matrix[5 + 12 * i][rowInOut] ^= state[4]; \
Matrix[6 + 12 * i][rowInOut] ^= state[5]; \
Matrix[7 + 12 * i][rowInOut] ^= state[6]; \
Matrix[8 + 12 * i][rowInOut] ^= state[7]; \
Matrix[9 + 12 * i][rowInOut] ^= state[8]; \
Matrix[10+ 12 * i][rowInOut] ^= state[9]; \
Matrix[11+ 12 * i][rowInOut] ^= state[10]; \
} \
}
#define reduceDuplexRow(rowIn, rowInOut, rowOut) { \ #define reduceDuplexRow(rowIn, rowInOut, rowOut) { \
for (int i = 0; i < 8; i++) { \ for (int i = 0; i < 8; i++) { \
for (int j = 0; j < 12; j++) \ for (int j = 0; j < 12; j++) \
state[j] ^= Matrix[12 * i + j][rowIn] + Matrix[12 * i + j][rowInOut]; \ state[j] ^= Matrix[12 * i + j][rowIn] + Matrix[12 * i + j][rowInOut]; \
round_lyra_v35(state); \ round_lyra(state); \
for (int j = 0; j < 12; j++) \ for (int j = 0; j < 12; j++) \
Matrix[j + 12 * i][rowOut] ^= state[j]; \ Matrix[j + 12 * i][rowOut] ^= state[j]; \
Matrix[0 + 12 * i][rowInOut] ^= state[11]; \ Matrix[0 + 12 * i][rowInOut] ^= state[11]; \
@ -73,291 +50,70 @@ static __constant__ uint2 blake2b_IV[8] = {
state[9] ^= Matrix[9][in]; \ state[9] ^= Matrix[9][in]; \
state[10] ^= Matrix[10][in]; \ state[10] ^= Matrix[10][in]; \
state[11] ^= Matrix[11][in]; \ state[11] ^= Matrix[11][in]; \
round_lyra_v35(state); \ round_lyra(state); \
round_lyra_v35(state); \ round_lyra(state); \
round_lyra_v35(state); \ round_lyra(state); \
round_lyra_v35(state); \ round_lyra(state); \
round_lyra_v35(state); \ round_lyra(state); \
round_lyra_v35(state); \ round_lyra(state); \
round_lyra_v35(state); \ round_lyra(state); \
round_lyra_v35(state); \ round_lyra(state); \
round_lyra_v35(state); \ round_lyra(state); \
round_lyra_v35(state); \ round_lyra(state); \
round_lyra_v35(state); \ round_lyra(state); \
round_lyra_v35(state); \ round_lyra(state); \
}
//// test version
#define reduceDuplexRowSetup_test(rowIn, rowInOut, rowOut) { \
for (int i = 0; i < 8; i++) { \
for (int j = 0; j < 12; j++) \
state[j] ^= Matrix[j][i][rowIn] + Matrix[j][i][rowInOut]; \
round_lyra_v35(state); \
for (int j = 0; j < 12; j++) \
Matrix[j][7-i][rowOut] = Matrix[j][i][rowIn] ^ state[j]; \
Matrix[0][i][rowInOut] ^= state[11]; \
Matrix[1][i][rowInOut] ^= state[0]; \
Matrix[2][i][rowInOut] ^= state[1]; \
Matrix[3][i][rowInOut] ^= state[2]; \
Matrix[4][i][rowInOut] ^= state[3]; \
Matrix[5][i][rowInOut] ^= state[4]; \
Matrix[6][i][rowInOut] ^= state[5]; \
Matrix[7][i][rowInOut] ^= state[6]; \
Matrix[8][i][rowInOut] ^= state[7]; \
Matrix[9][i][rowInOut] ^= state[8]; \
Matrix[10][i][rowInOut] ^= state[9]; \
Matrix[11][i][rowInOut] ^= state[10]; \
} \
}
#define reduceDuplexRow_test(rowIn, rowInOut, rowOut) { \
for (int i = 0; i < 8; i++) { \
for (int j = 0; j < 12; j++) \
state[j] ^= Matrix[j][i][rowIn] + Matrix[j][i][rowInOut]; \
round_lyra_v35(state); \
for (int j = 0; j < 12; j++) \
Matrix[j][i][rowOut] ^= state[j]; \
Matrix[0][i][rowInOut] ^= state[11]; \
Matrix[1][i][rowInOut] ^= state[0]; \
Matrix[2][i][rowInOut] ^= state[1]; \
Matrix[3][i][rowInOut] ^= state[2]; \
Matrix[4][i][rowInOut] ^= state[3]; \
Matrix[5][i][rowInOut] ^= state[4]; \
Matrix[6][i][rowInOut] ^= state[5]; \
Matrix[7][i][rowInOut] ^= state[6]; \
Matrix[8][i][rowInOut] ^= state[7]; \
Matrix[9][i][rowInOut] ^= state[8]; \
Matrix[10][i][rowInOut] ^= state[9]; \
Matrix[11][i][rowInOut] ^= state[10]; \
} \
}
#define absorbblock_test(in) { \
state[0] ^= Matrix[0][0][ in]; \
state[1] ^= Matrix[1][0][in]; \
state[2] ^= Matrix[2][0][in]; \
state[3] ^= Matrix[3][0][in]; \
state[4] ^= Matrix[4][0][in]; \
state[5] ^= Matrix[5][0][in]; \
state[6] ^= Matrix[6][0][in]; \
state[7] ^= Matrix[7][0][in]; \
state[8] ^= Matrix[8][0][in]; \
state[9] ^= Matrix[9][0][in]; \
state[10] ^= Matrix[10][0][in]; \
state[11] ^= Matrix[11][0][in]; \
round_lyra_v35(state); \
round_lyra_v35(state); \
round_lyra_v35(state); \
round_lyra_v35(state); \
round_lyra_v35(state); \
round_lyra_v35(state); \
round_lyra_v35(state); \
round_lyra_v35(state); \
round_lyra_v35(state); \
round_lyra_v35(state); \
round_lyra_v35(state); \
round_lyra_v35(state); \
}
//// compute 30 version
#define reduceDuplexRowSetup_v30(rowIn, rowInOut, rowOut) { \
for (int i = 0; i < 8; i++) { \
for (int j = 0; j < 12; j++) \
state[j] ^= Matrix[12 * i + j][rowIn] + Matrix[12 * i + j][rowInOut]; \
round_lyra_v30(state); \
for (int j = 0; j < 12; j++) \
Matrix[j + 84 - 12 * i][rowOut] = Matrix[12 * i + j][rowIn] ^ state[j]; \
Matrix[0 + 12 * i][rowInOut] ^= state[11]; \
Matrix[1 + 12 * i][rowInOut] ^= state[0]; \
Matrix[2 + 12 * i][rowInOut] ^= state[1]; \
Matrix[3 + 12 * i][rowInOut] ^= state[2]; \
Matrix[4 + 12 * i][rowInOut] ^= state[3]; \
Matrix[5 + 12 * i][rowInOut] ^= state[4]; \
Matrix[6 + 12 * i][rowInOut] ^= state[5]; \
Matrix[7 + 12 * i][rowInOut] ^= state[6]; \
Matrix[8 + 12 * i][rowInOut] ^= state[7]; \
Matrix[9 + 12 * i][rowInOut] ^= state[8]; \
Matrix[10 + 12 * i][rowInOut] ^= state[9]; \
Matrix[11 + 12 * i][rowInOut] ^= state[10]; \
} \
}
#define reduceDuplexRow_v30(rowIn, rowInOut, rowOut) { \
for (int i = 0; i < 8; i++) { \
for (int j = 0; j < 12; j++) \
state[j] ^= Matrix[12 * i + j][rowIn] + Matrix[12 * i + j][rowInOut]; \
round_lyra_v30(state); \
for (int j = 0; j < 12; j++) \
Matrix[j + 12 * i][rowOut] ^= state[j]; \
Matrix[0 + 12 * i][rowInOut] ^= state[11]; \
Matrix[1 + 12 * i][rowInOut] ^= state[0]; \
Matrix[2 + 12 * i][rowInOut] ^= state[1]; \
Matrix[3 + 12 * i][rowInOut] ^= state[2]; \
Matrix[4 + 12 * i][rowInOut] ^= state[3]; \
Matrix[5 + 12 * i][rowInOut] ^= state[4]; \
Matrix[6 + 12 * i][rowInOut] ^= state[5]; \
Matrix[7 + 12 * i][rowInOut] ^= state[6]; \
Matrix[8 + 12 * i][rowInOut] ^= state[7]; \
Matrix[9 + 12 * i][rowInOut] ^= state[8]; \
Matrix[10 + 12 * i][rowInOut] ^= state[9]; \
Matrix[11 + 12 * i][rowInOut] ^= state[10]; \
} \
}
#define absorbblock_v30(in) { \
state[0] ^= Matrix[0][in]; \
state[1] ^= Matrix[1][in]; \
state[2] ^= Matrix[2][in]; \
state[3] ^= Matrix[3][in]; \
state[4] ^= Matrix[4][in]; \
state[5] ^= Matrix[5][in]; \
state[6] ^= Matrix[6][in]; \
state[7] ^= Matrix[7][in]; \
state[8] ^= Matrix[8][in]; \
state[9] ^= Matrix[9][in]; \
state[10] ^= Matrix[10][in]; \
state[11] ^= Matrix[11][in]; \
round_lyra_v30(state); \
round_lyra_v30(state); \
round_lyra_v30(state); \
round_lyra_v30(state); \
round_lyra_v30(state); \
round_lyra_v30(state); \
round_lyra_v30(state); \
round_lyra_v30(state); \
round_lyra_v30(state); \
round_lyra_v30(state); \
round_lyra_v30(state); \
round_lyra_v30(state); \
} }
static __device__ __forceinline__ static __device__ __forceinline__
void Gfunc_v35(uint2 & a, uint2 &b, uint2 &c, uint2 &d) void Gfunc(uint2 & a, uint2 &b, uint2 &c, uint2 &d)
{ {
a += b; d ^= a; d = ROR2(d, 32); a += b; d ^= a; d = SWAPUINT2(d);
c += d; b ^= c; b = ROR2(b, 24); c += d; b ^= c; b = ROR2(b, 24);
a += b; d ^= a; d = ROR2(d, 16); a += b; d ^= a; d = ROR2(d, 16);
c += d; b ^= c; b = ROR2(b, 63); c += d; b ^= c; b = ROR2(b, 63);
} }
static __device__ __forceinline__ __device__ __forceinline__
void Gfunc_v30(uint64_t & a, uint64_t &b, uint64_t &c, uint64_t &d) static void round_lyra(uint2 *s)
{ {
a += b; d ^= a; d = ROTR64(d, 32); Gfunc(s[0], s[4], s[8], s[12]);
c += d; b ^= c; b = ROTR64(b, 24); Gfunc(s[1], s[5], s[9], s[13]);
a += b; d ^= a; d = ROTR64(d, 16); Gfunc(s[2], s[6], s[10], s[14]);
c += d; b ^= c; b = ROTR64(b, 63); Gfunc(s[3], s[7], s[11], s[15]);
Gfunc(s[0], s[5], s[10], s[15]);
Gfunc(s[1], s[6], s[11], s[12]);
Gfunc(s[2], s[7], s[8], s[13]);
Gfunc(s[3], s[4], s[9], s[14]);
} }
#define round_lyra_v35_new(state) { \ __device__ __forceinline__
Gfunc_v35(state[0], state[4], state[8], state[12]); \ void reduceDuplexRowSetup(const int rowIn, const int rowInOut, const int rowOut, uint2 state[16], uint2 Matrix[96][8])
Gfunc_v35(state[1], state[5], state[9], state[13]); \
Gfunc_v35(state[2], state[6], state[10], state[14]); \
Gfunc_v35(state[3], state[7], state[11], state[15]); \
Gfunc_v35(state[0], state[5], state[10], state[15]); \
Gfunc_v35(state[1], state[6], state[11], state[12]); \
Gfunc_v35(state[2], state[7], state[8], state[13]); \
Gfunc_v35(state[3], state[4], state[9], state[14]); \
}
static __device__ __forceinline__ void round_lyra_v35(uint2 *s)
{
Gfunc_v35(s[0], s[4], s[8], s[12]);
Gfunc_v35(s[1], s[5], s[9], s[13]);
Gfunc_v35(s[2], s[6], s[10], s[14]);
Gfunc_v35(s[3], s[7], s[11], s[15]);
Gfunc_v35(s[0], s[5], s[10], s[15]);
Gfunc_v35(s[1], s[6], s[11], s[12]);
Gfunc_v35(s[2], s[7], s[8], s[13]);
Gfunc_v35(s[3], s[4], s[9], s[14]);
}
static __device__ __forceinline__ void round_lyra_v30(uint64_t *s)
{
Gfunc_v30(s[0], s[4], s[8], s[12]);
Gfunc_v30(s[1], s[5], s[9], s[13]);
Gfunc_v30(s[2], s[6], s[10], s[14]);
Gfunc_v30(s[3], s[7], s[11], s[15]);
Gfunc_v30(s[0], s[5], s[10], s[15]);
Gfunc_v30(s[1], s[6], s[11], s[12]);
Gfunc_v30(s[2], s[7], s[8], s[13]);
Gfunc_v30(s[3], s[4], s[9], s[14]);
}
__global__ __launch_bounds__(TPB, 1)
void lyra2_gpu_hash_32_v30(int threads, uint32_t startNounce, uint64_t *outputHash)
{ {
int thread = (blockDim.x * blockIdx.x + threadIdx.x); for (int i = 0; i < 8; i++)
if (thread < threads)
{ {
uint64_t state[16];
#pragma unroll
for (int i = 0; i<4; i++) { state[i] = outputHash[threads*i + thread]; } //password
#pragma unroll
for (int i = 0; i<4; i++) { state[i + 4] = state[i]; } //salt
#pragma unroll #pragma unroll
for (int i = 0; i<8; i++) { state[i + 8] = devectorize(blake2b_IV[i]); } for (int j = 0; j < 12; j++)
state[j] ^= Matrix[12 * i + j][rowIn] + Matrix[12 * i + j][rowInOut];
// blake2blyra x2 round_lyra(state);
#pragma unroll 24
for (int i = 0; i<24; i++) { round_lyra_v30(state); } //because 12 is not enough
uint64_t Matrix[96][8]; // not cool
// reducedSqueezeRow0
#pragma unroll 8
for (int i = 0; i < 8; i++) {
int idx = 84-12*i;
#pragma unroll 12
for (int j = 0; j<12; j++) { Matrix[j + idx][0] = state[j]; }
round_lyra_v30(state);
}
// reducedSqueezeRow1
#pragma unroll 8
for (int i = 0; i < 8; i++)
{
int idx0= 12*i;
int idx1= 84-idx0;
#pragma unroll 12
for (int j = 0; j<12; j++) { state[j] ^= Matrix[j + idx0][0]; }
round_lyra_v30(state);
#pragma unroll 12
for (int j = 0; j<12; j++) { Matrix[j + idx1][1] = Matrix[j + idx0][0] ^ state[j]; }
}
reduceDuplexRowSetup_v30(1, 0, 2);
reduceDuplexRowSetup_v30(2, 1, 3);
reduceDuplexRowSetup_v30(3, 0, 4);
reduceDuplexRowSetup_v30(4, 3, 5);
reduceDuplexRowSetup_v30(5, 2, 6);
reduceDuplexRowSetup_v30(6, 1, 7);
uint64_t rowa;
rowa = state[0] & 7;
reduceDuplexRow_v30(7, rowa, 0);
rowa = state[0] & 7;
reduceDuplexRow_v30(0, rowa, 3);
rowa = state[0] & 7;
reduceDuplexRow_v30(3, rowa, 6);
rowa = state[0] & 7;
reduceDuplexRow_v30(6, rowa, 1);
rowa = state[0] & 7;
reduceDuplexRow_v30(1, rowa, 4);
rowa = state[0] & 7;
reduceDuplexRow_v30(4, rowa, 7);
rowa = state[0] & 7;
reduceDuplexRow_v30(7, rowa, 2);
rowa = state[0] & 7;
reduceDuplexRow_v30(2, rowa, 5);
absorbblock_v30(rowa);
#pragma unroll #pragma unroll
for (int i = 0; i<4; i++) { for (int j = 0; j < 12; j++)
outputHash[threads*i + thread] = state[i]; Matrix[j + 84 - 12 * i][rowOut] = Matrix[12 * i + j][rowIn] ^ state[j];
} //password
Matrix[0 + 12 * i][rowInOut] ^= state[11];
} //thread Matrix[1 + 12 * i][rowInOut] ^= state[0];
Matrix[2 + 12 * i][rowInOut] ^= state[1];
Matrix[3 + 12 * i][rowInOut] ^= state[2];
Matrix[4 + 12 * i][rowInOut] ^= state[3];
Matrix[5 + 12 * i][rowInOut] ^= state[4];
Matrix[6 + 12 * i][rowInOut] ^= state[5];
Matrix[7 + 12 * i][rowInOut] ^= state[6];
Matrix[8 + 12 * i][rowInOut] ^= state[7];
Matrix[9 + 12 * i][rowInOut] ^= state[8];
Matrix[10 + 12 * i][rowInOut] ^= state[9];
Matrix[11 + 12 * i][rowInOut] ^= state[10];
}
} }
__global__ __launch_bounds__(TPB, 1) __global__ __launch_bounds__(TPB, 1)
@ -367,16 +123,27 @@ void lyra2_gpu_hash_32(int threads, uint32_t startNounce, uint64_t *outputHash)
if (thread < threads) if (thread < threads)
{ {
uint2 state[16]; uint2 state[16];
#pragma unroll #pragma unroll
for (int i = 0; i<4; i++) { LOHI(state[i].x, state[i].y, outputHash[threads*i + thread]); } //password for (int i = 0; i<4; i++) {
LOHI(state[i].x, state[i].y, outputHash[threads*i + thread]);
} //password
#pragma unroll #pragma unroll
for (int i = 0; i<4; i++) { state[i + 4] = state[i]; } //salt for (int i = 0; i<4; i++) {
state[i + 4] = state[i];
} //salt
#pragma unroll #pragma unroll
for (int i = 0; i<8; i++) { state[i + 8] = blake2b_IV[i]; } for (int i = 0; i<8; i++) {
state[i + 8] = blake2b_IV[i];
}
// blake2blyra x2 // blake2blyra x2
//#pragma unroll 24 //#pragma unroll 24
for (int i = 0; i<24; i++) { round_lyra_v35(state); } //because 12 is not enough for (int i = 0; i<24; i++) {
round_lyra(state);
} //because 12 is not enough
uint2 Matrix[96][8]; // not cool uint2 Matrix[96][8]; // not cool
@ -385,8 +152,10 @@ void lyra2_gpu_hash_32(int threads, uint32_t startNounce, uint64_t *outputHash)
for (int i = 0; i < 8; i++) for (int i = 0; i < 8; i++)
{ {
#pragma unroll 12 #pragma unroll 12
for (int j = 0; j<12; j++) { Matrix[j + 84 - 12 * i][0] = state[j]; } for (int j = 0; j<12; j++) {
round_lyra_v35(state); Matrix[j + 84 - 12 * i][0] = state[j];
}
round_lyra(state);
} }
// reducedSqueezeRow1 // reducedSqueezeRow1
@ -394,18 +163,22 @@ void lyra2_gpu_hash_32(int threads, uint32_t startNounce, uint64_t *outputHash)
for (int i = 0; i < 8; i++) for (int i = 0; i < 8; i++)
{ {
#pragma unroll 12 #pragma unroll 12
for (int j = 0; j<12; j++) { state[j] ^= Matrix[j + 12 * i][0]; } for (int j = 0; j<12; j++) {
round_lyra_v35(state); state[j] ^= Matrix[j + 12 * i][0];
}
round_lyra(state);
#pragma unroll 12 #pragma unroll 12
for (int j = 0; j<12; j++) { Matrix[j + 84 - 12 * i][1] = Matrix[j + 12 * i][0] ^ state[j]; } for (int j = 0; j<12; j++) {
Matrix[j + 84 - 12 * i][1] = Matrix[j + 12 * i][0] ^ state[j];
}
} }
reduceDuplexRowSetup(1, 0, 2); reduceDuplexRowSetup(1, 0, 2,state, Matrix);
reduceDuplexRowSetup(2, 1, 3); reduceDuplexRowSetup(2, 1, 3, state, Matrix);
reduceDuplexRowSetup(3, 0, 4); reduceDuplexRowSetup(3, 0, 4, state, Matrix);
reduceDuplexRowSetup(4, 3, 5); reduceDuplexRowSetup(4, 3, 5, state, Matrix);
reduceDuplexRowSetup(5, 2, 6); reduceDuplexRowSetup(5, 2, 6, state, Matrix);
reduceDuplexRowSetup(6, 1, 7); reduceDuplexRowSetup(6, 1, 7, state, Matrix);
uint32_t rowa; uint32_t rowa;
rowa = state[0].x & 7; rowa = state[0].x & 7;
@ -435,12 +208,6 @@ void lyra2_gpu_hash_32(int threads, uint32_t startNounce, uint64_t *outputHash)
} //thread } //thread
} }
__host__
void lyra2_cpu_init(int thr_id, int threads)
{
//not used
}
__host__ __host__
void lyra2_cpu_hash_32(int thr_id, int threads, uint32_t startNounce, uint64_t *d_outputHash, int order) void lyra2_cpu_hash_32(int thr_id, int threads, uint32_t startNounce, uint64_t *d_outputHash, int order)
{ {
@ -449,14 +216,5 @@ void lyra2_cpu_hash_32(int thr_id, int threads, uint32_t startNounce, uint64_t *
dim3 grid((threads + threadsperblock - 1) / threadsperblock); dim3 grid((threads + threadsperblock - 1) / threadsperblock);
dim3 block(threadsperblock); dim3 block(threadsperblock);
if (device_sm[device_map[thr_id]] >= 320) { lyra2_gpu_hash_32 <<<grid, block>>> (threads, startNounce, d_outputHash);
lyra2_gpu_hash_32 <<<grid, block>>> (threads, startNounce, d_outputHash);
} else {
// kernel for compute30 card
lyra2_gpu_hash_32_v30 <<<grid, block >>> (threads, startNounce, d_outputHash);
}
MyStreamSynchronize(NULL, order, thr_id);
//cudaThreadSynchronize();
} }

2
lyra2/lyra2RE.cu
Unescape View File

@ -20,7 +20,6 @@ extern void skein256_cpu_hash_32(int thr_id, int threads, uint32_t startNonce, u
extern void skein256_cpu_init(int thr_id, int threads); extern void skein256_cpu_init(int thr_id, int threads);
extern void lyra2_cpu_hash_32(int thr_id, int threads, uint32_t startNonce, uint64_t *d_outputHash, int order); extern void lyra2_cpu_hash_32(int thr_id, int threads, uint32_t startNonce, uint64_t *d_outputHash, int order);
extern void lyra2_cpu_init(int thr_id, int threads);
extern void groestl256_setTarget(const void *ptarget); extern void groestl256_setTarget(const void *ptarget);
extern uint32_t groestl256_cpu_hash_32(int thr_id, int threads, uint32_t startNounce, uint64_t *d_outputHash, int order); extern uint32_t groestl256_cpu_hash_32(int thr_id, int threads, uint32_t startNounce, uint64_t *d_outputHash, int order);
@ -78,7 +77,6 @@ extern "C" int scanhash_lyra2(int thr_id, uint32_t *pdata,
keccak256_cpu_init(thr_id,throughput); keccak256_cpu_init(thr_id,throughput);
skein256_cpu_init(thr_id, throughput); skein256_cpu_init(thr_id, throughput);
groestl256_cpu_init(thr_id, throughput); groestl256_cpu_init(thr_id, throughput);
lyra2_cpu_init(thr_id, throughput);
CUDA_SAFE_CALL(cudaMalloc(&d_hash[thr_id], 16 * sizeof(uint32_t) * throughput)); CUDA_SAFE_CALL(cudaMalloc(&d_hash[thr_id], 16 * sizeof(uint32_t) * throughput));

Write Preview
Loading…
Cancel
Save