From 2d83f74a7eade7658e596d0ab3dece49befacd5d Mon Sep 17 00:00:00 2001 From: Tanguy Pruvot Date: Sat, 24 Oct 2015 15:13:35 +0200 Subject: [PATCH] vstudio: special ifdef for the constant (bmw) --- ccminer.vcxproj | 3 ++- ccminer.vcxproj.filters | 5 ++++- quark/cuda_bmw512.cu | 7 ++++++- quark/{cuda_bmw512_30.cu => cuda_bmw512_sm3.cuh} | 15 +++++++++++++-- 4 files changed, 25 insertions(+), 5 deletions(-) rename quark/{cuda_bmw512_30.cu => cuda_bmw512_sm3.cuh} (97%) diff --git a/ccminer.vcxproj b/ccminer.vcxproj index c0588b1..74c3aff 100644 --- a/ccminer.vcxproj +++ b/ccminer.vcxproj @@ -314,6 +314,7 @@ + @@ -527,4 +528,4 @@ - + \ 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 @@ Source Files\CUDA\quark + + Source Files\CUDA\quark + @@ -698,4 +701,4 @@ Ressources - + \ 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 #include +#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)