From b55acb719b933a5df9812f7af440cd3fc5c57da2 Mon Sep 17 00:00:00 2001
From: havenprotocol <havencurrency@gmail.com>
Date: Sun, 10 Jun 2018 21:35:24 +0200
Subject: [PATCH] Add support for CryptoNight Haven (small Heavy tweak)
- update pools.txt
- add new algorithm `cryptonight_haven`
- update all backends
---
.../backend/amd/amd_gpu/opencl/cryptonight.cl | 24 ++++++++---
.../backend/cpu/crypto/cryptonight_aesni.h | 43 +++++++++++++++++--
xmrstak/backend/cpu/minethd.cpp | 31 ++++++++++++-
xmrstak/backend/cryptonight.hpp | 20 +++++++--
xmrstak/backend/nvidia/nvcc_code/cuda_core.cu | 25 ++++++++---
.../backend/nvidia/nvcc_code/cuda_extra.cu | 25 ++++++++---
xmrstak/jconf.cpp | 3 +-
xmrstak/net/jpsock.cpp | 7 ++-
xmrstak/pools.tpl | 3 +-
9 files changed, 150 insertions(+), 31 deletions(-)
diff --git a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl
index 709a365..7bbc386 100644
--- a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl
+++ b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl
@@ -513,8 +513,8 @@ __kernel void JOIN(cn0,ALGO)(__global ulong *input, __global uint4 *Scratchpad,
mem_fence(CLK_LOCAL_MEM_FENCE);
-// cryptonight_heavy
-#if (ALGO == 4)
+// cryptonight_heavy or cryptonight_haven
+#if (ALGO == 4 || ALGO == 9)
__local uint4 xin[8][WORKSIZE];
/* Also left over threads perform this loop.
@@ -668,6 +668,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states
idx0 = a[0];
b_x = ((uint4 *)c)[0];
+
// cryptonight_heavy
#if (ALGO == 4)
long n = *((__global long*)(Scratchpad + (IDX((idx0 & MASK) >> 4))));
@@ -675,6 +676,14 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states
long q = n / (d | 0x5);
*((__global long*)(Scratchpad + (IDX((idx0 & MASK) >> 4)))) = n ^ q;
idx0 = d ^ q;
+#endif
+// cryptonight_haven
+#if (ALGO == 9)
+ long n = *((__global long*)(Scratchpad + (IDX((idx0 & MASK) >> 4))));
+ int d = ((__global int*)(Scratchpad + (IDX((idx0 & MASK) >> 4))))[2];
+ long q = n / (d | 0x5);
+ *((__global long*)(Scratchpad + (IDX((idx0 & MASK) >> 4)))) = n ^ q;
+ idx0 = (~d) ^ q;
#endif
}
}
@@ -734,8 +743,8 @@ __kernel void JOIN(cn2,ALGO) (__global uint4 *Scratchpad, __global ulong *states
}
barrier(CLK_LOCAL_MEM_FENCE);
-
-#if (ALGO == 4)
+// cryptonight_heavy or cryptonight_haven
+#if (ALGO == 4 || ALGO == 9)
__local uint4 xin[8][WORKSIZE];
#endif
@@ -744,7 +753,8 @@ __kernel void JOIN(cn2,ALGO) (__global uint4 *Scratchpad, __global ulong *states
if(gIdx < Threads)
#endif
{
-#if (ALGO == 4)
+// cryptonight_heavy or cryptonight_haven
+#if (ALGO == 4 || ALGO == 9)
#pragma unroll 2
for(int i = 0; i < (MEMORY >> 7); ++i)
{
@@ -790,8 +800,8 @@ __kernel void JOIN(cn2,ALGO) (__global uint4 *Scratchpad, __global ulong *states
#endif
}
-// cryptonight_heavy
-#if (ALGO == 4)
+// cryptonight_heavy or cryptonight_haven
+#if (ALGO == 4 || ALGO == 9)
/* Also left over threads perform this loop.
* The left over thread results will be ignored
*/
diff --git a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h
index 364e32d..e61381a 100644
--- a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h
+++ b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h
@@ -180,7 +180,7 @@ void cn_explode_scratchpad(const __m128i* input, __m128i* output)
xin6 = _mm_load_si128(input + 10);
xin7 = _mm_load_si128(input + 11);
- if(ALGO == cryptonight_heavy)
+ if(ALGO == cryptonight_heavy || ALGO == cryptonight_haven)
{
for(size_t i=0; i < 16; i++)
{
@@ -324,11 +324,11 @@ void cn_implode_scratchpad(const __m128i* input, __m128i* output)
aes_round(k9, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7);
}
- if(ALGO == cryptonight_heavy)
+ if(ALGO == cryptonight_heavy || ALGO == cryptonight_haven)
mix_and_propagate(xout0, xout1, xout2, xout3, xout4, xout5, xout6, xout7);
}
- if(ALGO == cryptonight_heavy)
+ if(ALGO == cryptonight_heavy || ALGO == cryptonight_haven)
{
for (size_t i = 0; i < MEM / sizeof(__m128i); i += 8)
{
@@ -375,7 +375,7 @@ void cn_implode_scratchpad(const __m128i* input, __m128i* output)
aes_round(k9, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7);
}
- if(ALGO == cryptonight_heavy)
+ if(ALGO == cryptonight_heavy || ALGO == cryptonight_haven)
mix_and_propagate(xout0, xout1, xout2, xout3, xout4, xout5, xout6, xout7);
}
@@ -540,6 +540,15 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c
((int64_t*)&l0[idx0 & MASK])[0] = n ^ q;
idx0 = d ^ q;
}
+ else if(ALGO == cryptonight_haven)
+ {
+ int64_t n = ((int64_t*)&l0[idx0 & MASK])[0];
+ int32_t d = ((int32_t*)&l0[idx0 & MASK])[2];
+ int64_t q = n / (d | 0x5);
+
+ ((int64_t*)&l0[idx0 & MASK])[0] = n ^ q;
+ idx0 = (~d) ^ q;
+ }
}
// Optim - 90% time boundary
@@ -671,6 +680,15 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
((int64_t*)&l0[idx0 & MASK])[0] = n ^ q;
idx0 = d ^ q;
}
+ else if(ALGO == cryptonight_haven)
+ {
+ int64_t n = ((int64_t*)&l0[idx0 & MASK])[0];
+ int32_t d = ((int32_t*)&l0[idx0 & MASK])[2];
+ int64_t q = n / (d | 0x5);
+
+ ((int64_t*)&l0[idx0 & MASK])[0] = n ^ q;
+ idx0 = (~d) ^ q;
+ }
if(PREFETCH)
_mm_prefetch((const char*)&l0[idx0 & MASK], _MM_HINT_T0);
@@ -707,6 +725,15 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
((int64_t*)&l1[idx1 & MASK])[0] = n ^ q;
idx1 = d ^ q;
}
+ else if(ALGO == cryptonight_haven)
+ {
+ int64_t n = ((int64_t*)&l1[idx1 & MASK])[0];
+ int32_t d = ((int32_t*)&l1[idx1 & MASK])[2];
+ int64_t q = n / (d | 0x5);
+
+ ((int64_t*)&l1[idx1 & MASK])[0] = n ^ q;
+ idx1 = (~d) ^ q;
+ }
if(PREFETCH)
_mm_prefetch((const char*)&l1[idx1 & MASK], _MM_HINT_T0);
@@ -768,6 +795,14 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
int64_t q = n / (d | 0x5); \
((int64_t*)&l[idx & MASK])[0] = n ^ q; \
idx = d ^ q; \
+ } \
+ else if(ALGO == cryptonight_haven) \
+ { \
+ int64_t n = ((int64_t*)&l[idx & MASK])[0]; \
+ int32_t d = ((int32_t*)&l[idx & MASK])[2]; \
+ int64_t q = n / (d | 0x5); \
+ ((int64_t*)&l[idx & MASK])[0] = n ^ q; \
+ idx = (~d) ^ q; \
}
#define CONST_INIT(ctx, n) \
diff --git a/xmrstak/backend/cpu/minethd.cpp b/xmrstak/backend/cpu/minethd.cpp
index b6a64fb..7e2a28b 100644
--- a/xmrstak/backend/cpu/minethd.cpp
+++ b/xmrstak/backend/cpu/minethd.cpp
@@ -383,6 +383,9 @@ minethd::cn_hash_fun minethd::func_selector(bool bHaveAes, bool bNoPrefetch, xmr
case cryptonight_masari:
algv = 7;
break;
+ case cryptonight_haven:
+ algv = 8;
+ break;
default:
algv = 2;
break;
@@ -420,7 +423,11 @@ minethd::cn_hash_fun minethd::func_selector(bool bHaveAes, bool bNoPrefetch, xmr
cryptonight_hash<cryptonight_masari, false, false>,
cryptonight_hash<cryptonight_masari, true, false>,
cryptonight_hash<cryptonight_masari, false, true>,
- cryptonight_hash<cryptonight_masari, true, true>
+ cryptonight_hash<cryptonight_masari, true, true>,
+ cryptonight_hash<cryptonight_haven, false, false>,
+ cryptonight_hash<cryptonight_haven, true, false>,
+ cryptonight_hash<cryptonight_haven, false, true>,
+ cryptonight_hash<cryptonight_haven, true, true>
};
std::bitset<2> digit;
@@ -568,6 +575,9 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes,
case cryptonight_masari:
algv = 7;
break;
+ case cryptonight_haven:
+ algv = 8;
+ break;
default:
algv = 2;
break;
@@ -693,7 +703,6 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes,
cryptonight_penta_hash<cryptonight_stellite, false, true>,
cryptonight_penta_hash<cryptonight_stellite, true, true>,
-
cryptonight_double_hash<cryptonight_masari, false, false>,
cryptonight_double_hash<cryptonight_masari, true, false>,
cryptonight_double_hash<cryptonight_masari, false, true>,
@@ -710,6 +719,24 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes,
cryptonight_penta_hash<cryptonight_masari, true, false>,
cryptonight_penta_hash<cryptonight_masari, false, true>,
cryptonight_penta_hash<cryptonight_masari, true, true>,
+
+ cryptonight_double_hash<cryptonight_haven, false, false>,
+ cryptonight_double_hash<cryptonight_haven, true, false>,
+ cryptonight_double_hash<cryptonight_haven, false, true>,
+ cryptonight_double_hash<cryptonight_haven, true, true>,
+ cryptonight_triple_hash<cryptonight_haven, false, false>,
+ cryptonight_triple_hash<cryptonight_haven, true, false>,
+ cryptonight_triple_hash<cryptonight_haven, false, true>,
+ cryptonight_triple_hash<cryptonight_haven, true, true>,
+ cryptonight_quad_hash<cryptonight_haven, false, false>,
+ cryptonight_quad_hash<cryptonight_haven, true, false>,
+ cryptonight_quad_hash<cryptonight_haven, false, true>,
+ cryptonight_quad_hash<cryptonight_haven, true, true>,
+ cryptonight_penta_hash<cryptonight_haven, false, false>,
+ cryptonight_penta_hash<cryptonight_haven, true, false>,
+ cryptonight_penta_hash<cryptonight_haven, false, true>,
+ cryptonight_penta_hash<cryptonight_haven, true, true>
+
};
std::bitset<2> digit;
diff --git a/xmrstak/backend/cryptonight.hpp b/xmrstak/backend/cryptonight.hpp
index 7065570..5953756 100644
--- a/xmrstak/backend/cryptonight.hpp
+++ b/xmrstak/backend/cryptonight.hpp
@@ -13,7 +13,8 @@ enum xmrstak_algo
cryptonight_aeon = 5,
cryptonight_ipbc = 6, // equal to cryptonight_aeon with a small tweak in the miner code
cryptonight_stellite = 7, //equal to cryptonight_monero but with one tiny change
- cryptonight_masari = 8 //equal to cryptonight_monero but with less iterations, used by masari
+ cryptonight_masari = 8, //equal to cryptonight_monero but with less iterations, used by masari
+ cryptonight_haven = 9 // // equal to cryptonight_heavy with a small tweak
};
// define aeon settings
@@ -58,6 +59,8 @@ inline constexpr size_t cn_select_memory<cryptonight_stellite>() { return CRYPTO
template<>
inline constexpr size_t cn_select_memory<cryptonight_masari>() { return CRYPTONIGHT_MEMORY; }
+template<>
+inline constexpr size_t cn_select_memory<cryptonight_haven>() { return CRYPTONIGHT_HEAVY_MEMORY; }
inline size_t cn_select_memory(xmrstak_algo algo)
{
@@ -73,6 +76,7 @@ inline size_t cn_select_memory(xmrstak_algo algo)
case cryptonight_lite:
return CRYPTONIGHT_LITE_MEMORY;
case cryptonight_heavy:
+ case cryptonight_haven:
return CRYPTONIGHT_HEAVY_MEMORY;
default:
return 0;
@@ -106,6 +110,10 @@ inline constexpr uint32_t cn_select_mask<cryptonight_stellite>() { return CRYPTO
template<>
inline constexpr uint32_t cn_select_mask<cryptonight_masari>() { return CRYPTONIGHT_MASK; }
+template<>
+inline constexpr uint32_t cn_select_mask<cryptonight_haven>() { return CRYPTONIGHT_HEAVY_MASK; }
+
+
inline size_t cn_select_mask(xmrstak_algo algo)
{
switch(algo)
@@ -120,6 +128,7 @@ inline size_t cn_select_mask(xmrstak_algo algo)
case cryptonight_lite:
return CRYPTONIGHT_LITE_MASK;
case cryptonight_heavy:
+ case cryptonight_haven:
return CRYPTONIGHT_HEAVY_MASK;
default:
return 0;
@@ -153,6 +162,10 @@ inline constexpr uint32_t cn_select_iter<cryptonight_stellite>() { return CRYPTO
template<>
inline constexpr uint32_t cn_select_iter<cryptonight_masari>() { return CRYPTONIGHT_MASARI_ITER; }
+template<>
+inline constexpr uint32_t cn_select_iter<cryptonight_haven>() { return CRYPTONIGHT_HEAVY_ITER; }
+
+
inline size_t cn_select_iter(xmrstak_algo algo)
{
switch(algo)
@@ -166,9 +179,10 @@ inline size_t cn_select_iter(xmrstak_algo algo)
case cryptonight_lite:
return CRYPTONIGHT_LITE_ITER;
case cryptonight_heavy:
+ case cryptonight_haven:
return CRYPTONIGHT_HEAVY_ITER;
- case cryptonight_masari:
- return CRYPTONIGHT_MASARI_ITER;
+ case cryptonight_masari:
+ return CRYPTONIGHT_MASARI_ITER;
default:
return 0;
}
diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
index 4d6d835..48243e3 100644
--- a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
+++ b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
@@ -242,7 +242,7 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti
a = (d_ctx_a + thread * 4)[sub];
idx0 = shuffle<4>(sPtr,sub, a, 0);
- if(ALGO == cryptonight_heavy)
+ if(ALGO == cryptonight_heavy || ALGO == cryptonight_haven)
{
if(partidx != 0)
{
@@ -341,6 +341,17 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti
idx0 = d ^ q;
}
+ else if(ALGO == cryptonight_haven)
+ {
+ int64_t n = loadGlobal64<uint64_t>( ( (uint64_t *) long_state ) + (( idx0 & MASK ) >> 3));
+ int32_t d = loadGlobal32<uint32_t>( (uint32_t*)(( (uint64_t *) long_state ) + (( idx0 & MASK) >> 3) + 1u ));
+ int64_t q = n / (d | 0x5);
+
+ if(sub&1)
+ storeGlobal64<uint64_t>( ( (uint64_t *) long_state ) + (( idx0 & MASK ) >> 3), n ^ q );
+
+ idx0 = (~d) ^ q;
+ }
}
}
@@ -348,7 +359,7 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti
{
(d_ctx_a + thread * 4)[sub] = a;
(d_ctx_b + thread * 4)[sub] = d[1];
- if(ALGO == cryptonight_heavy)
+ if(ALGO == cryptonight_heavy || ALGO == cryptonight_haven)
if(sub&1)
*(d_ctx_b + threads * 4 + thread) = idx0;
}
@@ -394,7 +405,7 @@ __global__ void cryptonight_core_gpu_phase3( int threads, int bfactor, int parti
cn_aes_pseudo_round_mut( sharedMemory, text, key );
- if(ALGO == cryptonight_heavy)
+ if(ALGO == cryptonight_heavy || ALGO == cryptonight_haven)
{
#pragma unroll
for ( int j = 0; j < 4; ++j )
@@ -431,7 +442,7 @@ void cryptonight_core_gpu_hash(nvid_ctx* ctx, uint32_t nonce)
CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_core_gpu_phase1<ITERATIONS,MEMORY><<< grid, block8 >>>( ctx->device_blocks*ctx->device_threads,
bfactorOneThree, i,
ctx->d_long_state,
- (ALGO == cryptonight_heavy ? ctx->d_ctx_state2 : ctx->d_ctx_state),
+ (ALGO == cryptonight_heavy || ALGO == cryptonight_haven ? ctx->d_ctx_state2 : ctx->d_ctx_state),
ctx->d_ctx_key1 ));
if ( partcount > 1 && ctx->device_bsleep > 0) compat_usleep( ctx->device_bsleep );
@@ -465,7 +476,7 @@ void cryptonight_core_gpu_hash(nvid_ctx* ctx, uint32_t nonce)
int roundsPhase3 = partcountOneThree;
- if(ALGO == cryptonight_heavy)
+ if(ALGO == cryptonight_heavy || ALGO == cryptonight_haven)
{
// cryptonight_heavy used two full rounds over the scratchpad memory
roundsPhase3 *= 2;
@@ -519,5 +530,9 @@ void cryptonight_core_cpu_hash(nvid_ctx* ctx, xmrstak_algo miner_algo, uint32_t
{
cryptonight_core_gpu_hash<CRYPTONIGHT_MASARI_ITER, CRYPTONIGHT_MASK, CRYPTONIGHT_MEMORY/4, cryptonight_masari>(ctx, startNonce);
}
+ else if(miner_algo == cryptonight_haven)
+ {
+ cryptonight_core_gpu_hash<CRYPTONIGHT_HEAVY_ITER, CRYPTONIGHT_HEAVY_MASK, CRYPTONIGHT_HEAVY_MEMORY/4, cryptonight_haven>(ctx, startNonce);
+ }
}
diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu
index 304997e..2cb3702 100644
--- a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu
+++ b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu
@@ -114,7 +114,7 @@ __global__ void cryptonight_extra_gpu_prepare( int threads, uint32_t * __restric
int thread = ( blockDim.x * blockIdx.x + threadIdx.x );
__shared__ uint32_t sharedMemory[1024];
- if(ALGO == cryptonight_heavy)
+ if(ALGO == cryptonight_heavy || ALGO == cryptonight_haven)
{
cn_aes_gpu_init( sharedMemory );
__syncthreads( );
@@ -148,7 +148,7 @@ __global__ void cryptonight_extra_gpu_prepare( int threads, uint32_t * __restric
memcpy( d_ctx_key2 + thread * 40, ctx_key2, 40 * 4 );
memcpy( d_ctx_state + thread * 50, ctx_state, 50 * 4 );
- if(ALGO == cryptonight_heavy)
+ if(ALGO == cryptonight_heavy || ALGO == cryptonight_haven)
{
for(int i=0; i < 16; i++)
@@ -172,7 +172,7 @@ __global__ void cryptonight_extra_gpu_final( int threads, uint64_t target, uint3
__shared__ uint32_t sharedMemory[1024];
- if(ALGO == cryptonight_heavy)
+ if(ALGO == cryptonight_heavy || ALGO == cryptonight_haven)
{
cn_aes_gpu_init( sharedMemory );
__syncthreads( );
@@ -189,7 +189,7 @@ __global__ void cryptonight_extra_gpu_final( int threads, uint64_t target, uint3
for ( i = 0; i < 50; i++ )
state[i] = ctx_state[i];
- if(ALGO == cryptonight_heavy)
+ if(ALGO == cryptonight_heavy || ALGO == cryptonight_haven)
{
uint32_t key[40];
@@ -287,7 +287,7 @@ extern "C" int cryptonight_extra_cpu_init(nvid_ctx* ctx)
size_t wsize = ctx->device_blocks * ctx->device_threads;
CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_ctx_state, 50 * sizeof(uint32_t) * wsize));
size_t ctx_b_size = 4 * sizeof(uint32_t) * wsize;
- if(cryptonight_heavy == ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo())
+ if(cryptonight_heavy == ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() || cryptonight_haven == ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo())
{
// extent ctx_b to hold the state of idx0
ctx_b_size += sizeof(uint32_t) * wsize;
@@ -326,6 +326,11 @@ extern "C" void cryptonight_extra_cpu_prepare(nvid_ctx* ctx, uint32_t startNonce
CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_extra_gpu_prepare<cryptonight_heavy><<<grid, block >>>( wsize, ctx->d_input, ctx->inputlen, startNonce,
ctx->d_ctx_state,ctx->d_ctx_state2, ctx->d_ctx_a, ctx->d_ctx_b, ctx->d_ctx_key1, ctx->d_ctx_key2 ));
}
+ else if(miner_algo == cryptonight_haven)
+ {
+ CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_extra_gpu_prepare<cryptonight_haven><<<grid, block >>>( wsize, ctx->d_input, ctx->inputlen, startNonce,
+ ctx->d_ctx_state,ctx->d_ctx_state2, ctx->d_ctx_a, ctx->d_ctx_b, ctx->d_ctx_key1, ctx->d_ctx_key2 ));
+ }
else
{
/* pass two times d_ctx_state because the second state is used later in phase1,
@@ -355,6 +360,14 @@ extern "C" void cryptonight_extra_cpu_final(nvid_ctx* ctx, uint32_t startNonce,
cryptonight_extra_gpu_final<cryptonight_heavy><<<grid, block >>>( wsize, target, ctx->d_result_count, ctx->d_result_nonce, ctx->d_ctx_state,ctx->d_ctx_key2 )
);
}
+ else if(miner_algo == cryptonight_haven)
+ {
+ CUDA_CHECK_MSG_KERNEL(
+ ctx->device_id,
+ "\n**suggestion: Try to increase the value of the attribute 'bfactor' in the NVIDIA config file.**",
+ cryptonight_extra_gpu_final<cryptonight_haven><<<grid, block >>>( wsize, target, ctx->d_result_count, ctx->d_result_nonce, ctx->d_ctx_state,ctx->d_ctx_key2 )
+ );
+ }
else
{
// fallback for all other algorithms
@@ -612,7 +625,7 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx)
// up to 16kibyte extra memory is used per thread for some kernel (lmem/local memory)
// 680bytes are extra meta data memory per hash
size_t perThread = hashMemSize + 16192u + 680u;
- if(cryptonight_heavy == ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo())
+ if(cryptonight_heavy == ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() || cryptonight_haven == ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo())
perThread += 50 * 4; // state double buffer
size_t max_intensity = limitedMemory / perThread;
diff --git a/xmrstak/jconf.cpp b/xmrstak/jconf.cpp
index c49a176..c35fb51 100644
--- a/xmrstak/jconf.cpp
+++ b/xmrstak/jconf.cpp
@@ -93,6 +93,7 @@ xmrstak::coin_selection coins[] = {
{ "croat", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr },
{ "cryptonight", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr },
{ "cryptonight_masari", {cryptonight_monero, cryptonight_masari, 255u}, {cryptonight_monero, cryptonight_monero, 0u},nullptr },
+ { "cryptonight_haven", {cryptonight_heavy, cryptonight_haven, 255u}, {cryptonight_heavy, cryptonight_heavy, 0u}, nullptr },
{ "cryptonight_heavy", {cryptonight_heavy, cryptonight_heavy, 0u}, {cryptonight_heavy, cryptonight_heavy, 0u}, nullptr },
{ "cryptonight_lite", {cryptonight_aeon, cryptonight_lite, 255u}, {cryptonight_aeon, cryptonight_lite, 7u}, nullptr },
{ "cryptonight_lite_v7", {cryptonight_lite, cryptonight_aeon, 255u}, {cryptonight_aeon, cryptonight_lite, 7u}, nullptr },
@@ -102,7 +103,7 @@ xmrstak::coin_selection coins[] = {
{ "edollar", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr },
{ "electroneum", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr },
{ "graft", {cryptonight_monero, cryptonight, 8u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr },
- { "haven", {cryptonight_heavy, cryptonight, 2u}, {cryptonight_heavy, cryptonight_heavy, 0u}, nullptr },
+ { "haven", {cryptonight_haven, cryptonight_heavy, 3u}, {cryptonight_heavy, cryptonight_heavy, 0u}, nullptr },
{ "intense", {cryptonight_monero, cryptonight, 4u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr },
{ "ipbc", {cryptonight_aeon, cryptonight_ipbc, 255u}, {cryptonight_aeon, cryptonight_aeon, 255u}, nullptr },
{ "karbo", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr },
diff --git a/xmrstak/net/jpsock.cpp b/xmrstak/net/jpsock.cpp
index cf04c86..9fce9b7 100644
--- a/xmrstak/net/jpsock.cpp
+++ b/xmrstak/net/jpsock.cpp
@@ -435,7 +435,7 @@ bool jpsock::process_pool_job(const opq_json_val* params, const uint64_t message
const uint32_t iWorkLen = blob->GetStringLength() / 2;
oPoolJob.iWorkLen = iWorkLen;
-
+
if (iWorkLen > sizeof(pool_job::bWorkBlob))
return set_socket_error("PARSE error: Invalid job length. Are you sure you are mining the correct coin?");
@@ -487,7 +487,7 @@ bool jpsock::process_pool_job(const opq_json_val* params, const uint64_t message
lck.unlock();
// send event after current job data are updated
executor::inst()->push_event(ex_event(oPoolJob, pool_id));
-
+
return true;
}
@@ -697,6 +697,9 @@ bool jpsock::cmd_submit(const char* sJobId, uint32_t iNonce, const uint8_t* bRes
case cryptonight_heavy:
algo_name = "cryptonight_heavy";
break;
+ case cryptonight_haven:
+ algo_name = "cryptonight_haven";
+ break;
case cryptonight_masari:
algo_name = "cryptonight_masari";
break;
diff --git a/xmrstak/pools.tpl b/xmrstak/pools.tpl
index 50299cd..6966d22 100644
--- a/xmrstak/pools.tpl
+++ b/xmrstak/pools.tpl
@@ -26,7 +26,7 @@ POOLCONF],
* edollar
* electroneum
* graft
- * haven
+ * haven (automatic switch with block version 3 to cryptonight_haven)
* intense
* ipbc
* karbo
@@ -45,6 +45,7 @@ POOLCONF],
* cryptonight
* cryptonight_v7
* # 4MiB scratchpad memory
+ * cryptonight_haven
* cryptonight_heavy
*/
--
GitLab