GOSTSec/ccminer
1
0
Fork 0
mirror of https://github.com/GOSTSec/ccminer synced 2025-01-26 14:34:22 +00:00
Code Issues Releases Wiki Activity

sia: cuda impl. and suprnova getwork over stratum

Browse Source 
Signed-off-by: Tanguy Pruvot <tanguy.pruvot@gmail.com>
This commit is contained in:
Tanguy Pruvot 2016-07-29 16:09:33 +02:00
parent 3d701944b2
commit b31fb5316b
10 changed files with 642 additions and 16 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
Makefile.am
View File

@ -46,6 +46,7 @@ ccminer_SOURCES = elist.h miner.h compat.h \
quark/quarkcoin.cu quark/cuda_quark_compactionTest.cu \
neoscrypt/neoscrypt.cpp neoscrypt/neoscrypt-cpu.c neoscrypt/cuda_neoscrypt.cu \
cuda_nist5.cu pentablake.cu skein.cu cuda_skeincoin.cu skein2.cpp zr5.cu \
sia.cu sph/blake2b.c \
sph/bmw.c sph/blake.c sph/groestl.c sph/jh.c sph/keccak.c sph/skein.c \
sph/cubehash.c sph/echo.c sph/luffa.c sph/sha2.c sph/shavite.c sph/simd.c \
sph/hamsi.c sph/hamsi_helper.c sph/streebog.c \

2
algos.h
View File

@ -32,6 +32,7 @@ enum sha_algos {
ALGO_QUBIT,
ALGO_SCRYPT,
ALGO_SCRYPT_JANE,
ALGO_SIA,
ALGO_SIB,
ALGO_SKEIN,
ALGO_SKEIN2,
@ -81,6 +82,7 @@ static const char *algo_names[] = {
"qubit",
"scrypt",
"scrypt-jane",
"sia",
"sib",
"skein",
"skein2",

2
bench.cpp
View File

@ -56,6 +56,7 @@ void algo_free_all(int thr_id)
free_groestlcoin(thr_id);
free_heavy(thr_id);
free_jackpot(thr_id);
free_lbry(thr_id);
free_luffa(thr_id);
free_lyra2(thr_id);
free_lyra2v2(thr_id);
@ -67,6 +68,7 @@ void algo_free_all(int thr_id)
free_qubit(thr_id);
free_skeincoin(thr_id);
free_skein2(thr_id);
free_sia(thr_id);
free_sib(thr_id);
free_s3(thr_id);
free_vanilla(thr_id);

88
ccminer.cpp
View File

@ -570,6 +570,8 @@ static void calc_network_diff(struct work *work)
uint32_t nbits = have_longpoll ? work->data[18] : swab32(work->data[18]);
if (opt_algo == ALGO_LBRY) nbits = swab32(work->data[26]);
if (opt_algo == ALGO_DECRED) nbits = work->data[29];
if (opt_algo == ALGO_SIA) nbits = work->data[11]; // unsure if correct
uint32_t bits = (nbits & 0xffffff);
int16_t shift = (swab32(nbits) & 0xff); // 0x1c = 28
@ -647,7 +649,7 @@ static bool work_decode(const json_t *val, struct work *work)
stratum_diff = work->targetdiff;
work->tx_count = use_pok = 0;
if (work->data[0] & POK_BOOL_MASK) {
if (opt_algo == ALGO_ZR5 && work->data[0] & POK_BOOL_MASK) {
use_pok = 1;
json_t *txs = json_object_get(val, "txs");
if (txs && json_is_array(txs)) {
@ -844,6 +846,10 @@ static bool submit_upstream_work(CURL *curl, struct work *work)
le32enc(&ntime, work->data[25]);
le32enc(&nonce, work->data[27]);
break;
case ALGO_SIA:
be32enc(&ntime, work->data[10]);
be32enc(&nonce, work->data[8]);
break;
case ALGO_ZR5:
check_dups = true;
be32enc(&ntime, work->data[17]);
@ -874,6 +880,9 @@ static bool submit_upstream_work(CURL *curl, struct work *work)
if (opt_algo == ALGO_DECRED) {
xnonce2str = bin2hex((const uchar*)&work->data[36], stratum.xnonce1_size);
} else if (opt_algo == ALGO_SIA) {
uint16_t high_nonce = swab32(work->data[9]) >> 16;
xnonce2str = bin2hex((unsigned char*)(&high_nonce), 2);
} else {
xnonce2str = bin2hex(work->xnonce2, work->xnonce2_len);
}
@ -1394,6 +1403,7 @@ static bool stratum_gen_work(struct stratum_ctx *sctx, struct work *work)
/* Generate merkle root */
switch (opt_algo) {
case ALGO_DECRED:
case ALGO_SIA:
// getwork over stratum, no merkle to generate
break;
case ALGO_HEAVY:
@ -1446,7 +1456,6 @@ static bool stratum_gen_work(struct stratum_ctx *sctx, struct work *work)
sctx->xnonce1_size = sizeof(work->data)-(36*4);
}
memcpy(&work->data[36], sctx->xnonce1, sctx->xnonce1_size);
// work->data[36] = swab32(vote); // alt vote submission method
work->data[37] = (rand()*4) << 8; // random work data
sctx->job.height = work->data[32];
//applog_hex(work->data, 180);
@ -1458,6 +1467,18 @@ static bool stratum_gen_work(struct stratum_ctx *sctx, struct work *work)
work->data[25] = le32dec(sctx->job.ntime);
work->data[26] = le32dec(sctx->job.nbits);
work->data[28] = 0x80000000;
} else if (opt_algo == ALGO_SIA) {
uint32_t extra = 0;
memcpy(&extra, &sctx->job.coinbase[32], 2);
for (i = 0; i < 8; i++) // reversed hash
work->data[i] = ((uint32_t*)sctx->job.prevhash)[7-i];
work->data[8] = 0; // nonce
work->data[9] = swab32(extra) | ((rand() << 8) & 0xffff);
work->data[10] = be32dec(sctx->job.ntime);
work->data[11] = be32dec(sctx->job.nbits);
memcpy(&work->data[12], sctx->job.coinbase, 32); // merkle_root
work->data[20] = 0x80000000;
if (opt_debug) applog_hex(work->data, 80);
} else {
for (i = 0; i < 8; i++)
work->data[9 + i] = be32dec((uint32_t *)merkle_root + i);
@ -1490,7 +1511,7 @@ static bool stratum_gen_work(struct stratum_ctx *sctx, struct work *work)
pthread_mutex_unlock(&stratum_work_lock);
if (opt_debug && opt_algo != ALGO_DECRED) {
if (opt_debug && opt_algo != ALGO_DECRED && opt_algo != ALGO_SIA) {
uint32_t utm = work->data[17];
if (opt_algo != ALGO_ZR5) utm = swab32(utm);
char *tm = atime2str(utm - sctx->srvtime_diff);
@ -1673,11 +1694,19 @@ static void *miner_thread(void *userdata)
uint32_t start_nonce;
uint32_t scan_time = have_longpoll ? LP_SCANTIME : opt_scantime;
uint64_t max64, minmax = 0x100000;
int nodata_check_oft = 0;
bool regen = false;
// &work.data[19]
int wcmplen = (opt_algo == ALGO_DECRED) ? 140 : 76;
if (opt_algo == ALGO_LBRY) wcmplen = 108;
int wcmpoft = 0;
if (opt_algo == ALGO_LBRY) wcmplen = 108;
else if (opt_algo == ALGO_SIA) {
wcmpoft = (32+16)/4;
wcmplen = 32;
}
uint32_t *nonceptr = (uint32_t*) (((char*)work.data) + wcmplen);
if (have_stratum) {
@ -1698,7 +1727,14 @@ static void *miner_thread(void *userdata)
nonceptr = (uint32_t*) (((char*)work.data) + wcmplen);
pthread_mutex_lock(&g_work_lock);
extrajob |= work_done;
if (nonceptr[0] >= end_nonce || extrajob) {
regen = (nonceptr[0] >= end_nonce);
if (opt_algo == ALGO_SIA) {
regen = (nonceptr[1] & 0xFF00 >= 0xF000);
}
regen = regen || extrajob;
if (regen) {
work_done = false;
extrajob = false;
if (stratum_gen_work(&stratum, &g_work))
@ -1774,6 +1810,20 @@ static void *miner_thread(void *userdata)
// and make an unique work (extradata)
nonceptr[1] += 1;
nonceptr[2] |= thr_id;
} else if (opt_algo == ALGO_SIA) {
// suprnova job_id check without data/target/height change...
check_stratum_jobs = true;
if (check_stratum_jobs && strcmp(work.job_id, g_work.job_id)) {
pthread_mutex_unlock(&g_work_lock);
work_done = true;
continue;
}
nonceptr[1] += opt_n_threads;
nonceptr[1] |= thr_id;
// range max
nonceptr[0] = 0;
end_nonce = UINT32_MAX;
} else if (opt_benchmark) {
// randomize work
nonceptr[-1] += 1;
@ -1796,11 +1846,14 @@ static void *miner_thread(void *userdata)
}
loopcnt++;
/* prevent gpu scans before a job is received */
//if (opt_algo != ALGO_DECRED) // uncomment to allow testnet
if (have_stratum && work.data[0] == 0 && !opt_benchmark) {
// prevent gpu scans before a job is received
if (opt_algo == ALGO_SIA) nodata_check_oft = 7; // no stratum version
else if (opt_algo == ALGO_DECRED) nodata_check_oft = 4; // testnet ver is 0
else nodata_check_oft = 0;
if (have_stratum && work.data[nodata_check_oft] == 0 && !opt_benchmark) {
sleep(1);
if (!thr_id) pools[cur_pooln].wait_time += 1;
gpulog(LOG_DEBUG, thr_id, "no data");
continue;
}
@ -1931,6 +1984,7 @@ static void *miner_thread(void *userdata)
case ALGO_KECCAK:
case ALGO_LBRY:
case ALGO_LUFFA:
case ALGO_SIA:
case ALGO_SKEIN:
case ALGO_SKEIN2:
minmax = 0x1000000;
@ -2096,6 +2150,9 @@ static void *miner_thread(void *userdata)
case ALGO_SKEIN2:
rc = scanhash_skein2(thr_id, &work, max_nonce, &hashes_done);
break;
case ALGO_SIA:
rc = scanhash_sia(thr_id, &work, max_nonce, &hashes_done);
break;
case ALGO_SIB:
rc = scanhash_sib(thr_id, &work, max_nonce, &hashes_done);
break;
@ -2152,8 +2209,9 @@ static void *miner_thread(void *userdata)
gettimeofday(&tv_end, NULL);
// todo: update all algos to use work->nonces
work.nonces[0] = nonceptr[0];
if (opt_algo != ALGO_DECRED && opt_algo != ALGO_BLAKE2S && opt_algo != ALGO_LBRY) {
if (opt_algo != ALGO_SIA) // reversed endian
work.nonces[0] = nonceptr[0];
if (opt_algo != ALGO_DECRED && opt_algo != ALGO_BLAKE2S && opt_algo != ALGO_LBRY && opt_algo != ALGO_SIA) {
work.nonces[1] = nonceptr[2];
}
@ -2203,7 +2261,7 @@ static void *miner_thread(void *userdata)
nonceptr[0] = UINT32_MAX;
}
if (check_dups && opt_algo != ALGO_DECRED)
if (check_dups && opt_algo != ALGO_DECRED && opt_algo != ALGO_SIA)
hashlog_remember_scan_range(&work);
/* output */
@ -2236,10 +2294,15 @@ static void *miner_thread(void *userdata)
/* if nonce found, submit work */
if (rc > 0 && !opt_benchmark) {
uint32_t curnonce = nonceptr[0]; // current scan position
if (opt_led_mode == LED_MODE_SHARES)
gpu_led_percent(dev_id, 50);
nonceptr[0] = work.nonces[0];
if (!submit_work(mythr, &work))
break;
nonceptr[0] = curnonce;
// prevent stale work in solo
// we can't submit twice a block!
@ -2261,6 +2324,7 @@ static void *miner_thread(void *userdata)
}
if (!submit_work(mythr, &work))
break;
nonceptr[0] = curnonce;
}
}
}
@ -3431,7 +3495,7 @@ int main(int argc, char *argv[])
cur_pooln = pool_get_first_valid(0);
pool_switch(-1, cur_pooln);
if (opt_algo == ALGO_DECRED) {
if (opt_algo == ALGO_DECRED || opt_algo == ALGO_SIA) {
allow_gbt = false;
allow_mininginfo = false;
}

6
ccminer.vcxproj
View File

@ -115,7 +115,7 @@
<MaxRegCount>80</MaxRegCount>
<PtxAsOptionV>true</PtxAsOptionV>
<Keep>true</Keep>
<CodeGeneration>compute_61,sm_61;compute_52,sm_52</CodeGeneration>
<CodeGeneration>compute_52,sm_52</CodeGeneration>
<Include>$(NVTOOLSEXT_PATH)\include;..\..\..\Common\C99</Include>
<TargetMachinePlatform>64</TargetMachinePlatform>
</CudaCompile>
@ -266,6 +266,7 @@
<ClCompile Include="skein2.cpp" />
<ClCompile Include="sph\aes_helper.c" />
<ClCompile Include="sph\blake.c" />
<ClCompile Include="sph\blake2b.c" />
<ClCompile Include="sph\blake2s.c" />
<ClCompile Include="sph\bmw.c" />
<ClCompile Include="sph\cubehash.c" />
@ -322,6 +323,8 @@
<ClInclude Include="quark\cuda_skein512_sp.cuh" />
<ClInclude Include="res\resource.h" />
<ClInclude Include="scrypt\salsa_kernel.h" />
<ClInclude Include="sph\blake2b.h" />
<ClInclude Include="sph\blake2s.h" />
<ClInclude Include="sph\sph_blake.h" />
<ClInclude Include="sph\sph_bmw.h" />
<ClInclude Include="sph\sph_cubehash.h" />
@ -469,6 +472,7 @@
<CudaCompile Include="lyra2\cuda_lyra2.cu" />
<CudaCompile Include="lyra2\lyra2REv2.cu" />
<CudaCompile Include="lyra2\cuda_lyra2v2.cu" />
<CudaCompile Include="sia.cu" />
<CudaCompile Include="skein.cu">
<MaxRegCount>64</MaxRegCount>
</CudaCompile>

18
ccminer.vcxproj.filters
View File

@ -129,6 +129,12 @@
<ClCompile Include="sph\blake.c">
<Filter>Source Files\sph</Filter>
</ClCompile>
<ClCompile Include="sph\blake2b.c">
<Filter>Source Files\sph</Filter>
</ClCompile>
<ClCompile Include="sph\blake2s.c">
<Filter>Source Files\sph</Filter>
</ClCompile>
<ClCompile Include="sph\bmw.c">
<Filter>Source Files\sph</Filter>
</ClCompile>
@ -264,9 +270,6 @@
<ClCompile Include="bignum.cpp">
<Filter>Source Files</Filter>
</ClCompile>
<ClCompile Include="sph\blake2s.c">
<Filter>Source Files\sph</Filter>
</ClCompile>
<ClCompile Include="nvapi.cpp">
<Filter>Source Files</Filter>
</ClCompile>
@ -314,6 +317,12 @@
<ClInclude Include="x11\cuda_x11.h">
<Filter>Header Files\CUDA</Filter>
</ClInclude>
<ClInclude Include="sph\blake2b.h">
<Filter>Header Files\sph</Filter>
</ClInclude>
<ClInclude Include="sph\blake2s.h">
<Filter>Header Files\sph</Filter>
</ClInclude>
<ClInclude Include="sph\sph_blake.h">
<Filter>Header Files\sph</Filter>
</ClInclude>
@ -664,6 +673,9 @@
<CudaCompile Include="zr5.cu">
<Filter>Source Files\CUDA</Filter>
</CudaCompile>
<CudaCompile Include="sia.cu">
<Filter>Source Files\CUDA</Filter>
</CudaCompile>
<CudaCompile Include="skein.cu">
<Filter>Source Files\CUDA</Filter>
</CudaCompile>

3
miner.h
View File

@ -285,6 +285,7 @@ extern int scanhash_nist5(int thr_id, struct work *work, uint32_t max_nonce, uns
extern int scanhash_pentablake(int thr_id, struct work *work, uint32_t max_nonce, unsigned long *hashes_done);
extern int scanhash_quark(int thr_id, struct work *work, uint32_t max_nonce, unsigned long *hashes_done);
extern int scanhash_qubit(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done);
extern int scanhash_sia(int thr_id, struct work *work, uint32_t max_nonce, unsigned long *hashes_done);
extern int scanhash_sib(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done);
extern int scanhash_skeincoin(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done);
extern int scanhash_skein2(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done);
@ -320,6 +321,7 @@ extern void free_fugue256(int thr_id);
extern void free_groestlcoin(int thr_id);
extern void free_heavy(int thr_id);
extern void free_jackpot(int thr_id);
extern void free_lbry(int thr_id);
extern void free_luffa(int thr_id);
extern void free_lyra2(int thr_id);
extern void free_lyra2v2(int thr_id);
@ -329,6 +331,7 @@ extern void free_nist5(int thr_id);
extern void free_pentablake(int thr_id);
extern void free_quark(int thr_id);
extern void free_qubit(int thr_id);
extern void free_sia(int thr_id);
extern void free_sib(int thr_id);
extern void free_skeincoin(int thr_id);
extern void free_skein2(int thr_id);

301
sia.cu Normal file
View File

@ -0,0 +1,301 @@
/**
* Blake2-B CUDA Implementation
*
* tpruvot@github July 2016
*
*/
#include <miner.h>
#include <string.h>
#include <stdint.h>
#include <sph/blake2b.h>
#include <cuda_helper.h>
#include <cuda_vector_uint2x4.h>
#define TPB 512
#define NBN 2
static uint32_t *d_resNonces[MAX_GPUS];
__device__ uint64_t d_data[10];
static __constant__ const int8_t blake2b_sigma[12][16] = {
{ 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 } ,
{ 14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 } ,
{ 11, 8, 12, 0, 5, 2, 15, 13, 10, 14, 3, 6, 7, 1, 9, 4 } ,
{ 7, 9, 3, 1, 13, 12, 11, 14, 2, 6, 5, 10, 4, 0, 15, 8 } ,
{ 9, 0, 5, 7, 2, 4, 10, 15, 14, 1, 11, 12, 6, 8, 3, 13 } ,
{ 2, 12, 6, 10, 0, 11, 8, 3, 4, 13, 7, 5, 15, 14, 1, 9 } ,
{ 12, 5, 1, 15, 14, 13, 4, 10, 0, 7, 6, 3, 9, 2, 8, 11 } ,
{ 13, 11, 7, 14, 12, 1, 3, 9, 5, 0, 15, 4, 8, 6, 2, 10 } ,
{ 6, 15, 14, 9, 11, 3, 0, 8, 12, 2, 13, 7, 1, 4, 10, 5 } ,
{ 10, 2, 8, 4, 7, 6, 1, 5, 15, 11, 9, 14, 3, 12, 13, 0 } ,
{ 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 } ,
{ 14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 }
};
// host mem align
#define A 64
extern "C" void blake2b_hash(void *output, const void *input)
{
uint8_t _ALIGN(A) hash[32];
blake2b_ctx ctx;
blake2b_init(&ctx, 32, NULL, 0);
blake2b_update(&ctx, input, 80);
blake2b_final(&ctx, hash);
memcpy(output, hash, 32);
}
// ----------------------------------------------------------------
__device__ __forceinline__
static void G(const int r, const int i, uint64_t &a, uint64_t &b, uint64_t &c, uint64_t &d, uint64_t const m[16])
{
a = a + b + m[ blake2b_sigma[r][2*i] ];
((uint2*)&d)[0] = SWAPUINT2( ((uint2*)&d)[0] ^ ((uint2*)&a)[0] );
c = c + d;
((uint2*)&b)[0] = ROR24( ((uint2*)&b)[0] ^ ((uint2*)&c)[0] );
a = a + b + m[ blake2b_sigma[r][2*i+1] ];
((uint2*)&d)[0] = ROR16( ((uint2*)&d)[0] ^ ((uint2*)&a)[0] );
c = c + d;
((uint2*)&b)[0] = ROR2( ((uint2*)&b)[0] ^ ((uint2*)&c)[0], 63U);
}
#define ROUND(r) \
G(r, 0, v[0], v[4], v[ 8], v[12], m); \
G(r, 1, v[1], v[5], v[ 9], v[13], m); \
G(r, 2, v[2], v[6], v[10], v[14], m); \
G(r, 3, v[3], v[7], v[11], v[15], m); \
G(r, 4, v[0], v[5], v[10], v[15], m); \
G(r, 5, v[1], v[6], v[11], v[12], m); \
G(r, 6, v[2], v[7], v[ 8], v[13], m); \
G(r, 7, v[3], v[4], v[ 9], v[14], m);
// simplified for the last round
__device__ __forceinline__
static void H(const int r, const int i, uint64_t &a, uint64_t &b, uint64_t &c, uint64_t &d, uint64_t const m[16])
{
a = a + b + m[ blake2b_sigma[r][2*i] ];
((uint2*)&d)[0] = SWAPUINT2( ((uint2*)&d)[0] ^ ((uint2*)&a)[0] );
c = c + d;
((uint2*)&b)[0] = ROR24( ((uint2*)&b)[0] ^ ((uint2*)&c)[0] );
a = a + b + m[ blake2b_sigma[r][2*i+1] ];
((uint2*)&d)[0] = ROR16( ((uint2*)&d)[0] ^ ((uint2*)&a)[0] );
c = c + d;
}
// we only check v[0] and v[8]
#define ROUND_F(r) \
G(r, 0, v[0], v[4], v[ 8], v[12], m); \
G(r, 1, v[1], v[5], v[ 9], v[13], m); \
G(r, 2, v[2], v[6], v[10], v[14], m); \
G(r, 3, v[3], v[7], v[11], v[15], m); \
G(r, 4, v[0], v[5], v[10], v[15], m); \
G(r, 5, v[1], v[6], v[11], v[12], m); \
H(r, 6, v[2], v[7], v[ 8], v[13], m);
__global__
//__launch_bounds__(128, 8) /* to force 64 regs */
void blake2b_gpu_hash(const uint32_t threads, const uint32_t startNonce, uint32_t *resNonce, const uint2 target2)
{
const uint32_t nonce = (blockDim.x * blockIdx.x + threadIdx.x) + startNonce;
__shared__ uint64_t s_target;
if (!threadIdx.x) s_target = devectorize(target2);
uint64_t m[16];
m[0] = d_data[0];
m[1] = d_data[1];
m[2] = d_data[2];
m[3] = d_data[3];
m[4] = d_data[4] | nonce;
m[5] = d_data[5];
m[6] = d_data[6];
m[7] = d_data[7];
m[8] = d_data[8];
m[9] = d_data[9];
m[10] = m[11] = 0;
m[12] = m[13] = m[14] = m[15] = 0;
uint64_t v[16] = {
0x6a09e667f2bdc928, 0xbb67ae8584caa73b, 0x3c6ef372fe94f82b, 0xa54ff53a5f1d36f1,
0x510e527fade682d1, 0x9b05688c2b3e6c1f, 0x1f83d9abfb41bd6b, 0x5be0cd19137e2179,
0x6a09e667f3bcc908, 0xbb67ae8584caa73b, 0x3c6ef372fe94f82b, 0xa54ff53a5f1d36f1,
0x510e527fade68281, 0x9b05688c2b3e6c1f, 0xe07c265404be4294, 0x5be0cd19137e2179
};
ROUND( 0 );
ROUND( 1 );
ROUND( 2 );
ROUND( 3 );
ROUND( 4 );
ROUND( 5 );
ROUND( 6 );
ROUND( 7 );
ROUND( 8 );
ROUND( 9 );
ROUND( 10 );
ROUND_F( 11 );
uint64_t h64 = cuda_swab64(0x6a09e667f2bdc928 ^ v[0] ^ v[8]);
if (h64 <= s_target) {
resNonce[1] = resNonce[0];
resNonce[0] = nonce;
s_target = h64;
}
// if (!nonce) printf("%016lx ", s_target);
}
__host__
uint32_t blake2b_hash_cuda(const int thr_id, const uint32_t threads, const uint32_t startNonce, const uint2 target2, uint32_t &secNonce)
{
uint32_t resNonces[NBN] = { UINT32_MAX, UINT32_MAX };
uint32_t result = UINT32_MAX;
dim3 grid((threads + TPB-1)/TPB);
dim3 block(TPB);
/* Check error on Ctrl+C or kill to prevent segfaults on exit */
if (cudaMemset(d_resNonces[thr_id], 0xff, NBN*sizeof(uint32_t)) != cudaSuccess)
return result;
blake2b_gpu_hash <<<grid, block, 8>>> (threads, startNonce, d_resNonces[thr_id], target2);
cudaThreadSynchronize();
if (cudaSuccess == cudaMemcpy(resNonces, d_resNonces[thr_id], NBN*sizeof(uint32_t), cudaMemcpyDeviceToHost)) {
result = resNonces[0];
secNonce = resNonces[1];
if (secNonce == result) secNonce = UINT32_MAX;
}
return result;
}
__host__
void blake2b_setBlock(uint32_t *data)
{
CUDA_SAFE_CALL(cudaMemcpyToSymbol(d_data, data, 80, 0, cudaMemcpyHostToDevice));
}
static bool init[MAX_GPUS] = { 0 };
int scanhash_sia(int thr_id, struct work *work, uint32_t max_nonce, unsigned long *hashes_done)
{
uint32_t _ALIGN(A) hash[8];
uint32_t _ALIGN(A) vhashcpu[8];
uint32_t _ALIGN(A) inputdata[20];
uint32_t *pdata = work->data;
uint32_t *ptarget = work->target;
const uint32_t Htarg = ptarget[7];
const uint32_t first_nonce = pdata[8];
int dev_id = device_map[thr_id];
int intensity = (device_sm[dev_id] >= 500 && !is_windows()) ? 28 : 25;
if (device_sm[dev_id] >= 520 && is_windows()) intensity = 26;
if (device_sm[dev_id] < 350) intensity = 22;
uint32_t throughput = cuda_default_throughput(thr_id, 1U << intensity);
if (init[thr_id]) throughput = min(throughput, max_nonce - first_nonce);
if (!init[thr_id])
{
cudaSetDevice(dev_id);
if (opt_cudaschedule == -1 && gpu_threads == 1) {
cudaDeviceReset();
// reduce cpu usage (linux)
cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
//cudaDeviceSetCacheConfig(cudaFuncCachePreferL1);
CUDA_LOG_ERROR();
}
CUDA_CALL_OR_RET_X(cudaMalloc(&d_resNonces[thr_id], NBN * sizeof(uint32_t)), -1);
init[thr_id] = true;
}
memcpy(inputdata, pdata, 80);
inputdata[11] = 0; // nbits
const uint2 target = make_uint2(ptarget[6], ptarget[7]);
blake2b_setBlock(inputdata);
do {
uint32_t secNonce = UINT32_MAX;
uint32_t foundNonce = blake2b_hash_cuda(thr_id, throughput, pdata[8], target, secNonce);
*hashes_done = pdata[8] - first_nonce + throughput;
if (foundNonce != UINT32_MAX)
{
int res = 0;
inputdata[8] = foundNonce;
blake2b_hash(hash, inputdata);
if (swab32(hash[0]) <= Htarg) {
// sia hash target is reversed (start of hash)
swab256(vhashcpu, hash);
// applog_hex(vhashcpu, 32);
if (fulltest(vhashcpu, ptarget)) {
work_set_target_ratio(work, vhashcpu);
work->nonces[0] = foundNonce;
res ++;
}
}
if (secNonce != UINT32_MAX) {
inputdata[8] = secNonce;
blake2b_hash(hash, inputdata);
if (swab32(hash[0]) <= Htarg) {
if (opt_debug)
gpulog(LOG_BLUE, thr_id, "found second nonce %08x", secNonce);
swab256(vhashcpu, hash);
if (fulltest(vhashcpu, ptarget)) {
work->nonces[1] = secNonce;
if (bn_hash_target_ratio(vhashcpu, ptarget) > work->shareratio) {
work_set_target_ratio(work, vhashcpu);
xchg(work->nonces[0], work->nonces[1]);
}
res++;
}
}
}
if (res) {
pdata[8] = max_nonce;
return res;
}
}
if ((uint64_t) throughput + pdata[8] >= max_nonce) {
pdata[8] = max_nonce;
break;
}
pdata[8] += throughput;
} while (!work_restart[thr_id].restart);
*hashes_done = pdata[8] - first_nonce;
return 0;
}
// cleanup
extern "C" void free_sia(int thr_id)
{
if (!init[thr_id])
return;
cudaThreadSynchronize();
cudaFree(d_resNonces[thr_id]);
init[thr_id] = false;
cudaDeviceSynchronize();
}

196
sph/blake2b.c Normal file
View File

@ -0,0 +1,196 @@
/*
* Copyright 2009 Colin Percival, 2014 savale
* 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 <stdlib.h>
#include <stdint.h>
#include <string.h>
#include "sph_types.h"
#include "blake2b.h"
// Cyclic right rotation.
#ifndef ROTR64
#define ROTR64(x, y) (((x) >> (y)) ^ ((x) << (64 - (y))))
#endif
// Little-endian byte access.
#define B2B_GET64(p) \
(((uint64_t) ((uint8_t *) (p))[0]) ^ \
(((uint64_t) ((uint8_t *) (p))[1]) << 8) ^ \
(((uint64_t) ((uint8_t *) (p))[2]) << 16) ^ \
(((uint64_t) ((uint8_t *) (p))[3]) << 24) ^ \
(((uint64_t) ((uint8_t *) (p))[4]) << 32) ^ \
(((uint64_t) ((uint8_t *) (p))[5]) << 40) ^ \
(((uint64_t) ((uint8_t *) (p))[6]) << 48) ^ \
(((uint64_t) ((uint8_t *) (p))[7]) << 56))
// G Mixing function.
#define B2B_G(a, b, c, d, x, y) { \
v[a] = v[a] + v[b] + x; \
v[d] = ROTR64(v[d] ^ v[a], 32); \
v[c] = v[c] + v[d]; \
v[b] = ROTR64(v[b] ^ v[c], 24); \
v[a] = v[a] + v[b] + y; \
v[d] = ROTR64(v[d] ^ v[a], 16); \
v[c] = v[c] + v[d]; \
v[b] = ROTR64(v[b] ^ v[c], 63); }
// Initialization Vector.
static const uint64_t blake2b_iv[8] = {
0x6A09E667F3BCC908, 0xBB67AE8584CAA73B,
0x3C6EF372FE94F82B, 0xA54FF53A5F1D36F1,
0x510E527FADE682D1, 0x9B05688C2B3E6C1F,
0x1F83D9ABFB41BD6B, 0x5BE0CD19137E2179
};
// Compression function. "last" flag indicates last block.
static void blake2b_compress(blake2b_ctx *ctx, int last)
{
const uint8_t sigma[12][16] = {
{ 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 },
{ 14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 },
{ 11, 8, 12, 0, 5, 2, 15, 13, 10, 14, 3, 6, 7, 1, 9, 4 },
{ 7, 9, 3, 1, 13, 12, 11, 14, 2, 6, 5, 10, 4, 0, 15, 8 },
{ 9, 0, 5, 7, 2, 4, 10, 15, 14, 1, 11, 12, 6, 8, 3, 13 },
{ 2, 12, 6, 10, 0, 11, 8, 3, 4, 13, 7, 5, 15, 14, 1, 9 },
{ 12, 5, 1, 15, 14, 13, 4, 10, 0, 7, 6, 3, 9, 2, 8, 11 },
{ 13, 11, 7, 14, 12, 1, 3, 9, 5, 0, 15, 4, 8, 6, 2, 10 },
{ 6, 15, 14, 9, 11, 3, 0, 8, 12, 2, 13, 7, 1, 4, 10, 5 },
{ 10, 2, 8, 4, 7, 6, 1, 5, 15, 11, 9, 14, 3, 12, 13, 0 },
{ 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 },
{ 14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 }
};
int i;
uint64_t v[16], m[16];
for (i = 0; i < 8; i++) { // init work variables
v[i] = ctx->h[i];
v[i + 8] = blake2b_iv[i];
}
v[12] ^= ctx->t[0]; // low 64 bits of offset
v[13] ^= ctx->t[1]; // high 64 bits
if (last) // last block flag set ?
v[14] = ~v[14];
for (i = 0; i < 16; i++) // get little-endian words
m[i] = B2B_GET64(&ctx->b[8 * i]);
for (i = 0; i < 12; i++) { // twelve rounds
B2B_G( 0, 4, 8, 12, m[sigma[i][ 0]], m[sigma[i][ 1]]);
B2B_G( 1, 5, 9, 13, m[sigma[i][ 2]], m[sigma[i][ 3]]);
B2B_G( 2, 6, 10, 14, m[sigma[i][ 4]], m[sigma[i][ 5]]);
B2B_G( 3, 7, 11, 15, m[sigma[i][ 6]], m[sigma[i][ 7]]);
B2B_G( 0, 5, 10, 15, m[sigma[i][ 8]], m[sigma[i][ 9]]);
B2B_G( 1, 6, 11, 12, m[sigma[i][10]], m[sigma[i][11]]);
B2B_G( 2, 7, 8, 13, m[sigma[i][12]], m[sigma[i][13]]);
B2B_G( 3, 4, 9, 14, m[sigma[i][14]], m[sigma[i][15]]);
}
for( i = 0; i < 8; ++i )
ctx->h[i] ^= v[i] ^ v[i + 8];
}
// Initialize the hashing context "ctx" with optional key "key".
// 1 <= outlen <= 64 gives the digest size in bytes.
// Secret key (also <= 64 bytes) is optional (keylen = 0).
int blake2b_init(blake2b_ctx *ctx, size_t outlen,
const void *key, size_t keylen) // (keylen=0: no key)
{
size_t i;
if (outlen == 0 || outlen > 64 || keylen > 64)
return -1; // illegal parameters
for (i = 0; i < 8; i++) // state, "param block"
ctx->h[i] = blake2b_iv[i];
ctx->h[0] ^= 0x01010000 ^ (keylen << 8) ^ outlen;
ctx->t[0] = 0; // input count low word
ctx->t[1] = 0; // input count high word
ctx->c = 0; // pointer within buffer
ctx->outlen = outlen;
for (i = keylen; i < 128; i++) // zero input block
ctx->b[i] = 0;
if (keylen > 0) {
blake2b_update(ctx, key, keylen);
ctx->c = 128; // at the end
}
return 0;
}
// Add "inlen" bytes from "in" into the hash.
void blake2b_update(blake2b_ctx *ctx,
const void *in, size_t inlen) // data bytes
{
size_t i;
for (i = 0; i < inlen; i++) {
if (ctx->c == 128) { // buffer full ?
ctx->t[0] += ctx->c; // add counters
if (ctx->t[0] < ctx->c) // carry overflow ?
ctx->t[1]++; // high word
blake2b_compress(ctx, 0); // compress (not last)
ctx->c = 0; // counter to zero
}
ctx->b[ctx->c++] = ((const uint8_t *) in)[i];
}
}
// Generate the message digest (size given in init).
// Result placed in "out".
void blake2b_final(blake2b_ctx *ctx, void *out)
{
size_t i;
ctx->t[0] += ctx->c; // mark last block offset
if (ctx->t[0] < ctx->c) // carry overflow
ctx->t[1]++; // high word
while (ctx->c < 128) // fill up with zeros
ctx->b[ctx->c++] = 0;
blake2b_compress(ctx, 1); // final block flag = 1
// little endian convert and store
for (i = 0; i < ctx->outlen; i++) {
((uint8_t *) out)[i] =
(ctx->h[i >> 3] >> (8 * (i & 7))) & 0xFF;
}
}

41
sph/blake2b.h Normal file
View File

@ -0,0 +1,41 @@
#pragma once
#ifndef __BLAKE2B_H__
#define __BLAKE2B_H__
#include <stddef.h>
#include <stdint.h>
#if defined(_MSC_VER)
#include <inttypes.h>
#define inline __inline
#define ALIGN(x) __declspec(align(x))
#else
#define ALIGN(x) __attribute__((aligned(x)))
#endif
#if defined(_MSC_VER) || defined(__x86_64__) || defined(__x86__)
#define NATIVE_LITTLE_ENDIAN
#endif
// state context
ALIGN(64) typedef struct {
uint8_t b[128]; // input buffer
uint64_t h[8]; // chained state
uint64_t t[2]; // total number of bytes
size_t c; // pointer for b[]
size_t outlen; // digest size
} blake2b_ctx;
#if defined(__cplusplus)
extern "C" {
#endif
int blake2b_init(blake2b_ctx *ctx, size_t outlen, const void *key, size_t keylen);
void blake2b_update(blake2b_ctx *ctx, const void *in, size_t inlen);
void blake2b_final(blake2b_ctx *ctx, void *out);
#if defined(__cplusplus)
}
#endif
#endif