Skip to content

Commit

Permalink
make some variables local instead of global
Browse files Browse the repository at this point in the history
  • Loading branch information
@KlausT
KlausT committed Aug 31, 2015
1 parent cad5548 commit 42b44b7
Show file tree
Hide file tree
Showing 25 changed files with 362 additions and 390 deletions.
7 changes: 3 additions & 4 deletions Algo256/blake256.cu
View file
Original file line number Original file line Diff line number Diff line change
Expand Up @@ -500,8 +500,6 @@ void blake256_cpu_setBlock_16(int thr_id, uint32_t *penddata, const uint32_t *mi
} }
#endif #endif


static volatile bool init[MAX_GPUS] = { false };

extern int scanhash_blake256(int thr_id, uint32_t *pdata, uint32_t *ptarget, extern int scanhash_blake256(int thr_id, uint32_t *pdata, uint32_t *ptarget,
uint32_t max_nonce, uint32_t *hashes_done, int8_t blakerounds=14) uint32_t max_nonce, uint32_t *hashes_done, int8_t blakerounds=14)
{ {
Expand Down Expand Up @@ -531,15 +529,16 @@ extern int scanhash_blake256(int thr_id, uint32_t *pdata, uint32_t *ptarget,
pdata[k] = swab32(pdata[k]); pdata[k] = swab32(pdata[k]);
} }


if (!init[thr_id]) static THREAD volatile bool init = false;
if(!init)
{ {
CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id])); CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id]));
cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
cudaDeviceSetCacheConfig(cudaFuncCachePreferL1); cudaDeviceSetCacheConfig(cudaFuncCachePreferL1);
CUDA_SAFE_CALL(cudaStreamCreate(&gpustream[thr_id])); CUDA_SAFE_CALL(cudaStreamCreate(&gpustream[thr_id]));
CUDA_SAFE_CALL(cudaMallocHost(&h_resNonce, NBN * sizeof(uint32_t))); CUDA_SAFE_CALL(cudaMallocHost(&h_resNonce, NBN * sizeof(uint32_t)));
CUDA_SAFE_CALL(cudaMalloc(&d_resNonce[thr_id], NBN * sizeof(uint32_t))); CUDA_SAFE_CALL(cudaMalloc(&d_resNonce[thr_id], NBN * sizeof(uint32_t)));
init[thr_id] = true; init = true;
} }


#if PRECALC64 #if PRECALC64
Expand Down
12 changes: 6 additions & 6 deletions Algo256/keccak256.cu
View file
Original file line number Original file line Diff line number Diff line change
Expand Up @@ -14,8 +14,6 @@ extern "C"


#include "cuda_helper.h" #include "cuda_helper.h"


static THREAD uint32_t *h_nounce;

extern void keccak256_cpu_init(int thr_id, uint32_t threads); extern void keccak256_cpu_init(int thr_id, uint32_t threads);
extern void keccak256_setBlock_80(int thr_id, void *pdata,const void *ptarget); extern void keccak256_setBlock_80(int thr_id, void *pdata,const void *ptarget);
extern void keccak256_cpu_hash_80(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *h_nounce); extern void keccak256_cpu_hash_80(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *h_nounce);
Expand All @@ -34,12 +32,12 @@ extern "C" void keccak256_hash(void *state, const void *input)
memcpy(state, hash, 32); memcpy(state, hash, 32);
} }


static volatile bool init[MAX_GPUS] = { false };

extern int scanhash_keccak256(int thr_id, uint32_t *pdata, extern int scanhash_keccak256(int thr_id, uint32_t *pdata,
uint32_t *ptarget, uint32_t max_nonce, uint32_t *ptarget, uint32_t max_nonce,
uint32_t *hashes_done) uint32_t *hashes_done)
{ {
static THREAD uint32_t *h_nounce = nullptr;

const uint32_t first_nonce = pdata[19]; const uint32_t first_nonce = pdata[19];
uint32_t intensity = (device_sm[device_map[thr_id]] > 500) ? 1 << 28 : 1 << 27;; uint32_t intensity = (device_sm[device_map[thr_id]] > 500) ? 1 << 28 : 1 << 27;;
uint32_t throughput = device_intensity(device_map[thr_id], __func__, intensity); // 256*4096 uint32_t throughput = device_intensity(device_map[thr_id], __func__, intensity); // 256*4096
Expand All @@ -49,15 +47,17 @@ extern int scanhash_keccak256(int thr_id, uint32_t *pdata,
if (opt_benchmark) if (opt_benchmark)
ptarget[7] = 0x0002; ptarget[7] = 0x0002;


if (!init[thr_id]) { static THREAD volatile bool init = false;
if(!init)
{
CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id])); CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id]));
cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
cudaDeviceSetCacheConfig(cudaFuncCachePreferL1); cudaDeviceSetCacheConfig(cudaFuncCachePreferL1);
CUDA_SAFE_CALL(cudaStreamCreate(&gpustream[thr_id])); CUDA_SAFE_CALL(cudaStreamCreate(&gpustream[thr_id]));
CUDA_SAFE_CALL(cudaMallocHost(&h_nounce, 2 * sizeof(uint32_t))); CUDA_SAFE_CALL(cudaMallocHost(&h_nounce, 2 * sizeof(uint32_t)));
keccak256_cpu_init(thr_id, (int)throughput); keccak256_cpu_init(thr_id, (int)throughput);
// CUDA_SAFE_CALL(cudaMallocHost(&h_nounce, 2 * sizeof(uint32_t))); // CUDA_SAFE_CALL(cudaMallocHost(&h_nounce, 2 * sizeof(uint32_t)));
init[thr_id] = true; init = true;
} }


uint32_t endiandata[20]; uint32_t endiandata[20];
Expand Down
95 changes: 46 additions & 49 deletions JHA/jackpotcoin.cu
View file
Original file line number Original file line Diff line number Diff line change
Expand Up @@ -10,8 +10,6 @@ extern "C"
#include "miner.h" #include "miner.h"
#include "cuda_helper.h" #include "cuda_helper.h"


static uint32_t *d_hash[MAX_GPUS];

extern void jackpot_keccak512_cpu_init(int thr_id, uint32_t threads); extern void jackpot_keccak512_cpu_init(int thr_id, uint32_t threads);
extern void jackpot_keccak512_cpu_setBlock(int thr_id, void *pdata, size_t inlen); extern void jackpot_keccak512_cpu_setBlock(int thr_id, void *pdata, size_t inlen);
extern void jackpot_keccak512_cpu_hash(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_hash); extern void jackpot_keccak512_cpu_hash(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_hash);
Expand All @@ -33,12 +31,6 @@ extern void jackpot_compactTest_cpu_hash_64(int thr_id, uint32_t threads, uint32


extern uint32_t cuda_check_hash_branch(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_inputHash); extern uint32_t cuda_check_hash_branch(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_inputHash);


// Speicher zur Generierung der Noncevektoren für die bedingten Hashes
static uint32_t *d_jackpotNonces[MAX_GPUS];
static uint32_t *d_branch1Nonces[MAX_GPUS];
static uint32_t *d_branch2Nonces[MAX_GPUS];
static uint32_t *d_branch3Nonces[MAX_GPUS];

// Original jackpothash Funktion aus einem miner Quelltext // Original jackpothash Funktion aus einem miner Quelltext
extern "C" unsigned int jackpothash(void *state, const void *input) extern "C" unsigned int jackpothash(void *state, const void *input)
{ {
Expand Down Expand Up @@ -82,8 +74,6 @@ extern "C" unsigned int jackpothash(void *state, const void *input)
return round; return round;
} }


static volatile bool init[MAX_GPUS] = { false };

extern int scanhash_jackpot(int thr_id, uint32_t *pdata, extern int scanhash_jackpot(int thr_id, uint32_t *pdata,
uint32_t *ptarget, uint32_t max_nonce, uint32_t *ptarget, uint32_t max_nonce,
uint32_t *hashes_done) uint32_t *hashes_done)
Expand All @@ -96,29 +86,36 @@ extern int scanhash_jackpot(int thr_id, uint32_t *pdata,
if (opt_benchmark) if (opt_benchmark)
ptarget[7] = 0x000f; ptarget[7] = 0x000f;


if (!init[thr_id]) static THREAD uint32_t *d_hash = nullptr;
static THREAD uint32_t *d_jackpotNonces = nullptr;
static THREAD uint32_t *d_branch1Nonces = nullptr;
static THREAD uint32_t *d_branch2Nonces = nullptr;
static THREAD uint32_t *d_branch3Nonces = nullptr;
static THREAD volatile bool init = false;

if (!init)
{ {
CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id])); CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id]));
cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
cudaDeviceSetCacheConfig(cudaFuncCachePreferL1); cudaDeviceSetCacheConfig(cudaFuncCachePreferL1);
CUDA_SAFE_CALL(cudaStreamCreate(&gpustream[thr_id])); CUDA_SAFE_CALL(cudaStreamCreate(&gpustream[thr_id]));
get_cuda_arch(&cuda_arch[thr_id]); get_cuda_arch(&cuda_arch[thr_id]);


CUDA_SAFE_CALL(cudaMalloc(&d_hash[thr_id], 16 * sizeof(uint32_t) * throughput)); CUDA_SAFE_CALL(cudaMalloc(&d_hash, 16 * sizeof(uint32_t) * throughput));


jackpot_keccak512_cpu_init(thr_id, throughput); jackpot_keccak512_cpu_init(thr_id, throughput);
jackpot_compactTest_cpu_init(thr_id, throughput); jackpot_compactTest_cpu_init(thr_id, throughput);
quark_groestl512_cpu_init(thr_id, throughput); quark_groestl512_cpu_init(thr_id, throughput);
quark_skein512_cpu_init(thr_id); quark_skein512_cpu_init(thr_id);
cuda_check_cpu_init(thr_id, throughput); cuda_check_cpu_init(thr_id, throughput);


cudaMalloc(&d_branch1Nonces[thr_id], sizeof(uint32_t)*throughput*2); cudaMalloc(&d_branch1Nonces, sizeof(uint32_t)*throughput*2);
cudaMalloc(&d_branch2Nonces[thr_id], sizeof(uint32_t)*throughput*2); cudaMalloc(&d_branch2Nonces, sizeof(uint32_t)*throughput*2);
cudaMalloc(&d_branch3Nonces[thr_id], sizeof(uint32_t)*throughput*2); cudaMalloc(&d_branch3Nonces, sizeof(uint32_t)*throughput*2);


CUDA_SAFE_CALL(cudaMalloc(&d_jackpotNonces[thr_id], sizeof(uint32_t)*throughput*2)); CUDA_SAFE_CALL(cudaMalloc(&d_jackpotNonces, sizeof(uint32_t)*throughput*2));


init[thr_id] = true; init = true;
} }


uint32_t endiandata[22]; uint32_t endiandata[22];
Expand All @@ -130,74 +127,74 @@ extern int scanhash_jackpot(int thr_id, uint32_t *pdata,


do { do {
// erstes Keccak512 Hash mit CUDA // erstes Keccak512 Hash mit CUDA
jackpot_keccak512_cpu_hash(thr_id, throughput, pdata[19], d_hash[thr_id]); jackpot_keccak512_cpu_hash(thr_id, throughput, pdata[19], d_hash);


uint32_t nrm1, nrm2, nrm3; uint32_t nrm1, nrm2, nrm3;


// Runde 1 (ohne Gröstl) // Runde 1 (ohne Gröstl)


jackpot_compactTest_cpu_hash_64(thr_id, throughput, pdata[19], d_hash[thr_id], NULL, jackpot_compactTest_cpu_hash_64(thr_id, throughput, pdata[19], d_hash, NULL,
d_branch1Nonces[thr_id], &nrm1, d_branch1Nonces, &nrm1,
d_branch3Nonces[thr_id], &nrm3); d_branch3Nonces, &nrm3);


// verfolge den skein-pfad weiter // verfolge den skein-pfad weiter
quark_skein512_cpu_hash_64(thr_id, nrm3, pdata[19], d_branch3Nonces[thr_id], d_hash[thr_id]); quark_skein512_cpu_hash_64(thr_id, nrm3, pdata[19], d_branch3Nonces, d_hash);


// noch schnell Blake & JH // noch schnell Blake & JH
jackpot_compactTest_cpu_hash_64(thr_id, nrm3, pdata[19], d_hash[thr_id], d_branch3Nonces[thr_id], jackpot_compactTest_cpu_hash_64(thr_id, nrm3, pdata[19], d_hash, d_branch3Nonces,
d_branch1Nonces[thr_id], &nrm1, d_branch1Nonces, &nrm1,
d_branch2Nonces[thr_id], &nrm2); d_branch2Nonces, &nrm2);


if (nrm1+nrm2 == nrm3) { if (nrm1+nrm2 == nrm3) {
quark_blake512_cpu_hash_64(thr_id, nrm1, pdata[19], d_branch1Nonces[thr_id], d_hash[thr_id]); quark_blake512_cpu_hash_64(thr_id, nrm1, pdata[19], d_branch1Nonces, d_hash);
quark_jh512_cpu_hash_64(thr_id, nrm2, pdata[19], d_branch2Nonces[thr_id], d_hash[thr_id]); quark_jh512_cpu_hash_64(thr_id, nrm2, pdata[19], d_branch2Nonces, d_hash);
} }


// Runde 3 (komplett) // Runde 3 (komplett)


// jackpotNonces in branch1/2 aufsplitten gemäss if (hash[0] & 0x01) // jackpotNonces in branch1/2 aufsplitten gemäss if (hash[0] & 0x01)
jackpot_compactTest_cpu_hash_64(thr_id, nrm3, pdata[19], d_hash[thr_id], d_branch3Nonces[thr_id], jackpot_compactTest_cpu_hash_64(thr_id, nrm3, pdata[19], d_hash, d_branch3Nonces,
d_branch1Nonces[thr_id], &nrm1, d_branch1Nonces, &nrm1,
d_branch2Nonces[thr_id], &nrm2); d_branch2Nonces, &nrm2);


if (nrm1+nrm2 == nrm3) { if (nrm1+nrm2 == nrm3) {
quark_groestl512_cpu_hash_64(thr_id, nrm1, pdata[19], d_branch1Nonces[thr_id], d_hash[thr_id]); quark_groestl512_cpu_hash_64(thr_id, nrm1, pdata[19], d_branch1Nonces, d_hash);
quark_skein512_cpu_hash_64(thr_id, nrm2, pdata[19], d_branch2Nonces[thr_id], d_hash[thr_id]); quark_skein512_cpu_hash_64(thr_id, nrm2, pdata[19], d_branch2Nonces, d_hash);
} }


// jackpotNonces in branch1/2 aufsplitten gemäss if (hash[0] & 0x01) // jackpotNonces in branch1/2 aufsplitten gemäss if (hash[0] & 0x01)
jackpot_compactTest_cpu_hash_64(thr_id, nrm3, pdata[19], d_hash[thr_id], d_branch3Nonces[thr_id], jackpot_compactTest_cpu_hash_64(thr_id, nrm3, pdata[19], d_hash, d_branch3Nonces,
d_branch1Nonces[thr_id], &nrm1, d_branch1Nonces, &nrm1,
d_branch2Nonces[thr_id], &nrm2); d_branch2Nonces, &nrm2);


if (nrm1+nrm2 == nrm3) { if (nrm1+nrm2 == nrm3) {
quark_blake512_cpu_hash_64(thr_id, nrm1, pdata[19], d_branch1Nonces[thr_id], d_hash[thr_id]); quark_blake512_cpu_hash_64(thr_id, nrm1, pdata[19], d_branch1Nonces, d_hash);
quark_jh512_cpu_hash_64(thr_id, nrm2, pdata[19], d_branch2Nonces[thr_id], d_hash[thr_id]); quark_jh512_cpu_hash_64(thr_id, nrm2, pdata[19], d_branch2Nonces, d_hash);
} }


// Runde 3 (komplett) // Runde 3 (komplett)


// jackpotNonces in branch1/2 aufsplitten gemäss if (hash[0] & 0x01) // jackpotNonces in branch1/2 aufsplitten gemäss if (hash[0] & 0x01)
jackpot_compactTest_cpu_hash_64(thr_id, nrm3, pdata[19], d_hash[thr_id], d_branch3Nonces[thr_id], jackpot_compactTest_cpu_hash_64(thr_id, nrm3, pdata[19], d_hash, d_branch3Nonces,
d_branch1Nonces[thr_id], &nrm1, d_branch1Nonces, &nrm1,
d_branch2Nonces[thr_id], &nrm2); d_branch2Nonces, &nrm2);


if (nrm1+nrm2 == nrm3) { if (nrm1+nrm2 == nrm3) {
quark_groestl512_cpu_hash_64(thr_id, nrm1, pdata[19], d_branch1Nonces[thr_id], d_hash[thr_id]); quark_groestl512_cpu_hash_64(thr_id, nrm1, pdata[19], d_branch1Nonces, d_hash);
quark_skein512_cpu_hash_64(thr_id, nrm2, pdata[19], d_branch2Nonces[thr_id], d_hash[thr_id]); quark_skein512_cpu_hash_64(thr_id, nrm2, pdata[19], d_branch2Nonces, d_hash);
} }


// jackpotNonces in branch1/2 aufsplitten gemäss if (hash[0] & 0x01) // jackpotNonces in branch1/2 aufsplitten gemäss if (hash[0] & 0x01)
jackpot_compactTest_cpu_hash_64(thr_id, nrm3, pdata[19], d_hash[thr_id], d_branch3Nonces[thr_id], jackpot_compactTest_cpu_hash_64(thr_id, nrm3, pdata[19], d_hash, d_branch3Nonces,
d_branch1Nonces[thr_id], &nrm1, d_branch1Nonces, &nrm1,
d_branch2Nonces[thr_id], &nrm2); d_branch2Nonces, &nrm2);


if (nrm1+nrm2 == nrm3) { if (nrm1+nrm2 == nrm3) {
quark_blake512_cpu_hash_64(thr_id, nrm1, pdata[19], d_branch1Nonces[thr_id], d_hash[thr_id]); quark_blake512_cpu_hash_64(thr_id, nrm1, pdata[19], d_branch1Nonces, d_hash);
quark_jh512_cpu_hash_64(thr_id, nrm2, pdata[19], d_branch2Nonces[thr_id], d_hash[thr_id]); quark_jh512_cpu_hash_64(thr_id, nrm2, pdata[19], d_branch2Nonces, d_hash);
} }


uint32_t foundNonce = cuda_check_hash_branch(thr_id, nrm3, pdata[19], d_branch3Nonces[thr_id], d_hash[thr_id]); uint32_t foundNonce = cuda_check_hash_branch(thr_id, nrm3, pdata[19], d_branch3Nonces, d_hash);
if(stop_mining) {mining_has_stopped[thr_id] = true; cudaStreamDestroy(gpustream[thr_id]); pthread_exit(nullptr);} if(stop_mining) {mining_has_stopped[thr_id] = true; cudaStreamDestroy(gpustream[thr_id]); pthread_exit(nullptr);}
if(foundNonce != 0xffffffff) if(foundNonce != 0xffffffff)
{ {
Expand All @@ -211,7 +208,7 @@ extern int scanhash_jackpot(int thr_id, uint32_t *pdata,


if (vhash64[7] <= Htarg && fulltest(vhash64, ptarget)) { if (vhash64[7] <= Htarg && fulltest(vhash64, ptarget)) {
int res = 1; int res = 1;
uint32_t secNonce = cuda_check_hash_suppl(thr_id, throughput, pdata[19], d_hash[thr_id], foundNonce); uint32_t secNonce = cuda_check_hash_suppl(thr_id, throughput, pdata[19], d_hash, foundNonce);
*hashes_done = pdata[19] - first_nonce + throughput; *hashes_done = pdata[19] - first_nonce + throughput;
if (secNonce != 0) if (secNonce != 0)
{ {
Expand Down
10 changes: 5 additions & 5 deletions bitcoin.cu
View file
Original file line number Original file line Diff line number Diff line change
@@ -1,8 +1,6 @@
#include "miner.h" #include "miner.h"
#include "cuda_helper.h" #include "cuda_helper.h"


static THREAD uint32_t *h_nounce;

extern void bitcoin_cpu_init(int thr_id); extern void bitcoin_cpu_init(int thr_id);
extern void bitcoin_cpu_hash(int thr_id, uint32_t threads, uint32_t startNounce, const uint32_t *const ms, uint32_t merkle, uint32_t time, uint32_t compacttarget, uint32_t *const h_nounce); extern void bitcoin_cpu_hash(int thr_id, uint32_t threads, uint32_t startNounce, const uint32_t *const ms, uint32_t merkle, uint32_t time, uint32_t compacttarget, uint32_t *const h_nounce);
extern void bitcoin_midstate(const uint32_t *data, uint32_t *midstate); extern void bitcoin_midstate(const uint32_t *data, uint32_t *midstate);
Expand Down Expand Up @@ -122,20 +120,22 @@ void bitcoin_hash(uint32_t *output, const uint32_t *data, uint32_t nonce, const
be32enc(&output[7], h + hc[7]); be32enc(&output[7], h + hc[7]);
} }


static volatile bool init[MAX_GPUS] = { false };


int scanhash_bitcoin(int thr_id, uint32_t *pdata, int scanhash_bitcoin(int thr_id, uint32_t *pdata,
uint32_t *ptarget, uint32_t max_nonce, uint32_t *ptarget, uint32_t max_nonce,
uint32_t *hashes_done) uint32_t *hashes_done)
{ {
static THREAD uint32_t *h_nounce = nullptr;

const uint32_t first_nonce = pdata[19]; const uint32_t first_nonce = pdata[19];
uint32_t throughput = device_intensity(device_map[thr_id], __func__, 1U << 28); uint32_t throughput = device_intensity(device_map[thr_id], __func__, 1U << 28);
throughput = min(throughput, (max_nonce - first_nonce)) & 0xfffffc00; throughput = min(throughput, (max_nonce - first_nonce)) & 0xfffffc00;


if (opt_benchmark) if (opt_benchmark)
ptarget[7] = 0x0005; ptarget[7] = 0x0005;


if (!init[thr_id]) static THREAD volatile bool init = false;
if(!init)
{ {
CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id])); CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id]));
cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
Expand All @@ -144,7 +144,7 @@ int scanhash_bitcoin(int thr_id, uint32_t *pdata,
CUDA_SAFE_CALL(cudaStreamCreate(&gpustream[thr_id])); CUDA_SAFE_CALL(cudaStreamCreate(&gpustream[thr_id]));
bitcoin_cpu_init(thr_id); bitcoin_cpu_init(thr_id);
CUDA_SAFE_CALL(cudaMallocHost(&h_nounce, 2 * sizeof(uint32_t))); CUDA_SAFE_CALL(cudaMallocHost(&h_nounce, 2 * sizeof(uint32_t)));
init[thr_id] = true; init = true;
} }


uint32_t ms[8]; uint32_t ms[8];
Expand Down
16 changes: 8 additions & 8 deletions bitcredit/bitcredit.cu
View file
Original file line number Original file line Diff line number Diff line change
Expand Up @@ -6,9 +6,6 @@ extern "C"
#include "miner.h" #include "miner.h"
#include "cuda_helper.h" #include "cuda_helper.h"


static uint32_t *d_hash[MAX_GPUS];
static uint32_t *foundNonce;

extern void bitcredit_setBlockTarget(int thr_id, uint32_t * data, const uint32_t * midstate, const void *ptarget); extern void bitcredit_setBlockTarget(int thr_id, uint32_t * data, const uint32_t * midstate, const void *ptarget);
extern void bitcredit_cpu_init(int thr_id, uint32_t threads, uint32_t* hash); extern void bitcredit_cpu_init(int thr_id, uint32_t threads, uint32_t* hash);
extern void bitcredit_cpu_hash(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *result); extern void bitcredit_cpu_hash(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *result);
Expand Down Expand Up @@ -43,17 +40,20 @@ int scanhash_bitcredit(int thr_id, uint32_t *pdata,
int intensity = 256 * 256 * 64 * 8; int intensity = 256 * 256 * 64 * 8;
const uint32_t throughput = min(device_intensity(device_map[thr_id], __func__, intensity), (max_nonce - first_nonce)) & 0xfffffc00; // 19=256*256*8; const uint32_t throughput = min(device_intensity(device_map[thr_id], __func__, intensity), (max_nonce - first_nonce)) & 0xfffffc00; // 19=256*256*8;


static bool init[MAX_GPUS] = { false }; static THREAD uint32_t *d_hash = nullptr;
if(!init[thr_id]) static THREAD uint32_t *foundNonce = nullptr;

static THREAD bool init = false;
if(!init)
{ {
CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id])); CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id]));
cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
cudaDeviceSetCacheConfig(cudaFuncCachePreferL1); cudaDeviceSetCacheConfig(cudaFuncCachePreferL1);


CUDA_SAFE_CALL(cudaMalloc(&d_hash[thr_id], 8 * sizeof(uint32_t) * throughput)); CUDA_SAFE_CALL(cudaMalloc(&d_hash, 8 * sizeof(uint32_t) * throughput));
CUDA_SAFE_CALL(cudaMallocHost(&foundNonce, 2 * sizeof(uint32_t))); CUDA_SAFE_CALL(cudaMallocHost(&foundNonce, 2 * sizeof(uint32_t)));
bitcredit_cpu_init(thr_id, throughput, d_hash[thr_id]); bitcredit_cpu_init(thr_id, throughput, d_hash);
init[thr_id] = true; init = true;
} }


uint32_t endiandata[42]; uint32_t endiandata[42];
Expand Down
Loading

0 comments on commit 42b44b7

Please sign in to comment.