diff --git a/Makefile.am b/Makefile.am
index cdc75d6..a8d19b5 100644
--- a/Makefile.am
+++ b/Makefile.am
@@ -41,7 +41,7 @@ ccminer_SOURCES	= elist.h miner.h compat.h \
 			  quark/cuda_bmw512.cu quark/cuda_quark_keccak512.cu \
 			  quark/quarkcoin.cu quark/animecoin.cu \
 			  quark/cuda_quark_compactionTest.cu  \
-			  neoscrypt.cpp neoscrypt/neoscrypt.c neoscrypt/cuda_neoscrypt.cu \
+			  neoscrypt/neoscrypt.cpp neoscrypt/neoscrypt-cpu.c neoscrypt/cuda_neoscrypt.cu \
 			  cuda_nist5.cu pentablake.cu skein.cu skein2.cu zr5.cu \
 			  sph/bmw.c sph/blake.c sph/groestl.c sph/jh.c sph/keccak.c sph/skein.c \
 			  sph/cubehash.c sph/echo.c sph/luffa.c sph/sha2.c sph/shavite.c sph/simd.c \
diff --git a/ccminer.vcxproj b/ccminer.vcxproj
index fa6c230..b39472b 100644
--- a/ccminer.vcxproj
+++ b/ccminer.vcxproj
@@ -265,8 +265,8 @@
     <ClCompile Include="myriadgroestl.cpp" />
     <ClCompile Include="lyra2\Lyra2.c" />
     <ClCompile Include="lyra2\Sponge.c" />
-    <ClCompile Include="neoscrypt\neoscrypt.c" />
-    <ClCompile Include="neoscrypt.cpp" />
+    <ClCompile Include="neoscrypt\neoscrypt-cpu.c" />
+    <ClCompile Include="neoscrypt\neoscrypt.cpp" />
     <ClCompile Include="sph\aes_helper.c" />
     <ClCompile Include="sph\blake.c" />
     <ClCompile Include="sph\bmw.c" />
diff --git a/ccminer.vcxproj.filters b/ccminer.vcxproj.filters
index 5e65542..184bc9e 100644
--- a/ccminer.vcxproj.filters
+++ b/ccminer.vcxproj.filters
@@ -120,9 +120,9 @@
     <ClCompile Include="groestlcoin.cpp">
       <Filter>Source Files</Filter>
     </ClCompile>
-    <ClCompile Include="neoscrypt.cpp">
+    <ClCompile Include="neoscrypt\neoscrypt.cpp">
       <Filter>Source Files</Filter>
-    </CudaCompile>
+    </ClCompile>
     <ClCompile Include="sph\aes_helper.c">
       <Filter>Source Files\sph</Filter>
     </ClCompile>
@@ -234,7 +234,7 @@
     <ClCompile Include="scrypt-jane.cpp">
       <Filter>Source Files\CUDA\scrypt</Filter>
     </ClCompile>
-    <ClCompile Include="neoscrypt\neoscrypt.c">
+    <ClCompile Include="neoscrypt\neoscrypt-cpu.c">
       <Filter>Source Files</Filter>
     </ClCompile>
   </ItemGroup>
diff --git a/neoscrypt/neoscrypt.c b/neoscrypt/neoscrypt-cpu.c
similarity index 99%
rename from neoscrypt/neoscrypt.c
rename to neoscrypt/neoscrypt-cpu.c
index 75f1abf..d1ef2ff 100644
--- a/neoscrypt/neoscrypt.c
+++ b/neoscrypt/neoscrypt-cpu.c
@@ -34,8 +34,7 @@
 
 #include "neoscrypt.h"
 
-
-#if (WINDOWS)
+#ifdef WIN32
 /* sizeof(unsigned long) = 4 for MinGW64 */
 typedef unsigned long long ulong;
 #else
@@ -45,11 +44,9 @@ typedef unsigned int  uint;
 typedef unsigned char uchar;
 typedef unsigned int  bool;
 
-
 #define MIN(a, b) ((a) < (b) ? a : b)
 #define MAX(a, b) ((a) > (b) ? a : b)
 
-
 /* SHA-256 */
 
 static const uint32_t sha256_constants[64] = {
diff --git a/neoscrypt.cpp b/neoscrypt/neoscrypt.cpp
similarity index 84%
rename from neoscrypt.cpp
rename to neoscrypt/neoscrypt.cpp
index c24a7e9..f176331 100644
--- a/neoscrypt.cpp
+++ b/neoscrypt/neoscrypt.cpp
@@ -1,10 +1,6 @@
-
-extern "C" {
-#include "neoscrypt/neoscrypt.h"
-}
-
-#include "cuda_helper.h"
+#include <cuda_runtime.h>
 #include "miner.h"
+#include "neoscrypt/neoscrypt.h"
 
 static uint32_t *d_hash[MAX_GPUS] ;
 extern void neoscrypt_setBlockTarget(uint32_t * data, const void *ptarget);
@@ -13,6 +9,8 @@ extern uint32_t neoscrypt_cpu_hash_k4(int stratum, int thr_id, uint32_t threads,
 
 #define SHIFT 130
 
+static bool init[MAX_GPUS] = { 0 };
+
 int scanhash_neoscrypt(int thr_id, uint32_t *pdata, const uint32_t *ptarget, uint32_t max_nonce, unsigned long *hashes_done)
 {
 	const uint32_t first_nonce = pdata[19];
@@ -25,17 +23,21 @@ int scanhash_neoscrypt(int thr_id, uint32_t *pdata, const uint32_t *ptarget, uin
 	throughput = throughput / 32; /* set for max intensity ~= 20 */
 	throughput = min(throughput, max_nonce - first_nonce + 1);
 
-	static bool init[MAX_GPUS] = { 0 };
 	if (!init[thr_id])
 	{
 		cudaSetDevice(device_map[thr_id]);
 		cudaDeviceSetCacheConfig(cudaFuncCachePreferL1);
 
-		CUDA_SAFE_CALL(cudaMalloc(&d_hash[thr_id], 32 * SHIFT * sizeof(uint64_t) * throughput));
+		cudaMalloc(&d_hash[thr_id], 32 * SHIFT * sizeof(uint64_t) * throughput);
 		neoscrypt_cpu_init(thr_id, throughput, d_hash[thr_id]);
 
 		applog(LOG_INFO, "Using %d cuda threads", throughput);
-
+		if (cudaGetLastError() != cudaSuccess) {
+			cudaError_t err = cudaGetLastError();
+		        fprintf(stderr, "Cuda error in func '%s' at line %i : %s.\n",
+			                 __FUNCTION__, __LINE__, cudaGetErrorString(err) );
+			proper_exit(EXIT_FAILURE);
+		}
 		init[thr_id] = true;
 	}