diff --git a/Makefile.am b/Makefile.am
index ac2ecdad..125b432b 100644
--- a/Makefile.am
+++ b/Makefile.am
@@ -73,7 +73,7 @@ sgminer_SOURCES += algorithm/whirlcoin.c algorithm/whirlcoin.h
sgminer_SOURCES += algorithm/neoscrypt.c algorithm/neoscrypt.h
sgminer_SOURCES += algorithm/whirlpoolx.c algorithm/whirlpoolx.h
sgminer_SOURCES += algorithm/lyra2re.c algorithm/lyra2re.h algorithm/lyra2.c algorithm/lyra2.h algorithm/sponge.c algorithm/sponge.h
-sgminer_SOURCES += algorithm/lyra2re_old.c algorithm/lyra2re_old.h
+sgminer_SOURCES += algorithm/lyra2rev2.c algorithm/lyra2rev2.h
sgminer_SOURCES += algorithm/pluck.c algorithm/pluck.h
sgminer_SOURCES += algorithm/credits.c algorithm/credits.h
sgminer_SOURCES += algorithm/yescrypt.h algorithm/yescrypt.c algorithm/yescrypt_core.h algorithm/yescrypt-opt.c algorithm/yescryptcommon.c algorithm/sysendian.h
diff --git a/algorithm.c b/algorithm.c
index 6acab924..3f6914f4 100644
--- a/algorithm.c
+++ b/algorithm.c
@@ -33,7 +33,7 @@
#include "algorithm/neoscrypt.h"
#include "algorithm/whirlpoolx.h"
#include "algorithm/lyra2re.h"
-#include "algorithm/lyra2re_old.h"
+#include "algorithm/lyra2rev2.h"
#include "algorithm/pluck.h"
#include "algorithm/yescrypt.h"
#include "algorithm/credits.h"
@@ -62,7 +62,7 @@ const char *algorithm_type_str[] = {
"Neoscrypt",
"WhirlpoolX",
"Lyra2RE",
- "Lyra2REv2"
+ "Lyra2REV2"
"Pluck"
"Yescrypt",
"Yescrypt-multi"
@@ -798,7 +798,7 @@ static cl_int queue_whirlpoolx_kernel(struct __clState *clState, struct _dev_blk
return status;
}
-static cl_int queue_lyra2RE_kernel(struct __clState *clState, struct _dev_blk_ctx *blk, __maybe_unused cl_uint threads)
+static cl_int queue_lyra2re_kernel(struct __clState *clState, struct _dev_blk_ctx *blk, __maybe_unused cl_uint threads)
{
cl_kernel *kernel;
unsigned int num;
@@ -842,7 +842,7 @@ static cl_int queue_lyra2RE_kernel(struct __clState *clState, struct _dev_blk_ct
return status;
}
-static cl_int queue_lyra2REv2_kernel(struct __clState *clState, struct _dev_blk_ctx *blk, __maybe_unused cl_uint threads)
+static cl_int queue_lyra2rev2_kernel(struct __clState *clState, struct _dev_blk_ctx *blk, __maybe_unused cl_uint threads)
{
cl_kernel *kernel;
unsigned int num;
@@ -992,10 +992,8 @@ static algorithm_settings_t algos[] = {
{ "fresh", ALGO_FRESH, "", 1, 256, 256, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 4, 4 * 16 * 4194304, 0, fresh_regenhash, queue_fresh_kernel, gen_hash, NULL },
- { "lyra2re", ALGO_LYRA2RE, "", 1, 128, 128, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 4, 2 * 8 * 4194304, 0, lyra2reold_regenhash, queue_lyra2RE_kernel, gen_hash, NULL },
-
- { "lyra2rev2", ALGO_LYRA2REv2, "", 1, 256, 256, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 6, -1, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, lyra2re_regenhash, queue_lyra2REv2_kernel, gen_hash, append_neoscrypt_compiler_options },
-
+ { "lyra2re", ALGO_LYRA2RE, "", 1, 128, 128, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 4, 2 * 8 * 4194304, 0, lyra2re_regenhash, queue_lyra2re_kernel, gen_hash, NULL },
+ { "lyra2rev2", ALGO_LYRA2REV2, "", 1, 256, 256, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 6, -1, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, lyra2rev2_regenhash, queue_lyra2rev2_kernel, gen_hash, append_neoscrypt_compiler_options },
// kernels starting from this will have difficulty calculated by using fuguecoin algorithm
#define A_FUGUE(a, b, c) \
diff --git a/algorithm.h b/algorithm.h
index 8b7185a4..9187eb53 100644
--- a/algorithm.h
+++ b/algorithm.h
@@ -30,7 +30,7 @@ typedef enum {
ALGO_NEOSCRYPT,
ALGO_WHIRLPOOLX,
ALGO_LYRA2RE,
- ALGO_LYRA2REv2,
+ ALGO_LYRA2REV2,
ALGO_PLUCK,
ALGO_YESCRYPT,
ALGO_YESCRYPT_MULTI,
diff --git a/algorithm/lyra2.c b/algorithm/lyra2.c
index 42640e76..865d8e17 100644
--- a/algorithm/lyra2.c
+++ b/algorithm/lyra2.c
@@ -61,16 +61,18 @@ int LYRA2(void *K, uint64_t kLen, const void *pwd, uint64_t pwdlen, const void *
const int64_t ROW_LEN_INT64 = BLOCK_LEN_INT64 * nCols;
const int64_t ROW_LEN_BYTES = ROW_LEN_INT64 * 8;
+ // for Lyra2REv2, nCols = 4, v1 was using 8
+ const int64_t BLOCK_LEN = (nCols == 4) ? BLOCK_LEN_BLAKE2_SAFE_INT64 : BLOCK_LEN_BLAKE2_SAFE_BYTES;
i = (int64_t) ((int64_t) nRows * (int64_t) ROW_LEN_BYTES);
- uint64_t *wholeMatrix = malloc(i);
+ uint64_t *wholeMatrix = (uint64_t*)malloc(i);
if (wholeMatrix == NULL) {
return -1;
}
memset(wholeMatrix, 0, i);
//Allocates pointers to each row of the matrix
- uint64_t **memMatrix = malloc(nRows * sizeof (uint64_t*));
+ uint64_t **memMatrix = (uint64_t**)malloc(nRows * sizeof (uint64_t*));
if (memMatrix == NULL) {
return -1;
}
@@ -122,7 +124,7 @@ int LYRA2(void *K, uint64_t kLen, const void *pwd, uint64_t pwdlen, const void *
//======================= Initializing the Sponge State ====================//
//Sponge state: 16 uint64_t, BLOCK_LEN_INT64 words of them for the bitrate (b) and the remainder for the capacity (c)
- uint64_t *state = malloc(16 * sizeof (uint64_t));
+ uint64_t *state = (uint64_t*)malloc(16 * sizeof (uint64_t));
if (state == NULL) {
return -1;
}
@@ -134,7 +136,7 @@ int LYRA2(void *K, uint64_t kLen, const void *pwd, uint64_t pwdlen, const void *
ptrWord = wholeMatrix;
for (i = 0; i < nBlocksInput; i++) {
absorbBlockBlake2Safe(state, ptrWord); //absorbs each block of pad(pwd || salt || basil)
- ptrWord += BLOCK_LEN_BLAKE2_SAFE_INT64; //goes to next block of pad(pwd || salt || basil)
+ ptrWord += BLOCK_LEN; //goes to next block of pad(pwd || salt || basil)
}
//Initializes M[0] and M[1]
@@ -196,7 +198,7 @@ int LYRA2(void *K, uint64_t kLen, const void *pwd, uint64_t pwdlen, const void *
absorbBlock(state, memMatrix[rowa]);
//Squeezes the key
- squeeze(state, K, kLen);
+ squeeze(state, (unsigned char*)K, kLen);
//==========================================================================/
//========================= Freeing the memory =============================//
diff --git a/algorithm/lyra2re.c b/algorithm/lyra2re.c
index cfc5adbb..61f2b34f 100644
--- a/algorithm/lyra2re.c
+++ b/algorithm/lyra2re.c
@@ -36,8 +36,6 @@
#include "sph/sph_groestl.h"
#include "sph/sph_skein.h"
#include "sph/sph_keccak.h"
-#include "sph/sph_bmw.h"
-#include "sph/sph_cubehash.h"
#include "lyra2.h"
/*
@@ -57,10 +55,9 @@ be32enc_vect(uint32_t *dst, const uint32_t *src, uint32_t len)
inline void lyra2rehash(void *state, const void *input)
{
sph_blake256_context ctx_blake;
- sph_bmw256_context ctx_bmw;
+ sph_groestl256_context ctx_groestl;
sph_keccak256_context ctx_keccak;
sph_skein256_context ctx_skein;
- sph_cubehash256_context ctx_cube;
uint32_t hashA[8], hashB[8];
@@ -72,23 +69,17 @@ inline void lyra2rehash(void *state, const void *input)
sph_keccak256 (&ctx_keccak,hashA, 32);
sph_keccak256_close(&ctx_keccak, hashB);
- sph_cubehash256_init(&ctx_cube);
- sph_cubehash256(&ctx_cube, hashB, 32);
- sph_cubehash256_close(&ctx_cube, hashA);
+ LYRA2(hashA, 32, hashB, 32, hashB, 32, 1, 8, 8);
- LYRA2(hashB, 32, hashA, 32, hashA, 32, 1, 4, 4);
sph_skein256_init(&ctx_skein);
- sph_skein256 (&ctx_skein, hashB, 32);
- sph_skein256_close(&ctx_skein, hashA);
+ sph_skein256 (&ctx_skein, hashA, 32);
+ sph_skein256_close(&ctx_skein, hashB);
- sph_cubehash256_init(&ctx_cube);
- sph_cubehash256(&ctx_cube, hashA, 32);
- sph_cubehash256_close(&ctx_cube, hashB);
- sph_bmw256_init(&ctx_bmw);
- sph_bmw256 (&ctx_bmw, hashB, 32);
- sph_bmw256_close(&ctx_bmw, hashA);
+ sph_groestl256_init(&ctx_groestl);
+ sph_groestl256 (&ctx_groestl, hashB, 32);
+ sph_groestl256_close(&ctx_groestl, hashA);
memcpy(state, hashA, 32);
}
diff --git a/algorithm/lyra2re.h b/algorithm/lyra2re.h
index 8bc52ac4..8a58e747 100644
--- a/algorithm/lyra2re.h
+++ b/algorithm/lyra2re.h
@@ -2,8 +2,6 @@
#define LYRA2RE_H
#include "miner.h"
-#define LYRA_SCRATCHBUF_SIZE (1536) // matrix size [12][4][4] uint64_t or equivalent
-#define LYRA_SECBUF_SIZE (4) // (not used)
extern int lyra2re_test(unsigned char *pdata, const unsigned char *ptarget,
uint32_t nonce);
diff --git a/algorithm/lyra2re_old.h b/algorithm/lyra2re_old.h
deleted file mode 100644
index 0788dfb3..00000000
--- a/algorithm/lyra2re_old.h
+++ /dev/null
@@ -1,10 +0,0 @@
-#ifndef LYRA2REOLD_H
-#define LYRA2REOLD_H
-
-#include "miner.h"
-
-extern int lyra2reold_test(unsigned char *pdata, const unsigned char *ptarget,
- uint32_t nonce);
-extern void lyra2reold_regenhash(struct work *work);
-
-#endif /* LYRA2RE_H */
diff --git a/algorithm/lyra2re_old.c b/algorithm/lyra2rev2.c
similarity index 77%
rename from algorithm/lyra2re_old.c
rename to algorithm/lyra2rev2.c
index 31a0a1ab..aea0082a 100644
--- a/algorithm/lyra2re_old.c
+++ b/algorithm/lyra2rev2.c
@@ -36,6 +36,8 @@
#include "sph/sph_groestl.h"
#include "sph/sph_skein.h"
#include "sph/sph_keccak.h"
+#include "sph/sph_bmw.h"
+#include "sph/sph_cubehash.h"
#include "lyra2.h"
/*
@@ -52,13 +54,13 @@ be32enc_vect(uint32_t *dst, const uint32_t *src, uint32_t len)
}
-inline void lyra2rehash_old(void *state, const void *input)
+inline void lyra2rev2hash(void *state, const void *input)
{
sph_blake256_context ctx_blake;
- sph_groestl256_context ctx_groestl;
+ sph_bmw256_context ctx_bmw;
sph_keccak256_context ctx_keccak;
sph_skein256_context ctx_skein;
-
+ sph_cubehash256_context ctx_cube;
uint32_t hashA[8], hashB[8];
sph_blake256_init(&ctx_blake);
@@ -69,32 +71,41 @@ inline void lyra2rehash_old(void *state, const void *input)
sph_keccak256 (&ctx_keccak,hashA, 32);
sph_keccak256_close(&ctx_keccak, hashB);
- LYRA2(hashA, 32, hashB, 32, hashB, 32, 1, 8, 8);
+ sph_cubehash256_init(&ctx_cube);
+ sph_cubehash256(&ctx_cube, hashB, 32);
+ sph_cubehash256_close(&ctx_cube, hashA);
+
+ LYRA2(hashB, 32, hashA, 32, hashA, 32, 1, 4, 4);
+
+ sph_skein256_init(&ctx_skein);
+ sph_skein256 (&ctx_skein, hashB, 32);
+ sph_skein256_close(&ctx_skein, hashA);
- sph_skein256_init(&ctx_skein);
- sph_skein256 (&ctx_skein, hashA, 32);
- sph_skein256_close(&ctx_skein, hashB);
+ sph_cubehash256_init(&ctx_cube);
+ sph_cubehash256(&ctx_cube, hashA, 32);
+ sph_cubehash256_close(&ctx_cube, hashB);
+ sph_bmw256_init(&ctx_bmw);
+ sph_bmw256 (&ctx_bmw, hashB, 32);
+ sph_bmw256_close(&ctx_bmw, hashA);
- sph_groestl256_init(&ctx_groestl);
- sph_groestl256 (&ctx_groestl, hashB, 32);
- sph_groestl256_close(&ctx_groestl, hashA);
+//printf("cpu hash %08x %08x %08x %08x\n",hashA[0],hashA[1],hashA[2],hashA[3]);
- memcpy(state, hashA, 32);
+ memcpy(state, hashA, 32);
}
static const uint32_t diff1targ = 0x0000ffff;
/* Used externally as confirmation of correct OCL code */
-int lyra2reold_test(unsigned char *pdata, const unsigned char *ptarget, uint32_t nonce)
+int lyra2rev2_test(unsigned char *pdata, const unsigned char *ptarget, uint32_t nonce)
{
uint32_t tmp_hash7, Htarg = le32toh(((const uint32_t *)ptarget)[7]);
uint32_t data[20], ohash[8];
be32enc_vect(data, (const uint32_t *)pdata, 19);
data[19] = htobe32(nonce);
- lyra2rehash_old(ohash, data);
+ lyra2rev2hash(ohash, data);
tmp_hash7 = be32toh(ohash[7]);
applog(LOG_DEBUG, "htarget %08lx diff1 %08lx hash %08lx",
@@ -108,7 +119,7 @@ int lyra2reold_test(unsigned char *pdata, const unsigned char *ptarget, uint32_t
return 1;
}
-void lyra2reold_regenhash(struct work *work)
+void lyra2rev2_regenhash(struct work *work)
{
uint32_t data[20];
uint32_t *nonce = (uint32_t *)(work->data + 76);
@@ -116,10 +127,10 @@ void lyra2reold_regenhash(struct work *work)
be32enc_vect(data, (const uint32_t *)work->data, 19);
data[19] = htobe32(*nonce);
- lyra2rehash_old(ohash, data);
+ lyra2rev2hash(ohash, data);
}
-bool scanhash_lyra2reold(struct thr_info *thr, const unsigned char __maybe_unused *pmidstate,
+bool scanhash_lyra2rev2(struct thr_info *thr, const unsigned char __maybe_unused *pmidstate,
unsigned char *pdata, unsigned char __maybe_unused *phash1,
unsigned char __maybe_unused *phash, const unsigned char *ptarget,
uint32_t max_nonce, uint32_t *last_nonce, uint32_t n)
@@ -137,7 +148,7 @@ bool scanhash_lyra2reold(struct thr_info *thr, const unsigned char __maybe_unuse
*nonce = ++n;
data[19] = (n);
- lyra2rehash_old(ostate, data);
+ lyra2rev2hash(ostate, data);
tmp_hash7 = (ostate[7]);
applog(LOG_INFO, "data7 %08lx",
diff --git a/algorithm/lyra2rev2.h b/algorithm/lyra2rev2.h
new file mode 100644
index 00000000..1a31f76f
--- /dev/null
+++ b/algorithm/lyra2rev2.h
@@ -0,0 +1,11 @@
+#ifndef LYRA2REV2_H
+#define LYRA2REV2_H
+
+#include "miner.h"
+#define LYRA_SCRATCHBUF_SIZE (1536) // matrix size [12][4][4] uint64_t or equivalent
+#define LYRA_SECBUF_SIZE (4) // (not used)
+extern int lyra2rev2_test(unsigned char *pdata, const unsigned char *ptarget,
+ uint32_t nonce);
+extern void lyra2rev2_regenhash(struct work *work);
+
+#endif /* LYRA2REV2_H */
diff --git a/driver-opencl.c b/driver-opencl.c
index 0e45e555..4a9d0693 100644
--- a/driver-opencl.c
+++ b/driver-opencl.c
@@ -1366,7 +1366,7 @@ static bool opencl_thread_init(struct thr_info *thr)
static bool opencl_prepare_work(struct thr_info __maybe_unused *thr, struct work *work)
{
- if (work->pool->algorithm.type == ALGO_LYRA2RE || work->pool->algorithm.type == ALGO_LYRA2REv2) {
+ if (work->pool->algorithm.type == ALGO_LYRA2RE || work->pool->algorithm.type == ALGO_LYRA2REV2) {
work->blk.work = work;
precalc_hash_blake256(&work->blk, 0, (uint32_t *)(work->data));
}
diff --git a/kernel/lyra2rev2.cl b/kernel/lyra2rev2.cl
index 0fe0440d..a165a751 100644
--- a/kernel/lyra2rev2.cl
+++ b/kernel/lyra2rev2.cl
@@ -31,8 +31,8 @@
// typedef unsigned int uint;
#pragma OPENCL EXTENSION cl_amd_printf : enable
-#ifndef LYRA2RE_CL
-#define LYRA2RE_CL
+#ifndef LYRA2REV2_CL
+#define LYRA2REV2_CL
#if __ENDIAN_LITTLE__
#define SPH_LITTLE_ENDIAN 1
@@ -522,4 +522,4 @@ __kernel void search6(__global uchar* hashes, __global uint* output, const ulong
}
-#endif // LYRA2RE_CL
\ No newline at end of file
+#endif // LYRA2REV2_CL
\ No newline at end of file
diff --git a/ocl.c b/ocl.c
index cb00790f..98f337bd 100644
--- a/ocl.c
+++ b/ocl.c
@@ -37,7 +37,7 @@
#include "algorithm/neoscrypt.h"
#include "algorithm/pluck.h"
#include "algorithm/yescrypt.h"
-#include "algorithm/lyra2re.h"
+#include "algorithm/lyra2rev2.h"
/* FIXME: only here for global config vars, replace with configuration.h
* or similar as soon as config is in a struct instead of littered all
@@ -586,7 +586,7 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize, algorithm_t *alg
}
// Lyra2re v2 TC
- else if (cgpu->algorithm.type == ALGO_LYRA2REv2 && !cgpu->opt_tc) {
+ else if (cgpu->algorithm.type == ALGO_LYRA2REV2 && !cgpu->opt_tc) {
size_t glob_thread_count;
long max_int;
unsigned char type = 0;
@@ -797,7 +797,7 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize, algorithm_t *alg
applog(LOG_DEBUG, "yescrypt buffer sizes: %lu RW, %lu R", (unsigned long)bufsize, (unsigned long)readbufsize);
// scrypt/n-scrypt
}
- else if (algorithm->type == ALGO_LYRA2REv2) {
+ else if (algorithm->type == ALGO_LYRA2REV2) {
/* The scratch/pad-buffer needs 32kBytes memory per thread. */
bufsize = LYRA_SCRATCHBUF_SIZE * cgpu->thread_concurrency;
buf1size = 4* 8 * cgpu->thread_concurrency; //matrix
@@ -855,7 +855,7 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize, algorithm_t *alg
return NULL;
}
}
- else if (algorithm->type == ALGO_LYRA2REv2) {
+ else if (algorithm->type == ALGO_LYRA2REV2) {
// need additionnal buffers
clState->buffer1 = clCreateBuffer(clState->context, CL_MEM_READ_WRITE, buf1size, NULL, &status);
if (status != CL_SUCCESS && !clState->buffer1) {
diff --git a/winbuild/sgminer.vcxproj b/winbuild/sgminer.vcxproj
index 6ec8a4cd..e71685e4 100644
--- a/winbuild/sgminer.vcxproj
+++ b/winbuild/sgminer.vcxproj
@@ -265,6 +265,7 @@
+
@@ -330,6 +331,7 @@
+
diff --git a/winbuild/sgminer.vcxproj.filters b/winbuild/sgminer.vcxproj.filters
index 02c26210..14c9ed2f 100644
--- a/winbuild/sgminer.vcxproj.filters
+++ b/winbuild/sgminer.vcxproj.filters
@@ -218,6 +218,9 @@
Source Files\algorithm
+
+ Source Files\algorithm
+
@@ -415,6 +418,9 @@
Header Files\algorithm
+
+ Header Files\algorithm
+