From c336a8399ae1d6340e8a7fae687daf800551fb40 Mon Sep 17 00:00:00 2001
From: phm <phm@inbox.com>
Date: Sat, 22 Feb 2014 16:58:00 +0100
Subject: [PATCH] Added support for INKcoin.

---
 driver-opencl.c   | 1 +
 inkcoin.c         | 6 +++---
 kernel/inkcoin.cl | 2 +-
 miner.h           | 2 +-
 ocl.c             | 1 +
 sgminer.c         | 4 +++-
 6 files changed, 10 insertions(+), 6 deletions(-)

diff --git a/driver-opencl.c b/driver-opencl.c
index 000c453a..c41a7ec8 100644
--- a/driver-opencl.c
+++ b/driver-opencl.c
@@ -1358,6 +1358,7 @@ static bool opencl_thread_prepare(struct thr_info *thr)
 				break;
 			case KL_INKCOIN:
 				cgpu->kname = INKCOIN_KERNNAME;
+				break;
 			case KL_QUARKCOIN:
 				cgpu->kname = QUARKCOIN_KERNNAME;
 				break;
diff --git a/inkcoin.c b/inkcoin.c
index 56652543..d4252f3d 100644
--- a/inkcoin.c
+++ b/inkcoin.c
@@ -51,7 +51,7 @@ inline void inkhash(void *state, const void *input)
 {
     uint32_t hash[16];
     sph_shavite512_context ctx_shavite;
-    
+
     sph_shavite512_init(&ctx_shavite);
     sph_shavite512 (&ctx_shavite, input, 80);
     sph_shavite512_close(&ctx_shavite, hash);
@@ -76,7 +76,7 @@ int inkcoin_test(unsigned char *pdata, const unsigned char *ptarget, uint32_t no
 	be32enc_vect(data, (const uint32_t *)pdata, 19);
 	data[19] = htobe32(nonce);
 	//scratchbuf = alloca(SCRATCHBUF_SIZE);
-	xhash(ohash, data);
+	inkhash(ohash, data);
 	tmp_hash7 = be32toh(ohash[7]);
 
 	applog(LOG_DEBUG, "htarget %08lx diff1 %08lx hash %08lx",
@@ -99,7 +99,7 @@ void inkcoin_regenhash(struct work *work)
 
         be32enc_vect(data, (const uint32_t *)work->data, 19);
         data[19] = htobe32(*nonce);
-        xhash(ohash, data);
+        inkhash(ohash, data);
 }
 
 bool scanhash_inkcoin(struct thr_info *thr, const unsigned char __maybe_unused *pmidstate,
diff --git a/kernel/inkcoin.cl b/kernel/inkcoin.cl
index a8459945..33dba240 100644
--- a/kernel/inkcoin.cl
+++ b/kernel/inkcoin.cl
@@ -148,7 +148,7 @@ __kernel void search(__global unsigned char* block, volatile __global uint* outp
     rk10 = DEC32LE(block + 16 * 4);;
     rk11 = DEC32LE(block + 17 * 4);;
     rk12 = DEC32LE(block + 18 * 4);;
-    rk13 = 72768; // gid;
+    rk13 = gid;
     rk14 = 0x80;
     rk15 = rk16 = rk17 = rk18 = rk19 = rk1A = 0;
     rk1B = 0x2800000;
diff --git a/miner.h b/miner.h
index 433d6891..19f720a4 100644
--- a/miner.h
+++ b/miner.h
@@ -382,8 +382,8 @@ enum cl_kernels {
 	KL_PSW,
 	KL_ZUIKKIS,
 	KL_QUARKCOIN,
-	KL_DARKCOIN,
 	KL_INKCOIN,
+	KL_DARKCOIN,
 };
 
 enum dev_reason {
diff --git a/ocl.c b/ocl.c
index 62d72e5c..a15471bf 100644
--- a/ocl.c
+++ b/ocl.c
@@ -463,6 +463,7 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize)
 			applog(LOG_WARNING, "Kernel inkcoin is experimental.");
 			strcpy(filename, INKCOIN_KERNNAME".cl");
 			strcpy(binaryfilename, INKCOIN_KERNNAME);
+			break;
 		case KL_QUARKCOIN:
 			applog(LOG_WARNING, "Kernel quarkcoin is experimental.");
 			strcpy(filename, QUARKCOIN_KERNNAME".cl");
diff --git a/sgminer.c b/sgminer.c
index 3f539560..9f9943c1 100644
--- a/sgminer.c
+++ b/sgminer.c
@@ -4232,6 +4232,7 @@ void write_config(FILE *fcfg)
 					break;
 				case KL_INKCOIN:
 					fprintf(fcfg, INKCOIN_KERNNAME);
+					break;
 				case KL_QUARKCOIN:
 					fprintf(fcfg, QUARKCOIN_KERNNAME);
 					break;
@@ -6044,6 +6045,7 @@ static void rebuild_nonce(struct work *work, uint32_t nonce)
 			break;
 		case KL_INKCOIN:
 			inkcoin_regenhash(work);
+			break;
 		case KL_QUARKCOIN:
 			quarkcoin_regenhash(work);
 			break;
@@ -6060,7 +6062,7 @@ bool test_nonce(struct work *work, uint32_t nonce)
 	uint32_t diff1targ;
 
 	rebuild_nonce(work, nonce);
-	diff1targ = 0x0000ffffUL;
+	diff1targ = 0x000000ffUL;
 
 	return (le32toh(*hash_32) <= diff1targ);
 }