From 7f85a0b1bb0ead0acb4a2e8d217653a1be01e8e0 Mon Sep 17 00:00:00 2001 From: Michael Allar Date: Tue, 2 Jan 2018 22:43:44 -0800 Subject: [PATCH] Added Dallar kernel --- algorithm.c | 2 + algorithm/dallar.c | 155 ++++++++++++++ algorithm/dallar.h | 10 + kernel/dallar.cl | 336 +++++++++++++++++++++++++++++++ sgminer.c | 22 +- util.c | 2 +- util.h | 1 + winbuild/packages.config | 5 + winbuild/sgminer.sln | 5 +- winbuild/sgminer.vcxproj | 29 ++- winbuild/sgminer.vcxproj.filters | 7 + 11 files changed, 555 insertions(+), 19 deletions(-) create mode 100644 algorithm/dallar.c create mode 100644 algorithm/dallar.h create mode 100644 kernel/dallar.cl create mode 100644 winbuild/packages.config diff --git a/algorithm.c b/algorithm.c index fc0ee4a97..5bd9d2b5f 100644 --- a/algorithm.c +++ b/algorithm.c @@ -22,6 +22,7 @@ #include "algorithm/myriadcoin-groestl.h" #include "algorithm/fuguecoin.h" #include "algorithm/groestlcoin.h" +#include "algorithm/dallar.h" #include "algorithm/twecoin.h" #include "algorithm/marucoin.h" #include "algorithm/maxcoin.h" @@ -800,6 +801,7 @@ static algorithm_settings_t algos[] = { { a, ALGO_FUGUE, "", 1, 256, 256, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 0, 0, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, b, queue_sph_kernel, c, NULL } A_FUGUE("fuguecoin", fuguecoin_regenhash, sha256), A_FUGUE("groestlcoin", groestlcoin_regenhash, sha256), + A_FUGUE("dallar", dallar_regenhash, sha256), A_FUGUE("diamond", groestlcoin_regenhash, gen_hash), #undef A_FUGUE diff --git a/algorithm/dallar.c b/algorithm/dallar.c new file mode 100644 index 000000000..fb2734552 --- /dev/null +++ b/algorithm/dallar.c @@ -0,0 +1,155 @@ +/*- + * Copyright 2009 Colin Percival, 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. + * + * This file was originally written by Colin Percival as part of the Tarsnap + * online backup system. + */ + +#include "config.h" +#include "miner.h" + +#include +#include +#include + +#include "sph/sph_groestl.h" +#include "sph/sph_sha2.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]); +} + + +#ifdef __APPLE_CC__ +static +#endif +inline void throestlhash(void *state, const void *input) +{ + sph_groestl512_context ctx_groestl; + + uint32_t hash[16]; + + sph_groestl512_init(&ctx_groestl); + sph_groestl512(&ctx_groestl, input, 80); + sph_groestl512_close(&ctx_groestl, (void*) hash); + + sph_groestl512_init(&ctx_groestl); + sph_groestl512(&ctx_groestl, hash, 64); + sph_groestl512_close(&ctx_groestl, (void*) hash); + + sph_groestl512_init(&ctx_groestl); + sph_groestl512(&ctx_groestl, hash, 64); + sph_groestl512_close(&ctx_groestl, (void*)hash); + + memcpy(state, hash, 32); +} + +static const uint32_t diff1targ = 0x0000ffff; + + +/* Used externally as confirmation of correct OCL code */ +int dallar_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]; + + be32enc_vect(data, (const uint32_t *)pdata, 19); + data[19] = htobe32(nonce); + throestlhash(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 dallar_regenhash(struct work *work) +{ + uint32_t data[20]; + 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); + throestlhash(ohash, data); +} + +bool scanhash_dallar(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); + 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); + throestlhash(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/algorithm/dallar.h b/algorithm/dallar.h new file mode 100644 index 000000000..43741d1cc --- /dev/null +++ b/algorithm/dallar.h @@ -0,0 +1,10 @@ +#ifndef DALLAR_H +#define DALLAR_H + +#include "miner.h" + +extern int dallar_test(unsigned char *pdata, const unsigned char *ptarget, + uint32_t nonce); +extern void dallar_regenhash(struct work *work); + +#endif /* DALLAR_H */ diff --git a/kernel/dallar.cl b/kernel/dallar.cl new file mode 100644 index 000000000..2a0d6989a --- /dev/null +++ b/kernel/dallar.cl @@ -0,0 +1,336 @@ +/* + * GroestlCoin 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 GROESTLCOIN_CL +#define GROESTLCOIN_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) (as_uint(x)) +#define SPH_ROTL32(x, n) rotate(as_uint(x), as_uint(n)) +#define SPH_ROTR32(x, n) SPH_ROTL32(x, (32 - (n))) + +#define SPH_C64(x) ((sph_u64)(x ## UL)) +#define SPH_T64(x) (as_ulong(x)) +#define SPH_ROTL64(x, n) rotate(as_ulong(x), (n) & 0xFFFFFFFFFFFFFFFFUL) +#define SPH_ROTR64(x, n) SPH_ROTL64(x, (64 - (n))) + +#define SPH_ECHO_64 1 +#define SPH_SIMD_NOCOPY 0 +#define SPH_CUBEHASH_UNROLL 0 + +#ifndef SPH_LUFFA_PARALLEL + #define SPH_LUFFA_PARALLEL 0 +#endif + +#include "groestl.cl" + +#define SWAP4(x) as_uint(as_uchar4(x).wzyx) +#define SWAP8(x) as_ulong(as_uchar8(x).s76543210) + +#if SPH_BIG_ENDIAN + #define ENC64E(x) SWAP8(x) + #define DEC64E(x) SWAP8(*(const __global sph_u64 *) (x)); +#else + #define ENC64E(x) (x) + #define DEC64E(x) (*(const __global sph_u64 *) (x)); +#endif + +#define ROL32(x, n) rotate(x, (uint) n) +#define SHR(x, n) ((x) >> n) +#define SWAP32(a) (as_uint(as_uchar4(a).wzyx)) + +#define S0(x) (ROL32(x, 25) ^ ROL32(x, 14) ^ SHR(x, 3)) +#define S1(x) (ROL32(x, 15) ^ ROL32(x, 13) ^ SHR(x, 10)) + +#define S2(x) (ROL32(x, 30) ^ ROL32(x, 19) ^ ROL32(x, 10)) +#define S3(x) (ROL32(x, 26) ^ ROL32(x, 21) ^ ROL32(x, 7)) + +#define P(a,b,c,d,e,f,g,h,x,K) \ +{ \ + temp1 = h + S3(e) + F1(e,f,g) + (K + x); \ + d += temp1; h = temp1 + S2(a) + F0(a,b,c); \ +} + +#define PLAST(a,b,c,d,e,f,g,h,x,K) \ +{ \ + d += h + S3(e) + F1(e,f,g) + (x + K); \ +} + +#define F0(y, x, z) bitselect(z, y, z ^ x) +#define F1(x, y, z) bitselect(z, y, x) + +#define R0 (W0 = S1(W14) + W9 + S0(W1) + W0) +#define R1 (W1 = S1(W15) + W10 + S0(W2) + W1) +#define R2 (W2 = S1(W0) + W11 + S0(W3) + W2) +#define R3 (W3 = S1(W1) + W12 + S0(W4) + W3) +#define R4 (W4 = S1(W2) + W13 + S0(W5) + W4) +#define R5 (W5 = S1(W3) + W14 + S0(W6) + W5) +#define R6 (W6 = S1(W4) + W15 + S0(W7) + W6) +#define R7 (W7 = S1(W5) + W0 + S0(W8) + W7) +#define R8 (W8 = S1(W6) + W1 + S0(W9) + W8) +#define R9 (W9 = S1(W7) + W2 + S0(W10) + W9) +#define R10 (W10 = S1(W8) + W3 + S0(W11) + W10) +#define R11 (W11 = S1(W9) + W4 + S0(W12) + W11) +#define R12 (W12 = S1(W10) + W5 + S0(W13) + W12) +#define R13 (W13 = S1(W11) + W6 + S0(W14) + W13) +#define R14 (W14 = S1(W12) + W7 + S0(W15) + W14) +#define R15 (W15 = S1(W13) + W8 + S0(W0) + W15) + +#define RD14 (S1(W12) + W7 + S0(W15) + W14) +#define RD15 (S1(W13) + W8 + S0(W0) + W15) + +__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; + +#if !SPH_SMALL_FOOTPRINT_GROESTL + __local sph_u64 T0_C[256], T1_C[256], T2_C[256], T3_C[256]; + __local sph_u64 T4_C[256], T5_C[256], T6_C[256], T7_C[256]; +#else + __local sph_u64 T0_C[256], T4_C[256]; +#endif + int init = get_local_id(0); + int step = get_local_size(0); + + for (int i = init; i < 256; i += step) + { + T0_C[i] = T0[i]; + T4_C[i] = T4[i]; +#if !SPH_SMALL_FOOTPRINT_GROESTL + T1_C[i] = T1[i]; + T2_C[i] = T2[i]; + T3_C[i] = T3[i]; + T5_C[i] = T5[i]; + T6_C[i] = T6[i]; + T7_C[i] = T7[i]; +#endif + } + barrier(CLK_LOCAL_MEM_FENCE); // groestl +#define T0 T0_C +#define T1 T1_C +#define T2 T2_C +#define T3 T3_C +#define T4 T4_C +#define T5 T5_C +#define T6 T6_C +#define T7 T7_C + + + // groestl + { + sph_u64 H[16]; +//#pragma unroll 15 + for (unsigned int u = 0; u < 15; u ++) + H[u] = 0; +#if USE_LE + H[15] = ((sph_u64)(512 & 0xFF) << 56) | ((sph_u64)(512 & 0xFF00) << 40); +#else + H[15] = (sph_u64)512; +#endif + + sph_u64 g[16], m[16]; + m[0] = DEC64E(block + 0 * 8); + m[1] = DEC64E(block + 1 * 8); + m[2] = DEC64E(block + 2 * 8); + m[3] = DEC64E(block + 3 * 8); + m[4] = DEC64E(block + 4 * 8); + m[5] = DEC64E(block + 5 * 8); + m[6] = DEC64E(block + 6 * 8); + m[7] = DEC64E(block + 7 * 8); + m[8] = DEC64E(block + 8 * 8); + m[9] = DEC64E(block + 9 * 8); + m[9] &= 0x00000000FFFFFFFF; + m[9] |= ((sph_u64) gid << 32); + +//#pragma unroll 16 + for (unsigned int u = 0; u < 16; u ++) + g[u] = m[u] ^ H[u]; + m[10] = 0x80; g[10] = m[10] ^ H[10]; + m[11] = 0; g[11] = m[11] ^ H[11]; + m[12] = 0; g[12] = m[12] ^ H[12]; + m[13] = 0; g[13] = m[13] ^ H[13]; + m[14] = 0; g[14] = m[14] ^ H[14]; + m[15] = 0x100000000000000; g[15] = m[15] ^ H[15]; + PERM_BIG_P(g); + PERM_BIG_Q(m); + +//#pragma unroll 16 + for (unsigned int u = 0; u < 16; u ++) + H[u] ^= g[u] ^ m[u]; + sph_u64 xH[16]; + +//#pragma unroll 16 + for (unsigned int u = 0; u < 16; u ++) + xH[u] = H[u]; + PERM_BIG_P(xH); + +//#pragma unroll 16 + for (unsigned int u = 0; u < 16; u ++) + H[u] ^= xH[u]; + +//#pragma unroll 8 + for (unsigned int u = 0; u < 8; u ++) + hash.h8[u] = ENC64E(H[u + 8]); + +//#pragma unroll 15 + for (unsigned int u = 0; u < 15; u ++) + H[u] = 0; + #if USE_LE + H[15] = ((sph_u64)(512 & 0xFF) << 56) | ((sph_u64)(512 & 0xFF00) << 40); + #else + H[15] = (sph_u64)512; + #endif + + m[0] = hash.h8[0]; + m[1] = hash.h8[1]; + m[2] = hash.h8[2]; + m[3] = hash.h8[3]; + m[4] = hash.h8[4]; + m[5] = hash.h8[5]; + m[6] = hash.h8[6]; + m[7] = hash.h8[7]; + +//#pragma unroll 16 + for (unsigned int u = 0; u < 16; u ++) + g[u] = m[u] ^ H[u]; + m[8] = 0x80; g[8] = m[8] ^ H[8]; + m[9] = 0; g[9] = m[9] ^ H[9]; + m[10] = 0; g[10] = m[10] ^ H[10]; + m[11] = 0; g[11] = m[11] ^ H[11]; + m[12] = 0; g[12] = m[12] ^ H[12]; + m[13] = 0; g[13] = m[13] ^ H[13]; + m[14] = 0; g[14] = m[14] ^ H[14]; + m[15] = 0x100000000000000; g[15] = m[15] ^ H[15]; + PERM_BIG_P(g); + PERM_BIG_Q(m); + +//#pragma unroll 16 + for (unsigned int u = 0; u < 16; u ++) + H[u] ^= g[u] ^ m[u]; + + //#pragma unroll 16 + for (unsigned int u = 0; u < 16; u ++) + xH[u] = H[u]; + PERM_BIG_P(xH); + + //#pragma unroll 16 + for (unsigned int u = 0; u < 16; u ++) + H[u] ^= xH[u]; + +//#pragma unroll 8 + for (unsigned int u = 0; u < 8; u ++) + hash.h8[u] = H[u + 8]; + +//#pragma unroll 15 + for (unsigned int u = 0; u < 15; u ++) + H[u] = 0; + #if USE_LE + H[15] = ((sph_u64)(512 & 0xFF) << 56) | ((sph_u64)(512 & 0xFF00) << 40); + #else + H[15] = (sph_u64)512; + #endif + + m[0] = hash.h8[0]; + m[1] = hash.h8[1]; + m[2] = hash.h8[2]; + m[3] = hash.h8[3]; + m[4] = hash.h8[4]; + m[5] = hash.h8[5]; + m[6] = hash.h8[6]; + m[7] = hash.h8[7]; + +//#pragma unroll 16 + for (unsigned int u = 0; u < 16; u ++) + g[u] = m[u] ^ H[u]; + m[8] = 0x80; g[8] = m[8] ^ H[8]; + m[9] = 0; g[9] = m[9] ^ H[9]; + m[10] = 0; g[10] = m[10] ^ H[10]; + m[11] = 0; g[11] = m[11] ^ H[11]; + m[12] = 0; g[12] = m[12] ^ H[12]; + m[13] = 0; g[13] = m[13] ^ H[13]; + m[14] = 0; g[14] = m[14] ^ H[14]; + m[15] = 0x100000000000000; g[15] = m[15] ^ H[15]; + PERM_BIG_P(g); + PERM_BIG_Q(m); + +//#pragma unroll 16 + for (unsigned int u = 0; u < 16; u ++) + H[u] ^= g[u] ^ m[u]; + + //#pragma unroll 16 + for (unsigned int u = 0; u < 16; u ++) + xH[u] = H[u]; + PERM_BIG_P(xH); + + //#pragma unroll 16 + for (unsigned int u = 0; u < 16; u ++) + H[u] ^= xH[u]; + +//#pragma unroll 8 + for (unsigned int u = 0; u < 8; u ++) + hash.h8[u] = H[u + 8]; + + barrier(CLK_GLOBAL_MEM_FENCE); + } + + bool result = (hash.h8[3] <= target); + if (result) + output[output[0xFF]++] = SWAP4(gid); +} + +#endif // GROESTLCOIN_CL diff --git a/sgminer.c b/sgminer.c index fe4039246..ce216aa9f 100644 --- a/sgminer.c +++ b/sgminer.c @@ -1571,7 +1571,7 @@ struct opt_table opt_config_table[] = { "Disable 'extranonce' stratum subscribe for pool"), OPT_WITH_ARG("--pass|--pool-pass|-p", set_pass, NULL, NULL, - "Password for groestlcoin JSON-RPC server"), + "Password for dallar JSON-RPC server"), OPT_WITHOUT_ARG("--per-device-stats", opt_set_bool, &want_per_device_stats, "Force verbose mode and output per-device statistics"), @@ -1802,10 +1802,10 @@ struct opt_table opt_config_table[] = { "Set GPU thread concurrency for scrypt mining, comma separated"), OPT_WITH_ARG("--url|--pool-url|-o", set_url, NULL, NULL, - "URL for groestlcoin JSON-RPC server"), + "URL for dallar JSON-RPC server"), OPT_WITH_ARG("--user|--pool-user|-u", set_user, NULL, NULL, - "Username for groestlcoin JSON-RPC server"), + "Username for dallar JSON-RPC server"), OPT_WITH_ARG("--vectors", set_vector, NULL, NULL, opt_hidden), @@ -1822,7 +1822,7 @@ struct opt_table opt_config_table[] = { "Override detected optimal worksize - one value or comma separated list"), OPT_WITH_ARG("--userpass|--pool-userpass|-O", set_userpass, NULL, NULL, - "Username:Password pair for groestlcoin JSON-RPC server"), + "Username:Password pair for dallar JSON-RPC server"), OPT_WITHOUT_ARG("--worktime", opt_set_bool, &opt_worktime, "Display extra work time debug information"), @@ -1840,7 +1840,7 @@ struct opt_table opt_config_table[] = { #ifdef HAVE_LIBCURL OPT_WITH_ARG("--grs-address", opt_set_charp, NULL, &opt_grs_address, - "Set groestlcoin target address when solo mining to groestlcoind (mandatory)"), + "Set dallar target address when solo mining to dallard (mandatory)"), OPT_WITH_ARG("--grs-sig", opt_set_charp, NULL, &opt_grs_sig, "Set signature to add to coinbase when solo mining (optional)"), @@ -2330,7 +2330,7 @@ static void gbt_merkle_bins(struct pool *pool, json_t *transaction_arr) pool->merkles = 0; pool->transactions = json_array_size(transaction_arr); binlen = pool->transactions * 32 + 32; - hashbin = alloca(binlen + 32); + hashbin = static_cast(alloca(binlen + 32)); memset(hashbin, 0, 32); binleft = binlen / 32; if (pool->transactions) { @@ -2348,7 +2348,7 @@ static void gbt_merkle_bins(struct pool *pool, json_t *transaction_arr) len += strlen(txn); } - pool->txn_data = cgmalloc(len + 1); + pool->txn_data = static_cast(cgmalloc(len + 1)); pool->txn_data[len] = '\0'; for (i = 0; i < pool->transactions; i++) { @@ -2366,7 +2366,7 @@ static void gbt_merkle_bins(struct pool *pool, json_t *transaction_arr) int txn_len; txn_len = len / 2; - txn_bin = cgmalloc(txn_len); + txn_bin = static_cast(cgmalloc(txn_len)); hex2bin(txn_bin, txn, txn_len); /* This is needed for pooled mining since only * transaction data and not hashes are sent */ @@ -2569,7 +2569,7 @@ static bool gbt_solo_decode(struct pool *pool, json_t *res_val) + 1 + 25 // txout + 4; // lock free(pool->coinbase); - pool->coinbase = cgcalloc(len, 1); + pool->coinbase = static_cast(cgcalloc(len, 1)); cg_memcpy(pool->coinbase + 41, pool->scriptsig_base, ofs); cg_memcpy(pool->coinbase + 41 + ofs, "\xff\xff\xff\xff", 4); pool->coinbase[41 + ofs + 4] = 1; @@ -6256,10 +6256,10 @@ static bool setup_gbt_solo(CURL *curl, struct pool *pool) if (!valid_val) goto out; if (!json_is_true(valid_val)) { - applog(LOG_ERR, "Groestlcoin address %s is NOT valid", opt_grs_address); + applog(LOG_ERR, "Dallar address %s is NOT valid", opt_grs_address); goto out; } - applog(LOG_NOTICE, "Solo mining to valid GRS address: %s", opt_grs_address); + applog(LOG_NOTICE, "Solo mining to valid DAL address: %s", opt_grs_address); ret = true; address_to_pubkeyhash(pool->script_pubkey, opt_grs_address); hex2bin(scriptsig_header_bin, scriptsig_header, 41); diff --git a/util.c b/util.c index 15323f223..3fb76b9ce 100644 --- a/util.c +++ b/util.c @@ -838,7 +838,7 @@ unsigned char *ser_string(char *s, int *slen) size_t len = strlen(s); unsigned char *ret; - ret = cgmalloc(1 + len + 8); // Leave room for largest size + ret = static_cast(cgmalloc(1 + len + 8)); // Leave room for largest size if (len < 253) { ret[0] = len; cg_memcpy(ret + 1, s, len); diff --git a/util.h b/util.h index a178307b9..9fbce80e0 100644 --- a/util.h +++ b/util.h @@ -163,6 +163,7 @@ void cgsem_reset(cgsem_t *cgsem); void cgsem_destroy(cgsem_t *cgsem); bool cg_completion_timeout(void *fn, void *fnarg, int timeout); void _cg_memcpy(void *dest, const void *src, unsigned int n, const char *file, const char *func, const int line); +void align_len(size_t *len); #define cgsem_init(_sem) _cgsem_init(_sem, __FILE__, __func__, __LINE__) #define cgsem_post(_sem) _cgsem_post(_sem, __FILE__, __func__, __LINE__) diff --git a/winbuild/packages.config b/winbuild/packages.config new file mode 100644 index 000000000..191bfe5c2 --- /dev/null +++ b/winbuild/packages.config @@ -0,0 +1,5 @@ + + + + + \ No newline at end of file diff --git a/winbuild/sgminer.sln b/winbuild/sgminer.sln index 32fde6c73..08384be42 100644 --- a/winbuild/sgminer.sln +++ b/winbuild/sgminer.sln @@ -1,9 +1,12 @@  Microsoft Visual Studio Solution File, Format Version 12.00 # Visual Studio 2013 -VisualStudioVersion = 12.0.31101.0 +VisualStudioVersion = 12.0.40629.0 MinimumVisualStudioVersion = 10.0.40219.1 Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "sgminer", "sgminer.vcxproj", "{CCA64DCD-6401-42A3-ABC3-89E48A36D239}" + ProjectSection(ProjectDependencies) = postProject + {AFE7D2AA-025C-4837-B4B2-81117E010B3B} = {AFE7D2AA-025C-4837-B4B2-81117E010B3B} + EndProjectSection EndProject Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "jansson", "jansson\jansson.vcxproj", "{AFE7D2AA-025C-4837-B4B2-81117E010B3B}" EndProject diff --git a/winbuild/sgminer.vcxproj b/winbuild/sgminer.vcxproj index 6ec8a4cdc..0ae7553bf 100644 --- a/winbuild/sgminer.vcxproj +++ b/winbuild/sgminer.vcxproj @@ -65,11 +65,15 @@ - + + d2c84da7 + false $(SolutionDir)output\x86\$(Configuration)\ $(SolutionDir)output\x86\obj\ + C:\Program Files %28x86%29\AMD APP SDK\3.0\include;$(IncludePath) + C:\Program Files %28x86%29\AMD APP SDK\3.0\lib\x86;$(LibraryPath) false @@ -80,6 +84,8 @@ false $(SolutionDir)output\x86\$(Configuration)\ $(SolutionDir)output\x86\obj\ + C:\Program Files %28x86%29\AMD APP SDK\3.0\include;$(IncludePath) + C:\Program Files %28x86%29\AMD APP SDK\3.0\lib\x86;$(LibraryPath) false @@ -106,7 +112,7 @@ Console true $(ProjectDir)dist\lib\x86;$(ProjectDir)jansson\Debug\x86 - %(AdditionalDependencies) + ws2_32.lib;%(AdditionalDependencies) false true @@ -114,9 +120,8 @@ LIBCMT - - xcopy /Y /E /I "$(ProjectDir)..\kernel" "$(OutDir)\kernel" - + xcopy /Y /E /I "$(ProjectDir)dist\bin\x86\*.dll" "$(OutDir)" +xcopy /Y /E /I "$(ProjectDir)..\kernel" "$(OutDir)\kernel" @@ -193,7 +198,7 @@ true false $(ProjectDir)dist\lib\x86;$(ProjectDir)jansson\Release\x86 - %(AdditionalDependencies) + ws2_32.lib;%(AdditionalDependencies) LIBCMT @@ -263,6 +268,7 @@ + @@ -328,6 +334,7 @@ + @@ -391,9 +398,19 @@ + + + + + + This project references NuGet package(s) that are missing on this computer. Enable NuGet Package Restore to download them. For more information, see http://go.microsoft.com/fwlink/?LinkID=322105. The missing file is {0}. + + + + \ No newline at end of file diff --git a/winbuild/sgminer.vcxproj.filters b/winbuild/sgminer.vcxproj.filters index 02c26210e..d9af9f566 100644 --- a/winbuild/sgminer.vcxproj.filters +++ b/winbuild/sgminer.vcxproj.filters @@ -218,6 +218,9 @@ Source Files\algorithm + + Source Files\algorithm + @@ -415,8 +418,12 @@ Header Files\algorithm + + Header Files\algorithm + + \ No newline at end of file