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/nvcc_code/cuda_extra.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) (limited to 'xmrstak/backend/nvidia/nvcc_code') 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/nvcc_code') 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 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/nvcc_code/cuda_extra.cu | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) (limited to 'xmrstak/backend/nvidia/nvcc_code') 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