add support for IPBC coin

- add algorithm `cryptonight_lite_v7_xor`
- update documentation

10 files changed, 123 insertions, 34 deletions

diff --git a/xmrstak/backend/amd/amd_gpu/gpu.cpp b/xmrstak/backend/amd/amd_gpu/gpu.cpp
index 03100d0..e44120b 100644
--- a/xmrstak/backend/amd/amd_gpu/gpu.cpp
+++ b/xmrstak/backend/amd/amd_gpu/gpu.cpp
@@ -1004,7 +1004,7 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar
 		return(ERR_OCL_API);
 	}
 
-	if(miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon )
+	if(miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon || miner_algo == cryptonight_ipbc)
 	{
 		// Input
 		if ((ret = clSetKernelArg(ctx->Kernels[kernel_storage][1], 3, sizeof(cl_mem), &ctx->InputBuffer)) != CL_SUCCESS)
diff --git a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl
index d2ae1a7..0738c04 100644
--- a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl
+++ b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl
@@ -553,8 +553,8 @@ __kernel void JOIN(cn0,ALGO)(__global ulong *input, __global uint4 *Scratchpad,
 
 __attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
 __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states, ulong Threads
-// cryptonight_monero || cryptonight_aeon
-#if(ALGO == 3 || ALGO == 5)
+// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc
+#if(ALGO == 3 || ALGO == 5 || ALGO == 6)
 , __global ulong *input
 #endif
 )
@@ -574,7 +574,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states
 	}
 
 	barrier(CLK_LOCAL_MEM_FENCE);
-#if(ALGO == 3 || ALGO == 5)
+#if(ALGO == 3 || ALGO == 5 || ALGO == 6)
     uint2 tweak1_2;
 #endif
 	uint4 b_x;
@@ -598,7 +598,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states
 		b[1] = states[3] ^ states[7];
 
 		b_x = ((uint4 *)b)[0];
-#if(ALGO == 3 || ALGO == 5)
+#if(ALGO == 3 || ALGO == 5 || ALGO == 6)
 		tweak1_2 = as_uint2(input[4]);
 		tweak1_2.s0 >>= 24;
 		tweak1_2.s0 |= tweak1_2.s1 << 8;
@@ -625,7 +625,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states
 			((uint4 *)c)[0] = AES_Round(AES0, AES1, AES2, AES3, ((uint4 *)c)[0], ((uint4 *)a)[0]);
 
 			b_x ^= ((uint4 *)c)[0];
-#if(ALGO == 3 || ALGO == 5)
+#if(ALGO == 3 || ALGO == 5 || ALGO == 6)
 			uint table = 0x75310U;
 			uint index = ((b_x.s2 >> 26) & 12) | ((b_x.s2 >> 23) & 2);
 			b_x.s2 ^= ((table >> index) & 0x30U) << 24;
@@ -639,10 +639,19 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states
 			a[0] += mul_hi(c[0], as_ulong2(tmp).s0);
 
 
-#if(ALGO == 3 || ALGO == 5)
+#if(ALGO == 3 || ALGO == 5 || ALGO == 6)
+
+#	if(ALGO == 6)
+			uint2 ipbc_tmp = tweak1_2 ^ ((uint2 *)&(a[0]))[0];
+			((uint2 *)&(a[1]))[0] ^= ipbc_tmp;
+			Scratchpad[IDX((c[0] & MASK) >> 4)] = ((uint4 *)a)[0];
+			((uint2 *)&(a[1]))[0] ^= ipbc_tmp;
+#	else
 			((uint2 *)&(a[1]))[0] ^= tweak1_2;
 			Scratchpad[IDX((c[0] & MASK) >> 4)] = ((uint4 *)a)[0];
 			((uint2 *)&(a[1]))[0] ^= tweak1_2;
+#	endif
+
 #else
 			Scratchpad[IDX((c[0] & MASK) >> 4)] = ((uint4 *)a)[0];
 #endif
diff --git a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h
index 7562de1..e2f4d88 100644
--- a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h
+++ b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h
@@ -444,7 +444,7 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c
 	constexpr size_t ITERATIONS = cn_select_iter<ALGO>();
 	constexpr size_t MEM = cn_select_memory<ALGO>();
 
-	if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon) && len < 43)
+	if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) && len < 43)
 	{
 		memset(output, 0, 32);
 		return;
@@ -453,7 +453,7 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c
 	keccak((const uint8_t *)input, len, ctx0->hash_state, 200);
 
 	uint64_t monero_const;
-	if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
+	if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc)
 	{
 		monero_const  =  *reinterpret_cast<const uint64_t*>(reinterpret_cast<const uint8_t*>(input) + 35);
 		monero_const ^=  *(reinterpret_cast<const uint64_t*>(ctx0->hash_state) + 24);
@@ -482,7 +482,7 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c
 		else
 			cx = _mm_aesenc_si128(cx, _mm_set_epi64x(ah0, al0));
 
-		if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
+		if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc)
 			cryptonight_monero_tweak((uint64_t*)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx));
 		else
 			_mm_store_si128((__m128i *)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx));
@@ -506,8 +506,13 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c
 			_mm_prefetch((const char*)&l0[al0 & MASK], _MM_HINT_T0);
 		ah0 += lo;
 
-		if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
-			((uint64_t*)&l0[idx0 & MASK])[1] = ah0 ^ monero_const;
+		if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc)
+		{
+			if(ALGO == cryptonight_ipbc)
+				((uint64_t*)&l0[idx0 & MASK])[1] = ah0 ^ monero_const ^ ((uint64_t*)&l0[idx0 & MASK])[0];
+			else
+				((uint64_t*)&l0[idx0 & MASK])[1] = ah0 ^ monero_const;
+		}
 		else
 			((uint64_t*)&l0[idx0 & MASK])[1] = ah0;
 		ah0 ^= ch;
@@ -544,7 +549,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
 	constexpr size_t ITERATIONS = cn_select_iter<ALGO>();
 	constexpr size_t MEM = cn_select_memory<ALGO>();
 
-	if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon) && len < 43)
+	if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) && len < 43)
 	{
 		memset(output, 0, 64);
 		return;
@@ -554,7 +559,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
 	keccak((const uint8_t *)input+len, len, ctx[1]->hash_state, 200);
 
 	uint64_t monero_const_0, monero_const_1;
-	if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
+	if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc)
 	{
 		monero_const_0  =  *reinterpret_cast<const uint64_t*>(reinterpret_cast<const uint8_t*>(input) + 35);
 		monero_const_0 ^=  *(reinterpret_cast<const uint64_t*>(ctx[0]->hash_state) + 24);
@@ -592,7 +597,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
 		else
 			cx = _mm_aesenc_si128(cx, _mm_set_epi64x(axh0, axl0));
 
-		if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
+		if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc)
 			cryptonight_monero_tweak((uint64_t*)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx));
 		else
 			_mm_store_si128((__m128i *)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx));
@@ -610,7 +615,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
 		else
 			cx = _mm_aesenc_si128(cx, _mm_set_epi64x(axh1, axl1));
 
-		if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
+		if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc)
 			cryptonight_monero_tweak((uint64_t*)&l1[idx1 & MASK], _mm_xor_si128(bx1, cx));
 		else
 			_mm_store_si128((__m128i *)&l1[idx1 & MASK], _mm_xor_si128(bx1, cx));
@@ -631,8 +636,13 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
 		axh0 += lo;
 		((uint64_t*)&l0[idx0 & MASK])[0] = axl0;
 
-		if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
-			((uint64_t*)&l0[idx0 & MASK])[1] = axh0 ^ monero_const_0;
+		if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc)
+		{
+			if(ALGO == cryptonight_ipbc)
+				((uint64_t*)&l0[idx0 & MASK])[1] = axh0 ^ monero_const_0 ^ ((uint64_t*)&l0[idx0 & MASK])[0];
+			else
+				((uint64_t*)&l0[idx0 & MASK])[1] = axh0 ^ monero_const_0;
+		}
 		else
 			((uint64_t*)&l0[idx0 & MASK])[1] = axh0;
 
@@ -662,8 +672,13 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
 		axh1 += lo;
 		((uint64_t*)&l1[idx1 & MASK])[0] = axl1;
 
-		if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
-			((uint64_t*)&l1[idx1 & MASK])[1] = axh1 ^ monero_const_1;
+		if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc)
+		{
+			if(ALGO == cryptonight_ipbc)
+				((uint64_t*)&l1[idx1 & MASK])[1] = axh1 ^ monero_const_1 ^ ((uint64_t*)&l1[idx1 & MASK])[0];
+			else
+				((uint64_t*)&l1[idx1 & MASK])[1] = axh1 ^ monero_const_1;
+		}
 		else
 			((uint64_t*)&l1[idx1 & MASK])[1] = axh1;
 
@@ -709,7 +724,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
 	else							\
 		c = _mm_aesenc_si128(c, a);			\
 	b = _mm_xor_si128(b, c);				\
-	if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) \
+	if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) \
 		cryptonight_monero_tweak((uint64_t*)ptr, b); \
 	else \
 		_mm_store_si128(ptr, b);\
@@ -724,8 +739,12 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
 #define CN_STEP4(a, b, c, l, mc, ptr, idx)				\
 	lo = _umul128(idx, _mm_cvtsi128_si64(b), &hi);		\
 	a = _mm_add_epi64(a, _mm_set_epi64x(lo, hi));		\
-	if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) \
+	if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) \
+	{ \
 		_mm_store_si128(ptr, _mm_xor_si128(a, mc)); \
+		if (ALGO == cryptonight_ipbc) \
+			((uint64_t*)ptr)[1] ^= ((uint64_t*)ptr)[0];\
+	} \
 	else \
 		_mm_store_si128(ptr, a);\
 	a = _mm_xor_si128(a, b); \
@@ -751,7 +770,7 @@ void cryptonight_triple_hash(const void* input, size_t len, void* output, crypto
 	constexpr size_t ITERATIONS = cn_select_iter<ALGO>();
 	constexpr size_t MEM = cn_select_memory<ALGO>();
 
-	if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon) && len < 43)
+	if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) && len < 43)
 	{
 		memset(output, 0, 32 * 3);
 		return;
@@ -845,7 +864,7 @@ void cryptonight_quad_hash(const void* input, size_t len, void* output, cryptoni
 	constexpr size_t ITERATIONS = cn_select_iter<ALGO>();
 	constexpr size_t MEM = cn_select_memory<ALGO>();
 
-	if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon) && len < 43)
+	if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) && len < 43)
 	{
 		memset(output, 0, 32 * 4);
 		return;
@@ -954,7 +973,7 @@ void cryptonight_penta_hash(const void* input, size_t len, void* output, crypton
 	constexpr size_t ITERATIONS = cn_select_iter<ALGO>();
 	constexpr size_t MEM = cn_select_memory<ALGO>();
 
-	if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon) && len < 43)
+	if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) && len < 43)
 	{
 		memset(output, 0, 32 * 5);
 		return;
diff --git a/xmrstak/backend/cpu/minethd.cpp b/xmrstak/backend/cpu/minethd.cpp
index dc89c50..bbb4eec 100644
--- a/xmrstak/backend/cpu/minethd.cpp
+++ b/xmrstak/backend/cpu/minethd.cpp
@@ -286,7 +286,9 @@ bool minethd::self_test()
 	else if(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() == cryptonight_aeon)
 	{
 	}
-
+	else if(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() == cryptonight_ipbc)
+	{
+	}
 	for (int i = 0; i < MAX_N; i++)
 		cryptonight_free_ctx(ctx[i]);
 
@@ -369,6 +371,9 @@ minethd::cn_hash_fun minethd::func_selector(bool bHaveAes, bool bNoPrefetch, xmr
 	case cryptonight_aeon:
 		algv = 4;
 		break;
+	case cryptonight_ipbc:
+		algv = 5;
+		break;
 	default:
 		algv = 2;
 		break;
@@ -394,7 +399,11 @@ minethd::cn_hash_fun minethd::func_selector(bool bHaveAes, bool bNoPrefetch, xmr
 		cryptonight_hash<cryptonight_aeon, false, false>,
 		cryptonight_hash<cryptonight_aeon, true, false>,
 		cryptonight_hash<cryptonight_aeon, false, true>,
-		cryptonight_hash<cryptonight_aeon, true, true>
+		cryptonight_hash<cryptonight_aeon, true, true>,
+		cryptonight_hash<cryptonight_ipbc, false, false>,
+		cryptonight_hash<cryptonight_ipbc, true, false>,
+		cryptonight_hash<cryptonight_ipbc, false, true>,
+		cryptonight_hash<cryptonight_ipbc, true, true>
 	};
 
 	std::bitset<2> digit;
@@ -531,6 +540,9 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes,
 	case cryptonight_aeon:
 		algv = 4;
 		break;
+	case cryptonight_ipbc:
+		algv = 5;
+		break;
 	default:
 		algv = 2;
 		break;
@@ -620,7 +632,24 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes,
 		cryptonight_penta_hash<cryptonight_aeon, false, false>,
 		cryptonight_penta_hash<cryptonight_aeon, true, false>,
 		cryptonight_penta_hash<cryptonight_aeon, false, true>,
-		cryptonight_penta_hash<cryptonight_aeon, true, true>
+		cryptonight_penta_hash<cryptonight_aeon, true, true>,
+
+		cryptonight_double_hash<cryptonight_ipbc, false, false>,
+		cryptonight_double_hash<cryptonight_ipbc, true, false>,
+		cryptonight_double_hash<cryptonight_ipbc, false, true>,
+		cryptonight_double_hash<cryptonight_ipbc, true, true>,
+		cryptonight_triple_hash<cryptonight_ipbc, false, false>,
+		cryptonight_triple_hash<cryptonight_ipbc, true, false>,
+		cryptonight_triple_hash<cryptonight_ipbc, false, true>,
+		cryptonight_triple_hash<cryptonight_ipbc, true, true>,
+		cryptonight_quad_hash<cryptonight_ipbc, false, false>,
+		cryptonight_quad_hash<cryptonight_ipbc, true, false>,
+		cryptonight_quad_hash<cryptonight_ipbc, false, true>,
+		cryptonight_quad_hash<cryptonight_ipbc, true, true>,
+		cryptonight_penta_hash<cryptonight_ipbc, false, false>,
+		cryptonight_penta_hash<cryptonight_ipbc, true, false>,
+		cryptonight_penta_hash<cryptonight_ipbc, false, true>,
+		cryptonight_penta_hash<cryptonight_ipbc, true, true>
 	};
 
 	std::bitset<2> digit;
diff --git a/xmrstak/backend/cryptonight.hpp b/xmrstak/backend/cryptonight.hpp
index 8a8e259..b175cda 100644
--- a/xmrstak/backend/cryptonight.hpp
+++ b/xmrstak/backend/cryptonight.hpp
@@ -10,7 +10,8 @@ enum xmrstak_algo
 	cryptonight_lite = 2,
 	cryptonight_monero = 3,
 	cryptonight_heavy = 4,
-	cryptonight_aeon = 5
+	cryptonight_aeon = 5,
+	cryptonight_ipbc = 6 // equal to cryptonight_aeon with a small tweak in the miner code
 };
 
 // define aeon settings
@@ -44,6 +45,9 @@ inline constexpr size_t cn_select_memory<cryptonight_heavy>() { return CRYPTONIG
 template<>
 inline constexpr size_t cn_select_memory<cryptonight_aeon>() { return CRYPTONIGHT_LITE_MEMORY; }
 
+template<>
+inline constexpr size_t cn_select_memory<cryptonight_ipbc>() { return CRYPTONIGHT_LITE_MEMORY; }
+
 
 inline size_t cn_select_memory(xmrstak_algo algo)
 {
@@ -51,7 +55,8 @@ inline size_t cn_select_memory(xmrstak_algo algo)
 	{
 	case cryptonight_monero:
 	case cryptonight:
-		return CRYPTONIGHT_MEMORY;	
+		return CRYPTONIGHT_MEMORY;
+	case cryptonight_ipbc:
 	case cryptonight_aeon:
 	case cryptonight_lite:
 		return CRYPTONIGHT_LITE_MEMORY;
@@ -80,6 +85,9 @@ inline constexpr uint32_t cn_select_mask<cryptonight_heavy>() { return CRYPTONIG
 template<>
 inline constexpr uint32_t cn_select_mask<cryptonight_aeon>() { return CRYPTONIGHT_LITE_MASK; }
 
+template<>
+inline constexpr uint32_t cn_select_mask<cryptonight_ipbc>() { return CRYPTONIGHT_LITE_MASK; }
+
 inline size_t cn_select_mask(xmrstak_algo algo)
 {
 	switch(algo)
@@ -87,6 +95,7 @@ inline size_t cn_select_mask(xmrstak_algo algo)
 	case cryptonight_monero:
 	case cryptonight:
 		return CRYPTONIGHT_MASK;
+	case cryptonight_ipbc:
 	case cryptonight_aeon:
 	case cryptonight_lite:
 		return CRYPTONIGHT_LITE_MASK;
@@ -115,6 +124,9 @@ inline constexpr uint32_t cn_select_iter<cryptonight_heavy>() { return CRYPTONIG
 template<>
 inline constexpr uint32_t cn_select_iter<cryptonight_aeon>() { return CRYPTONIGHT_LITE_ITER; }
 
+template<>
+inline constexpr uint32_t cn_select_iter<cryptonight_ipbc>() { return CRYPTONIGHT_LITE_ITER; }
+
 inline size_t cn_select_iter(xmrstak_algo algo)
 {
 	switch(algo)
@@ -122,6 +134,7 @@ inline size_t cn_select_iter(xmrstak_algo algo)
 	case cryptonight_monero:
 	case cryptonight:
 		return CRYPTONIGHT_ITER;
+	case cryptonight_ipbc:
 	case cryptonight_aeon:
 	case cryptonight_lite:
 		return CRYPTONIGHT_LITE_ITER;
diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
index 43740d2..3bb910e 100644
--- a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
+++ b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
@@ -231,7 +231,7 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti
 	uint32_t t1[2], t2[2], res;
 
 	uint32_t tweak1_2[2];
-	if (ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
+	if (ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc)
 	{
 		uint32_t * state = d_ctx_state + thread * 50;
 		tweak1_2[0] = (d_input[8] >> 24) | (d_input[9] << 8);
@@ -275,7 +275,7 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti
 			t1[0] = shuffle<4>(sPtr,sub, d[x], 0);
 
 			const uint32_t z = d[0] ^ d[1];
-			if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
+			if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc)
 			{
 				const uint32_t table = 0x75310U;
 				const uint32_t index = ((z >> 26) & 12) | ((z >> 23) & 2);
@@ -304,10 +304,17 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti
 			res = *( (uint64_t *) t2 )  >> ( sub & 1 ? 32 : 0 );
 
 			
-			if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
+			if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc)
 			{
 				const uint32_t tweaked_res = tweak1_2[sub & 1] ^ res;
-				const uint32_t long_state_update = sub2 ? tweaked_res : res;
+				uint32_t long_state_update = sub2 ? tweaked_res : res;
+				
+				if (ALGO == cryptonight_ipbc)
+				{
+					uint32_t value = shuffle<4>(sPtr,sub, long_state_update, sub & 1) ^ long_state_update;
+					long_state_update = sub >= 2 ? value : long_state_update;
+				}
+				
 				storeGlobal32( long_state + j, long_state_update );
 			}
 			else
@@ -492,5 +499,9 @@ void cryptonight_core_cpu_hash(nvid_ctx* ctx, xmrstak_algo miner_algo, uint32_t
 	{
 		cryptonight_core_gpu_hash<CRYPTONIGHT_LITE_ITER, CRYPTONIGHT_LITE_MASK, CRYPTONIGHT_LITE_MEMORY/4, cryptonight_aeon>(ctx, startNonce);
 	}
+	else if(miner_algo == cryptonight_ipbc)
+	{
+		cryptonight_core_gpu_hash<CRYPTONIGHT_LITE_ITER, CRYPTONIGHT_LITE_MASK, CRYPTONIGHT_LITE_MEMORY/4, cryptonight_ipbc>(ctx, startNonce);
+	}
 
 }
diff --git a/xmrstak/jconf.cpp b/xmrstak/jconf.cpp
index 091ee7a..8a81528 100644
--- a/xmrstak/jconf.cpp
+++ b/xmrstak/jconf.cpp
@@ -88,12 +88,14 @@ xmrstak::coin_selection coins[] = {
 	{ "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 },
+	{ "cryptonight_lite_v7_xor", {cryptonight_aeon, cryptonight_ipbc, 255u}, {cryptonight_aeon, cryptonight_aeon, 255u}, nullptr },
 	{ "cryptonight_v7",      {cryptonight_monero, cryptonight_monero, 0u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr },
 	{ "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 },
 	{ "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 },
 	{ "monero7",             {cryptonight_monero, cryptonight_monero, 0u}, {cryptonight_monero, cryptonight_monero, 0u}, "pool.usxmrpool.com:3333" },
 	{ "stellite",            {cryptonight_monero, cryptonight, 3u},        {cryptonight_monero, cryptonight_monero, 0u}, nullptr },
diff --git a/xmrstak/misc/executor.cpp b/xmrstak/misc/executor.cpp
index e7092b9..8808fd7 100644
--- a/xmrstak/misc/executor.cpp
+++ b/xmrstak/misc/executor.cpp
@@ -575,6 +575,7 @@ void executor::ex_main()
 			pools.emplace_front(0, "nervproxy1.ddns.net:4444", "donate", "", "x", 0.0, true, false, "", false);
 		break;
 
+	case cryptonight_ipbc:
 	case cryptonight_aeon:
 	case cryptonight_lite:
 		if(dev_tls)
diff --git a/xmrstak/net/jpsock.cpp b/xmrstak/net/jpsock.cpp
index 74d1c26..14b78cd 100644
--- a/xmrstak/net/jpsock.cpp
+++ b/xmrstak/net/jpsock.cpp
@@ -653,6 +653,9 @@ bool jpsock::cmd_submit(const char* sJobId, uint32_t iNonce, const uint8_t* bRes
 		case cryptonight_aeon:
 			algo_name = "cryptonight-aeonv7";
 			break;
+		case cryptonight_ipbc:
+			algo_name = "cryptonight-ipbc";
+			break;
 		case cryptonight_heavy:
 			algo_name = "cryptonight-heavy";
 			break;
diff --git a/xmrstak/pools.tpl b/xmrstak/pools.tpl
index f0667a1..d6433ee 100644
--- a/xmrstak/pools.tpl
+++ b/xmrstak/pools.tpl
@@ -28,6 +28,7 @@ POOLCONF],
  *    graft
  *    haven
  *    intense
+ *    ipbc
  *    karbo
  *    monero7 (use this for Monero's new PoW)
  *    sumokoin (automatic switch with block version 3 to cryptonight_heavy)
@@ -38,6 +39,7 @@ POOLCONF],
  *    # 1MiB scratchpad memory
  *    cryptonight_lite
  *    cryptonight_lite_v7
+ *    cryptonight_lite_v7_xor (algorithm used by ipbc)
  *    # 2MiB scratchpad memory
  *    cryptonight
  *    cryptonight_v7