Add MaxCoin support

Makefile.am

@@ -55,6 +55,7 @@ sgminer_SOURCES += groestlcoin.c groestlcoin.h
 sgminer_SOURCES += sifcoin.c sifcoin.h
 sgminer_SOURCES += twecoin.c twecoin.h
 sgminer_SOURCES += marucoin.c marucoin.h
+sgminer_SOURCES += maxcoin.c maxcoin.h
 sgminer_SOURCES += kernel/*.cl
 
 bin_SCRIPTS	= $(top_srcdir)/kernel/*.cl

algorithm.c

@@ -23,6 +23,7 @@
 #include "groestlcoin.h"
 #include "twecoin.h"
 #include "marucoin.h"
+#include "maxcoin.h"
 
 #include <inttypes.h>
 #include <string.h>
@@ -79,12 +80,32 @@ static cl_int queue_sph_kernel(struct __clState *clState, struct _dev_blk_ctx *b
     return status;
 }
 
+static cl_int queue_maxcoin_kernel(struct __clState *clState, struct _dev_blk_ctx *blk, __maybe_unused cl_uint threads)
+{
+    cl_kernel *kernel = &clState->kernel;
+    unsigned int num = 0;
+    cl_int status = 0;
+
+    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);
+
+    return status;
+}
+
 typedef struct _algorithm_settings_t {
     const char *name; /* Human-readable identifier */
     double   diff_multiplier1;
     double   diff_multiplier2;
+    double   share_diff_multiplier;
+    uint32_t xintensity_shift;
+    uint32_t intensity_shift;
+    uint32_t found_idx;
     unsigned long long   diff_nonce;
     unsigned long long   diff_numerator;
+    uint32_t diff1targ;
     void     (*regenhash)(struct work *);
     cl_int   (*queue_kernel)(struct __clState *, struct _dev_blk_ctx *, cl_uint);
     void     (*gen_hash)(const unsigned char *, unsigned int, unsigned char *);
@@ -93,7 +114,7 @@ typedef struct _algorithm_settings_t {
 static algorithm_settings_t algos[] = {
     // kernels starting from this will have difficulty calculated by using litecoin algorithm
 #define A_SCRYPT(a) \
-    { a, 1, 65536, 0x0000ffff00000000ULL, 0xFFFFFFFFULL, scrypt_regenhash, queue_scrypt_kernel, gen_hash}
+    { a, 1, 65536, 65536, 0, 0, 0xFF, 0x0000ffff00000000ULL, 0xFFFFFFFFULL, 0x0000ffffUL, scrypt_regenhash, queue_scrypt_kernel, gen_hash}
     A_SCRYPT( "ckolivas" ),
     A_SCRYPT( "alexkarnew" ),
     A_SCRYPT( "alexkarnold" ),
@@ -104,7 +125,7 @@ static algorithm_settings_t algos[] = {
 
     // kernels starting from this will have difficulty calculated by using quarkcoin algorithm
 #define A_QUARK(a, b) \
-    { a, 256, 256, 0x000000ffff000000ULL, 0xFFFFFFULL, b, queue_sph_kernel, gen_hash}
+    { a, 256, 256, 256, 0, 0, 0xFF, 0x000000ffff000000ULL, 0xFFFFFFULL, 0x0000ffffUL, b, queue_sph_kernel, gen_hash}
     A_QUARK( "quarkcoin", quarkcoin_regenhash),
     A_QUARK( "qubitcoin", qubitcoin_regenhash),
     A_QUARK( "animecoin", animecoin_regenhash),
@@ -113,24 +134,25 @@ static algorithm_settings_t algos[] = {
 
     // kernels starting from this will have difficulty calculated by using bitcoin algorithm
 #define A_DARK(a, b) \
-    { a, 1, 1, 0x00000000ffff0000ULL, 0xFFFFULL, b, queue_sph_kernel, gen_hash}
+    { a, 1, 1, 1, 0, 0, 0xFF, 0x00000000ffff0000ULL, 0xFFFFULL, 0x0000ffffUL, b, queue_sph_kernel, gen_hash}
     A_DARK( "darkcoin",           darkcoin_regenhash),
     A_DARK( "inkcoin",            inkcoin_regenhash),
     A_DARK( "myriadcoin-groestl", myriadcoin_groestl_regenhash),
     A_DARK( "marucoin",           marucoin_regenhash),
 #undef A_DARK
 
-    { "twecoin", 1, 1, 0x00000000ffff0000ULL, 0xFFFFULL, twecoin_regenhash, queue_sph_kernel, sha256},
+    { "twecoin", 1, 1, 1, 0, 0, 0xFF, 0x00000000ffff0000ULL, 0xFFFFULL, 0x0000ffffUL, twecoin_regenhash, queue_sph_kernel, sha256},
+    { "maxcoin", 1, 256, 1, 4, 15, 0x0F, 0x00000000ffff0000ULL, 0xFFFFULL, 0x000000ffUL, maxcoin_regenhash, queue_maxcoin_kernel, sha256},
 
     // kernels starting from this will have difficulty calculated by using fuguecoin algorithm
 #define A_FUGUE(a, b) \
-    { a, 1, 256, 0x00000000ffff0000ULL, 0xFFFFULL, b, queue_sph_kernel, sha256}
+    { a, 1, 256, 256, 0, 0, 0xFF, 0x00000000ffff0000ULL, 0xFFFFULL, 0x0000ffffUL, b, queue_sph_kernel, sha256}
     A_FUGUE( "fuguecoin",   fuguecoin_regenhash),
     A_FUGUE( "groestlcoin", groestlcoin_regenhash),
 #undef A_FUGUE
 
     // Terminator (do not remove)
-    { NULL, 0, 0, 0, 0, NULL, NULL, NULL}
+    { NULL, 0, 0, 0, 0, 0, 0, 0, 0, 0, NULL, NULL, NULL}
 };
 
 void copy_algorithm_settings(algorithm_t* dest, const char* algo) {
@@ -143,8 +165,13 @@ void copy_algorithm_settings(algorithm_t* dest, const char* algo) {
 
             dest->diff_multiplier1 = src->diff_multiplier1;
             dest->diff_multiplier2 = src->diff_multiplier2;
+            dest->share_diff_multiplier = src->share_diff_multiplier;
+            dest->xintensity_shift = src->xintensity_shift;
+            dest->intensity_shift = src->intensity_shift;
+            dest->found_idx = src->found_idx;
             dest->diff_nonce = src->diff_nonce;
             dest->diff_numerator = src->diff_numerator;
+            dest->diff1targ = src->diff1targ;
             dest->regenhash = src->regenhash;
             dest->queue_kernel = src->queue_kernel;
             dest->gen_hash = src->gen_hash;

algorithm.h

@@ -25,8 +25,13 @@ typedef struct _algorithm_t {
     uint8_t  nfactor;  /* Factor of N above (n = 2^nfactor) */
     double   diff_multiplier1;
     double   diff_multiplier2;
+    double   share_diff_multiplier;
+    uint32_t xintensity_shift;
+    uint32_t intensity_shift;
+    uint32_t found_idx;
     unsigned long long   diff_nonce;
     unsigned long long   diff_numerator;
+    uint32_t diff1targ;
     void     (*regenhash)(struct work *);
     cl_int   (*queue_kernel)(struct __clState *, struct _dev_blk_ctx *, cl_uint);
     void     (*gen_hash)(const unsigned char *, unsigned int, unsigned char *);

driver-opencl.c

@@ -964,16 +964,20 @@ void manage_gpu(void)
 static _clState *clStates[MAX_GPUDEVICES];
 
 static void set_threads_hashes(unsigned int vectors, unsigned int compute_shaders, int64_t *hashes, size_t *globalThreads,
-			       unsigned int minthreads, __maybe_unused int *intensity, __maybe_unused int *xintensity, __maybe_unused int *rawintensity)
+			       unsigned int minthreads, __maybe_unused int *intensity, __maybe_unused int *xintensity,
+             __maybe_unused int *rawintensity, algorithm_t *algorithm)
 {
 	unsigned int threads = 0;
 	while (threads < minthreads) {
 		if (*rawintensity > 0) {
 			threads = *rawintensity;
 		} else if (*xintensity > 0) {
-			threads = compute_shaders * *xintensity;
+      if (algorithm->xintensity_shift)
+        threads = compute_shaders * (1 << (algorithm->xintensity_shift + *xintensity));
+      else
+        threads = compute_shaders * *xintensity;
 		} else {
-			threads = 1 << *intensity;
+			threads = 1 << (algorithm->intensity_shift + *intensity);
 		}
 		if (threads < minthreads) {
 			if (likely(*intensity < MAX_INTENSITY))
@@ -1285,7 +1289,6 @@ static bool opencl_thread_init(struct thr_info *thr)
 	return true;
 }
 
-
 static bool opencl_prepare_work(struct thr_info __maybe_unused *thr, struct work *work)
 {
 	work->blk.work = work;
@@ -1308,7 +1311,7 @@ static int64_t opencl_scanhash(struct thr_info *thr, struct work *work,
 	size_t globalThreads[1];
 	size_t localThreads[1] = { clState->wsize };
 	int64_t hashes;
-	int found = FOUND;
+	int found = gpu->algorithm.found_idx;
 	int buffersize = BUFFERSIZE;
 
 	/* Windows' timer resolution is only 15ms so oversample 5x */
@@ -1330,7 +1333,7 @@ static int64_t opencl_scanhash(struct thr_info *thr, struct work *work,
 	}
 
 	set_threads_hashes(clState->vwidth, clState->compute_shaders, &hashes, globalThreads, localThreads[0],
-			   &gpu->intensity, &gpu->xintensity, &gpu->rawintensity);
+			   &gpu->intensity, &gpu->xintensity, &gpu->rawintensity, &gpu->algorithm);
 	if (hashes > gpu->max_hashes)
 		gpu->max_hashes = hashes;
 
@@ -1369,7 +1372,7 @@ static int64_t opencl_scanhash(struct thr_info *thr, struct work *work,
 	/* This finish flushes the readbuffer set with CL_FALSE in clEnqueueReadBuffer */
 	clFinish(clState->commandQueue);
 
-	/* FOUND entry is used as a counter to say how many nonces exist */
+	/* found entry is used as a counter to say how many nonces exist */
 	if (thrdata->res[found]) {
 		/* Clear the buffer again */
 		status = clEnqueueWriteBuffer(clState->commandQueue, clState->outputBuffer, CL_FALSE, 0,

findnonce.c

@@ -183,7 +183,7 @@ static void *postcalc_hash(void *userdata)
 	struct thr_info *thr = pcd->thr;
 	unsigned int entry = 0;
 
-	int found = FOUND;
+	int found = thr->cgpu->algorithm.found_idx;
 
 	pthread_detach(pthread_self());
 
@@ -199,6 +199,8 @@ static void *postcalc_hash(void *userdata)
 
 	for (entry = 0; entry < pcd->res[found]; entry++) {
 		uint32_t nonce = pcd->res[entry];
+		if (found == 0x0F)
+        nonce = swab32(nonce);
 
 		applog(LOG_DEBUG, "OCL NONCE %u found in slot %d", nonce, entry);
 		submit_nonce(thr, pcd->work, nonce);

findnonce.h

@@ -5,10 +5,8 @@
 #include "config.h"
 
 #define MAXTHREADS (0xFFFFFFFEULL)
-
 #define MAXBUFFERS (0x100)
 #define BUFFERSIZE (sizeof(uint32_t) * MAXBUFFERS)
-#define FOUND (0xFF)
 
 extern void precalc_hash(dev_blk_ctx *blk, uint32_t *state, uint32_t *data);
 extern void postcalc_hash_async(struct thr_info *thr, struct work *work, uint32_t *res);

kernel/maxcoin.cl

@@ -0,0 +1,131 @@
+#define ARGS_25(x) x ## 0, x ## 1, x ## 2, x ## 3, x ## 4, x ## 5, x ## 6, x ## 7, x ## 8, x ## 9, x ## 10, x ## 11, x ## 12, x ## 13, x ## 14, x ## 15, x ## 16, x ## 17, x ## 18, x ## 19, x ## 20, x ## 21, x ## 22, x ## 23, x ## 24
+
+__constant uint2 keccak_round_constants[24] = 
+{
+  (uint2)(0x00000001,0x00000000), (uint2)(0x00008082,0x00000000),
+  (uint2)(0x0000808a,0x80000000), (uint2)(0x80008000,0x80000000),
+  (uint2)(0x0000808b,0x00000000), (uint2)(0x80000001,0x00000000),
+  (uint2)(0x80008081,0x80000000), (uint2)(0x00008009,0x80000000),
+  (uint2)(0x0000008a,0x00000000), (uint2)(0x00000088,0x00000000),
+  (uint2)(0x80008009,0x00000000), (uint2)(0x8000000a,0x00000000),
+  (uint2)(0x8000808b,0x00000000), (uint2)(0x0000008b,0x80000000),
+  (uint2)(0x00008089,0x80000000), (uint2)(0x00008003,0x80000000),
+  (uint2)(0x00008002,0x80000000), (uint2)(0x00000080,0x80000000),
+  (uint2)(0x0000800a,0x00000000), (uint2)(0x8000000a,0x80000000),
+  (uint2)(0x80008081,0x80000000), (uint2)(0x00008080,0x80000000),
+  (uint2)(0x80000001,0x00000000), (uint2)(0x80008008,0x80000000)
+};
+
+uint2 ROTL64_1(const uint2 x, const uint y)
+{
+  return (uint2)((x.x<<y)^(x.y>>(32-y)),(x.y<<y)^(x.x>>(32-y)));
+}
+uint2 ROTL64_2(const uint2 x, const uint y)
+{
+  return (uint2)((x.y<<y)^(x.x>>(32-y)),(x.x<<y)^(x.y>>(32-y)));
+}
+
+#define RND(i) \
+    m0 = *s0 ^ *s5 ^ *s10 ^ *s15 ^ *s20 ^ ROTL64_1(*s2 ^ *s7 ^ *s12 ^ *s17 ^ *s22, 1);\
+    m1 = *s1 ^ *s6 ^ *s11 ^ *s16 ^ *s21 ^ ROTL64_1(*s3 ^ *s8 ^ *s13 ^ *s18 ^ *s23, 1);\
+    m2 = *s2 ^ *s7 ^ *s12 ^ *s17 ^ *s22 ^ ROTL64_1(*s4 ^ *s9 ^ *s14 ^ *s19 ^ *s24, 1);\
+    m3 = *s3 ^ *s8 ^ *s13 ^ *s18 ^ *s23 ^ ROTL64_1(*s0 ^ *s5 ^ *s10 ^ *s15 ^ *s20, 1);\
+    m4 = *s4 ^ *s9 ^ *s14 ^ *s19 ^ *s24 ^ ROTL64_1(*s1 ^ *s6 ^ *s11 ^ *s16 ^ *s21, 1);\
+\
+    m5 = *s1^m0;\
+\
+    *s0 ^= m4;\
+    *s1 = ROTL64_2(*s6^m0, 12);\
+    *s6 = ROTL64_1(*s9^m3, 20);\
+    *s9 = ROTL64_2(*s22^m1, 29);\
+    *s22 = ROTL64_2(*s14^m3, 7);\
+    *s14 = ROTL64_1(*s20^m4, 18);\
+    *s20 = ROTL64_2(*s2^m1, 30);\
+    *s2 = ROTL64_2(*s12^m1, 11);\
+    *s12 = ROTL64_1(*s13^m2, 25);\
+    *s13 = ROTL64_1(*s19^m3,  8);\
+    *s19 = ROTL64_2(*s23^m2, 24);\
+    *s23 = ROTL64_2(*s15^m4, 9);\
+    *s15 = ROTL64_1(*s4^m3, 27);\
+    *s4 = ROTL64_1(*s24^m3, 14);\
+    *s24 = ROTL64_1(*s21^m0,  2);\
+    *s21 = ROTL64_2(*s8^m2, 23);\
+    *s8 = ROTL64_2(*s16^m0, 13);\
+    *s16 = ROTL64_2(*s5^m4, 4);\
+    *s5 = ROTL64_1(*s3^m2, 28);\
+    *s3 = ROTL64_1(*s18^m2, 21);\
+    *s18 = ROTL64_1(*s17^m1, 15);\
+    *s17 = ROTL64_1(*s11^m0, 10);\
+    *s11 = ROTL64_1(*s7^m1,  6);\
+    *s7 = ROTL64_1(*s10^m4,  3);\
+    *s10 = ROTL64_1(      m5,  1);\
+    \
+    m5 = *s0; m6 = *s1; *s0 = bitselect(*s0^*s2,*s0,*s1); *s1 = bitselect(*s1^*s3,*s1,*s2); *s2 = bitselect(*s2^*s4,*s2,*s3); *s3 = bitselect(*s3^m5,*s3,*s4); *s4 = bitselect(*s4^m6,*s4,m5);\
+    m5 = *s5; m6 = *s6; *s5 = bitselect(*s5^*s7,*s5,*s6); *s6 = bitselect(*s6^*s8,*s6,*s7); *s7 = bitselect(*s7^*s9,*s7,*s8); *s8 = bitselect(*s8^m5,*s8,*s9); *s9 = bitselect(*s9^m6,*s9,m5);\
+    m5 = *s10; m6 = *s11; *s10 = bitselect(*s10^*s12,*s10,*s11); *s11 = bitselect(*s11^*s13,*s11,*s12); *s12 = bitselect(*s12^*s14,*s12,*s13); *s13 = bitselect(*s13^m5,*s13,*s14); *s14 = bitselect(*s14^m6,*s14,m5);\
+    m5 = *s15; m6 = *s16; *s15 = bitselect(*s15^*s17,*s15,*s16); *s16 = bitselect(*s16^*s18,*s16,*s17); *s17 = bitselect(*s17^*s19,*s17,*s18); *s18 = bitselect(*s18^m5,*s18,*s19); *s19 = bitselect(*s19^m6,*s19,m5);\
+    m5 = *s20; m6 = *s21; *s20 = bitselect(*s20^*s22,*s20,*s21); *s21 = bitselect(*s21^*s23,*s21,*s22); *s22 = bitselect(*s22^*s24,*s22,*s23); *s23 = bitselect(*s23^m5,*s23,*s24); *s24 = bitselect(*s24^m6,*s24,m5);\
+\
+    *s0 ^= keccak_round_constants[i];
+
+void keccak_block_noabsorb(ARGS_25(uint2* s))
+{
+  uint2 m0,m1,m2,m3,m4,m5,m6;
+  RND(0);
+  for (int i = 1; i < 22; ++i) 
+  {
+    RND(i);
+    ++i;
+    RND(i);
+    ++i;
+    RND(i);
+  }
+  RND(22);
+  RND(23);
+}
+
+__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
+__kernel void search(__global const uint2*restrict in, __global uint*restrict output)
+{
+  uint2 ARGS_25(state);
+  
+  state0 = in[0];
+  state1 = in[1];
+  state2 = in[2];
+  state3 = in[3];
+  state4 = in[4];
+  state5 = in[5];
+  state6 = in[6];
+  state7 = in[7];
+  state8 = in[8];
+  state9 = (uint2)(in[9].x,get_global_id(0));
+  state10 = (uint2)(1,0);
+  state11 = 0;
+  state12 = 0;
+  state13 = 0;
+  state14 = 0;
+  state15 = 0;
+  state16 = (uint2)(0,0x80000000U);
+  state17 = 0;
+  state18 = 0;
+  state19 = 0;
+  state20 = 0;
+  state21 = 0;
+  state22 = 0;
+  state23 = 0;
+  state24 = 0;
+  
+  keccak_block_noabsorb(ARGS_25(&state));
+  
+#define FOUND (0x0F)
+#define SETFOUND(Xnonce) output[output[FOUND]++] = Xnonce
+        
+        if ((state3.y & 0xFFFFFFF0U) == 0)
+        {
+                SETFOUND(get_global_id(0));
+        }
+}
+
+/*-
+ * Scrypt-jane public domain, OpenCL implementation of scrypt(keccak,chacha,SCRYPTN,1,1) 2013 mtrlt
+ */

maxcoin.c

@@ -0,0 +1,321 @@
+#include "config.h"
+#include "miner.h"
+
+#include <stdlib.h>
+#include <stdint.h>
+#include <string.h>
+
+#include <stdio.h>
+
+#define CL_SET_BLKARG(blkvar) status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->blkvar)
+#define CL_SET_ARG(var) status |= clSetKernelArg(*kernel, num++, sizeof(var), (void *)&var)
+
+struct uint256
+{
+  unsigned char v[32];
+};
+typedef struct uint256 uint256;
+
+typedef unsigned long long  UINT64;
+
+#define ROL(a, offset) ((a << offset) | (a >> (64-offset)))
+
+static const UINT64 MaxcoinF_RoundConstants[24] = 
+{
+  0x0000000000000001ULL,
+  0x0000000000008082ULL,
+  0x800000000000808aULL,
+  0x8000000080008000ULL,
+  0x000000000000808bULL,
+  0x0000000080000001ULL,
+  0x8000000080008081ULL,
+  0x8000000000008009ULL,
+  0x000000000000008aULL,
+  0x0000000000000088ULL,
+  0x0000000080008009ULL,
+  0x000000008000000aULL,
+  0x000000008000808bULL,
+  0x800000000000008bULL,
+  0x8000000000008089ULL,
+  0x8000000000008003ULL,
+  0x8000000000008002ULL,
+  0x8000000000000080ULL,
+  0x000000000000800aULL,
+  0x800000008000000aULL,
+  0x8000000080008081ULL,
+  0x8000000000008080ULL,
+  0x0000000080000001ULL,
+  0x8000000080008008ULL
+};
+
+struct bin32
+{
+  UINT64 v0;
+  UINT64 v1;
+  UINT64 v2;
+  UINT64 v3;
+};
+
+void maxcoin1(unsigned char *out, const unsigned char *inraw, unsigned inrawlen)
+{
+  unsigned char temp[136];
+  unsigned round;
+
+  UINT64 Aba, Abe, Abi, Abo, Abu;
+  UINT64 Aga, Age, Agi, Ago, Agu;
+  UINT64 Aka, Ake, Aki, Ako, Aku;
+  UINT64 Ama, Ame, Ami, Amo, Amu;
+  UINT64 Asa, Ase, Asi, Aso, Asu;
+  UINT64 BCa, BCe, BCi, BCo, BCu;
+  UINT64 Da, De, Di, Do, Du;
+  UINT64 Eba, Ebe, Ebi, Ebo, Ebu;
+  UINT64 Ega, Ege, Egi, Ego, Egu;
+  UINT64 Eka, Eke, Eki, Eko, Eku;
+  UINT64 Ema, Eme, Emi, Emo, Emu;
+  UINT64 Esa, Ese, Esi, Eso, Esu;
+
+  memcpy(temp, inraw, inrawlen);
+  temp[inrawlen++] = 1;
+  memset( temp+inrawlen, 0, 136 - inrawlen);
+  temp[136-1] |= 0x80;
+  const UINT64 *in = (const UINT64 *)temp;
+
+  //copyFromState(A, state)
+  Aba = in[ 0];
+  Abe = in[ 1];
+  Abi = in[ 2];
+  Abo = in[ 3];
+  Abu = in[ 4];
+  Aga = in[ 5];
+  Age = in[ 6];
+  Agi = in[ 7];
+  Ago = in[ 8];
+  Agu = in[ 9];
+  Aka = in[10];
+  Ake = in[11];
+  Aki = in[12];
+  Ako = in[13];
+  Aku = in[14];
+  Ama = in[15];
+  Ame = in[16];
+  Ami = 0;
+  Amo = 0;
+  Amu = 0;
+  Asa = 0;
+  Ase = 0;
+  Asi = 0;
+  Aso = 0;
+  Asu = 0;
+
+  for( round = 0; round < 24; round += 2 )
+  {
+    //    prepareTheta
+    BCa = Aba^Aga^Aka^Ama^Asa;
+    BCe = Abe^Age^Ake^Ame^Ase;
+    BCi = Abi^Agi^Aki^Ami^Asi;
+    BCo = Abo^Ago^Ako^Amo^Aso;
+    BCu = Abu^Agu^Aku^Amu^Asu;
+
+    //thetaRhoPiChiIotaPrepareTheta(round  , A, E)
+    Da = BCu^ROL(BCe, 1);
+    De = BCa^ROL(BCi, 1);
+    Di = BCe^ROL(BCo, 1);
+    Do = BCi^ROL(BCu, 1);
+    Du = BCo^ROL(BCa, 1);
+
+    Aba ^= Da;
+    BCa = Aba;
+    Age ^= De;
+    BCe = ROL(Age, 44);
+    Aki ^= Di;
+    BCi = ROL(Aki, 43);
+    Amo ^= Do;
+    BCo = ROL(Amo, 21);
+    Asu ^= Du;
+    BCu = ROL(Asu, 14);
+    Eba =   BCa ^((~BCe)&  BCi );
+    Eba ^= MaxcoinF_RoundConstants[round];
+    Ebe =   BCe ^((~BCi)&  BCo );
+    Ebi =   BCi ^((~BCo)&  BCu );
+    Ebo =   BCo ^((~BCu)&  BCa );
+    Ebu =   BCu ^((~BCa)&  BCe );
+
+    Abo ^= Do;
+    BCa = ROL(Abo, 28);
+    Agu ^= Du;
+    BCe = ROL(Agu, 20);
+    Aka ^= Da;
+    BCi = ROL(Aka,  3);
+    Ame ^= De;
+    BCo = ROL(Ame, 45);
+    Asi ^= Di;
+    BCu = ROL(Asi, 61);
+    Ega =   BCa ^((~BCe)&  BCi );
+    Ege =   BCe ^((~BCi)&  BCo );
+    Egi =   BCi ^((~BCo)&  BCu );
+    Ego =   BCo ^((~BCu)&  BCa );
+    Egu =   BCu ^((~BCa)&  BCe );
+
+    Abe ^= De;
+    BCa = ROL(Abe,  1);
+    Agi ^= Di;
+    BCe = ROL(Agi,  6);
+    Ako ^= Do;
+    BCi = ROL(Ako, 25);
+    Amu ^= Du;
+    BCo = ROL(Amu,  8);
+    Asa ^= Da;
+    BCu = ROL(Asa, 18);
+    Eka =   BCa ^((~BCe)&  BCi );
+    Eke =   BCe ^((~BCi)&  BCo );
+    Eki =   BCi ^((~BCo)&  BCu );
+    Eko =   BCo ^((~BCu)&  BCa );
+    Eku =   BCu ^((~BCa)&  BCe );
+
+    Abu ^= Du;
+    BCa = ROL(Abu, 27);
+    Aga ^= Da;
+    BCe = ROL(Aga, 36);
+    Ake ^= De;
+    BCi = ROL(Ake, 10);
+    Ami ^= Di;
+    BCo = ROL(Ami, 15);
+    Aso ^= Do;
+    BCu = ROL(Aso, 56);
+    Ema =   BCa ^((~BCe)&  BCi );
+    Eme =   BCe ^((~BCi)&  BCo );
+    Emi =   BCi ^((~BCo)&  BCu );
+    Emo =   BCo ^((~BCu)&  BCa );
+    Emu =   BCu ^((~BCa)&  BCe );
+
+    Abi ^= Di;
+    BCa = ROL(Abi, 62);
+    Ago ^= Do;
+    BCe = ROL(Ago, 55);
+    Aku ^= Du;
+    BCi = ROL(Aku, 39);
+    Ama ^= Da;
+    BCo = ROL(Ama, 41);
+    Ase ^= De;
+    BCu = ROL(Ase,  2);
+    Esa =   BCa ^((~BCe)&  BCi );
+    Ese =   BCe ^((~BCi)&  BCo );
+    Esi =   BCi ^((~BCo)&  BCu );
+    Eso =   BCo ^((~BCu)&  BCa );
+    Esu =   BCu ^((~BCa)&  BCe );
+
+    //    prepareTheta
+    BCa = Eba^Ega^Eka^Ema^Esa;
+    BCe = Ebe^Ege^Eke^Eme^Ese;
+    BCi = Ebi^Egi^Eki^Emi^Esi;
+    BCo = Ebo^Ego^Eko^Emo^Eso;
+    BCu = Ebu^Egu^Eku^Emu^Esu;
+
+    //thetaRhoPiChiIotaPrepareTheta(round+1, E, A)
+    Da = BCu^ROL(BCe, 1);
+    De = BCa^ROL(BCi, 1);
+    Di = BCe^ROL(BCo, 1);
+    Do = BCi^ROL(BCu, 1);
+    Du = BCo^ROL(BCa, 1);
+
+    Eba ^= Da;
+    BCa = Eba;
+    Ege ^= De;
+    BCe = ROL(Ege, 44);
+    Eki ^= Di;
+    BCi = ROL(Eki, 43);
+    Emo ^= Do;
+    BCo = ROL(Emo, 21);
+    Esu ^= Du;
+    BCu = ROL(Esu, 14);
+    Aba =   BCa ^((~BCe)&  BCi );
+    Aba ^= MaxcoinF_RoundConstants[round+1];
+    Abe =   BCe ^((~BCi)&  BCo );
+    Abi =   BCi ^((~BCo)&  BCu );
+    Abo =   BCo ^((~BCu)&  BCa );
+    Abu =   BCu ^((~BCa)&  BCe );
+
+    Ebo ^= Do;
+    BCa = ROL(Ebo, 28);
+    Egu ^= Du;
+    BCe = ROL(Egu, 20);
+    Eka ^= Da;
+    BCi = ROL(Eka, 3);
+    Eme ^= De;
+    BCo = ROL(Eme, 45);
+    Esi ^= Di;
+    BCu = ROL(Esi, 61);
+    Aga =   BCa ^((~BCe)&  BCi );
+    Age =   BCe ^((~BCi)&  BCo );
+    Agi =   BCi ^((~BCo)&  BCu );
+    Ago =   BCo ^((~BCu)&  BCa );
+    Agu =   BCu ^((~BCa)&  BCe );
+
+    Ebe ^= De;
+    BCa = ROL(Ebe, 1);
+    Egi ^= Di;
+    BCe = ROL(Egi, 6);
+    Eko ^= Do;
+    BCi = ROL(Eko, 25);
+    Emu ^= Du;
+    BCo = ROL(Emu, 8);
+    Esa ^= Da;
+    BCu = ROL(Esa, 18);
+    Aka =   BCa ^((~BCe)&  BCi );
+    Ake =   BCe ^((~BCi)&  BCo );
+    Aki =   BCi ^((~BCo)&  BCu );
+    Ako =   BCo ^((~BCu)&  BCa );
+    Aku =   BCu ^((~BCa)&  BCe );
+
+    Ebu ^= Du;
+    BCa = ROL(Ebu, 27);
+    Ega ^= Da;
+    BCe = ROL(Ega, 36);
+    Eke ^= De;
+    BCi = ROL(Eke, 10);
+    Emi ^= Di;
+    BCo = ROL(Emi, 15);
+    Eso ^= Do;
+    BCu = ROL(Eso, 56);
+    Ama =   BCa ^((~BCe)&  BCi );
+    Ame =   BCe ^((~BCi)&  BCo );
+    Ami =   BCi ^((~BCo)&  BCu );
+    Amo =   BCo ^((~BCu)&  BCa );
+    Amu =   BCu ^((~BCa)&  BCe );
+
+    Ebi ^= Di;
+    BCa = ROL(Ebi, 62);
+    Ego ^= Do;
+    BCe = ROL(Ego, 55);
+    Eku ^= Du;
+    BCi = ROL(Eku, 39);
+    Ema ^= Da;
+    BCo = ROL(Ema, 41);
+    Ese ^= De;
+    BCu = ROL(Ese, 2);
+    Asa =   BCa ^((~BCe)&  BCi );
+    Ase =   BCe ^((~BCi)&  BCo );
+    Asi =   BCi ^((~BCo)&  BCu );
+    Aso =   BCo ^((~BCu)&  BCa );
+    Asu =   BCu ^((~BCa)&  BCe );
+  }
+  {
+  UINT64 *out64 = (UINT64 *)out;
+  out64[ 0] = Aba;
+  out64[ 1] = Abe;
+  out64[ 2] = Abi;
+  out64[ 3] = Abo;
+  }
+}
+
+void maxcoin_regenhash(struct work *work)
+{
+  uint256 result; 
+
+    unsigned int data[20], datacopy[20]; // aligned for flip80
+    memcpy(datacopy, work->data, 80);
+    flip80(data, datacopy); 
+    maxcoin1((unsigned char*)&result, (unsigned char*)data, 80);
+
+  memcpy(work->hash, &result, 32);
+}

maxcoin.h

@@ -0,0 +1,9 @@
+#ifndef MAXCOIN_H
+#define MAXCOIN_H
+
+#include "miner.h"
+
+// extern int maxcoin_test(unsigned char *pdata, const unsigned char *ptarget, uint32_t nonce);
+extern void maxcoin_regenhash(struct work *work);
+
+#endif /* MAXCOIN_H */

sgminer.c

@@ -3719,7 +3719,7 @@ static double share_diff(const struct work *work)
 	double d64, s64;
 	double ret;
 
-	d64 = work->pool->algorithm.diff_multiplier2 * truediffone;
+	d64 = work->pool->algorithm.share_diff_multiplier * truediffone;
 	s64 = le256todouble(work->hash);
 	if (unlikely(!s64))
 		s64 = 0;
@@ -5947,7 +5947,6 @@ void set_target(unsigned char *dest_target, double diff, double diff_multiplier2
 		diff = 1.0;
 	}
 
-	// FIXME: is target set right?
 	d64 = diff_multiplier2 * truediffone;
 	d64 /= diff;
 
@@ -6223,7 +6222,7 @@ bool test_nonce(struct work *work, uint32_t nonce)
 	uint32_t diff1targ;
 
 	rebuild_nonce(work, nonce);
-	diff1targ = 0x0000ffffUL;
+	diff1targ = work->pool->algorithm.diff1targ;
 
 	return (le32toh(*hash_32) <= diff1targ);
 }
@@ -6243,11 +6242,11 @@ bool test_nonce_diff(struct work *work, uint32_t nonce, double diff)
 static void update_work_stats(struct thr_info *thr, struct work *work)
 {
 	double test_diff = current_diff;
-	test_diff *= work->pool->algorithm.diff_multiplier2;
+	test_diff *= work->pool->algorithm.share_diff_multiplier;
 
 	work->share_diff = share_diff(work);
 
-	test_diff *= work->pool->algorithm.diff_multiplier2;
+	test_diff *= work->pool->algorithm.share_diff_multiplier;
 
 	if (unlikely(work->share_diff >= test_diff)) {
 		work->block = true;