fix wrong memory detection

Free and total memory is only evaluated on the first device.
To detect the gpu memory the gpu must be selected.

- create context on the gpu before the memory is checked
- add smx to the auto detection
- change the result code of `cuda_get_deviceinfo()`

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" +

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;
 }