GOSTSec/ccminer
1
0
Fork 0
mirror of https://github.com/GOSTSec/ccminer synced 2025-01-10 14:57:53 +00:00
Code Issues Releases Wiki Activity

lyra2/v2: fixes for vstudio

Browse Source
This commit is contained in:
Tanguy Pruvot 2015-10-13 01:40:47 +02:00
parent 5bf1f98200
commit 32f212469b
6 changed files with 56 additions and 39 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
ccminer.vcxproj
View File

@ -246,6 +246,7 @@
<ClCompile Include="myriadgroestl.cpp" /> <ClCompile Include="myriadgroestl.cpp" />
<ClCompile Include="lyra2\Lyra2.c" /> <ClCompile Include="lyra2\Lyra2.c" />
<ClCompile Include="lyra2\Sponge.c" /> <ClCompile Include="lyra2\Sponge.c" />
<ClInclude Include="lyra2\cuda_lyra2_sm2.cuh" />
<ClInclude Include="neoscrypt\neoscrypt.h" /> <ClInclude Include="neoscrypt\neoscrypt.h" />
<ClCompile Include="neoscrypt\neoscrypt.cpp" /> <ClCompile Include="neoscrypt\neoscrypt.cpp" />
<ClCompile Include="neoscrypt\neoscrypt-cpu.c" /> <ClCompile Include="neoscrypt\neoscrypt-cpu.c" />

36
ccminer.vcxproj.filters
View File

@ -79,6 +79,9 @@
<Filter Include="Source Files\neoscrypt"> <Filter Include="Source Files\neoscrypt">
<UniqueIdentifier>{2037fd0d-e7a2-4da8-956c-150aec726a99}</UniqueIdentifier> <UniqueIdentifier>{2037fd0d-e7a2-4da8-956c-150aec726a99}</UniqueIdentifier>
</Filter> </Filter>
<Filter Include="Source Files\CUDA\lyra2">
<UniqueIdentifier>{1613763f-895c-4321-b58b-6f5849868956}</UniqueIdentifier>
</Filter>
</ItemGroup> </ItemGroup>
<ItemGroup> <ItemGroup>
<ClCompile Include="compat\jansson\dump.c"> <ClCompile Include="compat\jansson\dump.c">
@ -419,6 +422,12 @@
<ClInclude Include="bignum.hpp"> <ClInclude Include="bignum.hpp">
<Filter>Header Files</Filter> <Filter>Header Files</Filter>
</ClInclude> </ClInclude>
<ClInclude Include="lyra2\cuda_lyra2v2_sm3.cuh">
<Filter>Source Files\CUDA\lyra2</Filter>
</ClInclude>
<ClInclude Include="lyra2\cuda_lyra2_sm2.cuh">
<Filter>Source Files\CUDA\lyra2</Filter>
</ClInclude>
</ItemGroup> </ItemGroup>
<ItemGroup> <ItemGroup>
<CudaCompile Include="cuda.cpp"> <CudaCompile Include="cuda.cpp">
@ -607,21 +616,6 @@
<CudaCompile Include="Algo256\cuda_skein256.cu"> <CudaCompile Include="Algo256\cuda_skein256.cu">
<Filter>Source Files\CUDA\Algo256</Filter> <Filter>Source Files\CUDA\Algo256</Filter>
</CudaCompile> </CudaCompile>
<CudaCompile Include="lyra2\cuda_lyra2.cu">
<Filter>Source Files\CUDA</Filter>
</CudaCompile>
<CudaCompile Include="lyra2\lyra2RE.cu">
<Filter>Source Files\CUDA</Filter>
</CudaCompile>
<CudaCompile Include="lyra2\cuda_lyra2v2.cu">
<Filter>Source Files\CUDA</Filter>
</CudaCompile>
<CudaCompile Include="lyra2\lyra2REv2.cu">
<Filter>Source Files\CUDA</Filter>
</CudaCompile>
<ClInclude Include="lyra2\cuda_lyra2v2_sm3.cuh">
<Filter>Source Files\CUDA</Filter>
</ClInclude>
<CudaCompile Include="zr5.cu"> <CudaCompile Include="zr5.cu">
<Filter>Source Files\CUDA</Filter> <Filter>Source Files\CUDA</Filter>
</CudaCompile> </CudaCompile>
@ -670,6 +664,18 @@
<CudaCompile Include="Algo256\cuda_bmw.cu"> <CudaCompile Include="Algo256\cuda_bmw.cu">
<Filter>Source Files\CUDA\Algo256</Filter> <Filter>Source Files\CUDA\Algo256</Filter>
</CudaCompile> </CudaCompile>
<CudaCompile Include="lyra2\cuda_lyra2.cu">
<Filter>Source Files\CUDA\lyra2</Filter>
</CudaCompile>
<CudaCompile Include="lyra2\cuda_lyra2v2.cu">
<Filter>Source Files\CUDA\lyra2</Filter>
</CudaCompile>
<CudaCompile Include="lyra2\lyra2RE.cu">
<Filter>Source Files\CUDA\lyra2</Filter>
</CudaCompile>
<CudaCompile Include="lyra2\lyra2REv2.cu">
<Filter>Source Files\CUDA\lyra2</Filter>
</CudaCompile>
</ItemGroup> </ItemGroup>
<ItemGroup> <ItemGroup>
<Image Include="res\ccminer.ico"> <Image Include="res\ccminer.ico">

5
lyra2/cuda_lyra2.cu
View File

@ -11,6 +11,11 @@
#include "cuda_lyra2_sm2.cuh" #include "cuda_lyra2_sm2.cuh"
#ifdef __INTELLISENSE__
/* just for vstudio code colors */
#define __CUDA_ARCH__ 500
#endif
#if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 500 #if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 500
#include "cuda_lyra2_vectors.h" #include "cuda_lyra2_vectors.h"

8
lyra2/cuda_lyra2_sm2.cuh
View File

@ -10,9 +10,8 @@
#define TPB30 160 #define TPB30 160
#if __CUDA_ARCH__ >= 200 && __CUDA_ARCH__ <= 350 #if (__CUDA_ARCH__ >= 200 && __CUDA_ARCH__ <= 350) || !defined(__CUDA_ARCH__)
__constant__ static uint2 blake2b_IV[8] = {
static __constant__ uint2 blake2b_IV[8] = {
{ 0xf3bcc908, 0x6a09e667 }, { 0xf3bcc908, 0x6a09e667 },
{ 0x84caa73b, 0xbb67ae85 }, { 0x84caa73b, 0xbb67ae85 },
{ 0xfe94f82b, 0x3c6ef372 }, { 0xfe94f82b, 0x3c6ef372 },
@ -22,6 +21,9 @@ static __constant__ uint2 blake2b_IV[8] = {
{ 0xfb41bd6b, 0x1f83d9ab }, { 0xfb41bd6b, 0x1f83d9ab },
{ 0x137e2179, 0x5be0cd19 } { 0x137e2179, 0x5be0cd19 }
}; };
#endif
#if __CUDA_ARCH__ >= 200 && __CUDA_ARCH__ <= 350
#define reduceDuplexRow(rowIn, rowInOut, rowOut) { \ #define reduceDuplexRow(rowIn, rowInOut, rowOut) { \
for (int i = 0; i < 8; i++) { \ for (int i = 0; i < 8; i++) { \

39
lyra2/cuda_lyra2v2.cu
View File

@ -2,27 +2,26 @@
#include <stdint.h> #include <stdint.h>
#include <memory.h> #include <memory.h>
#ifdef __INTELLISENSE__
/* just for vstudio code colors */
#define __CUDA_ARCH__ 500
#endif
#define TPB52 10 #define TPB52 10
#define TPB50 16 #define TPB50 16
#include "cuda_lyra2v2_sm3.cuh" #include "cuda_lyra2v2_sm3.cuh"
#ifdef __INTELLISENSE__
/* just for vstudio code colors */
#define __CUDA_ARCH__ 500
#endif
#if __CUDA_ARCH__ >= 500 #if __CUDA_ARCH__ >= 500
#include "cuda_lyra2_vectors.h" #include "cuda_lyra2_vectors.h"
#define Nrow 4 #define Nrow 4
#define Ncol 4 #define Ncol 4
#define u64type uint2 #define uint4x2 uint28
#define vectype uint28
#define memshift 3 #define memshift 3
__device__ vectype *DMatrix; __device__ uint4x2 *DMatrix;
__device__ __forceinline__ __device__ __forceinline__
void Gfunc_v5(uint2 &a, uint2 &b, uint2 &c, uint2 &d) void Gfunc_v5(uint2 &a, uint2 &b, uint2 &c, uint2 &d)
@ -34,7 +33,7 @@ void Gfunc_v5(uint2 &a, uint2 &b, uint2 &c, uint2 &d)
} }
__device__ __forceinline__ __device__ __forceinline__
void round_lyra_v5(vectype* s) void round_lyra_v5(uint4x2* s)
{ {
Gfunc_v5(s[0].x, s[1].x, s[2].x, s[3].x); Gfunc_v5(s[0].x, s[1].x, s[2].x, s[3].x);
Gfunc_v5(s[0].y, s[1].y, s[2].y, s[3].y); Gfunc_v5(s[0].y, s[1].y, s[2].y, s[3].y);
@ -48,9 +47,9 @@ void round_lyra_v5(vectype* s)
} }
__device__ __forceinline__ __device__ __forceinline__
void reduceDuplex(vectype state[4], uint32_t thread) void reduceDuplex(uint4x2 state[4], uint32_t thread)
{ {
vectype state1[3]; uint4x2 state1[3];
uint32_t ps1 = (Nrow * Ncol * memshift * thread); uint32_t ps1 = (Nrow * Ncol * memshift * thread);
uint32_t ps2 = (memshift * (Ncol-1) + memshift * Ncol + Nrow * Ncol * memshift * thread); uint32_t ps2 = (memshift * (Ncol-1) + memshift * Ncol + Nrow * Ncol * memshift * thread);
@ -81,7 +80,7 @@ void reduceDuplex(vectype state[4], uint32_t thread)
} }
__device__ __forceinline__ __device__ __forceinline__
void reduceDuplex50(vectype state[4], uint32_t thread) void reduceDuplex50(uint4x2 state[4], uint32_t thread)
{ {
uint32_t ps1 = (Nrow * Ncol * memshift * thread); uint32_t ps1 = (Nrow * Ncol * memshift * thread);
uint32_t ps2 = (memshift * (Ncol - 1) + memshift * Ncol + Nrow * Ncol * memshift * thread); uint32_t ps2 = (memshift * (Ncol - 1) + memshift * Ncol + Nrow * Ncol * memshift * thread);
@ -105,9 +104,9 @@ void reduceDuplex50(vectype state[4], uint32_t thread)
} }
__device__ __forceinline__ __device__ __forceinline__
void reduceDuplexRowSetupV2(const int rowIn, const int rowInOut, const int rowOut, vectype state[4], uint32_t thread) void reduceDuplexRowSetupV2(const int rowIn, const int rowInOut, const int rowOut, uint4x2 state[4], uint32_t thread)
{ {
vectype state2[3], state1[3]; uint4x2 state2[3], state1[3];
uint32_t ps1 = (memshift * Ncol * rowIn + Nrow * Ncol * memshift * thread); uint32_t ps1 = (memshift * Ncol * rowIn + Nrow * Ncol * memshift * thread);
uint32_t ps2 = (memshift * Ncol * rowInOut + Nrow * Ncol * memshift * thread); uint32_t ps2 = (memshift * Ncol * rowInOut + Nrow * Ncol * memshift * thread);
@ -152,7 +151,7 @@ void reduceDuplexRowSetupV2(const int rowIn, const int rowInOut, const int rowOu
#pragma unroll #pragma unroll
for (int j = 0; j < 3; j++) for (int j = 0; j < 3; j++)
{ {
vectype tmp = state1[j] + state2[j]; uint4x2 tmp = state1[j] + state2[j];
state[j] ^= tmp; state[j] ^= tmp;
} }
@ -180,9 +179,9 @@ void reduceDuplexRowSetupV2(const int rowIn, const int rowInOut, const int rowOu
__device__ __forceinline__ __device__ __forceinline__
void reduceDuplexRowtV2(const int rowIn, const int rowInOut, const int rowOut, vectype* state, uint32_t thread) void reduceDuplexRowtV2(const int rowIn, const int rowInOut, const int rowOut, uint4x2* state, uint32_t thread)
{ {
vectype state1[3],state2[3]; uint4x2 state1[3],state2[3];
uint32_t ps1 = (memshift * Ncol * rowIn + Nrow * Ncol * memshift * thread); uint32_t ps1 = (memshift * Ncol * rowIn + Nrow * Ncol * memshift * thread);
uint32_t ps2 = (memshift * Ncol * rowInOut + Nrow * Ncol * memshift * thread); uint32_t ps2 = (memshift * Ncol * rowInOut + Nrow * Ncol * memshift * thread);
uint32_t ps3 = (memshift * Ncol * rowOut + Nrow * Ncol * memshift * thread); uint32_t ps3 = (memshift * Ncol * rowOut + Nrow * Ncol * memshift * thread);
@ -260,9 +259,9 @@ void lyra2v2_gpu_hash_32(uint32_t threads, uint32_t startNounce, uint2 *outputHa
{ {
const uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x); const uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x);
vectype state[4]; uint4x2 state[4];
uint28 blake2b_IV[2]; uint4x2 blake2b_IV[2];
if (threadIdx.x == 0) { if (threadIdx.x == 0) {
@ -344,7 +343,9 @@ void lyra2v2_gpu_hash_32(uint32_t threads, uint32_t startNounce, uint2 *outputHa
} }
#else #else
#include "cuda_helper.h" #include "cuda_helper.h"
#if __CUDA_ARCH__ < 300
__device__ void* DMatrix; __device__ void* DMatrix;
#endif
__global__ void lyra2v2_gpu_hash_32(uint32_t threads, uint32_t startNounce, uint2 *outputHash) {} __global__ void lyra2v2_gpu_hash_32(uint32_t threads, uint32_t startNounce, uint2 *outputHash) {}
#endif #endif

2
lyra2/cuda_lyra2v2_sm3.cuh
View File

@ -11,6 +11,8 @@
#if __CUDA_ARCH__ >= 300 && __CUDA_ARCH__ < 500 #if __CUDA_ARCH__ >= 300 && __CUDA_ARCH__ < 500
#include "cuda_lyra2_vectors.h"
#define Nrow 4 #define Nrow 4
#define Ncol 4 #define Ncol 4