diff options
Diffstat (limited to 'xmrstak')
| -rw-r--r-- | xmrstak/backend/amd/amd_gpu/gpu.cpp | 2 | ||||
| -rw-r--r-- | xmrstak/backend/amd/jconf.cpp | 7 | ||||
| -rw-r--r-- | xmrstak/backend/nvidia/nvcc_code/cuda_core.cu | 12 | ||||
| -rw-r--r-- | xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu | 8 |
4 files changed, 25 insertions, 4 deletions
diff --git a/xmrstak/backend/amd/amd_gpu/gpu.cpp b/xmrstak/backend/amd/amd_gpu/gpu.cpp index d9bc962..c39c567 100644 --- a/xmrstak/backend/amd/amd_gpu/gpu.cpp +++ b/xmrstak/backend/amd/amd_gpu/gpu.cpp @@ -518,13 +518,13 @@ std::vector<GpuContext> getAMDDevices(int index) printer::inst()->print_msg(L1,"WARNING: %s when calling clGetDeviceInfo to get CL_DEVICE_NAME for device %u.", err_to_str(clStatus), k); continue; } - printer::inst()->print_msg(L0,"Found OpenCL GPU %s.",ctx.name.c_str()); // if environment variable GPU_SINGLE_ALLOC_PERCENT is not set we can not allocate the full memory ctx.deviceIdx = k; ctx.freeMem = std::min(ctx.freeMem, maxMem); ctx.name = std::string(devNameVec.data()); ctx.DeviceID = device_list[k]; + printer::inst()->print_msg(L0,"Found OpenCL GPU %s.",ctx.name.c_str()); ctxVec.push_back(ctx); } } diff --git a/xmrstak/backend/amd/jconf.cpp b/xmrstak/backend/amd/jconf.cpp index 07afb19..f126342 100644 --- a/xmrstak/backend/amd/jconf.cpp +++ b/xmrstak/backend/amd/jconf.cpp @@ -56,9 +56,10 @@ struct configVal { Type iType; }; -//Same order as in configEnum, as per comment above +// Same order as in configEnum, as per comment above +// kNullType means any type configVal oConfigValues[] = { - { aGpuThreadsConf, "gpu_threads_conf", kArrayType }, + { aGpuThreadsConf, "gpu_threads_conf", kNullType }, { iPlatformIdx, "platform_index", kNumberType } }; @@ -68,6 +69,8 @@ inline bool checkType(Type have, Type want) { if(want == have) return true; + else if(want == kNullType) + return true; else if(want == kTrueType && have == kFalseType) return true; else if(want == kFalseType && have == kTrueType) diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu index 15a6f36..cc97274 100644 --- a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu +++ b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu @@ -74,24 +74,36 @@ __device__ __forceinline__ uint64_t cuda_mul128( uint64_t multiplier, uint64_t m template< typename T > __device__ __forceinline__ T loadGlobal64( T * const addr ) { +#if (__CUDA_ARCH__ < 700) T x; asm volatile( "ld.global.cg.u64 %0, [%1];" : "=l"( x ) : "l"( addr ) ); return x; +#else + return *addr; +#endif } template< typename T > __device__ __forceinline__ T loadGlobal32( T * const addr ) { +#if (__CUDA_ARCH__ < 700) T x; asm volatile( "ld.global.cg.u32 %0, [%1];" : "=r"( x ) : "l"( addr ) ); return x; +#else + return *addr; +#endif } template< typename T > __device__ __forceinline__ void storeGlobal32( T* addr, T const & val ) { +#if (__CUDA_ARCH__ < 700) asm volatile( "st.global.cg.u32 [%0], %1;" : : "l"( addr ), "r"( val ) ); +#else + *addr = val; +#endif } template<size_t ITERATIONS, uint32_t THREAD_SHIFT> diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu index d5d0039..92259db 100644 --- a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu +++ b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu @@ -204,7 +204,13 @@ extern "C" int cryptonight_extra_cpu_init(nvid_ctx* ctx) break; }; - CUDA_CHECK(ctx->device_id, cudaDeviceSetCacheConfig(cudaFuncCachePreferL1)); + const int gpuArch = ctx->device_arch[0] * 10 + ctx->device_arch[1]; + + /* Disable L1 cache for GPUs before Volta. + * L1 speed is increased and latency reduced with Volta. + */ + if(gpuArch < 70) + CUDA_CHECK(ctx->device_id, cudaDeviceSetCacheConfig(cudaFuncCachePreferL1)); size_t hashMemSize; if(::jconf::inst()->IsCurrencyMonero()) |
