From 75ae1dd60165141700c36ab699656e9e073e186d Mon Sep 17 00:00:00 2001 From: Tony Butler Date: Thu, 5 Apr 2018 23:19:48 -0600 Subject: Repair all 'namepsace' to 'namespace' (all within comments) --- xmrstak/backend/nvidia/autoAdjust.hpp | 2 +- xmrstak/backend/nvidia/jconf.hpp | 2 +- xmrstak/backend/nvidia/minethd.hpp | 2 +- 3 files changed, 3 insertions(+), 3 deletions(-) (limited to 'xmrstak/backend/nvidia') diff --git a/xmrstak/backend/nvidia/autoAdjust.hpp b/xmrstak/backend/nvidia/autoAdjust.hpp index be7d1ce..0b1da87 100644 --- a/xmrstak/backend/nvidia/autoAdjust.hpp +++ b/xmrstak/backend/nvidia/autoAdjust.hpp @@ -109,4 +109,4 @@ private: }; } // namespace nvidia -} // namepsace xmrstak +} // namespace xmrstak diff --git a/xmrstak/backend/nvidia/jconf.hpp b/xmrstak/backend/nvidia/jconf.hpp index 7f60f1d..b4ebaa0 100644 --- a/xmrstak/backend/nvidia/jconf.hpp +++ b/xmrstak/backend/nvidia/jconf.hpp @@ -49,4 +49,4 @@ private: }; } // namespace nvidia -} // namepsace xmrstak +} // namespace xmrstak diff --git a/xmrstak/backend/nvidia/minethd.hpp b/xmrstak/backend/nvidia/minethd.hpp index fcd24fa..89c2944 100644 --- a/xmrstak/backend/nvidia/minethd.hpp +++ b/xmrstak/backend/nvidia/minethd.hpp @@ -60,4 +60,4 @@ private: }; } // namespace nvidia -} // namepsace xmrstak +} // namespace xmrstak -- cgit v1.1 From 5fd4520afa4dd824676f9e24e3f3c771687a6a39 Mon Sep 17 00:00:00 2001 From: Tony Butler Date: Thu, 5 Apr 2018 23:30:19 -0600 Subject: Repair more typos in comments only --- xmrstak/backend/nvidia/autoAdjust.hpp | 2 +- xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu | 4 ++-- 2 files changed, 3 insertions(+), 3 deletions(-) (limited to 'xmrstak/backend/nvidia') diff --git a/xmrstak/backend/nvidia/autoAdjust.hpp b/xmrstak/backend/nvidia/autoAdjust.hpp index 0b1da87..d8bb621 100644 --- a/xmrstak/backend/nvidia/autoAdjust.hpp +++ b/xmrstak/backend/nvidia/autoAdjust.hpp @@ -50,7 +50,7 @@ public: ctx.device_blocks = -1; ctx.device_threads = -1; - // set all evice option those marked as auto (-1) to a valid value + // set all device option those marked as auto (-1) to a valid value #ifndef _WIN32 ctx.device_bfactor = 0; ctx.device_bsleep = 0; diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu index 02c157e..ead93c5 100644 --- a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu +++ b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu @@ -458,7 +458,7 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx) #undef XMRSTAK_PP_TOSTRING1 std::stringstream ss(archStringList); - //transform string list sperated with `+` into a vector of integers + //transform string list separated with `+` into a vector of integers int tmpArch; while ( ss >> tmpArch ) arch.push_back( tmpArch ); @@ -492,7 +492,7 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx) } } - // set all evice option those marked as auto (-1) to a valid value + // set all device option those marked as auto (-1) to a valid value if(ctx->device_blocks == -1) { /* good values based of my experience -- cgit v1.1 From 1b18f598aa1190a0e6126ed2c70e052e9403d180 Mon Sep 17 00:00:00 2001 From: psychocrypt Date: Sun, 8 Apr 2018 20:40:54 +0200 Subject: refactor scratchpad creation Use the maximum scratchpad size from before and after the fork. --- xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu | 10 ++++++++-- 1 file changed, 8 insertions(+), 2 deletions(-) (limited to 'xmrstak/backend/nvidia') diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu index ead93c5..f016ef4 100644 --- a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu +++ b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu @@ -279,7 +279,10 @@ extern "C" int cryptonight_extra_cpu_init(nvid_ctx* ctx) if(gpuArch < 70) CUDA_CHECK(ctx->device_id, cudaDeviceSetCacheConfig(cudaFuncCachePreferL1)); - size_t hashMemSize = cn_select_memory(::jconf::inst()->GetMiningAlgo()); + size_t hashMemSize = std::max( + cn_select_memory(::jconf::inst()->GetMiningAlgo()), + cn_select_memory(::jconf::inst()->GetMiningAlgoRoot()) + ); size_t wsize = ctx->device_blocks * ctx->device_threads; CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_ctx_state, 50 * sizeof(uint32_t) * wsize)); @@ -576,7 +579,10 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx) ctx->total_device_memory = totalMemory; ctx->free_device_memory = freeMemory; - size_t hashMemSize = cn_select_memory(::jconf::inst()->GetMiningAlgo()); + size_t hashMemSize = std::max( + cn_select_memory(::jconf::inst()->GetMiningAlgo()), + cn_select_memory(::jconf::inst()->GetMiningAlgoRoot()) + ); #ifdef WIN32 /* We use in windows bfactor (split slow kernel into smaller parts) to avoid -- cgit v1.1 From a5ddd040a6eeb0609e1d1b0f16c1d271d31c7377 Mon Sep 17 00:00:00 2001 From: psychocrypt Date: Mon, 16 Apr 2018 17:52:42 +0200 Subject: fix wrong algo selection In the case where the dev pool mines on a higher version than a monero fork coin the miner is not resetting the algorithm. This PR select the correct algorithm each time the block version hash changed. --- xmrstak/backend/nvidia/minethd.cpp | 11 +++++++++-- 1 file changed, 9 insertions(+), 2 deletions(-) (limited to 'xmrstak/backend/nvidia') diff --git a/xmrstak/backend/nvidia/minethd.cpp b/xmrstak/backend/nvidia/minethd.cpp index 804c06a..4593dab 100644 --- a/xmrstak/backend/nvidia/minethd.cpp +++ b/xmrstak/backend/nvidia/minethd.cpp @@ -247,6 +247,7 @@ void minethd::work_main() globalStates::inst().iConsumeCnt++; uint8_t version = 0; + size_t lastPoolId = 0; while (bQuit == 0) { @@ -264,13 +265,19 @@ void minethd::work_main() continue; } uint8_t new_version = oWork.getVersion(); - if(new_version != version) + if(new_version != version || oWork.iPoolId != lastPoolId) { - if(new_version >= ::jconf::inst()->GetMiningForkVersion()) + if(new_version >= ::jconf::inst()->GetMiningForkVersion() || oWork.iPoolId == 0) { miner_algo = ::jconf::inst()->GetMiningAlgo(); hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); } + else + { + miner_algo = ::jconf::inst()->GetMiningAlgoRoot(); + hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); + } + lastPoolId = oWork.iPoolId; version = new_version; } -- cgit v1.1 From 0877e2f654b203c5145bb8154fcfb1ad46ba8265 Mon Sep 17 00:00:00 2001 From: psychocrypt Date: Mon, 16 Apr 2018 21:46:33 +0200 Subject: add independent dev pool coin description - allow the dev pool to fork on a different block version than the user descriped coin All algorithm are centered around the user coin description. It is allowed to have two two different coin algorithms in the user coin description. It is only allowed to use algorithms for the dev pool coin description those are used in the user coin description. There are two ways to define a non forking coin. - set both user coin algorithm descriptions to the same algorithm and set version to zero - set the first algorithm in the user coin description to something you like to use in the dev pool and set the second algorithm to the correct representation of the coin. Set the version to 255. This will allow that the dev pool can mine on a different coin algorithm than the not forking user coin. Do not use an algorithm with different scratchpad size for the dev pool. --- xmrstak/backend/nvidia/minethd.cpp | 9 +++++---- xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu | 12 ++++++------ 2 files changed, 11 insertions(+), 10 deletions(-) (limited to 'xmrstak/backend/nvidia') diff --git a/xmrstak/backend/nvidia/minethd.cpp b/xmrstak/backend/nvidia/minethd.cpp index 4593dab..92f5f78 100644 --- a/xmrstak/backend/nvidia/minethd.cpp +++ b/xmrstak/backend/nvidia/minethd.cpp @@ -239,7 +239,7 @@ void minethd::work_main() cpu_ctx = cpu::minethd::minethd_alloc_ctx(); // start with root algorithm and switch later if fork version is reached - auto miner_algo = ::jconf::inst()->GetMiningAlgoRoot(); + auto miner_algo = ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot(); cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); uint32_t iNonce; @@ -267,14 +267,15 @@ void minethd::work_main() uint8_t new_version = oWork.getVersion(); if(new_version != version || oWork.iPoolId != lastPoolId) { - if(new_version >= ::jconf::inst()->GetMiningForkVersion() || oWork.iPoolId == 0) + coinDescription coinDesc = ::jconf::inst()->GetCurrentCoinSelection().GetDescription(oWork.iPoolId); + if(new_version >= coinDesc.GetMiningForkVersion()) { - miner_algo = ::jconf::inst()->GetMiningAlgo(); + miner_algo = coinDesc.GetMiningAlgo(); hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); } else { - miner_algo = ::jconf::inst()->GetMiningAlgoRoot(); + miner_algo = coinDesc.GetMiningAlgoRoot(); hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); } lastPoolId = oWork.iPoolId; diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu index f016ef4..f192f01 100644 --- a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu +++ b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu @@ -280,14 +280,14 @@ extern "C" int cryptonight_extra_cpu_init(nvid_ctx* ctx) CUDA_CHECK(ctx->device_id, cudaDeviceSetCacheConfig(cudaFuncCachePreferL1)); size_t hashMemSize = std::max( - cn_select_memory(::jconf::inst()->GetMiningAlgo()), - cn_select_memory(::jconf::inst()->GetMiningAlgoRoot()) + cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()), + cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot()) ); size_t wsize = ctx->device_blocks * ctx->device_threads; CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_ctx_state, 50 * sizeof(uint32_t) * wsize)); size_t ctx_b_size = 4 * sizeof(uint32_t) * wsize; - if(cryptonight_heavy == ::jconf::inst()->GetMiningAlgo()) + if(cryptonight_heavy == ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()) { // extent ctx_b to hold the state of idx0 ctx_b_size += sizeof(uint32_t) * wsize; @@ -580,8 +580,8 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx) ctx->free_device_memory = freeMemory; size_t hashMemSize = std::max( - cn_select_memory(::jconf::inst()->GetMiningAlgo()), - cn_select_memory(::jconf::inst()->GetMiningAlgoRoot()) + cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()), + cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot()) ); #ifdef WIN32 @@ -612,7 +612,7 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx) // up to 16kibyte extra memory is used per thread for some kernel (lmem/local memory) // 680bytes are extra meta data memory per hash size_t perThread = hashMemSize + 16192u + 680u; - if(cryptonight_heavy == ::jconf::inst()->GetMiningAlgo()) + if(cryptonight_heavy == ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()) perThread += 50 * 4; // state double buffer size_t max_intensity = limitedMemory / perThread; -- cgit v1.1