mirror of
https://github.com/GOSTSec/ccminer
synced 2025-01-09 14:28:15 +00:00
9eead77027
This will allow later more gpu candidates. Note: This is an unfinished work, we keep the previous behavior for now To finish this, all algos solutions should be migrated and submitted nonces attributes stored. Its required to handle the different share diff per nonce and fix the possible solved count error (if 1/2 nonces is solved).
233 lines
6.9 KiB
Plaintext
233 lines
6.9 KiB
Plaintext
extern "C" {
|
|
#include "sph/sph_blake.h"
|
|
#include "sph/sph_bmw.h"
|
|
#include "sph/sph_groestl.h"
|
|
#include "sph/sph_skein.h"
|
|
#include "sph/sph_jh.h"
|
|
#include "sph/sph_keccak.h"
|
|
#include "sph/sph_luffa.h"
|
|
#include "sph/sph_cubehash.h"
|
|
#include "sph/sph_shavite.h"
|
|
#include "sph/sph_simd.h"
|
|
#include "sph/sph_echo.h"
|
|
}
|
|
|
|
#include "miner.h"
|
|
#include "cuda_helper.h"
|
|
#include "cuda_x11.h"
|
|
|
|
#include <stdio.h>
|
|
#include <memory.h>
|
|
|
|
static uint32_t *d_hash[MAX_GPUS];
|
|
|
|
// X11 CPU Hash
|
|
extern "C" void x11hash(void *output, const void *input)
|
|
{
|
|
unsigned char _ALIGN(128) hash[128] = { 0 };
|
|
|
|
// blake1-bmw2-grs3-skein4-jh5-keccak6-luffa7-cubehash8-shavite9-simd10-echo11
|
|
|
|
sph_blake512_context ctx_blake;
|
|
sph_bmw512_context ctx_bmw;
|
|
sph_groestl512_context ctx_groestl;
|
|
sph_jh512_context ctx_jh;
|
|
sph_keccak512_context ctx_keccak;
|
|
sph_skein512_context ctx_skein;
|
|
sph_luffa512_context ctx_luffa;
|
|
sph_cubehash512_context ctx_cubehash;
|
|
sph_shavite512_context ctx_shavite;
|
|
sph_simd512_context ctx_simd;
|
|
sph_echo512_context ctx_echo;
|
|
|
|
sph_blake512_init(&ctx_blake);
|
|
sph_blake512 (&ctx_blake, input, 80);
|
|
sph_blake512_close(&ctx_blake, (void*) hash);
|
|
|
|
sph_bmw512_init(&ctx_bmw);
|
|
sph_bmw512 (&ctx_bmw, (const void*) hash, 64);
|
|
sph_bmw512_close(&ctx_bmw, (void*) hash);
|
|
|
|
sph_groestl512_init(&ctx_groestl);
|
|
sph_groestl512 (&ctx_groestl, (const void*) hash, 64);
|
|
sph_groestl512_close(&ctx_groestl, (void*) hash);
|
|
|
|
sph_skein512_init(&ctx_skein);
|
|
sph_skein512 (&ctx_skein, (const void*) hash, 64);
|
|
sph_skein512_close(&ctx_skein, (void*) hash);
|
|
|
|
sph_jh512_init(&ctx_jh);
|
|
sph_jh512 (&ctx_jh, (const void*) hash, 64);
|
|
sph_jh512_close(&ctx_jh, (void*) hash);
|
|
|
|
sph_keccak512_init(&ctx_keccak);
|
|
sph_keccak512 (&ctx_keccak, (const void*) hash, 64);
|
|
sph_keccak512_close(&ctx_keccak, (void*) hash);
|
|
|
|
sph_luffa512_init(&ctx_luffa);
|
|
sph_luffa512 (&ctx_luffa, (const void*) hash, 64);
|
|
sph_luffa512_close (&ctx_luffa, (void*) hash);
|
|
|
|
sph_cubehash512_init(&ctx_cubehash);
|
|
sph_cubehash512 (&ctx_cubehash, (const void*) hash, 64);
|
|
sph_cubehash512_close(&ctx_cubehash, (void*) hash);
|
|
|
|
sph_shavite512_init(&ctx_shavite);
|
|
sph_shavite512 (&ctx_shavite, (const void*) hash, 64);
|
|
sph_shavite512_close(&ctx_shavite, (void*) hash);
|
|
|
|
sph_simd512_init(&ctx_simd);
|
|
sph_simd512 (&ctx_simd, (const void*) hash, 64);
|
|
sph_simd512_close(&ctx_simd, (void*) hash);
|
|
|
|
sph_echo512_init(&ctx_echo);
|
|
sph_echo512 (&ctx_echo, (const void*) hash, 64);
|
|
sph_echo512_close(&ctx_echo, (void*) hash);
|
|
|
|
memcpy(output, hash, 32);
|
|
}
|
|
|
|
//#define _DEBUG
|
|
#define _DEBUG_PREFIX "x11"
|
|
#include "cuda_debug.cuh"
|
|
|
|
static bool init[MAX_GPUS] = { 0 };
|
|
|
|
extern "C" int scanhash_x11(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];
|
|
int intensity = (device_sm[device_map[thr_id]] >= 500 && !is_windows()) ? 20 : 19;
|
|
uint32_t throughput = cuda_default_throughput(thr_id, 1U << intensity); // 19=256*256*8;
|
|
//if (init[thr_id]) throughput = min(throughput, max_nonce - first_nonce);
|
|
|
|
if (opt_benchmark)
|
|
ptarget[7] = 0x5;
|
|
|
|
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);
|
|
|
|
quark_blake512_cpu_init(thr_id, throughput);
|
|
quark_bmw512_cpu_init(thr_id, throughput);
|
|
quark_groestl512_cpu_init(thr_id, throughput);
|
|
quark_skein512_cpu_init(thr_id, throughput);
|
|
quark_keccak512_cpu_init(thr_id, throughput);
|
|
quark_jh512_cpu_init(thr_id, throughput);
|
|
x11_luffaCubehash512_cpu_init(thr_id, throughput);
|
|
x11_shavite512_cpu_init(thr_id, throughput);
|
|
x11_echo512_cpu_init(thr_id, throughput);
|
|
if (x11_simd512_cpu_init(thr_id, throughput) != 0) {
|
|
return 0;
|
|
}
|
|
CUDA_CALL_OR_RET_X(cudaMalloc(&d_hash[thr_id], (size_t) 64 * throughput), 0);
|
|
|
|
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]);
|
|
|
|
quark_blake512_cpu_setBlock_80(thr_id, endiandata);
|
|
cuda_check_cpu_setTarget(ptarget);
|
|
|
|
do {
|
|
int order = 0;
|
|
uint32_t foundNonce;
|
|
|
|
// Hash with CUDA
|
|
quark_blake512_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id]); order++;
|
|
TRACE("blake :");
|
|
quark_bmw512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++);
|
|
TRACE("bmw :");
|
|
quark_groestl512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++);
|
|
TRACE("groestl:");
|
|
quark_skein512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++);
|
|
TRACE("skein :");
|
|
quark_jh512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++);
|
|
TRACE("jh512 :");
|
|
quark_keccak512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++);
|
|
TRACE("keccak :");
|
|
x11_luffaCubehash512_cpu_hash_64(thr_id, throughput, d_hash[thr_id], order++);
|
|
TRACE("luffa+c:");
|
|
x11_shavite512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++);
|
|
TRACE("shavite:");
|
|
x11_simd512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++);
|
|
TRACE("simd :");
|
|
x11_echo512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++);
|
|
TRACE("echo => ");
|
|
|
|
foundNonce = cuda_check_hash(thr_id, throughput, pdata[19], d_hash[thr_id]);
|
|
if (foundNonce != UINT32_MAX)
|
|
{
|
|
const uint32_t Htarg = ptarget[7];
|
|
uint32_t vhash64[8];
|
|
be32enc(&endiandata[19], foundNonce);
|
|
x11hash(vhash64, endiandata);
|
|
|
|
if (vhash64[7] <= Htarg && fulltest(vhash64, ptarget)) {
|
|
int res = 1;
|
|
// check if there was some other ones...
|
|
uint32_t secNonce = cuda_check_hash_suppl(thr_id, throughput, pdata[19], d_hash[thr_id], 1);
|
|
work_set_target_ratio(work, vhash64);
|
|
*hashes_done = pdata[19] - first_nonce + throughput;
|
|
if (secNonce != 0) {
|
|
be32enc(&endiandata[19], secNonce);
|
|
x11hash(vhash64, endiandata);
|
|
if (bn_hash_target_ratio(vhash64, ptarget) > work->shareratio[0])
|
|
work_set_target_ratio(work, vhash64);
|
|
pdata[21] = secNonce;
|
|
res++;
|
|
}
|
|
pdata[19] = foundNonce;
|
|
return res;
|
|
} else {
|
|
gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNonce);
|
|
pdata[19] = foundNonce + 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;
|
|
}
|
|
|
|
// cleanup
|
|
extern "C" void free_x11(int thr_id)
|
|
{
|
|
if (!init[thr_id])
|
|
return;
|
|
|
|
cudaThreadSynchronize();
|
|
|
|
cudaFree(d_hash[thr_id]);
|
|
|
|
quark_blake512_cpu_free(thr_id);
|
|
quark_groestl512_cpu_free(thr_id);
|
|
x11_simd512_cpu_free(thr_id);
|
|
|
|
cuda_check_cpu_free(thr_id);
|
|
init[thr_id] = false;
|
|
|
|
cudaDeviceSynchronize();
|
|
}
|