fix windows compile and broken aeon

- fix windows linker error during compile
- fix wrong parameter to call aeon (nvidia-backend)

2 files changed, 24 insertions, 8 deletions

diff --git a/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp b/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp
index 4e0ace7..0cfdaac 100644
--- a/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp
+++ b/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp
@@ -44,6 +44,4 @@ void cryptonight_extra_cpu_prepare(nvid_ctx* ctx, uint32_t startNonce);
 void cryptonight_extra_cpu_final(nvid_ctx* ctx, uint32_t startNonce, uint64_t target, uint32_t* rescount, uint32_t *resnonce);
 }
 
-template<size_t ITERATIONS, size_t THREAD_SHIFT, uint32_t MASK>
-void cryptonight_core_cpu_hash(nvid_ctx* ctx);
-
+void cryptonight_core_cpu_hash(nvid_ctx* ctx, bool mineMonero);
diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
index 5bea230..a92fa8c 100644
--- a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
+++ b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
@@ -1,3 +1,5 @@
+#include "xmrstak/backend/cryptonight.hpp"
+
 #include <stdio.h>
 #include <stdint.h>
 #include <string.h>
@@ -92,7 +94,7 @@ __device__ __forceinline__ void storeGlobal32( T* addr, T const & val )
 	asm volatile( "st.global.cg.u32 [%0], %1;" : : "l"( addr ), "r"( val ) );
 }
 
-template<size_t ITERATIONS, size_t THREAD_SHIFT>
+template<size_t ITERATIONS, uint32_t THREAD_SHIFT>
 __global__ void cryptonight_core_gpu_phase1( int threads, int bfactor, int partidx, uint32_t * __restrict__ long_state, uint32_t * __restrict__ ctx_state, uint32_t * __restrict__ ctx_key1 )
 {
 	__shared__ uint32_t sharedMemory[1024];
@@ -168,7 +170,7 @@ __forceinline__ __device__ uint32_t shuffle(volatile uint32_t* ptr,const uint32_
 #ifdef XMR_STAK_THREADS
 __launch_bounds__( XMR_STAK_THREADS * 4 )
 #endif
-template<size_t ITERATIONS, size_t THREAD_SHIFT, uint32_t MASK>
+template<size_t ITERATIONS, uint32_t THREAD_SHIFT, uint32_t MASK>
 __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int partidx, uint32_t * d_long_state, uint32_t * d_ctx_a, uint32_t * d_ctx_b )
 {
 	__shared__ uint32_t sharedMemory[1024];
@@ -257,7 +259,7 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti
 	}
 }
 
-template<size_t ITERATIONS, size_t THREAD_SHIFT>
+template<size_t ITERATIONS, uint32_t THREAD_SHIFT>
 __global__ void cryptonight_core_gpu_phase3( int threads, int bfactor, int partidx, const uint32_t * __restrict__ long_state, uint32_t * __restrict__ d_ctx_state, uint32_t * __restrict__ d_ctx_key2 )
 {
 	__shared__ uint32_t sharedMemory[1024];
@@ -292,8 +294,8 @@ __global__ void cryptonight_core_gpu_phase3( int threads, int bfactor, int parti
 	MEMCPY8( d_ctx_state + thread * 50 + sub + 16, text, 2 );
 }
 
-template<size_t ITERATIONS, size_t THREAD_SHIFT, uint32_t MASK>
-void cryptonight_core_cpu_hash(nvid_ctx* ctx)
+template<size_t ITERATIONS, uint32_t MASK, uint32_t THREAD_SHIFT>
+void cryptonight_core_gpu_hash(nvid_ctx* ctx)
 {
 	dim3 grid( ctx->device_blocks );
 	dim3 block( ctx->device_threads );
@@ -349,3 +351,19 @@ void cryptonight_core_cpu_hash(nvid_ctx* ctx)
 			ctx->d_ctx_state, ctx->d_ctx_key2 ));
 	}
 }
+
+void cryptonight_core_cpu_hash(nvid_ctx* ctx, bool mineMonero)
+{
+#ifndef CONF_NO_MONERO
+	if(mineMonero)
+	{
+		cryptonight_core_gpu_hash<MONERO_ITER, MONERO_MASK, 19u>(ctx);
+	}
+#endif
+#ifndef CONF_NO_AEON
+	if(!mineMonero)
+	{
+		cryptonight_core_gpu_hash<AEON_ITER, AEON_MASK, 18u>(ctx);
+	}
+#endif
+}