From 91e1d324c541fc2abe02db94b15122f6fdaba90c Mon Sep 17 00:00:00 2001 From: elbandi Date: Sun, 10 Jul 2016 15:13:54 +0000 Subject: [PATCH] Add Sia algo support --- Makefile.am | 1 + algorithm.c | 21 +++++ algorithm.h | 1 + algorithm/sia.c | 232 ++++++++++++++++++++++++++++++++++++++++++++++++ algorithm/sia.h | 8 ++ kernel/sia.cl | 120 +++++++++++++++++++++++++ sgminer.c | 18 +++- 7 files changed, 398 insertions(+), 3 deletions(-) create mode 100644 algorithm/sia.c create mode 100644 algorithm/sia.h create mode 100644 kernel/sia.cl diff --git a/Makefile.am b/Makefile.am index 137a5723..38f1bfed 100644 --- a/Makefile.am +++ b/Makefile.am @@ -75,6 +75,7 @@ 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/lyra2rev2.c algorithm/lyra2rev2.h sgminer_SOURCES += algorithm/pluck.c algorithm/pluck.h +sgminer_SOURCES += algorithm/sia.c algorithm/sia.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 sgminer_SOURCES += algorithm/blake256.c algorithm/blake256.h diff --git a/algorithm.c b/algorithm.c index a4637ba0..ca9ba551 100644 --- a/algorithm.c +++ b/algorithm.c @@ -39,6 +39,7 @@ #include "algorithm/credits.h" #include "algorithm/blake256.h" #include "algorithm/blakecoin.h" +#include "algorithm/sia.h" #include "algorithm/decred.h" #include "algorithm/lbry.h" @@ -72,6 +73,7 @@ const char *algorithm_type_str[] = { "Yescrypt-multi", "Blakecoin", "Blake", + "Sia", "Decred", "Vanilla", "Lbry" @@ -976,6 +978,24 @@ static cl_int queue_blake_kernel(_clState *clState, dev_blk_ctx *blk, __maybe_un return status; } +static cl_int queue_sia_kernel(struct __clState *clState, struct _dev_blk_ctx *blk, __maybe_unused cl_uint threads) +{ + cl_kernel *kernel = &clState->kernel; + unsigned int num = 0; + cl_ulong le_target; + cl_int status = 0; + + le_target = *(cl_ulong *)(blk->work->device_target + 24); + flip80(clState->cldata, blk->work->data); + status = clEnqueueWriteBuffer(clState->commandQueue, clState->CLbuffer0, true, 0, 80, clState->cldata, 0, NULL, NULL); + + CL_SET_ARG(clState->CLbuffer0); + CL_SET_ARG(clState->outputBuffer); + CL_SET_ARG(le_target); + + return status; +} + static cl_int queue_decred_kernel(_clState *clState, dev_blk_ctx *blk, __maybe_unused cl_uint threads) { cl_kernel *kernel = &clState->kernel; @@ -1129,6 +1149,7 @@ static algorithm_settings_t algos[] = { { "blake256r8", ALGO_BLAKECOIN, "", 1, 1, 1, 0, 0, 0xFF, 0xFFFFULL, 0x000000ffUL, 0, 128, 0, blakecoin_regenhash, blakecoin_midstate, blakecoin_prepare_work, queue_blake_kernel, sha256, NULL }, { "blake256r14", ALGO_BLAKE, "", 1, 1, 1, 0, 0, 0xFF, 0xFFFFULL, 0x00000000UL, 0, 128, 0, blake256_regenhash, blake256_midstate, blake256_prepare_work, queue_blake_kernel, gen_hash, NULL }, + { "sia", ALGO_SIA, "", 1, 1, 1, 0, 0, 0xFF, 0xFFFFULL, 0x0000FFFFUL, 0, 128, 0, sia_regenhash, NULL, NULL, queue_sia_kernel, NULL, NULL }, { "vanilla", ALGO_VANILLA, "", 1, 1, 1, 0, 0, 0xFF, 0xFFFFULL, 0x000000ffUL, 0, 128, 0, blakecoin_regenhash, blakecoin_midstate, blakecoin_prepare_work, queue_blake_kernel, gen_hash, NULL }, { "lbry", ALGO_LBRY, "", 1, 256, 256, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 2, 4 * 8 * 4194304, 0, lbry_regenhash, NULL, NULL, queue_lbry_kernel, gen_hash, NULL }, diff --git a/algorithm.h b/algorithm.h index c02da8a6..56056295 100644 --- a/algorithm.h +++ b/algorithm.h @@ -36,6 +36,7 @@ typedef enum { ALGO_YESCRYPT_MULTI, ALGO_BLAKECOIN, ALGO_BLAKE, + ALGO_SIA, ALGO_DECRED, ALGO_VANILLA, ALGO_LBRY diff --git a/algorithm/sia.c b/algorithm/sia.c new file mode 100644 index 00000000..c04ead43 --- /dev/null +++ b/algorithm/sia.c @@ -0,0 +1,232 @@ +/*- + * Copyright 2009 Colin Percival, 2014 savale + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions + * are met: + * 1. Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * 2. Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in the + * documentation and/or other materials provided with the distribution. + * + * THIS SOFTWARE IS PROVIDED BY THE AUTHOR AND CONTRIBUTORS ``AS IS'' AND + * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE + * ARE DISCLAIMED. IN NO EVENT SHALL THE AUTHOR OR CONTRIBUTORS BE LIABLE + * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL + * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS + * OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) + * HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT + * LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY + * OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF + * SUCH DAMAGE. + * + * This file was originally written by Colin Percival as part of the Tarsnap + * online backup system. + */ + +#include "config.h" +#include "miner.h" + +#include +#include +#include + +// Cyclic right rotation. + +#ifndef ROTR64 +#define ROTR64(x, y) (((x) >> (y)) ^ ((x) << (64 - (y)))) +#endif + +// Little-endian byte access. + +#define B2B_GET64(p) \ + (((uint64_t) ((uint8_t *) (p))[0]) ^ \ + (((uint64_t) ((uint8_t *) (p))[1]) << 8) ^ \ + (((uint64_t) ((uint8_t *) (p))[2]) << 16) ^ \ + (((uint64_t) ((uint8_t *) (p))[3]) << 24) ^ \ + (((uint64_t) ((uint8_t *) (p))[4]) << 32) ^ \ + (((uint64_t) ((uint8_t *) (p))[5]) << 40) ^ \ + (((uint64_t) ((uint8_t *) (p))[6]) << 48) ^ \ + (((uint64_t) ((uint8_t *) (p))[7]) << 56)) + +// G Mixing function. + +#define B2B_G(a, b, c, d, x, y) { \ + v[a] = v[a] + v[b] + x; \ + v[d] = ROTR64(v[d] ^ v[a], 32); \ + v[c] = v[c] + v[d]; \ + v[b] = ROTR64(v[b] ^ v[c], 24); \ + v[a] = v[a] + v[b] + y; \ + v[d] = ROTR64(v[d] ^ v[a], 16); \ + v[c] = v[c] + v[d]; \ + v[b] = ROTR64(v[b] ^ v[c], 63); } + +// Initialization Vector. + +static const uint64_t blake2b_iv[8] = { + 0x6A09E667F3BCC908, 0xBB67AE8584CAA73B, + 0x3C6EF372FE94F82B, 0xA54FF53A5F1D36F1, + 0x510E527FADE682D1, 0x9B05688C2B3E6C1F, + 0x1F83D9ABFB41BD6B, 0x5BE0CD19137E2179 +}; + +// state context +typedef struct { + uint8_t b[128]; // input buffer + uint64_t h[8]; // chained state + uint64_t t[2]; // total number of bytes + size_t c; // pointer for b[] + size_t outlen; // digest size +} blake2b_ctx; + +void blake2b_update(blake2b_ctx *ctx, // context + const void *in, size_t inlen); // data to be hashed + +// Compression function. "last" flag indicates last block. + +static void blake2b_compress(blake2b_ctx *ctx, int last) +{ + const uint8_t sigma[12][16] = { + { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 }, + { 14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 }, + { 11, 8, 12, 0, 5, 2, 15, 13, 10, 14, 3, 6, 7, 1, 9, 4 }, + { 7, 9, 3, 1, 13, 12, 11, 14, 2, 6, 5, 10, 4, 0, 15, 8 }, + { 9, 0, 5, 7, 2, 4, 10, 15, 14, 1, 11, 12, 6, 8, 3, 13 }, + { 2, 12, 6, 10, 0, 11, 8, 3, 4, 13, 7, 5, 15, 14, 1, 9 }, + { 12, 5, 1, 15, 14, 13, 4, 10, 0, 7, 6, 3, 9, 2, 8, 11 }, + { 13, 11, 7, 14, 12, 1, 3, 9, 5, 0, 15, 4, 8, 6, 2, 10 }, + { 6, 15, 14, 9, 11, 3, 0, 8, 12, 2, 13, 7, 1, 4, 10, 5 }, + { 10, 2, 8, 4, 7, 6, 1, 5, 15, 11, 9, 14, 3, 12, 13, 0 }, + { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 }, + { 14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 } + }; + int i; + uint64_t v[16], m[16]; + + for (i = 0; i < 8; i++) { // init work variables + v[i] = ctx->h[i]; + v[i + 8] = blake2b_iv[i]; + } + + v[12] ^= ctx->t[0]; // low 64 bits of offset + v[13] ^= ctx->t[1]; // high 64 bits + if (last) // last block flag set ? + v[14] = ~v[14]; + + for (i = 0; i < 16; i++) // get little-endian words + m[i] = B2B_GET64(&ctx->b[8 * i]); + + for (i = 0; i < 12; i++) { // twelve rounds + B2B_G( 0, 4, 8, 12, m[sigma[i][ 0]], m[sigma[i][ 1]]); + B2B_G( 1, 5, 9, 13, m[sigma[i][ 2]], m[sigma[i][ 3]]); + B2B_G( 2, 6, 10, 14, m[sigma[i][ 4]], m[sigma[i][ 5]]); + B2B_G( 3, 7, 11, 15, m[sigma[i][ 6]], m[sigma[i][ 7]]); + B2B_G( 0, 5, 10, 15, m[sigma[i][ 8]], m[sigma[i][ 9]]); + B2B_G( 1, 6, 11, 12, m[sigma[i][10]], m[sigma[i][11]]); + B2B_G( 2, 7, 8, 13, m[sigma[i][12]], m[sigma[i][13]]); + B2B_G( 3, 4, 9, 14, m[sigma[i][14]], m[sigma[i][15]]); + } + + for( i = 0; i < 8; ++i ) + ctx->h[i] ^= v[i] ^ v[i + 8]; +} + +// Initialize the hashing context "ctx" with optional key "key". +// 1 <= outlen <= 64 gives the digest size in bytes. +// Secret key (also <= 64 bytes) is optional (keylen = 0). + +int blake2b_init(blake2b_ctx *ctx, size_t outlen, + const void *key, size_t keylen) // (keylen=0: no key) +{ + size_t i; + + if (outlen == 0 || outlen > 64 || keylen > 64) + return -1; // illegal parameters + + for (i = 0; i < 8; i++) // state, "param block" + ctx->h[i] = blake2b_iv[i]; + ctx->h[0] ^= 0x01010000 ^ (keylen << 8) ^ outlen; + + ctx->t[0] = 0; // input count low word + ctx->t[1] = 0; // input count high word + ctx->c = 0; // pointer within buffer + ctx->outlen = outlen; + + for (i = keylen; i < 128; i++) // zero input block + ctx->b[i] = 0; + if (keylen > 0) { + blake2b_update(ctx, key, keylen); + ctx->c = 128; // at the end + } + + return 0; +} + +// Add "inlen" bytes from "in" into the hash. + +void blake2b_update(blake2b_ctx *ctx, + const void *in, size_t inlen) // data bytes +{ + size_t i; + + for (i = 0; i < inlen; i++) { + if (ctx->c == 128) { // buffer full ? + ctx->t[0] += ctx->c; // add counters + if (ctx->t[0] < ctx->c) // carry overflow ? + ctx->t[1]++; // high word + blake2b_compress(ctx, 0); // compress (not last) + ctx->c = 0; // counter to zero + } + ctx->b[ctx->c++] = ((const uint8_t *) in)[i]; + } +} + +// Generate the message digest (size given in init). +// Result placed in "out". + +void blake2b_final(blake2b_ctx *ctx, void *out) +{ + size_t i; + + ctx->t[0] += ctx->c; // mark last block offset + if (ctx->t[0] < ctx->c) // carry overflow + ctx->t[1]++; // high word + + while (ctx->c < 128) // fill up with zeros + ctx->b[ctx->c++] = 0; + blake2b_compress(ctx, 1); // final block flag = 1 + + // little endian convert and store + for (i = 0; i < ctx->outlen; i++) { + ((uint8_t *) out)[i] = + (ctx->h[i >> 3] >> (8 * (i & 7))) & 0xFF; + } +} + +#ifdef __APPLE_CC__ +static +#endif +void siaHash(void *state, const void *input) +{ + blake2b_ctx ctx; + blake2b_init(&ctx, 32, NULL, 0); + blake2b_update(&ctx, input, 80); + blake2b_final(&ctx, state); +} + +void sia_regenhash(struct work *work) +{ + uint32_t data[20]; + uint32_t hash[16]; + char *scratchbuf; + uint32_t *nonce = (uint32_t *)(work->data + 32); + uint32_t *ohash = (uint32_t *)(work->hash); + + be32enc_vect(data, (const uint32_t *)work->data, 20); + data[8] = htobe32(*nonce); + siaHash(hash, data); + swab256(ohash, hash); +} diff --git a/algorithm/sia.h b/algorithm/sia.h new file mode 100644 index 00000000..413a76c7 --- /dev/null +++ b/algorithm/sia.h @@ -0,0 +1,8 @@ +#ifndef SIAH_H +#define SIAH_H + +#include "miner.h" + +extern void sia_regenhash(struct work *work); + +#endif /* FRESHH_H */ \ No newline at end of file diff --git a/kernel/sia.cl b/kernel/sia.cl new file mode 100644 index 00000000..79a1354a --- /dev/null +++ b/kernel/sia.cl @@ -0,0 +1,120 @@ + +#if __ENDIAN_LITTLE__ + #define SPH_LITTLE_ENDIAN 1 +#else + #define SPH_BIG_ENDIAN 1 +#endif + +#define SPH_UPTR sph_u64 + +typedef unsigned int sph_u32; +typedef int sph_s32; +#ifndef __OPENCL_VERSION__ + typedef unsigned long long sph_u64; + typedef long long sph_s64; +#else + typedef unsigned long sph_u64; + typedef long sph_s64; +#endif + +#define SPH_64 1 +#define SPH_64_TRUE 1 + +#define SWAP4(x) as_uint(as_uchar4(x).wzyx) +#define SWAP8(x) as_ulong(as_uchar8(x).s76543210) + +#if SPH_BIG_ENDIAN + #define DEC64E(x) (x) + #define DEC64BE(x) (*(const __global sph_u64 *) (x)); + #define DEC32LE(x) SWAP4(*(const __global sph_u32 *) (x)); +#else + #define DEC64E(x) SWAP8(x) + #define DEC64BE(x) SWAP8(*(const __global sph_u64 *) (x)); + #define DEC64LE(x) (*(const __global sph_u64 *) (x)); + #define DEC32LE(x) (*(const __global sph_u32 *) (x)); +#endif + +inline static uint2 ror64(const uint2 x, const uint y) +{ + return (uint2)(((x).x>>y)^((x).y<<(32-y)),((x).y>>y)^((x).x<<(32-y))); +} +inline static uint2 ror64_2(const uint2 x, const uint y) +{ + return (uint2)(((x).y>>(y-32))^((x).x<<(64-y)),((x).x>>(y-32))^((x).y<<(64-y))); +} +__constant static const uchar blake2b_sigma[12][16] = { + { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 } , + { 14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 } , + { 11, 8, 12, 0, 5, 2, 15, 13, 10, 14, 3, 6, 7, 1, 9, 4 } , + { 7, 9, 3, 1, 13, 12, 11, 14, 2, 6, 5, 10, 4, 0, 15, 8 } , + { 9, 0, 5, 7, 2, 4, 10, 15, 14, 1, 11, 12, 6, 8, 3, 13 } , + { 2, 12, 6, 10, 0, 11, 8, 3, 4, 13, 7, 5, 15, 14, 1, 9 } , + { 12, 5, 1, 15, 14, 13, 4, 10, 0, 7, 6, 3, 9, 2, 8, 11 } , + { 13, 11, 7, 14, 12, 1, 3, 9, 5, 0, 15, 4, 8, 6, 2, 10 } , + { 6, 15, 14, 9, 11, 3, 0, 8, 12, 2, 13, 7, 1, 4, 10, 5 } , + { 10, 2, 8, 4, 7, 6, 1, 5, 15, 11, 9, 14, 3, 12, 13, 0 } , + { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 } , + { 14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 } }; + +__kernel void search(__global unsigned char* block, volatile __global uint* output, const ulong target) { + sph_u32 gid = get_global_id(0); + + ulong m[16]; + m[0] = DEC64LE(block + 0); + m[1] = DEC64LE(block + 8); + m[2] = DEC64LE(block + 16); + m[3] = DEC64LE(block + 24); + m[4] = DEC64LE(block + 32); + m[4] &= 0xFFFFFFFF00000000; + m[4] ^= (gid); + m[5] = DEC64LE(block + 40); + m[6] = DEC64LE(block + 48); + m[7] = DEC64LE(block + 56); + m[8] = DEC64LE(block + 64); + m[9] = DEC64LE(block + 72); + m[10] = m[11] = m[12] = m[13] = m[14] = m[15] = 0; + + ulong v[16] = { 0x6a09e667f2bdc928, 0xbb67ae8584caa73b, 0x3c6ef372fe94f82b, 0xa54ff53a5f1d36f1, + 0x510e527fade682d1, 0x9b05688c2b3e6c1f, 0x1f83d9abfb41bd6b, 0x5be0cd19137e2179, + 0x6a09e667f3bcc908, 0xbb67ae8584caa73b, 0x3c6ef372fe94f82b, 0xa54ff53a5f1d36f1, + 0x510e527fade68281, 0x9b05688c2b3e6c1f, 0xe07c265404be4294, 0x5be0cd19137e2179 }; + +#define G(r,i,a,b,c,d) \ + a = a + b + m[ blake2b_sigma[r][2*i] ]; \ + ((uint2*)&d)[0] = ((uint2*)&d)[0].yx ^ ((uint2*)&a)[0].yx; \ + c = c + d; \ + ((uint2*)&b)[0] = ror64( ((uint2*)&b)[0] ^ ((uint2*)&c)[0], 24U); \ + a = a + b + m[ blake2b_sigma[r][2*i+1] ]; \ + ((uint2*)&d)[0] = ror64( ((uint2*)&d)[0] ^ ((uint2*)&a)[0], 16U); \ + c = c + d; \ + ((uint2*)&b)[0] = ror64_2( ((uint2*)&b)[0] ^ ((uint2*)&c)[0], 63U); + +#define ROUND(r) \ + G(r,0,v[ 0],v[ 4],v[ 8],v[12]); \ + G(r,1,v[ 1],v[ 5],v[ 9],v[13]); \ + G(r,2,v[ 2],v[ 6],v[10],v[14]); \ + G(r,3,v[ 3],v[ 7],v[11],v[15]); \ + G(r,4,v[ 0],v[ 5],v[10],v[15]); \ + G(r,5,v[ 1],v[ 6],v[11],v[12]); \ + G(r,6,v[ 2],v[ 7],v[ 8],v[13]); \ + G(r,7,v[ 3],v[ 4],v[ 9],v[14]); + ROUND( 0 ); + ROUND( 1 ); + ROUND( 2 ); + ROUND( 3 ); + ROUND( 4 ); + ROUND( 5 ); + ROUND( 6 ); + ROUND( 7 ); + ROUND( 8 ); + ROUND( 9 ); + ROUND( 10 ); + ROUND( 11 ); + +#undef G +#undef ROUND + + bool result = (SWAP8(0x6a09e667f2bdc928 ^ v[0] ^ v[8]) <= target); + if (result) + output[output[0xFF]++] = SWAP4(gid); +} diff --git a/sgminer.c b/sgminer.c index 896f7abc..3ebc8ee8 100644 --- a/sgminer.c +++ b/sgminer.c @@ -5630,6 +5630,9 @@ static void *stratum_sthread(void *userdata) else if (pool->algorithm.type == ALGO_LBRY) { nonce = *((uint32_t *)(work->data + 108)); } + else if (pool->algorithm.type == ALGO_SIA) { + nonce = *((uint32_t *)(work->data + 32)); + } else { nonce = *((uint32_t *)(work->data + 76)); } @@ -6116,7 +6119,7 @@ static void gen_stratum_work(struct pool *pool, struct work *work) cg_wlock(&pool->data_lock); nonce2le = htole64(pool->nonce2); - if (pool->algorithm.type != ALGO_DECRED) { + if (pool->algorithm.type != ALGO_DECRED && pool->algorithm.type != ALGO_SIA) { /* Update coinbase. Always use an LE encoded nonce2 to fill in values * from left to right and prevent overflow errors with small n2sizes */ memcpy(pool->coinbase + pool->nonce2_offset, &nonce2le, pool->n2size); @@ -6127,7 +6130,7 @@ static void gen_stratum_work(struct pool *pool, struct work *work) /* Downgrade to a read lock to read off the pool variables */ cg_dwlock(&pool->data_lock); - if (pool->algorithm.type != ALGO_DECRED) { + if (pool->algorithm.type != ALGO_DECRED && pool->algorithm.type != ALGO_SIA) { /* Generate merkle root */ pool->algorithm.gen_hash(pool->coinbase, pool->swork.cb_len, merkle_root); memcpy(merkle_sha, merkle_root, 32); @@ -6181,6 +6184,14 @@ static void gen_stratum_work(struct pool *pool, struct work *work) memcpy(work->data + 144, pool->nonce1bin, nonce2_offset); memcpy(work->data + 144 + nonce2_offset, &nonce2le, pool->n2size); } + else if (pool->algorithm.type == ALGO_SIA) { + size_t nonce2_offset = MIN(pool->n1_len, 4); + swab256(work->data, pool->header_bin + 4); // prevhash + memcpy(work->data + 32 + 4, pool->nonce1bin, nonce2_offset); + memcpy(work->data + 32 + 4 + nonce2_offset, &nonce2le, pool->n2size); + memcpy(work->data + 32 + 8, pool->header_bin + 68, 4); // timestamp + flip32(work->data + 32 + 8 + 8, pool->coinbase); // merkleroot + } else { data32 = (uint32_t *)merkle_sha; swap32 = (uint32_t *)merkle_root; @@ -7153,6 +7164,7 @@ static void rebuild_nonce(struct work *work, uint32_t nonce) if (work->pool->algorithm.type == ALGO_CRE) nonce_pos = 140; else if (work->pool->algorithm.type == ALGO_DECRED) nonce_pos = 140; else if (work->pool->algorithm.type == ALGO_LBRY) nonce_pos = 108; + else if (work->pool->algorithm.type == ALGO_SIA) nonce_pos = 32; uint32_t *work_nonce = (uint32_t *)(work->data + nonce_pos); @@ -7189,7 +7201,7 @@ static void update_work_stats(struct thr_info *thr, struct work *work) test_diff *= work->pool->algorithm.share_diff_multiplier; - if (unlikely(work->share_diff >= test_diff)) { + if (unlikely(work->pool->algorithm.type != ALGO_SIA && work->share_diff >= test_diff)) { work->block = true; work->pool->solved++; found_blocks++;