From e92d5ebf6e3b9ac6b1ef9847db0b20e19797e396 Mon Sep 17 00:00:00 2001 From: phm Date: Sat, 22 Feb 2014 12:46:46 +0100 Subject: [PATCH 1/2] Added support for InkCoin --- Makefile.am | 1 + configure.ac | 1 + driver-opencl.c | 6 ++ inkcoin.c | 147 ++++++++++++++++++++++++++++ inkcoin.h | 10 ++ kernel/inkcoin.cl | 242 ++++++++++++++++++++++++++++++++++++++++++++++ miner.h | 1 + ocl.c | 5 + sgminer.c | 6 ++ 9 files changed, 419 insertions(+) create mode 100644 inkcoin.c create mode 100644 inkcoin.h create mode 100644 kernel/inkcoin.cl diff --git a/Makefile.am b/Makefile.am index dda3b399..a6623be0 100644 --- a/Makefile.am +++ b/Makefile.am @@ -43,6 +43,7 @@ sgminer_SOURCES += findnonce.c findnonce.h sgminer_SOURCES += adl.c adl.h adl_functions.h sgminer_SOURCES += scrypt.c scrypt.h sgminer_SOURCES += darkcoin.c darkcoin.h +sgminer_SOURCES += inkcoin.c inkcoin.h sgminer_SOURCES += kernel/*.cl bin_SCRIPTS = $(top_srcdir)/kernel/*.cl diff --git a/configure.ac b/configure.ac index a7d5656e..15fd9740 100644 --- a/configure.ac +++ b/configure.ac @@ -346,6 +346,7 @@ AC_DEFINE_UNQUOTED([CKOLIVAS_KERNNAME], ["ckolivas"], [Filename for original scr AC_DEFINE_UNQUOTED([ZUIKKIS_KERNNAME], ["zuikkis"], [Filename for Zuikkis' optimised kernel]) AC_DEFINE_UNQUOTED([PSW_KERNNAME], ["psw"], [Filename for psw's experimental kernel]) AC_DEFINE_UNQUOTED([DARKCOIN_KERNNAME], ["darkcoin"], [Filename for DarkCoin optimised kernel]) +AC_DEFINE_UNQUOTED([INKCOIN_KERNNAME], ["inkcoin"], [Filename for InkCoin optimised kernel]) AC_SUBST(OPENCL_LIBS) AC_SUBST(OPENCL_FLAGS) diff --git a/driver-opencl.c b/driver-opencl.c index 6680686e..88504bf9 100644 --- a/driver-opencl.c +++ b/driver-opencl.c @@ -209,6 +209,8 @@ static enum cl_kernels select_kernel(char *arg) return KL_PSW; if (!strcmp(arg, DARKCOIN_KERNNAME)) return KL_DARKCOIN; + if (!strcmp(arg, INKCOIN_KERNNAME)) + return KL_INKCOIN; return KL_NONE; } @@ -1348,6 +1350,9 @@ static bool opencl_thread_prepare(struct thr_info *thr) case KL_DARKCOIN: cgpu->kname = DARKCOIN_KERNNAME; break; + case KL_INKCOIN: + cgpu->kname = INKCOIN_KERNNAME; + break; default: break; } @@ -1384,6 +1389,7 @@ static bool opencl_thread_init(struct thr_info *thr) thrdata->queue_kernel_parameters = &queue_scrypt_kernel; break; case KL_DARKCOIN: + case KL_INKCOIN: thrdata->queue_kernel_parameters = &queue_sph_kernel; break; default: diff --git a/inkcoin.c b/inkcoin.c new file mode 100644 index 00000000..56652543 --- /dev/null +++ b/inkcoin.c @@ -0,0 +1,147 @@ +/*- + * Copyright 2014 phm + * 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. + */ + +#include "config.h" +#include "miner.h" + +#include +#include +#include + + +#include "sph/sph_shavite.h" + +/* + * Encode a length len/4 vector of (uint32_t) into a length len vector of + * (unsigned char) in big-endian form. Assumes len is a multiple of 4. + */ +static inline void +be32enc_vect(uint32_t *dst, const uint32_t *src, uint32_t len) +{ + uint32_t i; + + for (i = 0; i < len; i++) + dst[i] = htobe32(src[i]); +} + +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); + + sph_shavite512_init(&ctx_shavite); + sph_shavite512(&ctx_shavite, hash, 64); + sph_shavite512_close(&ctx_shavite, hash); + + memcpy(state, hash, 32); +} + +static const uint32_t diff1targ = 0x0000ffff; + + +/* Used externally as confirmation of correct OCL code */ +int inkcoin_test(unsigned char *pdata, const unsigned char *ptarget, uint32_t nonce) +{ + uint32_t tmp_hash7, Htarg = le32toh(((const uint32_t *)ptarget)[7]); + uint32_t data[20], ohash[8]; + //char *scratchbuf; + + be32enc_vect(data, (const uint32_t *)pdata, 19); + data[19] = htobe32(nonce); + //scratchbuf = alloca(SCRATCHBUF_SIZE); + xhash(ohash, data); + tmp_hash7 = be32toh(ohash[7]); + + applog(LOG_DEBUG, "htarget %08lx diff1 %08lx hash %08lx", + (long unsigned int)Htarg, + (long unsigned int)diff1targ, + (long unsigned int)tmp_hash7); + if (tmp_hash7 > diff1targ) + return -1; + if (tmp_hash7 > Htarg) + return 0; + return 1; +} + +void inkcoin_regenhash(struct work *work) +{ + uint32_t data[20]; + char *scratchbuf; + uint32_t *nonce = (uint32_t *)(work->data + 76); + uint32_t *ohash = (uint32_t *)(work->hash); + + be32enc_vect(data, (const uint32_t *)work->data, 19); + data[19] = htobe32(*nonce); + xhash(ohash, data); +} + +bool scanhash_inkcoin(struct thr_info *thr, const unsigned char __maybe_unused *pmidstate, + unsigned char *pdata, unsigned char __maybe_unused *phash1, + unsigned char __maybe_unused *phash, const unsigned char *ptarget, + uint32_t max_nonce, uint32_t *last_nonce, uint32_t n) +{ + uint32_t *nonce = (uint32_t *)(pdata + 76); + char *scratchbuf; + uint32_t data[20]; + uint32_t tmp_hash7; + uint32_t Htarg = le32toh(((const uint32_t *)ptarget)[7]); + bool ret = false; + + be32enc_vect(data, (const uint32_t *)pdata, 19); + + while(1) { + uint32_t ostate[8]; + + *nonce = ++n; + data[19] = (n); + xhash(ostate, data); + tmp_hash7 = (ostate[7]); + + applog(LOG_INFO, "data7 %08lx", + (long unsigned int)data[7]); + + if (unlikely(tmp_hash7 <= Htarg)) { + ((uint32_t *)pdata)[19] = htobe32(n); + *last_nonce = n; + ret = true; + break; + } + + if (unlikely((n >= max_nonce) || thr->work_restart)) { + *last_nonce = n; + break; + } + } + + return ret; +} + + + diff --git a/inkcoin.h b/inkcoin.h new file mode 100644 index 00000000..986ed6f3 --- /dev/null +++ b/inkcoin.h @@ -0,0 +1,10 @@ +#ifndef INKCOIN_H +#define INKCOIN_H + +#include "miner.h" + +extern int inkcoin_test(unsigned char *pdata, const unsigned char *ptarget, + uint32_t nonce); +extern void inkcoin_regenhash(struct work *work); + +#endif /* INKCOIN_H */ diff --git a/kernel/inkcoin.cl b/kernel/inkcoin.cl new file mode 100644 index 00000000..a8459945 --- /dev/null +++ b/kernel/inkcoin.cl @@ -0,0 +1,242 @@ +/* + * InkCoin kernel implementation. + * + * ==========================(LICENSE BEGIN)============================ + * + * Copyright (c) 2014 phm + * + * Permission is hereby granted, free of charge, to any person obtaining + * a copy of this software and associated documentation files (the + * "Software"), to deal in the Software without restriction, including + * without limitation the rights to use, copy, modify, merge, publish, + * distribute, sublicense, and/or sell copies of the Software, and to + * permit persons to whom the Software is furnished to do so, subject to + * the following conditions: + * + * The above copyright notice and this permission notice shall be + * included in all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, + * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF + * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. + * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY + * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, + * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE + * SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. + * + * ===========================(LICENSE END)============================= + * + * @author phm + */ + +#ifndef DARKCOIN_CL +#define DARKCOIN_CL + +#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 SPH_C32(x) ((sph_u32)(x ## U)) +#define SPH_T32(x) ((x) & SPH_C32(0xFFFFFFFF)) +#define SPH_ROTL32(x, n) SPH_T32(((x) << (n)) | ((x) >> (32 - (n)))) +#define SPH_ROTR32(x, n) SPH_ROTL32(x, (32 - (n))) + +#define SPH_C64(x) ((sph_u64)(x ## UL)) +#define SPH_T64(x) ((x) & SPH_C64(0xFFFFFFFFFFFFFFFF)) +#define SPH_ROTL64(x, n) SPH_T64(((x) << (n)) | ((x) >> (64 - (n)))) +#define SPH_ROTR64(x, n) SPH_ROTL64(x, (64 - (n))) + +#define SPH_ECHO_64 1 +#define SPH_KECCAK_64 1 +#define SPH_JH_64 1 +#define SPH_SIMD_NOCOPY 0 +#define SPH_KECCAK_NOCOPY 0 +#define SPH_COMPACT_BLAKE_64 0 +#define SPH_LUFFA_PARALLEL 0 +#define SPH_SMALL_FOOTPRINT_GROESTL 0 +#define SPH_GROESTL_BIG_ENDIAN 0 + +#define SPH_CUBEHASH_UNROLL 0 +#define SPH_KECCAK_UNROLL 0 + +#include "shavite.cl" + +#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 DEC32LE(x) (*(const __global sph_u32 *) (x)); +#endif + +// __attribute__((reqd_work_group_size(WORKSIZE, 1, 1))) +__kernel void search(__global unsigned char* block, volatile __global uint* output, const ulong target) +{ + uint gid = get_global_id(0); + union { + unsigned char h1[64]; + uint h4[16]; + ulong h8[8]; + } hash; + + __local sph_u32 AES0[256], AES1[256], AES2[256], AES3[256]; + int init = get_local_id(0); + int step = get_local_size(0); + for (int i = init; i < 256; i += step) + { + AES0[i] = AES0_C[i]; + AES1[i] = AES1_C[i]; + AES2[i] = AES2_C[i]; + AES3[i] = AES3_C[i]; + } + barrier(CLK_LOCAL_MEM_FENCE); + + // shavite + { + // IV + sph_u32 h0 = SPH_C32(0x72FCCDD8), h1 = SPH_C32(0x79CA4727), h2 = SPH_C32(0x128A077B), h3 = SPH_C32(0x40D55AEC); + sph_u32 h4 = SPH_C32(0xD1901A06), h5 = SPH_C32(0x430AE307), h6 = SPH_C32(0xB29F5CD1), h7 = SPH_C32(0xDF07FBFC); + sph_u32 h8 = SPH_C32(0x8E45D73D), h9 = SPH_C32(0x681AB538), hA = SPH_C32(0xBDE86578), hB = SPH_C32(0xDD577E47); + sph_u32 hC = SPH_C32(0xE275EADE), hD = SPH_C32(0x502D9FCD), hE = SPH_C32(0xB9357178), hF = SPH_C32(0x022A4B9A); + + // state + sph_u32 rk00, rk01, rk02, rk03, rk04, rk05, rk06, rk07; + sph_u32 rk08, rk09, rk0A, rk0B, rk0C, rk0D, rk0E, rk0F; + sph_u32 rk10, rk11, rk12, rk13, rk14, rk15, rk16, rk17; + sph_u32 rk18, rk19, rk1A, rk1B, rk1C, rk1D, rk1E, rk1F; + + sph_u32 sc_count0 = (80 << 3), sc_count1 = 0, sc_count2 = 0, sc_count3 = 0; + + rk00 = DEC32LE(block + 0 * 4);; + rk01 = DEC32LE(block + 1 * 4);; + rk02 = DEC32LE(block + 2 * 4);; + rk03 = DEC32LE(block + 3 * 4);; + rk04 = DEC32LE(block + 4 * 4);; + rk05 = DEC32LE(block + 5 * 4);; + rk06 = DEC32LE(block + 6 * 4);; + rk07 = DEC32LE(block + 7 * 4);; + rk08 = DEC32LE(block + 8 * 4);; + rk09 = DEC32LE(block + 9 * 4);; + rk0A = DEC32LE(block + 10 * 4);; + rk0B = DEC32LE(block + 11 * 4);; + rk0C = DEC32LE(block + 12 * 4);; + rk0D = DEC32LE(block + 13 * 4);; + rk0E = DEC32LE(block + 14 * 4);; + rk0F = DEC32LE(block + 15 * 4);; + rk10 = DEC32LE(block + 16 * 4);; + rk11 = DEC32LE(block + 17 * 4);; + rk12 = DEC32LE(block + 18 * 4);; + rk13 = 72768; // gid; + rk14 = 0x80; + rk15 = rk16 = rk17 = rk18 = rk19 = rk1A = 0; + rk1B = 0x2800000; + rk1C = rk1D = rk1E = 0; + rk1F = 0x2000000; + + c512(buf); + + hash.h4[0] = h0; + hash.h4[1] = h1; + hash.h4[2] = h2; + hash.h4[3] = h3; + hash.h4[4] = h4; + hash.h4[5] = h5; + hash.h4[6] = h6; + hash.h4[7] = h7; + hash.h4[8] = h8; + hash.h4[9] = h9; + hash.h4[10] = hA; + hash.h4[11] = hB; + hash.h4[12] = hC; + hash.h4[13] = hD; + hash.h4[14] = hE; + hash.h4[15] = hF; + } + + + // shavite + { + // IV + sph_u32 h0 = SPH_C32(0x72FCCDD8), h1 = SPH_C32(0x79CA4727), h2 = SPH_C32(0x128A077B), h3 = SPH_C32(0x40D55AEC); + sph_u32 h4 = SPH_C32(0xD1901A06), h5 = SPH_C32(0x430AE307), h6 = SPH_C32(0xB29F5CD1), h7 = SPH_C32(0xDF07FBFC); + sph_u32 h8 = SPH_C32(0x8E45D73D), h9 = SPH_C32(0x681AB538), hA = SPH_C32(0xBDE86578), hB = SPH_C32(0xDD577E47); + sph_u32 hC = SPH_C32(0xE275EADE), hD = SPH_C32(0x502D9FCD), hE = SPH_C32(0xB9357178), hF = SPH_C32(0x022A4B9A); + + // state + sph_u32 rk00, rk01, rk02, rk03, rk04, rk05, rk06, rk07; + sph_u32 rk08, rk09, rk0A, rk0B, rk0C, rk0D, rk0E, rk0F; + sph_u32 rk10, rk11, rk12, rk13, rk14, rk15, rk16, rk17; + sph_u32 rk18, rk19, rk1A, rk1B, rk1C, rk1D, rk1E, rk1F; + + sph_u32 sc_count0 = (64 << 3), sc_count1 = 0, sc_count2 = 0, sc_count3 = 0; + + rk00 = hash.h4[0]; + rk01 = hash.h4[1]; + rk02 = hash.h4[2]; + rk03 = hash.h4[3]; + rk04 = hash.h4[4]; + rk05 = hash.h4[5]; + rk06 = hash.h4[6]; + rk07 = hash.h4[7]; + rk08 = hash.h4[8]; + rk09 = hash.h4[9]; + rk0A = hash.h4[10]; + rk0B = hash.h4[11]; + rk0C = hash.h4[12]; + rk0D = hash.h4[13]; + rk0E = hash.h4[14]; + rk0F = hash.h4[15]; + rk10 = 0x80; + rk11 = rk12 = rk13 = rk14 = rk15 = rk16 = rk17 = rk18 = rk19 = rk1A = 0; + rk1B = 0x2000000; + rk1C = rk1D = rk1E = 0; + rk1F = 0x2000000; + + c512(buf); + + hash.h4[0] = h0; + hash.h4[1] = h1; + hash.h4[2] = h2; + hash.h4[3] = h3; + hash.h4[4] = h4; + hash.h4[5] = h5; + hash.h4[6] = h6; + hash.h4[7] = h7; + hash.h4[8] = h8; + hash.h4[9] = h9; + hash.h4[10] = hA; + hash.h4[11] = hB; + hash.h4[12] = hC; + hash.h4[13] = hD; + hash.h4[14] = hE; + hash.h4[15] = hF; + } + + bool result = (hash.h8[3] <= target); + if (result) + output[output[0xFF]++] = SWAP4(gid); +} + +#endif // DARKCOIN_CL diff --git a/miner.h b/miner.h index f9360d16..a1e7ec9c 100644 --- a/miner.h +++ b/miner.h @@ -382,6 +382,7 @@ enum cl_kernels { KL_PSW, KL_ZUIKKIS, KL_DARKCOIN, + KL_INKCOIN, }; enum dev_reason { diff --git a/ocl.c b/ocl.c index c05dd68e..e42cbd7a 100644 --- a/ocl.c +++ b/ocl.c @@ -459,6 +459,11 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize) strcpy(filename, DARKCOIN_KERNNAME".cl"); strcpy(binaryfilename, DARKCOIN_KERNNAME); break; + case KL_INKCOIN: + applog(LOG_WARNING, "Kernel inkcoin is experimental."); + strcpy(filename, INKCOIN_KERNNAME".cl"); + strcpy(binaryfilename, INKCOIN_KERNNAME); + break; case KL_NONE: /* Shouldn't happen */ break; } diff --git a/sgminer.c b/sgminer.c index dadb7e31..18875ea0 100644 --- a/sgminer.c +++ b/sgminer.c @@ -4227,6 +4227,9 @@ void write_config(FILE *fcfg) case KL_DARKCOIN: fprintf(fcfg, DARKCOIN_KERNNAME); break; + case KL_INKCOIN: + fprintf(fcfg, INKCOIN_KERNNAME); + break; } } @@ -6034,6 +6037,9 @@ static void rebuild_nonce(struct work *work, uint32_t nonce) case KL_DARKCOIN: darkcoin_regenhash(work); break; + case KL_INKCOIN: + inkcoin_regenhash(work); + break; default: scrypt_regenhash(work); break; From c336a8399ae1d6340e8a7fae687daf800551fb40 Mon Sep 17 00:00:00 2001 From: phm Date: Sat, 22 Feb 2014 16:58:00 +0100 Subject: [PATCH 2/2] 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); }