GOSTSec
/
ccminer-gostd-lite
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity
Browse Source

groestl: use sp bitslice enhancement, prepare SM 2.x variant 

todo: simd512 SM 2.x variant (shfl op), and groestl/myriad functions
master
Tanguy Pruvot 10 years ago
parent
29c3db5dc6
commit
b521acb480
9 changed files with 532 additions and 468 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Split View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 6
      Makefile.am
  2. 507
      bitslice_transformations_quad.cu
  3. 2
      configure.ac
  4. 6
      cpuminer-config.h
  5. 13
      cuda_groestlcoin.cu
  6. 23
      cuda_myriadgroestl.cu
  7. 28
      quark/cuda_quark_groestl512.cu
  8. 330
      quark/cuda_quark_groestl512_sm20.cu
  9. 7
      x11/cuda_x11_simd512.cu

6
Makefile.am
Unescape View File

@ -71,8 +71,10 @@ ccminer_CPPFLAGS = @LIBCURL_CPPFLAGS@ @OPENMP_CFLAGS@ $(CPPFLAGS) $(PTHREAD_FLAG @@ -71,8 +71,10 @@ ccminer_CPPFLAGS = @LIBCURL_CPPFLAGS@ @OPENMP_CFLAGS@ $(CPPFLAGS) $(PTHREAD_FLAG
nvcc_ARCH = -gencode=arch=compute_50,code=\"sm_50,compute_50\"
#nvcc_ARCH += -gencode=arch=compute_52,code=\"sm_52,compute_52\"
#nvcc_ARCH += -gencode=arch=compute_35,code=\"sm_35,compute_35\"
#nvcc_ARCH += -gencode=arch=compute_30,code=\"sm_30,compute_30\"
#nvcc_ARCH = -gencode=arch=compute_35,code=\"sm_35,compute_35\"
#nvcc_ARCH = -gencode=arch=compute_30,code=\"sm_30,compute_30\"
#nvcc_ARCH = -gencode=arch=compute_20,code=\"sm_21,compute_20\"
nvcc_FLAGS = $(nvcc_ARCH) @CUDA_INCLUDES@ -I. @CUDA_CFLAGS@
nvcc_FLAGS += $(JANSSON_INCLUDES) --ptxas-options="-v"

507
bitslice_transformations_quad.cu
Unescape View File

@ -1,24 +1,49 @@ @@ -1,24 +1,49 @@
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 300
/**
* __shfl() returns the value of var held by the thread whose ID is given by srcLane.
* If srcLane is outside the range 0..width-1, the thread's own value of var is returned.
*/
#undef __shfl
#define __shfl(var, srcLane, width) (uint32_t)(var)
#endif
/* File included in quark/groestl (quark/jha,nist5/X11+) and groest/myriad coins for SM 3+ */
#define merge8(z,x,y)\
z=__byte_perm(x, y, 0x5140); \
#define SWAP8(x,y)\
x=__byte_perm(x, y, 0x5410); \
y=__byte_perm(x, y, 0x7632);
#define SWAP4(x,y)\
t = (y<<4); \
t = (x ^ t); \
t = 0xf0f0f0f0UL & t; \
x = (x ^ t); \
t= t>>4;\
y= y ^ t;
#define SWAP2(x,y)\
t = (y<<2); \
t = (x ^ t); \
t = 0xccccccccUL & t; \
x = (x ^ t); \
t= t>>2;\
y= y ^ t;
#define SWAP1(x,y)\
t = (y+y); \
t = (x ^ t); \
t = 0xaaaaaaaaUL & t; \
x = (x ^ t); \
t= t>>1;\
y= y ^ t;
#ifdef __CUDA_ARCH__
__device__ __forceinline__
void to_bitslice_quad(uint32_t *input, uint32_t *output)
void to_bitslice_quad(uint32_t *const __restrict__ input, uint32_t *const __restrict__ output)
{
uint32_t other[8];
const int n = threadIdx.x % 4;
uint32_t d[8];
uint32_t t;
const unsigned int n = threadIdx.x & 3;
#pragma unroll
for (int i = 0; i < 8; i++) {
input[i] = __shfl((int)input[i], n ^ (3*(n >=1 && n <=2)), 4);
other[i] = __shfl((int)input[i], (threadIdx.x + 1) % 4, 4);
other[i] = __shfl((int)input[i], (threadIdx.x + 1) & 3, 4);
input[i] = __shfl((int)input[i], threadIdx.x & 2, 4);
other[i] = __shfl((int)other[i], threadIdx.x & 2, 4);
if (threadIdx.x & 1) {
@ -27,402 +52,76 @@ void to_bitslice_quad(uint32_t *input, uint32_t *output) @@ -27,402 +52,76 @@ void to_bitslice_quad(uint32_t *input, uint32_t *output)
}
}
output[ 0] = (input[ 0] & 0x00000001);
output[ 0] |= ((other[ 0] & 0x00000001) << 1);
output[ 0] |= ((input[ 1] & 0x00000001) << 2);
output[ 0] |= ((other[ 1] & 0x00000001) << 3);
output[ 0] |= ((input[ 2] & 0x00000001) << 4);
output[ 0] |= ((other[ 2] & 0x00000001) << 5);
output[ 0] |= ((input[ 3] & 0x00000001) << 6);
output[ 0] |= ((other[ 3] & 0x00000001) << 7);
output[ 0] |= ((input[ 4] & 0x00000001) << 8);
output[ 0] |= ((other[ 4] & 0x00000001) << 9);
output[ 0] |= ((input[ 5] & 0x00000001) <<10);
output[ 0] |= ((other[ 5] & 0x00000001) <<11);
output[ 0] |= ((input[ 6] & 0x00000001) <<12);
output[ 0] |= ((other[ 6] & 0x00000001) <<13);
output[ 0] |= ((input[ 7] & 0x00000001) <<14);
output[ 0] |= ((other[ 7] & 0x00000001) <<15);
output[ 0] |= ((input[ 0] & 0x00000100) << 8);
output[ 0] |= ((other[ 0] & 0x00000100) << 9);
output[ 0] |= ((input[ 1] & 0x00000100) <<10);
output[ 0] |= ((other[ 1] & 0x00000100) <<11);
output[ 0] |= ((input[ 2] & 0x00000100) <<12);
output[ 0] |= ((other[ 2] & 0x00000100) <<13);
output[ 0] |= ((input[ 3] & 0x00000100) <<14);
output[ 0] |= ((other[ 3] & 0x00000100) <<15);
output[ 0] |= ((input[ 4] & 0x00000100) <<16);
output[ 0] |= ((other[ 4] & 0x00000100) <<17);
output[ 0] |= ((input[ 5] & 0x00000100) <<18);
output[ 0] |= ((other[ 5] & 0x00000100) <<19);
output[ 0] |= ((input[ 6] & 0x00000100) <<20);
output[ 0] |= ((other[ 6] & 0x00000100) <<21);
output[ 0] |= ((input[ 7] & 0x00000100) <<22);
output[ 0] |= ((other[ 7] & 0x00000100) <<23);
merge8(d[0], input[0], input[4]);
merge8(d[1], other[0], other[4]);
merge8(d[2], input[1], input[5]);
merge8(d[3], other[1], other[5]);
merge8(d[4], input[2], input[6]);
merge8(d[5], other[2], other[6]);
merge8(d[6], input[3], input[7]);
merge8(d[7], other[3], other[7]);
SWAP1(d[0], d[1]);
SWAP1(d[2], d[3]);
SWAP1(d[4], d[5]);
SWAP1(d[6], d[7]);
SWAP2(d[0], d[2]);
SWAP2(d[1], d[3]);
SWAP2(d[4], d[6]);
SWAP2(d[5], d[7]);
SWAP4(d[0], d[4]);
SWAP4(d[1], d[5]);
SWAP4(d[2], d[6]);
SWAP4(d[3], d[7]);
output[0] = d[0];
output[1] = d[1];
output[2] = d[2];
output[3] = d[3];
output[4] = d[4];
output[5] = d[5];
output[6] = d[6];
output[7] = d[7];
}
output[ 1] = ((input[ 0] & 0x00000002) >> 1);
output[ 1] |= (other[ 0] & 0x00000002);
output[ 1] |= ((input[ 1] & 0x00000002) << 1);
output[ 1] |= ((other[ 1] & 0x00000002) << 2);
output[ 1] |= ((input[ 2] & 0x00000002) << 3);
output[ 1] |= ((other[ 2] & 0x00000002) << 4);
output[ 1] |= ((input[ 3] & 0x00000002) << 5);
output[ 1] |= ((other[ 3] & 0x00000002) << 6);
output[ 1] |= ((input[ 4] & 0x00000002) << 7);
output[ 1] |= ((other[ 4] & 0x00000002) << 8);
output[ 1] |= ((input[ 5] & 0x00000002) << 9);
output[ 1] |= ((other[ 5] & 0x00000002) <<10);
output[ 1] |= ((input[ 6] & 0x00000002) <<11);
output[ 1] |= ((other[ 6] & 0x00000002) <<12);
output[ 1] |= ((input[ 7] & 0x00000002) <<13);
output[ 1] |= ((other[ 7] & 0x00000002) <<14);
output[ 1] |= ((input[ 0] & 0x00000200) << 7);
output[ 1] |= ((other[ 0] & 0x00000200) << 8);
output[ 1] |= ((input[ 1] & 0x00000200) << 9);
output[ 1] |= ((other[ 1] & 0x00000200) <<10);
output[ 1] |= ((input[ 2] & 0x00000200) <<11);
output[ 1] |= ((other[ 2] & 0x00000200) <<12);
output[ 1] |= ((input[ 3] & 0x00000200) <<13);
output[ 1] |= ((other[ 3] & 0x00000200) <<14);
output[ 1] |= ((input[ 4] & 0x00000200) <<15);
output[ 1] |= ((other[ 4] & 0x00000200) <<16);
output[ 1] |= ((input[ 5] & 0x00000200) <<17);
output[ 1] |= ((other[ 5] & 0x00000200) <<18);
output[ 1] |= ((input[ 6] & 0x00000200) <<19);
output[ 1] |= ((other[ 6] & 0x00000200) <<20);
output[ 1] |= ((input[ 7] & 0x00000200) <<21);
output[ 1] |= ((other[ 7] & 0x00000200) <<22);
__device__ __forceinline__
void from_bitslice_quad(const uint32_t *const __restrict__ input, uint32_t *const __restrict__ output)
{
uint32_t d[8];
uint32_t t;
output[ 2] = ((input[ 0] & 0x00000004) >> 2);
output[ 2] |= ((other[ 0] & 0x00000004) >> 1);
output[ 2] |= (input[ 1] & 0x00000004);
output[ 2] |= ((other[ 1] & 0x00000004) << 1);
output[ 2] |= ((input[ 2] & 0x00000004) << 2);
output[ 2] |= ((other[ 2] & 0x00000004) << 3);
output[ 2] |= ((input[ 3] & 0x00000004) << 4);
output[ 2] |= ((other[ 3] & 0x00000004) << 5);
output[ 2] |= ((input[ 4] & 0x00000004) << 6);
output[ 2] |= ((other[ 4] & 0x00000004) << 7);
output[ 2] |= ((input[ 5] & 0x00000004) << 8);
output[ 2] |= ((other[ 5] & 0x00000004) << 9);
output[ 2] |= ((input[ 6] & 0x00000004) <<10);
output[ 2] |= ((other[ 6] & 0x00000004) <<11);
output[ 2] |= ((input[ 7] & 0x00000004) <<12);
output[ 2] |= ((other[ 7] & 0x00000004) <<13);
output[ 2] |= ((input[ 0] & 0x00000400) << 6);
output[ 2] |= ((other[ 0] & 0x00000400) << 7);
output[ 2] |= ((input[ 1] & 0x00000400) << 8);
output[ 2] |= ((other[ 1] & 0x00000400) << 9);
output[ 2] |= ((input[ 2] & 0x00000400) <<10);
output[ 2] |= ((other[ 2] & 0x00000400) <<11);
output[ 2] |= ((input[ 3] & 0x00000400) <<12);
output[ 2] |= ((other[ 3] & 0x00000400) <<13);
output[ 2] |= ((input[ 4] & 0x00000400) <<14);
output[ 2] |= ((other[ 4] & 0x00000400) <<15);
output[ 2] |= ((input[ 5] & 0x00000400) <<16);
output[ 2] |= ((other[ 5] & 0x00000400) <<17);
output[ 2] |= ((input[ 6] & 0x00000400) <<18);
output[ 2] |= ((other[ 6] & 0x00000400) <<19);
output[ 2] |= ((input[ 7] & 0x00000400) <<20);
output[ 2] |= ((other[ 7] & 0x00000400) <<21);
d[0] = __byte_perm(input[0], input[4], 0x7531);
d[1] = __byte_perm(input[1], input[5], 0x7531);
d[2] = __byte_perm(input[2], input[6], 0x7531);
d[3] = __byte_perm(input[3], input[7], 0x7531);
output[ 3] = ((input[ 0] & 0x00000008) >> 3);
output[ 3] |= ((other[ 0] & 0x00000008) >> 2);
output[ 3] |= ((input[ 1] & 0x00000008) >> 1);
output[ 3] |= (other[ 1] & 0x00000008);
output[ 3] |= ((input[ 2] & 0x00000008) << 1);
output[ 3] |= ((other[ 2] & 0x00000008) << 2);
output[ 3] |= ((input[ 3] & 0x00000008) << 3);
output[ 3] |= ((other[ 3] & 0x00000008) << 4);
output[ 3] |= ((input[ 4] & 0x00000008) << 5);
output[ 3] |= ((other[ 4] & 0x00000008) << 6);
output[ 3] |= ((input[ 5] & 0x00000008) << 7);
output[ 3] |= ((other[ 5] & 0x00000008) << 8);
output[ 3] |= ((input[ 6] & 0x00000008) << 9);
output[ 3] |= ((other[ 6] & 0x00000008) <<10);
output[ 3] |= ((input[ 7] & 0x00000008) <<11);
output[ 3] |= ((other[ 7] & 0x00000008) <<12);
output[ 3] |= ((input[ 0] & 0x00000800) << 5);
output[ 3] |= ((other[ 0] & 0x00000800) << 6);
output[ 3] |= ((input[ 1] & 0x00000800) << 7);
output[ 3] |= ((other[ 1] & 0x00000800) << 8);
output[ 3] |= ((input[ 2] & 0x00000800) << 9);
output[ 3] |= ((other[ 2] & 0x00000800) <<10);
output[ 3] |= ((input[ 3] & 0x00000800) <<11);
output[ 3] |= ((other[ 3] & 0x00000800) <<12);
output[ 3] |= ((input[ 4] & 0x00000800) <<13);
output[ 3] |= ((other[ 4] & 0x00000800) <<14);
output[ 3] |= ((input[ 5] & 0x00000800) <<15);
output[ 3] |= ((other[ 5] & 0x00000800) <<16);
output[ 3] |= ((input[ 6] & 0x00000800) <<17);
output[ 3] |= ((other[ 6] & 0x00000800) <<18);
output[ 3] |= ((input[ 7] & 0x00000800) <<19);
output[ 3] |= ((other[ 7] & 0x00000800) <<20);
SWAP1(d[0], d[1]);
SWAP1(d[2], d[3]);
output[ 4] = ((input[ 0] & 0x00000010) >> 4);
output[ 4] |= ((other[ 0] & 0x00000010) >> 3);
output[ 4] |= ((input[ 1] & 0x00000010) >> 2);
output[ 4] |= ((other[ 1] & 0x00000010) >> 1);
output[ 4] |= (input[ 2] & 0x00000010);
output[ 4] |= ((other[ 2] & 0x00000010) << 1);
output[ 4] |= ((input[ 3] & 0x00000010) << 2);
output[ 4] |= ((other[ 3] & 0x00000010) << 3);
output[ 4] |= ((input[ 4] & 0x00000010) << 4);
output[ 4] |= ((other[ 4] & 0x00000010) << 5);
output[ 4] |= ((input[ 5] & 0x00000010) << 6);
output[ 4] |= ((other[ 5] & 0x00000010) << 7);
output[ 4] |= ((input[ 6] & 0x00000010) << 8);
output[ 4] |= ((other[ 6] & 0x00000010) << 9);
output[ 4] |= ((input[ 7] & 0x00000010) <<10);
output[ 4] |= ((other[ 7] & 0x00000010) <<11);
output[ 4] |= ((input[ 0] & 0x00001000) << 4);
output[ 4] |= ((other[ 0] & 0x00001000) << 5);
output[ 4] |= ((input[ 1] & 0x00001000) << 6);
output[ 4] |= ((other[ 1] & 0x00001000) << 7);
output[ 4] |= ((input[ 2] & 0x00001000) << 8);
output[ 4] |= ((other[ 2] & 0x00001000) << 9);
output[ 4] |= ((input[ 3] & 0x00001000) <<10);
output[ 4] |= ((other[ 3] & 0x00001000) <<11);
output[ 4] |= ((input[ 4] & 0x00001000) <<12);
output[ 4] |= ((other[ 4] & 0x00001000) <<13);
output[ 4] |= ((input[ 5] & 0x00001000) <<14);
output[ 4] |= ((other[ 5] & 0x00001000) <<15);
output[ 4] |= ((input[ 6] & 0x00001000) <<16);
output[ 4] |= ((other[ 6] & 0x00001000) <<17);
output[ 4] |= ((input[ 7] & 0x00001000) <<18);
output[ 4] |= ((other[ 7] & 0x00001000) <<19);
SWAP2(d[0], d[2]);
SWAP2(d[1], d[3]);
output[ 5] = ((input[ 0] & 0x00000020) >> 5);
output[ 5] |= ((other[ 0] & 0x00000020) >> 4);
output[ 5] |= ((input[ 1] & 0x00000020) >> 3);
output[ 5] |= ((other[ 1] & 0x00000020) >> 2);
output[ 5] |= ((input[ 2] & 0x00000020) >> 1);
output[ 5] |= (other[ 2] & 0x00000020);
output[ 5] |= ((input[ 3] & 0x00000020) << 1);
output[ 5] |= ((other[ 3] & 0x00000020) << 2);
output[ 5] |= ((input[ 4] & 0x00000020) << 3);
output[ 5] |= ((other[ 4] & 0x00000020) << 4);
output[ 5] |= ((input[ 5] & 0x00000020) << 5);
output[ 5] |= ((other[ 5] & 0x00000020) << 6);
output[ 5] |= ((input[ 6] & 0x00000020) << 7);
output[ 5] |= ((other[ 6] & 0x00000020) << 8);
output[ 5] |= ((input[ 7] & 0x00000020) << 9);
output[ 5] |= ((other[ 7] & 0x00000020) <<10);
output[ 5] |= ((input[ 0] & 0x00002000) << 3);
output[ 5] |= ((other[ 0] & 0x00002000) << 4);
output[ 5] |= ((input[ 1] & 0x00002000) << 5);
output[ 5] |= ((other[ 1] & 0x00002000) << 6);
output[ 5] |= ((input[ 2] & 0x00002000) << 7);
output[ 5] |= ((other[ 2] & 0x00002000) << 8);
output[ 5] |= ((input[ 3] & 0x00002000) << 9);
output[ 5] |= ((other[ 3] & 0x00002000) <<10);
output[ 5] |= ((input[ 4] & 0x00002000) <<11);
output[ 5] |= ((other[ 4] & 0x00002000) <<12);
output[ 5] |= ((input[ 5] & 0x00002000) <<13);
output[ 5] |= ((other[ 5] & 0x00002000) <<14);
output[ 5] |= ((input[ 6] & 0x00002000) <<15);
output[ 5] |= ((other[ 6] & 0x00002000) <<16);
output[ 5] |= ((input[ 7] & 0x00002000) <<17);
output[ 5] |= ((other[ 7] & 0x00002000) <<18);
t = __byte_perm(d[0], d[2], 0x5410);
d[2] = __byte_perm(d[0], d[2], 0x7632);
d[0] = t;
output[ 6] = ((input[ 0] & 0x00000040) >> 6);
output[ 6] |= ((other[ 0] & 0x00000040) >> 5);
output[ 6] |= ((input[ 1] & 0x00000040) >> 4);
output[ 6] |= ((other[ 1] & 0x00000040) >> 3);
output[ 6] |= ((input[ 2] & 0x00000040) >> 2);
output[ 6] |= ((other[ 2] & 0x00000040) >> 1);
output[ 6] |= (input[ 3] & 0x00000040);
output[ 6] |= ((other[ 3] & 0x00000040) << 1);
output[ 6] |= ((input[ 4] & 0x00000040) << 2);
output[ 6] |= ((other[ 4] & 0x00000040) << 3);
output[ 6] |= ((input[ 5] & 0x00000040) << 4);
output[ 6] |= ((other[ 5] & 0x00000040) << 5);
output[ 6] |= ((input[ 6] & 0x00000040) << 6);
output[ 6] |= ((other[ 6] & 0x00000040) << 7);
output[ 6] |= ((input[ 7] & 0x00000040) << 8);
output[ 6] |= ((other[ 7] & 0x00000040) << 9);
output[ 6] |= ((input[ 0] & 0x00004000) << 2);
output[ 6] |= ((other[ 0] & 0x00004000) << 3);
output[ 6] |= ((input[ 1] & 0x00004000) << 4);
output[ 6] |= ((other[ 1] & 0x00004000) << 5);
output[ 6] |= ((input[ 2] & 0x00004000) << 6);
output[ 6] |= ((other[ 2] & 0x00004000) << 7);
output[ 6] |= ((input[ 3] & 0x00004000) << 8);
output[ 6] |= ((other[ 3] & 0x00004000) << 9);
output[ 6] |= ((input[ 4] & 0x00004000) <<10);
output[ 6] |= ((other[ 4] & 0x00004000) <<11);
output[ 6] |= ((input[ 5] & 0x00004000) <<12);
output[ 6] |= ((other[ 5] & 0x00004000) <<13);
output[ 6] |= ((input[ 6] & 0x00004000) <<14);
output[ 6] |= ((other[ 6] & 0x00004000) <<15);
output[ 6] |= ((input[ 7] & 0x00004000) <<16);
output[ 6] |= ((other[ 7] & 0x00004000) <<17);
t = __byte_perm(d[1], d[3], 0x5410);
d[3] = __byte_perm(d[1], d[3], 0x7632);
d[1] = t;
output[ 7] = ((input[ 0] & 0x00000080) >> 7);
output[ 7] |= ((other[ 0] & 0x00000080) >> 6);
output[ 7] |= ((input[ 1] & 0x00000080) >> 5);
output[ 7] |= ((other[ 1] & 0x00000080) >> 4);
output[ 7] |= ((input[ 2] & 0x00000080) >> 3);
output[ 7] |= ((other[ 2] & 0x00000080) >> 2);
output[ 7] |= ((input[ 3] & 0x00000080) >> 1);
output[ 7] |= (other[ 3] & 0x00000080);
output[ 7] |= ((input[ 4] & 0x00000080) << 1);
output[ 7] |= ((other[ 4] & 0x00000080) << 2);
output[ 7] |= ((input[ 5] & 0x00000080) << 3);
output[ 7] |= ((other[ 5] & 0x00000080) << 4);
output[ 7] |= ((input[ 6] & 0x00000080) << 5);
output[ 7] |= ((other[ 6] & 0x00000080) << 6);
output[ 7] |= ((input[ 7] & 0x00000080) << 7);
output[ 7] |= ((other[ 7] & 0x00000080) << 8);
output[ 7] |= ((input[ 0] & 0x00008000) << 1);
output[ 7] |= ((other[ 0] & 0x00008000) << 2);
output[ 7] |= ((input[ 1] & 0x00008000) << 3);
output[ 7] |= ((other[ 1] & 0x00008000) << 4);
output[ 7] |= ((input[ 2] & 0x00008000) << 5);
output[ 7] |= ((other[ 2] & 0x00008000) << 6);
output[ 7] |= ((input[ 3] & 0x00008000) << 7);
output[ 7] |= ((other[ 3] & 0x00008000) << 8);
output[ 7] |= ((input[ 4] & 0x00008000) << 9);
output[ 7] |= ((other[ 4] & 0x00008000) <<10);
output[ 7] |= ((input[ 5] & 0x00008000) <<11);
output[ 7] |= ((other[ 5] & 0x00008000) <<12);
output[ 7] |= ((input[ 6] & 0x00008000) <<13);
output[ 7] |= ((other[ 6] & 0x00008000) <<14);
output[ 7] |= ((input[ 7] & 0x00008000) <<15);
output[ 7] |= ((other[ 7] & 0x00008000) <<16);
}
SWAP4(d[0], d[2]);
SWAP4(d[1], d[3]);
__device__ __forceinline__
void from_bitslice_quad(uint32_t *input, uint32_t *output)
{
output[ 0] = ((input[ 0] & 0x00000100) >> 8);
output[ 0] |= ((input[ 1] & 0x00000100) >> 7);
output[ 0] |= ((input[ 2] & 0x00000100) >> 6);
output[ 0] |= ((input[ 3] & 0x00000100) >> 5);
output[ 0] |= ((input[ 4] & 0x00000100) >> 4);
output[ 0] |= ((input[ 5] & 0x00000100) >> 3);
output[ 0] |= ((input[ 6] & 0x00000100) >> 2);
output[ 0] |= ((input[ 7] & 0x00000100) >> 1);
output[ 0] |= ((input[ 0] & 0x01000000) >>16);
output[ 0] |= ((input[ 1] & 0x01000000) >>15);
output[ 0] |= ((input[ 2] & 0x01000000) >>14);
output[ 0] |= ((input[ 3] & 0x01000000) >>13);
output[ 0] |= ((input[ 4] & 0x01000000) >>12);
output[ 0] |= ((input[ 5] & 0x01000000) >>11);
output[ 0] |= ((input[ 6] & 0x01000000) >>10);
output[ 0] |= ((input[ 7] & 0x01000000) >> 9);
output[ 2] = ((input[ 0] & 0x00000200) >> 9);
output[ 2] |= ((input[ 1] & 0x00000200) >> 8);
output[ 2] |= ((input[ 2] & 0x00000200) >> 7);
output[ 2] |= ((input[ 3] & 0x00000200) >> 6);
output[ 2] |= ((input[ 4] & 0x00000200) >> 5);
output[ 2] |= ((input[ 5] & 0x00000200) >> 4);
output[ 2] |= ((input[ 6] & 0x00000200) >> 3);
output[ 2] |= ((input[ 7] & 0x00000200) >> 2);
output[ 2] |= ((input[ 0] & 0x02000000) >>17);
output[ 2] |= ((input[ 1] & 0x02000000) >>16);
output[ 2] |= ((input[ 2] & 0x02000000) >>15);
output[ 2] |= ((input[ 3] & 0x02000000) >>14);
output[ 2] |= ((input[ 4] & 0x02000000) >>13);
output[ 2] |= ((input[ 5] & 0x02000000) >>12);
output[ 2] |= ((input[ 6] & 0x02000000) >>11);
output[ 2] |= ((input[ 7] & 0x02000000) >>10);
output[ 4] = ((input[ 0] & 0x00000400) >>10);
output[ 4] |= ((input[ 1] & 0x00000400) >> 9);
output[ 4] |= ((input[ 2] & 0x00000400) >> 8);
output[ 4] |= ((input[ 3] & 0x00000400) >> 7);
output[ 4] |= ((input[ 4] & 0x00000400) >> 6);
output[ 4] |= ((input[ 5] & 0x00000400) >> 5);
output[ 4] |= ((input[ 6] & 0x00000400) >> 4);
output[ 4] |= ((input[ 7] & 0x00000400) >> 3);
output[ 4] |= ((input[ 0] & 0x04000000) >>18);
output[ 4] |= ((input[ 1] & 0x04000000) >>17);
output[ 4] |= ((input[ 2] & 0x04000000) >>16);
output[ 4] |= ((input[ 3] & 0x04000000) >>15);
output[ 4] |= ((input[ 4] & 0x04000000) >>14);
output[ 4] |= ((input[ 5] & 0x04000000) >>13);
output[ 4] |= ((input[ 6] & 0x04000000) >>12);
output[ 4] |= ((input[ 7] & 0x04000000) >>11);
output[ 6] = ((input[ 0] & 0x00000800) >>11);
output[ 6] |= ((input[ 1] & 0x00000800) >>10);
output[ 6] |= ((input[ 2] & 0x00000800) >> 9);
output[ 6] |= ((input[ 3] & 0x00000800) >> 8);
output[ 6] |= ((input[ 4] & 0x00000800) >> 7);
output[ 6] |= ((input[ 5] & 0x00000800) >> 6);
output[ 6] |= ((input[ 6] & 0x00000800) >> 5);
output[ 6] |= ((input[ 7] & 0x00000800) >> 4);
output[ 6] |= ((input[ 0] & 0x08000000) >>19);
output[ 6] |= ((input[ 1] & 0x08000000) >>18);
output[ 6] |= ((input[ 2] & 0x08000000) >>17);
output[ 6] |= ((input[ 3] & 0x08000000) >>16);
output[ 6] |= ((input[ 4] & 0x08000000) >>15);
output[ 6] |= ((input[ 5] & 0x08000000) >>14);
output[ 6] |= ((input[ 6] & 0x08000000) >>13);
output[ 6] |= ((input[ 7] & 0x08000000) >>12);
output[ 8] = ((input[ 0] & 0x00001000) >>12);
output[ 8] |= ((input[ 1] & 0x00001000) >>11);
output[ 8] |= ((input[ 2] & 0x00001000) >>10);
output[ 8] |= ((input[ 3] & 0x00001000) >> 9);
output[ 8] |= ((input[ 4] & 0x00001000) >> 8);
output[ 8] |= ((input[ 5] & 0x00001000) >> 7);
output[ 8] |= ((input[ 6] & 0x00001000) >> 6);
output[ 8] |= ((input[ 7] & 0x00001000) >> 5);
output[ 8] |= ((input[ 0] & 0x10000000) >>20);
output[ 8] |= ((input[ 1] & 0x10000000) >>19);
output[ 8] |= ((input[ 2] & 0x10000000) >>18);
output[ 8] |= ((input[ 3] & 0x10000000) >>17);
output[ 8] |= ((input[ 4] & 0x10000000) >>16);
output[ 8] |= ((input[ 5] & 0x10000000) >>15);
output[ 8] |= ((input[ 6] & 0x10000000) >>14);
output[ 8] |= ((input[ 7] & 0x10000000) >>13);
output[10] = ((input[ 0] & 0x00002000) >>13);
output[10] |= ((input[ 1] & 0x00002000) >>12);
output[10] |= ((input[ 2] & 0x00002000) >>11);
output[10] |= ((input[ 3] & 0x00002000) >>10);
output[10] |= ((input[ 4] & 0x00002000) >> 9);
output[10] |= ((input[ 5] & 0x00002000) >> 8);
output[10] |= ((input[ 6] & 0x00002000) >> 7);
output[10] |= ((input[ 7] & 0x00002000) >> 6);
output[10] |= ((input[ 0] & 0x20000000) >>21);
output[10] |= ((input[ 1] & 0x20000000) >>20);
output[10] |= ((input[ 2] & 0x20000000) >>19);
output[10] |= ((input[ 3] & 0x20000000) >>18);
output[10] |= ((input[ 4] & 0x20000000) >>17);
output[10] |= ((input[ 5] & 0x20000000) >>16);
output[10] |= ((input[ 6] & 0x20000000) >>15);
output[10] |= ((input[ 7] & 0x20000000) >>14);
output[12] = ((input[ 0] & 0x00004000) >>14);
output[12] |= ((input[ 1] & 0x00004000) >>13);
output[12] |= ((input[ 2] & 0x00004000) >>12);
output[12] |= ((input[ 3] & 0x00004000) >>11);
output[12] |= ((input[ 4] & 0x00004000) >>10);
output[12] |= ((input[ 5] & 0x00004000) >> 9);
output[12] |= ((input[ 6] & 0x00004000) >> 8);
output[12] |= ((input[ 7] & 0x00004000) >> 7);
output[12] |= ((input[ 0] & 0x40000000) >>22);
output[12] |= ((input[ 1] & 0x40000000) >>21);
output[12] |= ((input[ 2] & 0x40000000) >>20);
output[12] |= ((input[ 3] & 0x40000000) >>19);
output[12] |= ((input[ 4] & 0x40000000) >>18);
output[12] |= ((input[ 5] & 0x40000000) >>17);
output[12] |= ((input[ 6] & 0x40000000) >>16);
output[12] |= ((input[ 7] & 0x40000000) >>15);
output[14] = ((input[ 0] & 0x00008000) >>15);
output[14] |= ((input[ 1] & 0x00008000) >>14);
output[14] |= ((input[ 2] & 0x00008000) >>13);
output[14] |= ((input[ 3] & 0x00008000) >>12);
output[14] |= ((input[ 4] & 0x00008000) >>11);
output[14] |= ((input[ 5] & 0x00008000) >>10);
output[14] |= ((input[ 6] & 0x00008000) >> 9);
output[14] |= ((input[ 7] & 0x00008000) >> 8);
output[14] |= ((input[ 0] & 0x80000000) >>23);
output[14] |= ((input[ 1] & 0x80000000) >>22);
output[14] |= ((input[ 2] & 0x80000000) >>21);
output[14] |= ((input[ 3] & 0x80000000) >>20);
output[14] |= ((input[ 4] & 0x80000000) >>19);
output[14] |= ((input[ 5] & 0x80000000) >>18);
output[14] |= ((input[ 6] & 0x80000000) >>17);
output[14] |= ((input[ 7] & 0x80000000) >>16);
output[0] = d[0];
output[2] = d[1];
output[4] = d[0] >> 16;
output[6] = d[1] >> 16;
output[8] = d[2];
output[10] = d[3];
output[12] = d[2] >> 16;
output[14] = d[3] >> 16;
#pragma unroll 8
for (int i = 0; i < 16; i+=2) {
@ -432,11 +131,3 @@ void from_bitslice_quad(uint32_t *input, uint32_t *output) @@ -432,11 +131,3 @@ void from_bitslice_quad(uint32_t *input, uint32_t *output)
if (threadIdx.x & 3) output[i] = output[i+1] = 0;
}
}
#else
/* host "fake" functions */
#define from_bitslice_quad(st, out)
#define to_bitslice_quad(in, msg) in[0] = (uint32_t) in[0];
#endif /* device only code */

2
configure.ac
Unescape View File

@ -1,4 +1,4 @@ @@ -1,4 +1,4 @@
AC_INIT([ccminer], [1.5.2-git])
AC_INIT([ccminer], [1.5.2])
AC_PREREQ([2.59c])
AC_CANONICAL_SYSTEM

6
cpuminer-config.h
Unescape View File

@ -156,7 +156,7 @@ @@ -156,7 +156,7 @@
#define PACKAGE_NAME "ccminer"
/* Define to the full name and version of this package. */
#define PACKAGE_STRING "ccminer 1.5.2-git"
#define PACKAGE_STRING "ccminer 1.5.2"
/* Define to the one symbol short name of this package. */
#define PACKAGE_TARNAME "ccminer"
@ -165,7 +165,7 @@ @@ -165,7 +165,7 @@
#define PACKAGE_URL ""
/* Define to the version of this package. */
#define PACKAGE_VERSION "1.5.2-git"
#define PACKAGE_VERSION "1.5.2"
/* If using the C implementation of alloca, define if you know the
direction of stack growth for your system; otherwise it will be
@ -188,7 +188,7 @@ @@ -188,7 +188,7 @@
#define USE_XOP 1
/* Version number of package */
#define VERSION "1.5.2-git"
#define VERSION "1.5.2"
/* Define curl_free() as free() if our version of curl lacks curl_free. */
/* #undef curl_free */

13
cuda_groestlcoin.cu
Unescape View File

@ -12,15 +12,20 @@ extern uint32_t *d_resultNonce[8]; @@ -12,15 +12,20 @@ extern uint32_t *d_resultNonce[8];
__constant__ uint32_t groestlcoin_gpu_msg[32];
// 64 Register Variante für Compute 3.0
#if __CUDA_ARCH__ >= 300
// 64 Registers Variant for Compute 3.0
#include "groestl_functions_quad.cu"
#include "bitslice_transformations_quad.cu"
#endif
#define SWAB32(x) cuda_swab32(x)
__global__ __launch_bounds__(256, 4)
void groestlcoin_gpu_hash_quad(int threads, uint32_t startNounce, uint32_t *resNounce)
{
#if __CUDA_ARCH__ >= 300
// durch 4 dividieren, weil jeweils 4 Threads zusammen ein Hash berechnen
int thread = (blockDim.x * blockIdx.x + threadIdx.x) / 4;
if (thread < threads)
@ -86,6 +91,7 @@ void groestlcoin_gpu_hash_quad(int threads, uint32_t startNounce, uint32_t *resN @@ -86,6 +91,7 @@ void groestlcoin_gpu_hash_quad(int threads, uint32_t startNounce, uint32_t *resN
resNounce[0] = nounce;
}
}
#endif
}
// Setup-Funktionen
@ -139,6 +145,11 @@ __host__ void groestlcoin_cpu_hash(int thr_id, int threads, uint32_t startNounce @@ -139,6 +145,11 @@ __host__ void groestlcoin_cpu_hash(int thr_id, int threads, uint32_t startNounce
// Größe des dynamischen Shared Memory Bereichs
size_t shared_size = 0;
if (device_sm[device_map[thr_id]] < 300) {
printf("Sorry, This algo is not supported by this GPU arch (SM 3.0 required)");
return;
}
cudaMemset(d_resultNonce[thr_id], 0xFF, sizeof(uint32_t));
groestlcoin_gpu_hash_quad<<<grid, block, shared_size>>>(threads, startNounce, d_resultNonce[thr_id]);

23
cuda_myriadgroestl.cu
Unescape View File

@ -5,6 +5,14 @@ @@ -5,6 +5,14 @@
#include "cuda_helper.h"
#if __CUDA_ARCH__ >= 300
// 64 Registers Variant for Compute 3.0
#include "groestl_functions_quad.cu"
#include "bitslice_transformations_quad.cu"
#endif
// globaler Speicher für alle HeftyHashes aller Threads
__constant__ uint32_t pTarget[8]; // Single GPU
uint32_t *d_outputHashes[8];
@ -40,13 +48,9 @@ uint32_t myr_sha256_cpu_w2Table[] = { @@ -40,13 +48,9 @@ uint32_t myr_sha256_cpu_w2Table[] = {
0x69bc7ac4, 0xbd11375b, 0xe3ba71e5, 0x3b209ff2, 0x18feee17, 0xe25ad9e7, 0x13375046, 0x0515089d,
0x4f0d0f04, 0x2627484e, 0x310128d2, 0xc668b434, 0x420841cc, 0x62d311b8, 0xe59ba771, 0x85a7a484 };
// 64 Register Variante für Compute 3.0
#include "groestl_functions_quad.cu"
#include "bitslice_transformations_quad.cu"
#define SWAB32(x) ( ((x & 0x000000FF) << 24) | ((x & 0x0000FF00) << 8) | ((x & 0x00FF0000) >> 8) | ((x & 0xFF000000) >> 24) )
#if __CUDA_ARCH__ < 350
#if __CUDA_ARCH__ < 320
// Kepler (Compute 3.0)
#define ROTR32(x, n) (((x) >> (n)) | ((x) << (32 - (n))))
#else
@ -220,6 +224,7 @@ __device__ void myriadgroestl_gpu_sha256(uint32_t *message) @@ -220,6 +224,7 @@ __device__ void myriadgroestl_gpu_sha256(uint32_t *message)
__global__ void __launch_bounds__(256, 4)
myriadgroestl_gpu_hash_quad(int threads, uint32_t startNounce, uint32_t *hashBuffer)
{
#if __CUDA_ARCH__ >= 300
// durch 4 dividieren, weil jeweils 4 Threads zusammen ein Hash berechnen
int thread = (blockDim.x * blockIdx.x + threadIdx.x) / 4;
if (thread < threads)
@ -250,11 +255,13 @@ __global__ void __launch_bounds__(256, 4) @@ -250,11 +255,13 @@ __global__ void __launch_bounds__(256, 4)
for(int k=0;k<16;k++) outpHash[k] = out_state[k];
}
}
#endif
}
__global__ void
myriadgroestl_gpu_hash_quad2(int threads, uint32_t startNounce, uint32_t *resNounce, uint32_t *hashBuffer)
{
#if __CUDA_ARCH__ >= 300
int thread = (blockDim.x * blockIdx.x + threadIdx.x);
if (thread < threads)
{
@ -291,6 +298,7 @@ __global__ void @@ -291,6 +298,7 @@ __global__ void
if(resNounce[0] > nounce)
resNounce[0] = nounce;
}
#endif
}
// Setup-Funktionen
@ -365,6 +373,11 @@ __host__ void myriadgroestl_cpu_hash(int thr_id, int threads, uint32_t startNoun @@ -365,6 +373,11 @@ __host__ void myriadgroestl_cpu_hash(int thr_id, int threads, uint32_t startNoun
dim3 grid(factor*((threads + threadsperblock-1)/threadsperblock));
dim3 block(threadsperblock);
if (device_sm[device_map[thr_id]] < 300) {
printf("Sorry, This algo is not supported by this GPU arch (SM 3.0 required)");
return;
}
myriadgroestl_gpu_hash_quad<<<grid, block, shared_size>>>(threads, startNounce, d_outputHashes[thr_id]);
dim3 grid2((threads + threadsperblock-1)/threadsperblock);
myriadgroestl_gpu_hash_quad2<<<grid2, block, shared_size>>>(threads, startNounce, d_resultNonce[thr_id], d_outputHashes[thr_id]);

28
quark/cuda_quark_groestl512.cu
Unescape View File

@ -8,13 +8,17 @@ @@ -8,13 +8,17 @@
#define TPB 256
#define THF 4
// 64 Register Variante für Compute 3.0
#if __CUDA_ARCH__ >= 300
#include "groestl_functions_quad.cu"
#include "bitslice_transformations_quad.cu"
#endif
#include "quark/cuda_quark_groestl512_sm20.cu"
__global__ __launch_bounds__(TPB, THF)
void quark_groestl512_gpu_hash_64_quad(int threads, uint32_t startNounce, uint32_t * __restrict g_hash, uint32_t * __restrict g_nonceVector)
{
#if __CUDA_ARCH__ >= 300
// durch 4 dividieren, weil jeweils 4 Threads zusammen ein Hash berechnen
int thread = (blockDim.x * blockIdx.x + threadIdx.x) >> 2;
if (thread < threads)
@ -54,11 +58,13 @@ void quark_groestl512_gpu_hash_64_quad(int threads, uint32_t startNounce, uint32 @@ -54,11 +58,13 @@ void quark_groestl512_gpu_hash_64_quad(int threads, uint32_t startNounce, uint32
for(int k=0;k<16;k++) outpHash[k] = hash[k];
}
}
#endif
}
__global__ void __launch_bounds__(TPB, THF)
quark_doublegroestl512_gpu_hash_64_quad(int threads, uint32_t startNounce, uint32_t *g_hash, uint32_t *g_nonceVector)
{
#if __CUDA_ARCH__ >= 300
int thread = (blockDim.x * blockIdx.x + threadIdx.x)>>2;
if (thread < threads)
{
@ -113,11 +119,15 @@ __global__ void __launch_bounds__(TPB, THF) @@ -113,11 +119,15 @@ __global__ void __launch_bounds__(TPB, THF)
for(int k=0;k<16;k++) outpHash[k] = hash[k];
}
}
#endif
}
// Setup-Funktionen
__host__ void quark_groestl512_cpu_init(int thr_id, int threads)
{
if (device_sm[device_map[thr_id]] < 300)
quark_groestl512_sm20_init(thr_id, threads);
}
__host__ void quark_groestl512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order)
@ -135,7 +145,10 @@ __host__ void quark_groestl512_cpu_hash_64(int thr_id, int threads, uint32_t sta @@ -135,7 +145,10 @@ __host__ void quark_groestl512_cpu_hash_64(int thr_id, int threads, uint32_t sta
// Größe des dynamischen Shared Memory Bereichs
size_t shared_size = 0;
if (device_sm[device_map[thr_id]] >= 300)
quark_groestl512_gpu_hash_64_quad<<<grid, block, shared_size>>>(threads, startNounce, d_hash, d_nonceVector);
else
quark_groestl512_sm20_hash_64(thr_id, threads, startNounce, d_nonceVector, d_hash, order);
// Strategisches Sleep Kommando zur Senkung der CPU Last
MyStreamSynchronize(NULL, order, thr_id);
@ -143,21 +156,18 @@ __host__ void quark_groestl512_cpu_hash_64(int thr_id, int threads, uint32_t sta @@ -143,21 +156,18 @@ __host__ void quark_groestl512_cpu_hash_64(int thr_id, int threads, uint32_t sta
__host__ void quark_doublegroestl512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order)
{
int threadsperblock = TPB;
// Compute 3.0 benutzt die registeroptimierte Quad Variante mit Warp Shuffle
// mit den Quad Funktionen brauchen wir jetzt 4 threads pro Hash, daher Faktor 4 bei der Blockzahl
const int factor = THF;
int threadsperblock = TPB;
// berechne wie viele Thread Blocks wir brauchen
dim3 grid(factor*((threads + threadsperblock-1)/threadsperblock));
dim3 block(threadsperblock);
// Größe des dynamischen Shared Memory Bereichs
size_t shared_size = 0;
if (device_sm[device_map[thr_id]] >= 300)
quark_doublegroestl512_gpu_hash_64_quad<<<grid, block, shared_size>>>(threads, startNounce, d_hash, d_nonceVector);
else
quark_doublegroestl512_sm20_hash_64(thr_id, threads, startNounce, d_nonceVector, d_hash, order);
// Strategisches Sleep Kommando zur Senkung der CPU Last
MyStreamSynchronize(NULL, order, thr_id);
}

330
quark/cuda_quark_groestl512_sm20.cu
Unescape View File

@ -0,0 +1,330 @@ @@ -0,0 +1,330 @@
// SM 2.1 variant
// #include "cuda_helper.h"
#define MAXWELL_OR_FERMI 0
#define USE_SHARED 1
#define SPH_C32(x) ((uint32_t)(x ## U))
#define SPH_T32(x) ((x) & SPH_C32(0xFFFFFFFF))
#define PC32up(j, r) ((uint32_t)((j) + (r)))
#define PC32dn(j, r) 0
#define QC32up(j, r) 0xFFFFFFFF
#define QC32dn(j, r) (((uint32_t)(r) << 24) ^ SPH_T32(~((uint32_t)(j) << 24)))
#define B32_0(x) __byte_perm(x, 0, 0x4440)
//((x) & 0xFF)
#define B32_1(x) __byte_perm(x, 0, 0x4441)
//(((x) >> 8) & 0xFF)
#define B32_2(x) __byte_perm(x, 0, 0x4442)
//(((x) >> 16) & 0xFF)
#define B32_3(x) __byte_perm(x, 0, 0x4443)
//((x) >> 24)
// a healthy mix between shared and textured access provides the highest speed on Compute 3.0 and 3.5!
#define T0up(x) (*((uint32_t*)mixtabs + ( (x))))
#define T0dn(x) tex1Dfetch(t0dn1, x)
#define T1up(x) tex1Dfetch(t1up1, x)
#define T1dn(x) (*((uint32_t*)mixtabs + (768+(x))))
#define T2up(x) tex1Dfetch(t2up1, x)
#define T2dn(x) (*((uint32_t*)mixtabs + (1280+(x))))
#define T3up(x) (*((uint32_t*)mixtabs + (1536+(x))))
#define T3dn(x) tex1Dfetch(t3dn1, x)
texture<unsigned int, 1, cudaReadModeElementType> t0up1;
texture<unsigned int, 1, cudaReadModeElementType> t0dn1;
texture<unsigned int, 1, cudaReadModeElementType> t1up1;
texture<unsigned int, 1, cudaReadModeElementType> t1dn1;
texture<unsigned int, 1, cudaReadModeElementType> t2up1;
texture<unsigned int, 1, cudaReadModeElementType> t2dn1;
texture<unsigned int, 1, cudaReadModeElementType> t3up1;
texture<unsigned int, 1, cudaReadModeElementType> t3dn1;
extern uint32_t T0up_cpu[];
extern uint32_t T0dn_cpu[];
extern uint32_t T1up_cpu[];
extern uint32_t T1dn_cpu[];
extern uint32_t T2up_cpu[];
extern uint32_t T2dn_cpu[];
extern uint32_t T3up_cpu[];
extern uint32_t T3dn_cpu[];
#if __CUDA_ARCH__ < 300
__device__ __forceinline__
void quark_groestl512_perm_P(uint32_t *a, char *mixtabs)
{
uint32_t t[32];
for(int r=0; r<14; r++)
{
switch(r) {
case 0:
#pragma unroll 16
for(int k=0;k<16;k++) a[(k*2)+0] ^= PC32up(k<< 4, 0); break;
case 1:
#pragma unroll 16
for(int k=0;k<16;k++) a[(k*2)+0] ^= PC32up(k<< 4, 1); break;
case 2:
#pragma unroll 16
for(int k=0;k<16;k++) a[(k*2)+0] ^= PC32up(k<< 4, 2); break;
case 3:
#pragma unroll 16
for(int k=0;k<16;k++) a[(k*2)+0] ^= PC32up(k<< 4, 3); break;
case 4:
#pragma unroll 16
for(int k=0;k<16;k++) a[(k*2)+0] ^= PC32up(k<< 4, 4); break;
case 5:
#pragma unroll 16
for(int k=0;k<16;k++) a[(k*2)+0] ^= PC32up(k<< 4, 5); break;
case 6:
#pragma unroll 16
for(int k=0;k<16;k++) a[(k*2)+0] ^= PC32up(k<< 4, 6); break;
case 7:
#pragma unroll 16
for(int k=0;k<16;k++) a[(k*2)+0] ^= PC32up(k<< 4, 7); break;
case 8:
#pragma unroll 16
for(int k=0;k<16;k++) a[(k*2)+0] ^= PC32up(k<< 4, 8); break;
case 9:
#pragma unroll 16
for(int k=0;k<16;k++) a[(k*2)+0] ^= PC32up(k<< 4, 9); break;
case 10:
#pragma unroll 16
for(int k=0;k<16;k++) a[(k*2)+0] ^= PC32up(k<< 4, 10); break;
case 11:
#pragma unroll 16
for(int k=0;k<16;k++) a[(k*2)+0] ^= PC32up(k<< 4, 11); break;
case 12:
#pragma unroll 16
for(int k=0;k<16;k++) a[(k*2)+0] ^= PC32up(k<< 4, 12); break;
case 13:
#pragma unroll 16
for(int k=0;k<16;k++) a[(k*2)+0] ^= PC32up(k<< 4, 13); break;
}
// RBTT
#pragma unroll 16
for(int k=0;k<32;k+=2) {
uint32_t t0_0 = B32_0(a[(k ) & 0x1f]), t9_0 = B32_0(a[(k + 9) & 0x1f]);
uint32_t t2_1 = B32_1(a[(k + 2) & 0x1f]), t11_1 = B32_1(a[(k + 11) & 0x1f]);
uint32_t t4_2 = B32_2(a[(k + 4) & 0x1f]), t13_2 = B32_2(a[(k + 13) & 0x1f]);
uint32_t t6_3 = B32_3(a[(k + 6) & 0x1f]), t23_3 = B32_3(a[(k + 23) & 0x1f]);
t[k + 0] = T0up( t0_0 ) ^ T1up( t2_1 ) ^ T2up( t4_2 ) ^ T3up( t6_3 ) ^
T0dn( t9_0 ) ^ T1dn( t11_1 ) ^ T2dn( t13_2 ) ^ T3dn( t23_3 );
t[k + 1] = T0dn( t0_0 ) ^ T1dn( t2_1 ) ^ T2dn( t4_2 ) ^ T3dn( t6_3 ) ^
T0up( t9_0 ) ^ T1up( t11_1 ) ^ T2up( t13_2 ) ^ T3up( t23_3 );
}
#pragma unroll 32
for(int k=0; k<32; k++) {
a[k] = t[k];
}
}
}
__device__ __forceinline__
void quark_groestl512_perm_Q(uint32_t *a, char *mixtabs)
{
for(int r=0; r<14; r++)
{
uint32_t t[32];
switch(r) {
case 0:
#pragma unroll 16
for(int k=0;k<16;k++) { a[(k*2)+0] ^= QC32up(k<< 4, 0); a[(k*2)+1] ^= QC32dn(k<< 4, 0);} break;
case 1:
#pragma unroll 16
for(int k=0;k<16;k++) { a[(k*2)+0] ^= QC32up(k<< 4, 1); a[(k*2)+1] ^= QC32dn(k<< 4, 1);} break;
case 2:
#pragma unroll 16
for(int k=0;k<16;k++) { a[(k*2)+0] ^= QC32up(k<< 4, 2); a[(k*2)+1] ^= QC32dn(k<< 4, 2);} break;
case 3:
#pragma unroll 16
for(int k=0;k<16;k++) { a[(k*2)+0] ^= QC32up(k<< 4, 3); a[(k*2)+1] ^= QC32dn(k<< 4, 3);} break;
case 4:
#pragma unroll 16
for(int k=0;k<16;k++) { a[(k*2)+0] ^= QC32up(k<< 4, 4); a[(k*2)+1] ^= QC32dn(k<< 4, 4);} break;
case 5:
#pragma unroll 16
for(int k=0;k<16;k++) { a[(k*2)+0] ^= QC32up(k<< 4, 5); a[(k*2)+1] ^= QC32dn(k<< 4, 5);} break;
case 6:
#pragma unroll 16
for(int k=0;k<16;k++) { a[(k*2)+0] ^= QC32up(k<< 4, 6); a[(k*2)+1] ^= QC32dn(k<< 4, 6);} break;
case 7:
#pragma unroll 16
for(int k=0;k<16;k++) { a[(k*2)+0] ^= QC32up(k<< 4, 7); a[(k*2)+1] ^= QC32dn(k<< 4, 7);} break;
case 8:
#pragma unroll 16
for(int k=0;k<16;k++) { a[(k*2)+0] ^= QC32up(k<< 4, 8); a[(k*2)+1] ^= QC32dn(k<< 4, 8);} break;
case 9:
#pragma unroll 16
for(int k=0;k<16;k++) { a[(k*2)+0] ^= QC32up(k<< 4, 9); a[(k*2)+1] ^= QC32dn(k<< 4, 9);} break;
case 10:
#pragma unroll 16
for(int k=0;k<16;k++) { a[(k*2)+0] ^= QC32up(k<< 4, 10); a[(k*2)+1] ^= QC32dn(k<< 4, 10);} break;
case 11:
#pragma unroll 16
for(int k=0;k<16;k++) { a[(k*2)+0] ^= QC32up(k<< 4, 11); a[(k*2)+1] ^= QC32dn(k<< 4, 11);} break;
case 12:
#pragma unroll 16
for(int k=0;k<16;k++) { a[(k*2)+0] ^= QC32up(k<< 4, 12); a[(k*2)+1] ^= QC32dn(k<< 4, 12);} break;
case 13:
#pragma unroll 16
for(int k=0;k<16;k++) { a[(k*2)+0] ^= QC32up(k<< 4, 13); a[(k*2)+1] ^= QC32dn(k<< 4, 13);} break;
}
// RBTT
#pragma unroll 16
for(int k=0;k<32;k+=2)
{
uint32_t t2_0 = B32_0(a[(k + 2) & 0x1f]), t1_0 = B32_0(a[(k + 1) & 0x1f]);
uint32_t t6_1 = B32_1(a[(k + 6) & 0x1f]), t5_1 = B32_1(a[(k + 5) & 0x1f]);
uint32_t t10_2 = B32_2(a[(k + 10) & 0x1f]), t9_2 = B32_2(a[(k + 9) & 0x1f]);
uint32_t t22_3 = B32_3(a[(k + 22) & 0x1f]), t13_3 = B32_3(a[(k + 13) & 0x1f]);
t[k + 0] = T0up( t2_0 ) ^ T1up( t6_1 ) ^ T2up( t10_2 ) ^ T3up( t22_3 ) ^
T0dn( t1_0 ) ^ T1dn( t5_1 ) ^ T2dn( t9_2 ) ^ T3dn( t13_3 );
t[k + 1] = T0dn( t2_0 ) ^ T1dn( t6_1 ) ^ T2dn( t10_2 ) ^ T3dn( t22_3 ) ^
T0up( t1_0 ) ^ T1up( t5_1 ) ^ T2up( t9_2 ) ^ T3up( t13_3 );
}
#pragma unroll 32
for(int k=0;k<32;k++)
a[k] = t[k];
}
}
#endif
__global__
void quark_groestl512_gpu_hash_64(int threads, uint32_t startNounce, uint32_t *g_hash, uint32_t *g_nonceVector)
{
#if __CUDA_ARCH__ < 300
extern __shared__ char mixtabs[];
if (threadIdx.x < 256)
{
*((uint32_t*)mixtabs + ( threadIdx.x)) = tex1Dfetch(t0up1, threadIdx.x);
*((uint32_t*)mixtabs + (256+threadIdx.x)) = tex1Dfetch(t0dn1, threadIdx.x);
*((uint32_t*)mixtabs + (512+threadIdx.x)) = tex1Dfetch(t1up1, threadIdx.x);
*((uint32_t*)mixtabs + (768+threadIdx.x)) = tex1Dfetch(t1dn1, threadIdx.x);
*((uint32_t*)mixtabs + (1024+threadIdx.x)) = tex1Dfetch(t2up1, threadIdx.x);
*((uint32_t*)mixtabs + (1280+threadIdx.x)) = tex1Dfetch(t2dn1, threadIdx.x);
*((uint32_t*)mixtabs + (1536+threadIdx.x)) = tex1Dfetch(t3up1, threadIdx.x);
*((uint32_t*)mixtabs + (1792+threadIdx.x)) = tex1Dfetch(t3dn1, threadIdx.x);
}
__syncthreads();
int thread = (blockDim.x * blockIdx.x + threadIdx.x);
if (thread < threads)
{
// GROESTL
uint32_t message[32];
uint32_t state[32];
uint32_t nounce = (g_nonceVector != NULL) ? g_nonceVector[thread] : (startNounce + thread);
int hashPosition = nounce - startNounce;
uint32_t *inpHash = &g_hash[16 * hashPosition];
#pragma unroll 16
for(int k=0; k<16; k++)
message[k] = inpHash[k];
#pragma unroll 14
for(int k=1; k<15; k++)
message[k+16] = 0;
message[16] = 0x80;
message[31] = 0x01000000;
#pragma unroll 32
for(int u=0; u<32; u++)
state[u] = message[u];
state[31] ^= 0x20000;
// Perm
quark_groestl512_perm_P(state, mixtabs);
state[31] ^= 0x20000;
quark_groestl512_perm_Q(message, mixtabs);
#pragma unroll 32
for(int u=0;u<32;u++) state[u] ^= message[u];
#pragma unroll 32
for(int u=0;u<32;u++) message[u] = state[u];
quark_groestl512_perm_P(message, mixtabs);
#pragma unroll 32
for(int u=0;u<32;u++) state[u] ^= message[u];
// Erzeugten Hash rausschreiben
uint32_t *outpHash = &g_hash[16 * hashPosition];
#pragma unroll 16
for(int k=0;k<16;k++) outpHash[k] = state[k+16];
}
#endif
}
#define texDef(texname, texmem, texsource, texsize) \
unsigned int *texmem; \
cudaMalloc(&texmem, texsize); \
cudaMemcpy(texmem, texsource, texsize, cudaMemcpyHostToDevice); \
texname.normalized = 0; \
texname.filterMode = cudaFilterModePoint; \
texname.addressMode[0] = cudaAddressModeClamp; \
{ cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<unsigned int>(); \
cudaBindTexture(NULL, &texname, texmem, &channelDesc, texsize ); } \
__host__
void quark_groestl512_sm20_init(int thr_id, int threads)
{
// Texturen mit obigem Makro initialisieren
texDef(t0up1, d_T0up, T0up_cpu, sizeof(uint32_t)*256);
texDef(t0dn1, d_T0dn, T0dn_cpu, sizeof(uint32_t)*256);
texDef(t1up1, d_T1up, T1up_cpu, sizeof(uint32_t)*256);
texDef(t1dn1, d_T1dn, T1dn_cpu, sizeof(uint32_t)*256);
texDef(t2up1, d_T2up, T2up_cpu, sizeof(uint32_t)*256);
texDef(t2dn1, d_T2dn, T2dn_cpu, sizeof(uint32_t)*256);
texDef(t3up1, d_T3up, T3up_cpu, sizeof(uint32_t)*256);
texDef(t3dn1, d_T3dn, T3dn_cpu, sizeof(uint32_t)*256);
}
__host__
void quark_groestl512_sm20_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order)
{
int threadsperblock = 512;
dim3 grid((threads + threadsperblock-1)/threadsperblock);
dim3 block(threadsperblock);
size_t shared_size = 8 * 256 * sizeof(uint32_t);
quark_groestl512_gpu_hash_64<<<grid, block, shared_size>>>(threads, startNounce, d_hash, d_nonceVector);
// MyStreamSynchronize(NULL, order, thr_id);
}
__host__
void quark_doublegroestl512_sm20_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order)
{
int threadsperblock = 512;
dim3 grid((threads + threadsperblock-1)/threadsperblock);
dim3 block(threadsperblock);
size_t shared_size = 8 * 256 * sizeof(uint32_t);
quark_groestl512_gpu_hash_64<<<grid, block, shared_size>>>(threads, startNounce, d_hash, d_nonceVector);
quark_groestl512_gpu_hash_64<<<grid, block, shared_size>>>(threads, startNounce, d_hash, d_nonceVector);
// MyStreamSynchronize(NULL, order, thr_id);
}

7
x11/cuda_x11_simd512.cu
Unescape View File

@ -181,6 +181,13 @@ do { \ @@ -181,6 +181,13 @@ do { \
#undef BUTTERFLY
}
#if defined(__CUDA_ARCH__)
#if __CUDA_ARCH__ < 300
#define __shfl(var, srcLane, width) (uint32_t)(var)
// #error __shfl() not supported by SM 2.x
#endif
#endif
/**
* FFT_16 using w=2 as 16th root of unity
* Unrolled decimation in frequency (DIF) radix-2 NTT.

Write Preview
Loading…
Cancel
Save