github-jwinterm/xmr-stak
1
0
Fork 0
mirror of https://github.com/jwinterm/xmr-stak.git synced 2026-10-18 05:51:45 +02:00
Code Issues Packages Projects Releases Wiki Activity

Merge pull request #1014 from psychocrypt/topic-speedupVolta

Browse Source 
speedup Volta
This commit is contained in:
fireice-uk
2018-02-03 10:30:16 +00:00
committed by GitHub
parent 4d7fdd2dea c528f51a76
commit 8bef63c913
2 changed files with 19 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files
xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
+12
View File
@@ -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>
xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu
+7 -1
View File
@@ -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())