"
"
"
diff --git a/xmrstak/misc/executor.cpp b/xmrstak/misc/executor.cpp
index 5eb9f07ef..18da71224 100644
--- a/xmrstak/misc/executor.cpp
+++ b/xmrstak/misc/executor.cpp
@@ -121,7 +121,9 @@ void executor::sched_reconnect()
rt, int_port(iReconnectAttempts));
auto work = xmrstak::miner_work();
- xmrstak::globalStates::inst().switch_work(work);
+ xmrstak::pool_data dat;
+
+ xmrstak::globalStates::inst().switch_work(work, dat);
push_timed_event(ex_event(EV_RECONNECT, usr_pool_id), rt);
}
@@ -166,7 +168,8 @@ void executor::log_result_ok(uint64_t iActualDiff)
jpsock* executor::pick_pool_by_id(size_t pool_id)
{
- assert(pool_id != invalid_pool_id);
+ if(pool_id == invalid_pool_id)
+ return nullptr;
if(pool_id == dev_pool_id)
return dev_pool;
@@ -233,14 +236,23 @@ void executor::on_pool_have_job(size_t pool_id, pool_job& oPoolJob)
jpsock* pool = pick_pool_by_id(pool_id);
- xmrstak::miner_work oWork(oPoolJob.sJobID, oPoolJob.bWorkBlob,
- oPoolJob.iWorkLen, oPoolJob.iResumeCnt, oPoolJob.iTarget,
- pool_id != dev_pool_id && ::jconf::inst()->NiceHashMode(),
- pool_id);
+ xmrstak::miner_work oWork(oPoolJob.sJobID, oPoolJob.bWorkBlob, oPoolJob.iWorkLen, oPoolJob.iTarget,
+ pool_id != dev_pool_id && ::jconf::inst()->NiceHashMode(), pool_id);
oWork.iTarget32 = oPoolJob.iTarget32;
- xmrstak::globalStates::inst().switch_work(oWork);
+ xmrstak::pool_data dat;
+ dat.iSavedNonce = oPoolJob.iSavedNonce;
+ dat.pool_id = pool_id;
+
+ xmrstak::globalStates::inst().switch_work(oWork, dat);
+
+ if(dat.pool_id != pool_id)
+ {
+ jpsock* prev_pool;
+ if((prev_pool = pick_pool_by_id(dat.pool_id)) != nullptr)
+ prev_pool->save_nonce(dat.iSavedNonce);
+ }
if(pool_id == dev_pool_id)
return;
@@ -251,7 +263,10 @@ void executor::on_pool_have_job(size_t pool_id, pool_job& oPoolJob)
printer::inst()->print_msg(L2, "Difficulty changed. Now: %llu.", int_port(iPoolDiff));
}
- printer::inst()->print_msg(L3, "New block detected.");
+ if(dat.pool_id == pool_id)
+ printer::inst()->print_msg(L3, "New block detected.");
+ else
+ printer::inst()->print_msg(L3, "Pool switched.");
}
void executor::on_miner_result(size_t pool_id, job_result& oResult)
@@ -356,13 +371,7 @@ void executor::on_switch_pool(size_t pool_id)
return;
}
- xmrstak::miner_work oWork(oPoolJob.sJobID, oPoolJob.bWorkBlob,
- oPoolJob.iWorkLen, oPoolJob.iResumeCnt, oPoolJob.iTarget,
- ::jconf::inst()->NiceHashMode(), pool_id);
-
- oWork.iTarget32 = oPoolJob.iTarget32;
-
- xmrstak::globalStates::inst().switch_work(oWork);
+ on_pool_have_job(current_pool_id, oPoolJob);
if(dev_pool->is_running())
push_timed_event(ex_event(EV_DEV_POOL_EXIT), 5);
diff --git a/xmrstak/net/jpsock.cpp b/xmrstak/net/jpsock.cpp
index ae20e5221..773592177 100644
--- a/xmrstak/net/jpsock.cpp
+++ b/xmrstak/net/jpsock.cpp
@@ -566,6 +566,12 @@ bool jpsock::cmd_submit(const char* sJobId, uint32_t iNonce, const uint8_t* bRes
return cmd_ret_wait(cmd_buffer, oResult);
}
+void jpsock::save_nonce(uint32_t nonce)
+{
+ std::unique_lock(job_mutex);
+ oCurrentJob.iSavedNonce = nonce;
+}
+
bool jpsock::get_current_job(pool_job& job)
{
std::unique_lock(job_mutex);
@@ -573,7 +579,6 @@ bool jpsock::get_current_job(pool_job& job)
if(oCurrentJob.iWorkLen == 0)
return false;
- oCurrentJob.iResumeCnt++;
job = oCurrentJob;
return true;
}
diff --git a/xmrstak/net/jpsock.hpp b/xmrstak/net/jpsock.hpp
index c2194ad0d..e94eaad22 100644
--- a/xmrstak/net/jpsock.hpp
+++ b/xmrstak/net/jpsock.hpp
@@ -50,6 +50,7 @@ class jpsock
inline uint64_t get_current_diff() { return iJobDiff; }
+ void save_nonce(uint32_t nonce);
bool get_current_job(pool_job& job);
size_t pool_id;
diff --git a/xmrstak/net/msgstruct.hpp b/xmrstak/net/msgstruct.hpp
index 05d15fe24..8e8254bdc 100644
--- a/xmrstak/net/msgstruct.hpp
+++ b/xmrstak/net/msgstruct.hpp
@@ -15,11 +15,11 @@ struct pool_job
// \todo remove workaround needed for amd
uint32_t iTarget32;
uint32_t iWorkLen;
- uint32_t iResumeCnt;
+ uint32_t iSavedNonce;
- pool_job() : iWorkLen(0), iResumeCnt(0) {}
+ pool_job() : iWorkLen(0), iSavedNonce(0) {}
pool_job(const char* sJobID, uint64_t iTarget, const uint8_t* bWorkBlob, uint32_t iWorkLen) :
- iTarget(iTarget), iWorkLen(iWorkLen), iResumeCnt(0)
+ iTarget(iTarget), iWorkLen(iWorkLen), iSavedNonce(0)
{
assert(iWorkLen <= sizeof(pool_job::bWorkBlob));
memcpy(this->sJobID, sJobID, sizeof(pool_job::sJobID));
diff --git a/doc/tuning.md b/doc/tuning.md
index 806f05915..645923559 100644
--- a/doc/tuning.md
+++ b/doc/tuning.md
@@ -53,15 +53,16 @@ In the most cases a `worksize` of `16` or `8` is optimal.
### Add More GPUs
-To add a new GPU you need to add a new config set to `gpu_threads_conf` and increase `gpu_thread_num"` to the number of gpus (entries in `gpu_threads_conf`).
-`index` is the number of the gpu.
+To add a new GPU you need to add a new config set to `gpu_threads_conf`. `index` is the OpenCL index of the gpu.
+`platform_index`is the index of the OpenCL platform (Intel / AMD / Nvidia).
+If you are unsure of either GPU or platform index value, you can use `clinfo` tool that comes with AMD APP SDK to dump the values.
```
-"gpu_thread_num" : 2,
-
"gpu_threads_conf" :
[
{ "index" : 0, "intensity" : 1000, "worksize" : 8, "affine_to_cpu" : false },
{ "index" : 1, "intensity" : 1000, "worksize" : 8, "affine_to_cpu" : false },
],
+
+"platform_index" : 0,
```
diff --git a/xmrstak/backend/amd/amd_gpu/gpu.cpp b/xmrstak/backend/amd/amd_gpu/gpu.cpp
index 392903aad..69ef18b84 100644
--- a/xmrstak/backend/amd/amd_gpu/gpu.cpp
+++ b/xmrstak/backend/amd/amd_gpu/gpu.cpp
@@ -844,7 +844,8 @@ size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput)
}
}*/
- if((ret = clEnqueueNDRangeKernel(ctx->CommandQueues, ctx->Kernels[1], 1, &ctx->Nonce, &g_thd, &w_size, 0, NULL, NULL)) != CL_SUCCESS)
+ size_t tmpNonce = ctx->Nonce;
+ if((ret = clEnqueueNDRangeKernel(ctx->CommandQueues, ctx->Kernels[1], 1, &tmpNonce, &g_thd, &w_size, 0, NULL, NULL)) != CL_SUCCESS)
{
printer::inst()->print_msg(L1,"Error %s when calling clEnqueueNDRangeKernel for kernel %d.", err_to_str(ret), 1);
return ERR_OCL_API;
@@ -896,8 +897,9 @@ size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput)
// round up to next multiple of w_size
BranchNonces[i] = ((BranchNonces[i] + w_size - 1u) / w_size) * w_size;
// number of global threads must be a multiple of the work group size (w_size)
- assert(BranchNonces%w_size == 0);
- if((ret = clEnqueueNDRangeKernel(ctx->CommandQueues, ctx->Kernels[i + 3], 1, &ctx->Nonce, BranchNonces + i, &w_size, 0, NULL, NULL)) != CL_SUCCESS)
+ assert(BranchNonces[i]%w_size == 0);
+ size_t tmpNonce = ctx->Nonce;
+ if((ret = clEnqueueNDRangeKernel(ctx->CommandQueues, ctx->Kernels[i + 3], 1, &tmpNonce, BranchNonces + i, &w_size, 0, NULL, NULL)) != CL_SUCCESS)
{
printer::inst()->print_msg(L1,"Error %s when calling clEnqueueNDRangeKernel for kernel %d.", err_to_str(ret), i + 3);
return ERR_OCL_API;
diff --git a/xmrstak/backend/amd/amd_gpu/gpu.hpp b/xmrstak/backend/amd/amd_gpu/gpu.hpp
index 5ff7ea17b..c2d708d90 100644
--- a/xmrstak/backend/amd/amd_gpu/gpu.hpp
+++ b/xmrstak/backend/amd/amd_gpu/gpu.hpp
@@ -36,7 +36,7 @@ struct GpuContext
int computeUnits;
std::string name;
- size_t Nonce;
+ uint32_t Nonce;
};
diff --git a/xmrstak/backend/amd/jconf.cpp b/xmrstak/backend/amd/jconf.cpp
index c2bf1fa19..0617aeb2f 100644
--- a/xmrstak/backend/amd/jconf.cpp
+++ b/xmrstak/backend/amd/jconf.cpp
@@ -48,7 +48,7 @@ using namespace rapidjson;
/*
* This enum needs to match index in oConfigValues, otherwise we will get a runtime error
*/
-enum configEnum { iGpuThreadNum, aGpuThreadsConf, iPlatformIdx };
+enum configEnum { aGpuThreadsConf, iPlatformIdx };
struct configVal {
configEnum iName;
@@ -58,7 +58,6 @@ struct configVal {
//Same order as in configEnum, as per comment above
configVal oConfigValues[] = {
- { iGpuThreadNum, "gpu_thread_num", kNumberType },
{ aGpuThreadsConf, "gpu_threads_conf", kArrayType },
{ iPlatformIdx, "platform_index", kNumberType }
};
@@ -235,14 +234,6 @@ bool jconf::parse_config(const char* sFilename)
}
size_t n_thd = prv->configValues[aGpuThreadsConf]->Size();
- if(prv->configValues[iGpuThreadNum]->GetUint64() != n_thd)
- {
- printer::inst()->print_msg(L0,
- "Invalid config file. Your GPU config array has %llu members, while you want to use %llu threads.",
- int_port(n_thd), int_port(prv->configValues[iGpuThreadNum]->GetUint64()));
- return false;
- }
-
thd_cfg c;
for(size_t i=0; i < n_thd; i++)
{
diff --git a/xmrstak/backend/amd/minethd.cpp b/xmrstak/backend/amd/minethd.cpp
index 5ca10d3c0..9d18860f1 100644
--- a/xmrstak/backend/amd/minethd.cpp
+++ b/xmrstak/backend/amd/minethd.cpp
@@ -167,13 +167,10 @@ void minethd::consume_work()
void minethd::work_main()
{
uint64_t iCount = 0;
-
cryptonight_ctx* cpu_ctx;
cpu_ctx = cpu::minethd::minethd_alloc_ctx();
cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/);
-
globalStates::inst().iConsumeCnt++;
- uint32_t* piNonce = (uint32_t*)(oWork.bWorkBlob + 39);
while (bQuit == 0)
{
@@ -190,17 +187,24 @@ void minethd::work_main()
continue;
}
- if(oWork.bNiceHash)
- pGpuCtx->Nonce = calc_nicehash_nonce(*piNonce, oWork.iResumeCnt);
- else
- pGpuCtx->Nonce = calc_start_nonce(oWork.iResumeCnt);
+ uint32_t h_per_round = pGpuCtx->rawIntensity;
+ size_t round_ctr = 0;
assert(sizeof(job_result::sJobID) == sizeof(pool_job::sJobID));
uint32_t target = oWork.iTarget32;
XMRSetJob(pGpuCtx, oWork.bWorkBlob, oWork.iWorkSize, target);
+ if(oWork.bNiceHash)
+ pGpuCtx->Nonce = *(uint32_t*)(oWork.bWorkBlob + 39);
+
while(globalStates::inst().iGlobalJobNo.load(std::memory_order_relaxed) == iJobNo)
{
+ //Allocate a new nonce every 16 rounds
+ if((round_ctr++ & 0xF) == 0)
+ {
+ globalStates::inst().calc_start_nonce(pGpuCtx->Nonce, oWork.bNiceHash, h_per_round * 16);
+ }
+
cl_uint results[0x100];
memset(results,0,sizeof(cl_uint)*(0x100));
diff --git a/xmrstak/backend/cpu/autoAdjust.hpp b/xmrstak/backend/cpu/autoAdjust.hpp
index e1d0b2586..e88d272e2 100644
--- a/xmrstak/backend/cpu/autoAdjust.hpp
+++ b/xmrstak/backend/cpu/autoAdjust.hpp
@@ -74,7 +74,7 @@ class autoAdjust
if(L3KB_size <= 0)
break;
- double_mode = L3KB_size / 2048 > (int32_t)(corecnt-i);
+ double_mode = L3KB_size / 1024 > (int32_t)(corecnt-i);
conf += std::string(" { \"low_power_mode\" : ");
conf += std::string(double_mode ? "true" : "false");
@@ -93,10 +93,9 @@ class autoAdjust
aff_id++;
if(double_mode)
-
L3KB_size -= 2048;
- else
- L3KB_size -= 1024;
+ else
+ L3KB_size -= 1024;
}
}
diff --git a/xmrstak/backend/cpu/crypto/cryptonight.h b/xmrstak/backend/cpu/crypto/cryptonight.h
index 3b6c0c56a..d47b79f87 100644
--- a/xmrstak/backend/cpu/crypto/cryptonight.h
+++ b/xmrstak/backend/cpu/crypto/cryptonight.h
@@ -8,7 +8,7 @@ extern "C" {
#include