Added support for INKcoin.

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;

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,

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;

miner.h

@@ -382,8 +382,8 @@ enum cl_kernels {
 	KL_PSW,
 	KL_ZUIKKIS,
 	KL_QUARKCOIN,
-	KL_DARKCOIN,
 	KL_INKCOIN,
+	KL_DARKCOIN,
 };
 
 enum dev_reason {

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");

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