diff --git a/ccminer.vcxproj b/ccminer.vcxproj
index c0588b1..74c3aff 100644
--- a/ccminer.vcxproj
+++ b/ccminer.vcxproj
@@ -314,6 +314,7 @@
     <ClInclude Include="algos.h" />
     <ClInclude Include="miner.h" />
     <ClInclude Include="nvml.h" />
+    <ClInclude Include="quark\cuda_bmw512_sm3.cuh" />
     <ClInclude Include="quark\cuda_quark_blake512_sp.cuh" />
     <ClInclude Include="quark\cuda_skein512_sp.cuh" />
     <ClInclude Include="res\resource.h" />
@@ -527,4 +528,4 @@
   <Target Name="AfterClean">
     <Delete Files="@(FilesToCopy->'$(OutDir)%(Filename)%(Extension)')" TreatErrorsAsWarnings="true" />
   </Target>
-</Project>
+</Project>
\ No newline at end of file
diff --git a/ccminer.vcxproj.filters b/ccminer.vcxproj.filters
index b123622..f27161a 100644
--- a/ccminer.vcxproj.filters
+++ b/ccminer.vcxproj.filters
@@ -434,6 +434,9 @@
     <ClInclude Include="quark\cuda_skein512_sp.cuh">
       <Filter>Source Files\CUDA\quark</Filter>
     </ClInclude>
+    <ClInclude Include="quark\cuda_bmw512_sm3.cuh">
+      <Filter>Source Files\CUDA\quark</Filter>
+    </ClInclude>
   </ItemGroup>
   <ItemGroup>
     <CudaCompile Include="cuda.cpp">
@@ -698,4 +701,4 @@
       <Filter>Ressources</Filter>
     </Text>
   </ItemGroup>
-</Project>
+</Project>
\ No newline at end of file
diff --git a/quark/cuda_bmw512.cu b/quark/cuda_bmw512.cu
index 1c7180c..8910ee8 100644
--- a/quark/cuda_bmw512.cu
+++ b/quark/cuda_bmw512.cu
@@ -9,7 +9,12 @@
 __constant__ uint64_t c_PaddedMessage80[16]; // padded message (80 bytes + padding)
 #endif
 
-#include "cuda_bmw512_30.cu"
+#include "cuda_bmw512_sm3.cuh"
+
+#ifdef __INTELLISENSE__
+/* just for vstudio code colors */
+#define __CUDA_ARCH__ 500
+#endif
 
 #undef SHL
 #undef SHR
diff --git a/quark/cuda_bmw512_30.cu b/quark/cuda_bmw512_sm3.cuh
similarity index 97%
rename from quark/cuda_bmw512_30.cu
rename to quark/cuda_bmw512_sm3.cuh
index 5b204e3..247c9f0 100644
--- a/quark/cuda_bmw512_30.cu
+++ b/quark/cuda_bmw512_sm3.cuh
@@ -1,6 +1,11 @@
 #include <stdio.h>
 #include <memory.h>
 
+#ifdef __INTELLISENSE__
+/* for vstudio code colors */
+#define __CUDA_ARCH__ 300
+#endif
+
 #include "cuda_helper.h"
 
 #define SHL(x, n)            ((x) << (n))
@@ -12,8 +17,7 @@
 	q[i+8] + ROTL64(q[i+9], 37) + q[i+10] + ROTL64(q[i+11], 43) + \
 	q[i+12] + ROTL64(q[i+13], 53) + (SHR(q[i+14],1) ^ q[i+14]) + (SHR(q[i+15],2) ^ q[i+15])
 
-#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 500
-
+#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 200 && __CUDA_ARCH__ < 500)
 __constant__ uint64_t d_constMem[16] = {
 	SPH_C64(0x8081828384858687),
 	SPH_C64(0x88898A8B8C8D8E8F),
@@ -32,6 +36,13 @@ __constant__ uint64_t d_constMem[16] = {
 	SPH_C64(0xF0F1F2F3F4F5F6F7),
 	SPH_C64(0xF8F9FAFBFCFDFEFF)
 };
+# ifdef __GNUC__
+// windows and linux doesnt require the same ifdef for __constant__
+#  pragma GCC diagnostic ignored "-Wunused-variable"
+# endif
+#endif
+
+#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 500
 
 __device__
 void Compression512_30(uint64_t *msg, uint64_t *hash)