diff options
author | fireice-uk <fireice-uk@users.noreply.github.com> | 2018-05-30 21:18:45 +0100 |
---|---|---|
committer | GitHub <noreply@github.com> | 2018-05-30 21:18:45 +0100 |
commit | c0ab1734332d6472225d8ac7394f6fcba71aabc9 (patch) | |
tree | b53a4c37905a0cb5dfa6a66f514cf3dc1ea94a21 /xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu | |
parent | 26a5d65f12b2f19a0a3ece39a2bc64718796367b (diff) | |
parent | 4f34bd18024fa71a8cab81d5a0b86cf5c7d9370e (diff) | |
download | xmr-stak-2.4.4.zip xmr-stak-2.4.4.tar.gz |
Merge pull request #1610 from fireice-uk/dev2.4.4
release 2.4.4
Diffstat (limited to 'xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu')
-rw-r--r-- | xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu | 14 |
1 files changed, 7 insertions, 7 deletions
diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu index f192f01..304997e 100644 --- a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu +++ b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu @@ -184,7 +184,7 @@ __global__ void cryptonight_extra_gpu_final( int threads, uint64_t target, uint3 uint32_t * __restrict__ ctx_state = d_ctx_state + thread * 50; uint64_t hash[4]; uint32_t state[50]; - + #pragma unroll for ( i = 0; i < 50; i++ ) state[i] = ctx_state[i]; @@ -296,7 +296,7 @@ extern "C" int cryptonight_extra_cpu_init(nvid_ctx* ctx) } else ctx->d_ctx_state2 = ctx->d_ctx_state; - + CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_ctx_key1, 40 * sizeof(uint32_t) * wsize)); CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_ctx_key2, 40 * sizeof(uint32_t) * wsize)); CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_ctx_text, 32 * sizeof(uint32_t) * wsize)); @@ -472,7 +472,7 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx) std::vector<int>::iterator it = std::find(arch.begin(), arch.end(), 20); if(it == arch.end()) { - printf("WARNING: NVIDIA GPU %d: miner not compiled for the gpu architecture %d.\n", ctx->device_id, gpuArch); + printf("WARNING: NVIDIA GPU %d: miner not compiled for CUDA architecture %d.\n", ctx->device_id, gpuArch); return 5; } } @@ -490,7 +490,7 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx) minSupportedArch = arch[i]; if(minSupportedArch < 30 || gpuArch < minSupportedArch) { - printf("WARNING: NVIDIA GPU %d: miner not compiled for the gpu architecture %d.\n", ctx->device_id, gpuArch); + printf("WARNING: NVIDIA GPU %d: miner not compiled for CUDA architecture %d.\n", ctx->device_id, gpuArch); return 5; } } @@ -517,7 +517,7 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx) */ ctx->device_threads = 64; constexpr size_t byteToMiB = 1024u * 1024u; - + // no limit by default 1TiB size_t maxMemUsage = byteToMiB * byteToMiB; if(props.major == 6) @@ -575,7 +575,7 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx) CUDA_CHECK(ctx->device_id, cudaFree(tmp)); // delete created context on the gpu CUDA_CHECK(ctx->device_id, cudaDeviceReset()); - + ctx->total_device_memory = totalMemory; ctx->free_device_memory = freeMemory; @@ -614,7 +614,7 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx) size_t perThread = hashMemSize + 16192u + 680u; if(cryptonight_heavy == ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()) perThread += 50 * 4; // state double buffer - + size_t max_intensity = limitedMemory / perThread; ctx->device_threads = max_intensity / ctx->device_blocks; // use only odd number of threads |