summaryrefslogtreecommitdiffstats
path: root/xmrstak/backend/nvidia
diff options
context:
space:
mode:
Diffstat (limited to 'xmrstak/backend/nvidia')
-rw-r--r--xmrstak/backend/nvidia/nvcc_code/cuda_core.cu2
-rw-r--r--xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu4
2 files changed, 5 insertions, 1 deletions
diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
index a92fa8c..dba6676 100644
--- a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
+++ b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
@@ -167,10 +167,10 @@ __forceinline__ __device__ uint32_t shuffle(volatile uint32_t* ptr,const uint32_
#endif
}
+template<size_t ITERATIONS, uint32_t THREAD_SHIFT, uint32_t MASK>
#ifdef XMR_STAK_THREADS
__launch_bounds__( XMR_STAK_THREADS * 4 )
#endif
-template<size_t ITERATIONS, uint32_t THREAD_SHIFT, uint32_t MASK>
__global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int partidx, uint32_t * d_long_state, uint32_t * d_ctx_a, uint32_t * d_ctx_b )
{
__shared__ uint32_t sharedMemory[1024];
diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu
index 5501d8d..333ae73 100644
--- a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu
+++ b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu
@@ -380,6 +380,10 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx)
*/
ctx->device_blocks = props.multiProcessorCount *
( props.major < 3 ? 2 : 3 );
+
+ // increase bfactor for low end devices to avoid that the miner is killed by the OS
+ if(props.multiProcessorCount < 6)
+ ctx->device_bfactor += 2;
}
if(ctx->device_threads == -1)
{
OpenPOWER on IntegriCloud