diff options
Diffstat (limited to 'xmrstak/backend')
| -rw-r--r-- | xmrstak/backend/amd/amd_gpu/gpu.cpp | 4 | ||||
| -rw-r--r-- | xmrstak/backend/amd/amd_gpu/gpu.hpp | 1 | ||||
| -rw-r--r-- | xmrstak/backend/nvidia/nvcc_code/cuda_core.cu | 2 | ||||
| -rw-r--r-- | xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu | 4 |
4 files changed, 8 insertions, 3 deletions
diff --git a/xmrstak/backend/amd/amd_gpu/gpu.cpp b/xmrstak/backend/amd/amd_gpu/gpu.cpp index 15b8457..879a2e4 100644 --- a/xmrstak/backend/amd/amd_gpu/gpu.cpp +++ b/xmrstak/backend/amd/amd_gpu/gpu.cpp @@ -476,7 +476,7 @@ std::vector<GpuContext> getAMDDevices(int index) if(clStatus == CL_SUCCESS) { std::string devVendor(devVendorVec.data()); - if( devVendor.find("Advanced Micro Devices") != std::string::npos) + if( devVendor.find("Advanced Micro Devices") != std::string::npos || devVendor.find("AMD") != std::string::npos) { GpuContext ctx; ctx.deviceIdx = k; @@ -541,7 +541,7 @@ int getAMDPlatformIdx() clGetPlatformInfo(platforms[i], CL_PLATFORM_VENDOR, infoSize, platformNameVec.data(), NULL); std::string platformName(platformNameVec.data()); - if( platformName.find("Advanced Micro Devices") != std::string::npos) + if( platformName.find("Advanced Micro Devices") != std::string::npos || platformName.find("Apple") != std::string::npos) { platformIndex = i; printer::inst()->print_msg(L0,"Found AMD platform index id = %i, name = %s",i , platformName.c_str()); diff --git a/xmrstak/backend/amd/amd_gpu/gpu.hpp b/xmrstak/backend/amd/amd_gpu/gpu.hpp index 123de01..abbd08d 100644 --- a/xmrstak/backend/amd/amd_gpu/gpu.hpp +++ b/xmrstak/backend/amd/amd_gpu/gpu.hpp @@ -9,6 +9,7 @@ #endif #include <stdint.h> +#include <string> #include <vector> #define ERR_SUCCESS (0) 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) { |
