summaryrefslogtreecommitdiffstats
path: root/xmrstak
diff options
context:
space:
mode:
authorpsychocrypt <psychocryptHPC@gmail.com>2018-04-04 22:37:13 +0200
committerTimothy Pearson <tpearson@raptorengineering.com>2018-06-04 21:07:11 +0000
commit567b867132d8de886fb890949ed8f61de31c777c (patch)
tree90c37dcd2dced215cd066d8424e816f28f15c3b0 /xmrstak
parenta91df76170f72dd9d391b2e2ed0c5f8f6ffba61e (diff)
downloadxmr-stak-567b867132d8de886fb890949ed8f61de31c777c.zip
xmr-stak-567b867132d8de886fb890949ed8f61de31c777c.tar.gz
NVIDIA: fix sumokoin
sumokoin is broken if `bfactor >= 5` is used (default for windows) sumokoin for `sm_20` is broken due to the missing extern shared memory - call phase3 kernel two times if sumokoin is enabled - create extern shared memory for phase3 kernel
Diffstat (limited to 'xmrstak')
-rw-r--r--xmrstak/backend/nvidia/nvcc_code/cuda_core.cu46
1 files changed, 20 insertions, 26 deletions
diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
index 7aa44e8..43740d2 100644
--- a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
+++ b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
@@ -352,7 +352,7 @@ __global__ void cryptonight_core_gpu_phase3( int threads, int bfactor, int parti
int sub = subv << 2;
const int batchsize = MEMORY >> bfactor;
- const int start = partidx * batchsize;
+ const int start = (partidx % (1 << bfactor)) * batchsize;
const int end = start + batchsize;
if ( thread >= threads )
@@ -365,15 +365,15 @@ __global__ void cryptonight_core_gpu_phase3( int threads, int bfactor, int parti
__syncthreads( );
#if( __CUDA_ARCH__ < 300 )
- extern __shared__ uint32_t shuffleMem[];
- volatile uint32_t* sPtr = (volatile uint32_t*)(shuffleMem + (threadIdx.x& 0xFFFFFFFC));
+ extern __shared__ uint32_t shuffleMem[];
+ volatile uint32_t* sPtr = (volatile uint32_t*)(shuffleMem + (threadIdx.x& 0xFFFFFFF8));
#else
- volatile uint32_t* sPtr = NULL;
+ volatile uint32_t* sPtr = NULL;
#endif
for ( int i = start; i < end; i += 32 )
{
-#pragma unroll
+ #pragma unroll
for ( int j = 0; j < 4; ++j )
text[j] ^= long_state[((IndexType) thread * MEMORY) + ( sub + i + j)];
@@ -381,25 +381,7 @@ __global__ void cryptonight_core_gpu_phase3( int threads, int bfactor, int parti
if(ALGO == cryptonight_heavy)
{
-#pragma unroll
- for ( int j = 0; j < 4; ++j )
- text[j] ^= shuffle<8>(sPtr, subv, text[j], (subv+1)&7);
- }
- }
-
- if(ALGO == cryptonight_heavy)
- {
- __syncthreads( );
-
- for ( int i = start; i < end; i += 32 )
- {
-#pragma unroll
- for ( int j = 0; j < 4; ++j )
- text[j] ^= long_state[((IndexType) thread * MEMORY) + ( sub + i + j)];
-
- cn_aes_pseudo_round_mut( sharedMemory, text, key );
-
-#pragma unroll
+ #pragma unroll
for ( int j = 0; j < 4; ++j )
text[j] ^= shuffle<8>(sPtr, subv, text[j], (subv+1)&7);
}
@@ -466,9 +448,21 @@ void cryptonight_core_gpu_hash(nvid_ctx* ctx, uint32_t nonce)
if ( partcount > 1 && ctx->device_bsleep > 0) compat_usleep( ctx->device_bsleep );
}
- for ( int i = 0; i < partcountOneThree; i++ )
+ int roundsPhase3 = partcountOneThree;
+
+ if(ALGO == cryptonight_heavy)
+ {
+ // cryptonight_heavy used two full rounds over the scratchpad memory
+ roundsPhase3 *= 2;
+ }
+
+ for ( int i = 0; i < roundsPhase3; i++ )
{
- CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_core_gpu_phase3<ITERATIONS,MEMORY, ALGO><<< grid, block8 >>>( ctx->device_blocks*ctx->device_threads,
+ CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_core_gpu_phase3<ITERATIONS,MEMORY, ALGO><<<
+ grid,
+ block8,
+ block8.x * sizeof(uint32_t) * static_cast< int >( ctx->device_arch[0] < 3 )
+ >>>( ctx->device_blocks*ctx->device_threads,
bfactorOneThree, i,
ctx->d_long_state,
ctx->d_ctx_state, ctx->d_ctx_key2 ));
OpenPOWER on IntegriCloud