summaryrefslogtreecommitdiffstats
path: root/xmrstak/backend/nvidia
diff options
context:
space:
mode:
Diffstat (limited to 'xmrstak/backend/nvidia')
-rw-r--r--xmrstak/backend/nvidia/autoAdjust.hpp13
-rw-r--r--xmrstak/backend/nvidia/minethd.cpp7
-rw-r--r--xmrstak/backend/nvidia/minethd.hpp1
-rw-r--r--xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu43
4 files changed, 50 insertions, 14 deletions
diff --git a/xmrstak/backend/nvidia/autoAdjust.hpp b/xmrstak/backend/nvidia/autoAdjust.hpp
index bf04518..d36a46a 100644
--- a/xmrstak/backend/nvidia/autoAdjust.hpp
+++ b/xmrstak/backend/nvidia/autoAdjust.hpp
@@ -60,17 +60,15 @@ public:
ctx.device_bfactor = 6;
ctx.device_bsleep = 25;
#endif
- if( cuda_get_deviceinfo(&ctx) != 1 )
- {
- printer::inst()->print_msg(L0, "Setup failed for GPU %d. Exitting.\n", i);
- std::exit(0);
- }
- nvidCtxVec.push_back(ctx);
+ if(cuda_get_deviceinfo(&ctx) == 0)
+ nvidCtxVec.push_back(ctx);
+ else
+ printer::inst()->print_msg(L0, "WARNING: NVIDIA setup failed for GPU %d.\n", i);
}
generateThreadConfig();
- return true;
+ return true;
}
@@ -94,6 +92,7 @@ private:
{
conf += std::string(" // gpu: ") + ctx.name + " architecture: " + std::to_string(ctx.device_arch[0] * 10 + ctx.device_arch[1]) + "\n";
conf += std::string(" // memory: ") + std::to_string(ctx.free_device_memory / byte2mib) + "/" + std::to_string(ctx.total_device_memory / byte2mib) + " MiB\n";
+ conf += std::string(" // smx: ") + std::to_string(ctx.device_mpcount) + "\n";
conf += std::string(" { \"index\" : ") + std::to_string(ctx.device_id) + ",\n" +
" \"threads\" : " + std::to_string(ctx.device_threads) + ", \"blocks\" : " + std::to_string(ctx.device_blocks) + ",\n" +
" \"bfactor\" : " + std::to_string(ctx.device_bfactor) + ", \"bsleep\" : " + std::to_string(ctx.device_bsleep) + ",\n" +
diff --git a/xmrstak/backend/nvidia/minethd.cpp b/xmrstak/backend/nvidia/minethd.cpp
index 0cf9a42..5170305 100644
--- a/xmrstak/backend/nvidia/minethd.cpp
+++ b/xmrstak/backend/nvidia/minethd.cpp
@@ -66,6 +66,7 @@ namespace nvidia
minethd::minethd(miner_work& pWork, size_t iNo, const jconf::thd_cfg& cfg)
{
+ this->backendType = iBackend::NVIDIA;
oWork = pWork;
bQuit = 0;
iThreadNo = (uint8_t)iNo;
@@ -78,6 +79,7 @@ minethd::minethd(miner_work& pWork, size_t iNo, const jconf::thd_cfg& cfg)
ctx.device_bsleep = (int)cfg.bsleep;
this->affinity = cfg.cpu_aff;
+ std::unique_lock<std::mutex> lck(thd_aff_set);
std::future<void> order_guard = order_fix.get_future();
oWorkThd = std::thread(&minethd::work_main, this);
@@ -206,6 +208,9 @@ void minethd::work_main()
bindMemoryToNUMANode(affinity);
order_fix.set_value();
+ std::unique_lock<std::mutex> lck(thd_aff_set);
+ lck.release();
+ std::this_thread::yield();
uint64_t iCount = 0;
cryptonight_ctx* cpu_ctx;
@@ -280,7 +285,7 @@ void minethd::work_main()
if ( (*((uint64_t*)(bResult + 24))) < oWork.iTarget)
executor::inst()->push_event(ex_event(job_result(oWork.sJobID, foundNonce[i], bResult), oWork.iPoolId));
else
- executor::inst()->log_result_error("NVIDIA Invalid Result");
+ executor::inst()->push_event(ex_event("NVIDIA Invalid Result", oWork.iPoolId));
}
iCount += h_per_round;
diff --git a/xmrstak/backend/nvidia/minethd.hpp b/xmrstak/backend/nvidia/minethd.hpp
index f6d989c..ffcdab1 100644
--- a/xmrstak/backend/nvidia/minethd.hpp
+++ b/xmrstak/backend/nvidia/minethd.hpp
@@ -45,6 +45,7 @@ private:
miner_work oWork;
std::promise<void> order_fix;
+ std::mutex thd_aff_set;
std::thread oWorkThd;
int64_t affinity;
diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu
index e18532f..61d45ed 100644
--- a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu
+++ b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu
@@ -270,6 +270,14 @@ extern "C" int cuda_get_devicecount( int* deviceCount)
return 1;
}
+/** get device information
+ *
+ * @return 0 = all OK,
+ * 1 = something went wrong,
+ * 2 = gpu cannot be selected,
+ * 3 = context cannot be created
+ * 4 = not enough memory
+ */
extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx)
{
cudaError_t err;
@@ -279,25 +287,25 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx)
if(err != cudaSuccess)
{
printf("Unable to query CUDA driver version! Is an nVidia driver installed?\n");
- return 0;
+ return 1;
}
if(version < CUDART_VERSION)
{
printf("Driver does not support CUDA %d.%d API! Update your nVidia driver!\n", CUDART_VERSION / 1000, (CUDART_VERSION % 1000) / 10);
- return 0;
+ return 1;
}
int GPU_N;
if(cuda_get_devicecount(&GPU_N) == 0)
{
- return 0;
+ return 1;
}
if(ctx->device_id >= GPU_N)
{
printf("Invalid device ID!\n");
- return 0;
+ return 1;
}
cudaDeviceProp props;
@@ -305,7 +313,7 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx)
if(err != cudaSuccess)
{
printf("\nGPU %d: %s\n%s line %d\n", ctx->device_id, cudaGetErrorString(err), __FILE__, __LINE__);
- return 0;
+ return 1;
}
ctx->device_name = strdup(props.name);
@@ -347,9 +355,31 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx)
maxMemUsage = size_t(1024u) * byteToMiB;
}
+ int* tmp;
+ cudaError_t err;
+ // a device must be selected to get the right memory usage later on
+ err = cudaSetDevice(ctx->device_id);
+ if(err != cudaSuccess)
+ {
+ printf("WARNING: NVIDIA GPU %d: cannot be selected.\n", ctx->device_id);
+ return 2;
+ }
+ // trigger that a context on the gpu will be allocated
+ err = cudaMalloc(&tmp, 256);
+ if(err != cudaSuccess)
+ {
+ printf("WARNING: NVIDIA GPU %d: context cannot be created.\n", ctx->device_id);
+ return 3;
+ }
+
+
size_t freeMemory = 0;
size_t totalMemory = 0;
CUDA_CHECK(ctx->device_id, cudaMemGetInfo(&freeMemory, &totalMemory));
+
+ cudaFree(tmp);
+ // delete created context on the gpu
+ cudaDeviceReset();
ctx->total_device_memory = totalMemory;
ctx->free_device_memory = freeMemory;
@@ -379,6 +409,7 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx)
printf("WARNING: NVIDIA GPU %d: already %s MiB memory in use, skip GPU.\n",
ctx->device_id,
std::to_string(usedMem/byteToMiB).c_str());
+ return 4;
}
else
maxMemUsage -= usedMem;
@@ -404,5 +435,5 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx)
}
- return 1;
+ return 0;
}
OpenPOWER on IntegriCloud