summaryrefslogtreecommitdiffstats
path: root/xmrstak/backend
diff options
context:
space:
mode:
Diffstat (limited to 'xmrstak/backend')
-rw-r--r--xmrstak/backend/amd/amd_gpu/gpu.cpp6
-rw-r--r--xmrstak/backend/amd/minethd.cpp29
-rw-r--r--xmrstak/backend/cpu/minethd.cpp60
-rw-r--r--xmrstak/backend/cpu/minethd.hpp2
-rw-r--r--xmrstak/backend/nvidia/minethd.cpp33
-rw-r--r--xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp6
-rw-r--r--xmrstak/backend/nvidia/nvcc_code/cuda_core.cu19
-rw-r--r--xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu8
8 files changed, 73 insertions, 90 deletions
diff --git a/xmrstak/backend/amd/amd_gpu/gpu.cpp b/xmrstak/backend/amd/amd_gpu/gpu.cpp
index 4a82fab..2973db4 100644
--- a/xmrstak/backend/amd/amd_gpu/gpu.cpp
+++ b/xmrstak/backend/amd/amd_gpu/gpu.cpp
@@ -940,7 +940,7 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar
/// @todo only activate if currency is monero
int cn_kernel_offset = 0;
- if((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon ) && version >= 7)
+ if(miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon)
{
cn_kernel_offset = 6;
}
@@ -966,7 +966,7 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar
return(ERR_OCL_API);
}
- if((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon ) && version >= 7)
+ if(miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon )
{
// Input
if ((ret = clSetKernelArg(ctx->Kernels[1 + cn_kernel_offset], 3, sizeof(cl_mem), &ctx->InputBuffer)) != CL_SUCCESS)
@@ -1134,7 +1134,7 @@ size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput, xmrstak_algo miner_algo,
size_t tmpNonce = ctx->Nonce;
/// @todo only activate if currency is monero
int cn_kernel_offset = 0;
- if((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) && version >= 7)
+ if(miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon)
{
cn_kernel_offset = 6;
}
diff --git a/xmrstak/backend/amd/minethd.cpp b/xmrstak/backend/amd/minethd.cpp
index 775aa08..9357f9a 100644
--- a/xmrstak/backend/amd/minethd.cpp
+++ b/xmrstak/backend/amd/minethd.cpp
@@ -191,10 +191,10 @@ void minethd::work_main()
uint64_t iCount = 0;
cryptonight_ctx* cpu_ctx;
cpu_ctx = cpu::minethd::minethd_alloc_ctx();
- auto miner_algo = ::jconf::inst()->GetMiningAlgo();
// start with root algorithm and switch later if fork version is reached
- cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, ::jconf::inst()->GetMiningAlgoRoot());
+ auto miner_algo = ::jconf::inst()->GetMiningAlgoRoot();
+ cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
globalStates::inst().iConsumeCnt++;
@@ -213,16 +213,18 @@ void minethd::work_main()
std::this_thread::sleep_for(std::chrono::milliseconds(100));
consume_work();
- uint8_t new_version = oWork.getVersion();
- const bool time_to_fork =
- ((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) && version < 7 && new_version >= 7) ||
- (miner_algo == cryptonight_heavy && version < 3 && new_version >= 3);
- if(time_to_fork)
+ continue;
+ }
+
+ uint8_t new_version = oWork.getVersion();
+ if(new_version != version)
+ {
+ if(new_version >= ::jconf::inst()->GetMiningForkHeight())
{
+ miner_algo = ::jconf::inst()->GetMiningAlgo();
hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
}
version = new_version;
- continue;
}
uint32_t h_per_round = pGpuCtx->rawIntensity;
@@ -230,7 +232,7 @@ void minethd::work_main()
assert(sizeof(job_result::sJobID) == sizeof(pool_job::sJobID));
uint64_t target = oWork.iTarget;
- /// \todo add monero hard for version
+
XMRSetJob(pGpuCtx, oWork.bWorkBlob, oWork.iWorkSize, target, miner_algo, version);
if(oWork.bNiceHash)
@@ -274,15 +276,6 @@ void minethd::work_main()
}
consume_work();
- uint8_t new_version = oWork.getVersion();
- const bool time_to_fork =
- ((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) && version < 7 && new_version >= 7) ||
- (miner_algo == cryptonight_heavy && version < 3 && new_version >= 3);
- if(time_to_fork)
- {
- hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
- }
- version = new_version;
}
}
diff --git a/xmrstak/backend/cpu/minethd.cpp b/xmrstak/backend/cpu/minethd.cpp
index b25d91e..dd303b5 100644
--- a/xmrstak/backend/cpu/minethd.cpp
+++ b/xmrstak/backend/cpu/minethd.cpp
@@ -413,14 +413,15 @@ void minethd::work_main()
lck.release();
std::this_thread::yield();
- cn_hash_fun hash_fun;
cryptonight_ctx* ctx;
uint64_t iCount = 0;
uint64_t* piHashVal;
uint32_t* piNonce;
job_result result;
- hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->GetMiningAlgo());
+ // start with root algorithm and switch later if fork version is reached
+ auto miner_algo = ::jconf::inst()->GetMiningAlgoRoot();
+ cn_hash_fun hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo);
ctx = minethd_alloc_ctx();
piHashVal = (uint64_t*)(result.bResult + 24);
@@ -428,6 +429,8 @@ void minethd::work_main()
globalStates::inst().inst().iConsumeCnt++;
result.iThreadId = iThreadNo;
+ uint8_t version = 0;
+
while (bQuit == 0)
{
if (oWork.bStall)
@@ -453,15 +456,16 @@ void minethd::work_main()
if(oWork.bNiceHash)
result.iNonce = *piNonce;
- auto miner_algo = ::jconf::inst()->GetMiningAlgo();
- const bool time_to_fork =
- ((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) && oWork.bWorkBlob[0] >= 7) ||
- (miner_algo == cryptonight_heavy && oWork.bWorkBlob[0] >= 3);
-
- if(time_to_fork)
- hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo);
- else
- hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->GetMiningAlgoRoot());
+ uint8_t new_version = oWork.getVersion();
+ if(new_version != version)
+ {
+ if(new_version >= ::jconf::inst()->GetMiningForkHeight())
+ {
+ miner_algo = ::jconf::inst()->GetMiningAlgo();
+ hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo);
+ }
+ version = new_version;
+ }
while(globalStates::inst().iGlobalJobNo.load(std::memory_order_relaxed) == iJobNo)
{
@@ -619,22 +623,22 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes,
void minethd::double_work_main()
{
- multiway_work_main<2>(func_multi_selector(2, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->GetMiningAlgo()));
+ multiway_work_main<2>();
}
void minethd::triple_work_main()
{
- multiway_work_main<3>(func_multi_selector(3, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->GetMiningAlgo()));
+ multiway_work_main<3>();
}
void minethd::quad_work_main()
{
- multiway_work_main<4>(func_multi_selector(4, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->GetMiningAlgo()));
+ multiway_work_main<4>();
}
void minethd::penta_work_main()
{
- multiway_work_main<5>(func_multi_selector(5, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->GetMiningAlgo()));
+ multiway_work_main<5>();
}
template<size_t N>
@@ -649,7 +653,7 @@ void minethd::prep_multiway_work(uint8_t *bWorkBlob, uint32_t **piNonce)
}
template<size_t N>
-void minethd::multiway_work_main(cn_hash_fun_multi hash_fun_multi)
+void minethd::multiway_work_main()
{
if(affinity >= 0) //-1 means no affinity
bindMemoryToNUMANode(affinity);
@@ -680,6 +684,11 @@ void minethd::multiway_work_main(cn_hash_fun_multi hash_fun_multi)
globalStates::inst().iConsumeCnt++;
+ // start with root algorithm and switch later if fork version is reached
+ auto miner_algo = ::jconf::inst()->GetMiningAlgoRoot();
+ cn_hash_fun_multi hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo);
+ uint8_t version = 0;
+
while (bQuit == 0)
{
if (oWork.bStall)
@@ -704,15 +713,16 @@ void minethd::multiway_work_main(cn_hash_fun_multi hash_fun_multi)
if(oWork.bNiceHash)
iNonce = *piNonce[0];
- auto miner_algo = ::jconf::inst()->GetMiningAlgo();
- const bool time_to_fork =
- ((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) && oWork.bWorkBlob[0] >= 7) ||
- (miner_algo == cryptonight_heavy && oWork.bWorkBlob[0] >= 3);
-
- if(time_to_fork)
- hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo);
- else
- hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->GetMiningAlgoRoot());
+ uint8_t new_version = oWork.getVersion();
+ if(new_version != version)
+ {
+ if(new_version >= ::jconf::inst()->GetMiningForkHeight())
+ {
+ miner_algo = ::jconf::inst()->GetMiningAlgo();
+ hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo);
+ }
+ version = new_version;
+ }
while (globalStates::inst().iGlobalJobNo.load(std::memory_order_relaxed) == iJobNo)
{
diff --git a/xmrstak/backend/cpu/minethd.hpp b/xmrstak/backend/cpu/minethd.hpp
index ef1bbd2..cd78343 100644
--- a/xmrstak/backend/cpu/minethd.hpp
+++ b/xmrstak/backend/cpu/minethd.hpp
@@ -36,7 +36,7 @@ private:
minethd(miner_work& pWork, size_t iNo, int iMultiway, bool no_prefetch, int64_t affinity);
template<size_t N>
- void multiway_work_main(cn_hash_fun_multi hash_fun_multi);
+ void multiway_work_main();
template<size_t N>
void prep_multiway_work(uint8_t *bWorkBlob, uint32_t **piNonce);
diff --git a/xmrstak/backend/nvidia/minethd.cpp b/xmrstak/backend/nvidia/minethd.cpp
index 79297b0..2e29d64 100644
--- a/xmrstak/backend/nvidia/minethd.cpp
+++ b/xmrstak/backend/nvidia/minethd.cpp
@@ -237,10 +237,10 @@ void minethd::work_main()
uint64_t iCount = 0;
cryptonight_ctx* cpu_ctx;
cpu_ctx = cpu::minethd::minethd_alloc_ctx();
- auto miner_algo = ::jconf::inst()->GetMiningAlgo();
// start with root algorithm and switch later if fork version is reached
- cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, ::jconf::inst()->GetMiningAlgoRoot());
+ auto miner_algo = ::jconf::inst()->GetMiningAlgoRoot();
+ cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
uint32_t iNonce;
@@ -261,16 +261,17 @@ void minethd::work_main()
std::this_thread::sleep_for(std::chrono::milliseconds(100));
consume_work();
- uint8_t new_version = oWork.getVersion();
- const bool time_to_fork =
- ((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) && version < 7 && new_version >= 7) ||
- (miner_algo == cryptonight_heavy && version < 3 && new_version >= 3);
- if(time_to_fork)
+ continue;
+ }
+ uint8_t new_version = oWork.getVersion();
+ if(new_version != version)
+ {
+ if(new_version >= ::jconf::inst()->GetMiningForkHeight())
{
+ miner_algo = ::jconf::inst()->GetMiningAlgo();
hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
}
version = new_version;
- continue;
}
cryptonight_extra_cpu_set_data(&ctx, oWork.bWorkBlob, oWork.iWorkSize);
@@ -294,11 +295,11 @@ void minethd::work_main()
uint32_t foundNonce[10];
uint32_t foundCount;
- cryptonight_extra_cpu_prepare(&ctx, iNonce, miner_algo, version);
+ cryptonight_extra_cpu_prepare(&ctx, iNonce, miner_algo);
- cryptonight_core_cpu_hash(&ctx, miner_algo, iNonce, version);
+ cryptonight_core_cpu_hash(&ctx, miner_algo, iNonce);
- cryptonight_extra_cpu_final(&ctx, iNonce, oWork.iTarget, &foundCount, foundNonce, miner_algo, version);
+ cryptonight_extra_cpu_final(&ctx, iNonce, oWork.iTarget, &foundCount, foundNonce, miner_algo);
for(size_t i = 0; i < foundCount; i++)
{
@@ -329,18 +330,8 @@ void minethd::work_main()
}
consume_work();
- uint8_t new_version = oWork.getVersion();
- const bool time_to_fork =
- ((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) && version < 7 && new_version >= 7) ||
- (miner_algo == cryptonight_heavy && version < 3 && new_version >= 3);
- if(time_to_fork)
- {
- hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
- }
- version = new_version;
}
}
} // namespace xmrstak
-
} //namespace nvidia
diff --git a/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp b/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp
index 29a3523..c1e31b9 100644
--- a/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp
+++ b/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp
@@ -45,8 +45,8 @@ int cuda_get_devicecount( int* deviceCount);
int cuda_get_deviceinfo(nvid_ctx *ctx);
int cryptonight_extra_cpu_init(nvid_ctx *ctx);
void cryptonight_extra_cpu_set_data( nvid_ctx* ctx, const void *data, uint32_t len);
-void cryptonight_extra_cpu_prepare(nvid_ctx* ctx, uint32_t startNonce, xmrstak_algo miner_algo, uint8_t version);
-void cryptonight_extra_cpu_final(nvid_ctx* ctx, uint32_t startNonce, uint64_t target, uint32_t* rescount, uint32_t *resnonce,xmrstak_algo miner_algo, uint8_t version);
+void cryptonight_extra_cpu_prepare(nvid_ctx* ctx, uint32_t startNonce, xmrstak_algo miner_algo);
+void cryptonight_extra_cpu_final(nvid_ctx* ctx, uint32_t startNonce, uint64_t target, uint32_t* rescount, uint32_t *resnonce,xmrstak_algo miner_algo);
}
-void cryptonight_core_cpu_hash(nvid_ctx* ctx, xmrstak_algo miner_algo, uint32_t startNonce, uint8_t version);
+void cryptonight_core_cpu_hash(nvid_ctx* ctx, xmrstak_algo miner_algo, uint32_t startNonce);
diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
index a1a9e15..7aa44e8 100644
--- a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
+++ b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
@@ -475,24 +475,16 @@ void cryptonight_core_gpu_hash(nvid_ctx* ctx, uint32_t nonce)
}
}
-void cryptonight_core_cpu_hash(nvid_ctx* ctx, xmrstak_algo miner_algo, uint32_t startNonce, uint8_t version)
+void cryptonight_core_cpu_hash(nvid_ctx* ctx, xmrstak_algo miner_algo, uint32_t startNonce)
{
if(miner_algo == cryptonight_monero)
{
- if(version >= 7)
- cryptonight_core_gpu_hash<CRYPTONIGHT_ITER, CRYPTONIGHT_MASK, CRYPTONIGHT_MEMORY/4, cryptonight_monero>(ctx, startNonce);
- else
- cryptonight_core_gpu_hash<CRYPTONIGHT_ITER, CRYPTONIGHT_MASK, CRYPTONIGHT_MEMORY/4, cryptonight>(ctx, startNonce);
+ cryptonight_core_gpu_hash<CRYPTONIGHT_ITER, CRYPTONIGHT_MASK, CRYPTONIGHT_MEMORY/4, cryptonight_monero>(ctx, startNonce);
}
else if(miner_algo == cryptonight_heavy)
{
- if(version >= 3)
- cryptonight_core_gpu_hash<CRYPTONIGHT_HEAVY_ITER, CRYPTONIGHT_HEAVY_MASK, CRYPTONIGHT_HEAVY_MEMORY/4, cryptonight_heavy>(ctx, startNonce);
- else
- {
- cryptonight_core_gpu_hash<CRYPTONIGHT_ITER, CRYPTONIGHT_MASK, CRYPTONIGHT_MEMORY/4, cryptonight>(ctx, startNonce);
- }
+ cryptonight_core_gpu_hash<CRYPTONIGHT_HEAVY_ITER, CRYPTONIGHT_HEAVY_MASK, CRYPTONIGHT_HEAVY_MEMORY/4, cryptonight_heavy>(ctx, startNonce);
}
else if(miner_algo == cryptonight)
{
@@ -504,10 +496,7 @@ void cryptonight_core_cpu_hash(nvid_ctx* ctx, xmrstak_algo miner_algo, uint32_t
}
else if(miner_algo == cryptonight_aeon)
{
- if(version >= 7)
- cryptonight_core_gpu_hash<CRYPTONIGHT_LITE_ITER, CRYPTONIGHT_LITE_MASK, CRYPTONIGHT_LITE_MEMORY/4, cryptonight_aeon>(ctx, startNonce);
- else
- cryptonight_core_gpu_hash<CRYPTONIGHT_LITE_ITER, CRYPTONIGHT_LITE_MASK, CRYPTONIGHT_LITE_MEMORY/4, cryptonight_lite>(ctx, startNonce);
+ cryptonight_core_gpu_hash<CRYPTONIGHT_LITE_ITER, CRYPTONIGHT_LITE_MASK, CRYPTONIGHT_LITE_MEMORY/4, cryptonight_aeon>(ctx, startNonce);
}
}
diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu
index c2a1f87..e2f0b2d 100644
--- a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu
+++ b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu
@@ -310,7 +310,7 @@ extern "C" int cryptonight_extra_cpu_init(nvid_ctx* ctx)
return 1;
}
-extern "C" void cryptonight_extra_cpu_prepare(nvid_ctx* ctx, uint32_t startNonce, xmrstak_algo miner_algo, uint8_t version)
+extern "C" void cryptonight_extra_cpu_prepare(nvid_ctx* ctx, uint32_t startNonce, xmrstak_algo miner_algo)
{
int threadsperblock = 128;
uint32_t wsize = ctx->device_blocks * ctx->device_threads;
@@ -318,7 +318,7 @@ extern "C" void cryptonight_extra_cpu_prepare(nvid_ctx* ctx, uint32_t startNonce
dim3 grid( ( wsize + threadsperblock - 1 ) / threadsperblock );
dim3 block( threadsperblock );
- if(miner_algo == cryptonight_heavy && version >= 3)
+ if(miner_algo == cryptonight_heavy)
{
CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_extra_gpu_prepare<cryptonight_heavy><<<grid, block >>>( wsize, ctx->d_input, ctx->inputlen, startNonce,
ctx->d_ctx_state,ctx->d_ctx_state2, ctx->d_ctx_a, ctx->d_ctx_b, ctx->d_ctx_key1, ctx->d_ctx_key2 ));
@@ -333,7 +333,7 @@ extern "C" void cryptonight_extra_cpu_prepare(nvid_ctx* ctx, uint32_t startNonce
}
}
-extern "C" void cryptonight_extra_cpu_final(nvid_ctx* ctx, uint32_t startNonce, uint64_t target, uint32_t* rescount, uint32_t *resnonce,xmrstak_algo miner_algo, uint8_t version)
+extern "C" void cryptonight_extra_cpu_final(nvid_ctx* ctx, uint32_t startNonce, uint64_t target, uint32_t* rescount, uint32_t *resnonce,xmrstak_algo miner_algo)
{
int threadsperblock = 128;
uint32_t wsize = ctx->device_blocks * ctx->device_threads;
@@ -344,7 +344,7 @@ extern "C" void cryptonight_extra_cpu_final(nvid_ctx* ctx, uint32_t startNonce,
CUDA_CHECK(ctx->device_id, cudaMemset( ctx->d_result_nonce, 0xFF, 10 * sizeof (uint32_t ) ));
CUDA_CHECK(ctx->device_id, cudaMemset( ctx->d_result_count, 0, sizeof (uint32_t ) ));
- if(miner_algo == cryptonight_heavy && version >= 3)
+ if(miner_algo == cryptonight_heavy)
{
CUDA_CHECK_MSG_KERNEL(
ctx->device_id,
OpenPOWER on IntegriCloud