mirror of
https://github.com/GOSTSec/ccminer
synced 2025-01-26 22:44:15 +00:00
invoke gost
This commit is contained in:
parent
1042762dc4
commit
aa4dcdf58d
2
algos.h
2
algos.h
@ -18,6 +18,7 @@ enum sha_algos {
|
||||
ALGO_DMD_GR,
|
||||
ALGO_FRESH,
|
||||
ALGO_FUGUE256, /* Fugue256 */
|
||||
ALGO_GOST,
|
||||
ALGO_GROESTL,
|
||||
ALGO_HEAVY, /* Heavycoin hash */
|
||||
ALGO_HMQ1725,
|
||||
@ -78,6 +79,7 @@ static const char *algo_names[] = {
|
||||
"dmd-gr",
|
||||
"fresh",
|
||||
"fugue256",
|
||||
"gost",
|
||||
"groestl",
|
||||
"heavy",
|
||||
"hmq1725",
|
||||
|
@ -236,6 +236,7 @@ Options:\n\
|
||||
dmd-gr Diamond-Groestl\n\
|
||||
fresh Freshcoin (shavite 80)\n\
|
||||
fugue256 Fuguecoin\n\
|
||||
gost GOST\n\
|
||||
groestl Groestlcoin\n\
|
||||
heavy Heavycoin\n\
|
||||
hmq1725 Doubloons / Espers\n\
|
||||
@ -2323,6 +2324,9 @@ static void *miner_thread(void *userdata)
|
||||
case ALGO_SIB:
|
||||
rc = scanhash_sib(thr_id, &work, max_nonce, &hashes_done);
|
||||
break;
|
||||
case ALGO_GOST:
|
||||
rc = scanhash_gost(thr_id, &work, max_nonce, &hashes_done);
|
||||
break;
|
||||
case ALGO_S3:
|
||||
rc = scanhash_s3(thr_id, &work, max_nonce, &hashes_done);
|
||||
break;
|
||||
|
95
gost/gost.cu
95
gost/gost.cu
@ -32,3 +32,98 @@ extern "C" void gosthash(void *output, const void *input)
|
||||
memcpy(output, hash, 32);
|
||||
}
|
||||
|
||||
//#define _DEBUG
|
||||
#define _DEBUG_PREFIX "sib"
|
||||
#include "cuda_debug.cuh"
|
||||
|
||||
static bool init[MAX_GPUS] = { 0 };
|
||||
|
||||
extern "C" int scanhash_gost(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done)
|
||||
{
|
||||
uint32_t *pdata = work->data;
|
||||
uint32_t *ptarget = work->target;
|
||||
const uint32_t first_nonce = pdata[19];
|
||||
const int dev_id = device_map[thr_id];
|
||||
int intensity = (device_sm[dev_id] >= 500 && !is_windows()) ? 19 : 18; // 2^18 = 262144 cuda threads
|
||||
if (device_sm[dev_id] >= 600) intensity = 20;
|
||||
uint32_t throughput = cuda_default_throughput(thr_id, 1U << intensity);
|
||||
if (init[thr_id]) throughput = min(throughput, max_nonce - first_nonce);
|
||||
|
||||
if (opt_benchmark)
|
||||
ptarget[7] = 0xf;
|
||||
|
||||
if (!init[thr_id])
|
||||
{
|
||||
cudaSetDevice(device_map[thr_id]);
|
||||
if (opt_cudaschedule == -1 && gpu_threads == 1) {
|
||||
cudaDeviceReset();
|
||||
// reduce cpu usage
|
||||
cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
|
||||
CUDA_LOG_ERROR();
|
||||
}
|
||||
gpulog(LOG_INFO, thr_id, "Intensity set to %g, %u cuda threads", throughput2intensity(throughput), throughput);
|
||||
|
||||
CUDA_CALL_OR_RET_X(cudaMalloc(&d_hash[thr_id], (size_t) 64 * throughput), -1);
|
||||
|
||||
cuda_check_cpu_init(thr_id, throughput);
|
||||
|
||||
init[thr_id] = true;
|
||||
}
|
||||
|
||||
uint32_t endiandata[20];
|
||||
for (int k=0; k < 20; k++)
|
||||
be32enc(&endiandata[k], pdata[k]);
|
||||
|
||||
do {
|
||||
int order = 0;
|
||||
|
||||
// Hash with CUDA
|
||||
//streebog_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);
|
||||
//TRACE("gost :");
|
||||
|
||||
work->nonces[0] = cuda_check_hash(thr_id, throughput, pdata[19], d_hash[thr_id]);
|
||||
if (work->nonces[0] != UINT32_MAX)
|
||||
{
|
||||
const uint32_t Htarg = ptarget[7];
|
||||
uint32_t _ALIGN(64) vhash[8];
|
||||
be32enc(&endiandata[19], work->nonces[0]);
|
||||
sibhash(vhash, endiandata);
|
||||
|
||||
if (vhash[7] <= Htarg && fulltest(vhash, ptarget)) {
|
||||
work->valid_nonces = 1;
|
||||
work_set_target_ratio(work, vhash);
|
||||
work->nonces[1] =cuda_check_hash_suppl(thr_id, throughput, pdata[19], d_hash[thr_id], 1);
|
||||
*hashes_done = pdata[19] - first_nonce + throughput;
|
||||
if (work->nonces[1] != 0) {
|
||||
be32enc(&endiandata[19], work->nonces[1]);
|
||||
sibhash(vhash, endiandata);
|
||||
bn_set_target_ratio(work, vhash, 1);
|
||||
work->valid_nonces++;
|
||||
pdata[19] = max(work->nonces[0], work->nonces[1]) + 1;
|
||||
} else {
|
||||
pdata[19] = work->nonces[0] + 1; // cursor
|
||||
}
|
||||
return work->valid_nonces;
|
||||
}
|
||||
else if (vhash[7] > Htarg) {
|
||||
gpu_increment_reject(thr_id);
|
||||
if (!opt_quiet)
|
||||
gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", work->nonces[0]);
|
||||
pdata[19] = work->nonces[0] + 1;
|
||||
continue;
|
||||
}
|
||||
}
|
||||
|
||||
if ((uint64_t) throughput + pdata[19] >= max_nonce) {
|
||||
pdata[19] = max_nonce;
|
||||
break;
|
||||
}
|
||||
pdata[19] += throughput;
|
||||
|
||||
} while (!work_restart[thr_id].restart);
|
||||
|
||||
*hashes_done = pdata[19] - first_nonce;
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
|
1
miner.h
1
miner.h
@ -304,6 +304,7 @@ extern int scanhash_sha256d(int thr_id, struct work *work, uint32_t max_nonce, u
|
||||
extern int scanhash_sha256t(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_gost(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);
|
||||
extern int scanhash_s3(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done);
|
||||
|
Loading…
x
Reference in New Issue
Block a user