Add Sia algo support

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

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 },

algorithm.h

@@ -36,6 +36,7 @@ typedef enum {
   ALGO_YESCRYPT_MULTI,
   ALGO_BLAKECOIN,
   ALGO_BLAKE,
+  ALGO_SIA,
   ALGO_DECRED,
   ALGO_VANILLA,
   ALGO_LBRY

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 <stdlib.h>
+#include <stdint.h>
+#include <string.h>
+
+// 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);
+}

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 */

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);
+}

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++;