Topic: NVIDIA Kepler (K20) from 134MHash/s to 330MHash/s with CUDA - page 11. (Read 73308 times)

diff -Naur ./cmake-rpcminer/CMakeLists.txt ../../rpcminer-cuda_svn//cmake-rpcminer/CMakeLists.txt
--- ./cmake-rpcminer/CMakeLists.txt 2013-01-28 19:27:46.000000000 +0100
+++ ../../rpcminer-cuda_svn//cmake-rpcminer/CMakeLists.txt 2013-04-01 11:22:19.000000000 +0200
@@ -32,6 +32,10 @@
 
 IF(BITCOIN_ENABLE_CUDA)
  ADD_DEFINITIONS(-D_BITCOIN_MINER_CUDA_)
+ #SET(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} -gencode=arch=compute_20,code=sm_20 -gencode=arch=compute_35,code=sm_35)
+ SET(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} -gencode=arch=compute_35,code=sm_35  -Xptxas=-v -Xopencc=-LIST:source=on)
+ #SET(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} -gencode=arch=compute_13,code=sm_13 -Xptxas=-v -Xopencc=-LIST:source=on)
+
  CUDA_ADD_EXECUTABLE(rpcminer ${BITCOIN_RPC_MINER_SRC} ${BITCOIN_RPC_MINER_CUDA_SRC})
  # Install generated PTX CUDA module
  INSTALL(FILES "${CMAKE_CURRENT_BINARY_DIR}/${generated_file_basename}.ptx" DESTINATION "${CMAKE_CURRENT_BINARY_DIR}/rpcminer-cuda" RENAME "bitcoinminercuda.ptx")
diff -Naur ./patch.txt ../../rpcminer-cuda_svn//patch.txt
--- ./patch.txt 2013-04-01 11:30:47.000000000 +0200
+++ ../../rpcminer-cuda_svn//patch.txt 1970-01-01 01:00:00.000000000 +0100
@@ -1,14 +0,0 @@
-diff -Naur ./cmake-rpcminer/CMakeLists.txt ../../rpcminer-cuda_svn//cmake-rpcminer/CMakeLists.txt
---- ./cmake-rpcminer/CMakeLists.txt 2013-01-28 19:27:46.000000000 +0100
-+++ ../../rpcminer-cuda_svn//cmake-rpcminer/CMakeLists.txt 2013-04-01 11:22:19.000000000 +0200
-@@ -32,6 +32,10 @@
-
- IF(BITCOIN_ENABLE_CUDA)
- ADD_DEFINITIONS(-D_BITCOIN_MINER_CUDA_)
-+ #SET(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} -gencode=arch=compute_20,code=sm_20 -gencode=arch=compute_35,code=sm_35)
-+ SET(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} -gencode=arch=compute_35,code=sm_35  -Xptxas=-v -Xopencc=-LIST:source=on)
-+ #SET(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} -gencode=arch=compute_13,code=sm_13 -Xptxas=-v -Xopencc=-LIST:source=on)
-+
- CUDA_ADD_EXECUTABLE(rpcminer ${BITCOIN_RPC_MINER_SRC} ${BITCOIN_RPC_MINER_CUDA_SRC})
- # Install generated PTX CUDA module
- INSTALL(FILES "${CMAKE_CURRENT_BINARY_DIR}/${generated_file_basename}.ptx" DESTINATION "${CMAKE_CURRENT_BINARY_DIR}/rpcminer-cuda" RENAME "bitcoinminercuda.ptx")
diff -Naur ./src/cuda/bitcoinminercuda.cpp ../../rpcminer-cuda_svn//src/cuda/bitcoinminercuda.cpp
--- ./src/cuda/bitcoinminercuda.cpp 2013-01-28 19:27:46.000000000 +0100
+++ ../../rpcminer-cuda_svn//src/cuda/bitcoinminercuda.cpp 2013-04-01 11:22:30.000000000 +0200
@@ -290,6 +290,8 @@
  {
  AllocateResources(m_numb,m_numt);
  }
+ m_out[0].m_bestnonce=0;
+ cuMemcpyHtoD(m_devout,m_out,/*m_numb*m_numt*/sizeof(cuda_out));
 
  cuMemcpyHtoD(m_devin,m_in,sizeof(cuda_in));
 
@@ -319,11 +321,11 @@
  cuFuncSetBlockShape(m_function,m_numt,1,1);
  cuLaunchGrid(m_function,m_numb,1);
 
- cuMemcpyDtoH(m_out,m_devout,m_numb*m_numt*sizeof(cuda_out));
+ cuMemcpyDtoH(m_out,m_devout,/*m_numb*m_numt*/sizeof(cuda_out));
 
  // very unlikely that we will find more than 1 hash with H=0
  // so we'll just return the first one and not even worry about G
- for(int i=0; i+ for(int i=0; i<1/*m_numb*m_numt*/; i++)
  {
  if(m_out[i].m_bestnonce!=0)// && m_out[i].m_bestg  {
diff -Naur ./src/cuda/bitcoinminercuda.cu ../../rpcminer-cuda_svn//src/cuda/bitcoinminercuda.cu
--- ./src/cuda/bitcoinminercuda.cu 2013-01-28 19:27:46.000000000 +0100
+++ ../../rpcminer-cuda_svn//src/cuda/bitcoinminercuda.cu 2013-04-01 11:22:30.000000000 +0200
@@ -18,20 +18,31 @@
 
 #include "cudashared.h"
 
-#define byteswap(x) (((x>>24) & 0x000000ff) | ((x>>8) & 0x0000ff00) | ((x<<8) & 0x00ff0000) | ((x<<24) & 0xff000000))
-#define rotateright(x,bits) (((x & 0xffffffff) >> bits) | (x << (32 - bits)))
-#define R(x) (work[x] = (rotateright(work[x-2],17)^rotateright(work[x-2],19)^((work[x-2]&0xffffffff)>>10)) + work[x -  7] + (rotateright(work[x-15],7)^rotateright(work[x-15],18)^((work[x-15]&0xffffffff)>>3)) + work[x - 16])
+#define rotateright(x,bits) (((x ) >> bits) | (x << (32 - bits)))
+
+#define R(x) (work[x] = (rotateright(work[x-2],17)^rotateright(work[x-2],19)^((work[x-2])>>10)) + work[x -  7] + (rotateright(work[x-15],7)^rotateright(work[x-15],18)^((work[x-15])>>3)) + work[x - 16])
+
 #define sharound(a,b,c,d,e,f,g,h,x,K) {t1=h+(rotateright(e,6)^rotateright(e,11)^rotateright(e,25))+(g^(e&(f^g)))+K+x; t2=(rotateright(a,2)^rotateright(a,13)^rotateright(a,22))+((a&b)|(c&(a|b))); d+=t1; h=t1+t2;}
 
-extern "C" __global__ void cuda_process(cuda_in *in, cuda_out *out, const unsigned int loops, const unsigned int bits)
+extern "C" __global__ void cuda_process(cuda_in __restrict__ *in, cuda_out __restrict__ *out, const unsigned int loops, const unsigned int bits)
 {
+  /*variable to check if any other block has a solution*/
+  __shared__ unsigned int canExit;
+  if(threadIdx.x==0)
+ canExit=out[0].m_bestnonce;
+  __syncthreads();
+  /*exit as fast as posible if one block has finished with solution*/
+  if(canExit!=0) return;
+
 
   unsigned int work[64];
   unsigned int A,B,C,D,E,F,G,H;
  const unsigned int myid=(blockIdx.x*blockDim.x+threadIdx.x);
  const unsigned int nonce=in->m_nonce + (myid << bits);
  unsigned int t1,t2;
- unsigned int bestnonce=0;
+ //unsigned int bestnonce=0;
+
+
  //unsigned int bestg=~0;
 
  unsigned int* in_m_AH = in->m_AH;
@@ -40,46 +51,42 @@
   unsigned int in_m_nbits = in->m_nbits;
 
  // the first 3 rounds we can do outside the loop because they depend on work[0] through work[2] which won't change
- unsigned int A1,B1,C1,D1,E1,F1,G1,H1;
- A1=in_m_AH[0];
- B1=in_m_AH[1];
- C1=in_m_AH[2];
- D1=in_m_AH[3];
- E1=in_m_AH[4];
- F1=in_m_AH[5];
- G1=in_m_AH[6];
- H1=in_m_AH[7];
- sharound(A1,B1,C1,D1,E1,F1,G1,H1,in_m_merkle,0x428A2F98);
- sharound(H1,A1,B1,C1,D1,E1,F1,G1,in_m_ntime,0x71374491);
- sharound(G1,H1,A1,B1,C1,D1,E1,F1,in_m_nbits,0xB5C0FBCF);
-    
-    #pragma unroll 1
- for(unsigned int it=0; it+ /* move old A1, ... H1 to shared to solve registers
+ * can also calculated on host and give to kernel, because its se same for all threads and blocks
+ */
+ __shared__ unsigned int AH[8];
+ __shared__ unsigned int AH2[8]; //cache for second round
+ if(threadIdx.x<8)
+ {
+  AH2[threadIdx.x]=AH[threadIdx.x]=in_m_AH[threadIdx.x];
+ }
+ __syncthreads();
+ if(threadIdx.x==0)
+ {
+ sharound(AH[0],AH[1],AH[2],AH[3],AH[4],AH[5],AH[6],AH[7],in_m_merkle,0x428A2F98);
+ sharound(AH[7],AH[0],AH[1],AH[2],AH[3],AH[4],AH[5],AH[6],in_m_ntime,0x71374491);
+ sharound(AH[6],AH[7],AH[0],AH[1],AH[2],AH[3],AH[4],AH[5],in_m_nbits,0xB5C0FBCF);
+ }
+ __syncthreads();
+
+        #pragma unroll 1
+ for(unsigned int it=0; it  {
- /*
- A=in_m_AH[0];
- B=in_m_AH[1];
- C=in_m_AH[2];
- D=in_m_AH[3];
- E=in_m_AH[4];
- F=in_m_AH[5];
- G=in_m_AH[6];
- H=in_m_AH[7];
- */
- A=A1;
- B=B1;
- C=C1;
- D=D1;
- E=E1;
- F=F1;
- G=G1;
- H=H1;
+ if(out[0].m_bestnonce!=0) return;
 
+ A=AH[0];
+ B=AH[1];
+ C=AH[2];
+ D=AH[3];
+ E=AH[4];
+ F=AH[5];
+ G=AH[6];
+ H=AH[7];
  work[0]=in_m_merkle;
  work[1]=in_m_ntime;
  work[2]=in_m_nbits;
  //work[3]=byteswap(nonce+it);
- work[3]=nonce + it;
+ work[3]=nonce +it;
  work[4]=0x80000000;
  work[5]=0x00000000;
  work[6]=0x00000000;
@@ -160,14 +167,14 @@
 
  // hash the hash now
 
- work[0]=in_m_AH[0]+A;
- work[1]=in_m_AH[1]+B;
- work[2]=in_m_AH[2]+C;
- work[3]=in_m_AH[3]+D;
- work[4]=in_m_AH[4]+E;
- work[5]=in_m_AH[5]+F;
- work[6]=in_m_AH[6]+G;
- work[7]=in_m_AH[7]+H;
+ work[0]=AH2[0]+A;
+ work[1]=AH2[1]+B;
+ work[2]=AH2[2]+C;
+ work[3]=AH2[3]+D;
+ work[4]=AH2[4]+E;
+ work[5]=AH2[5]+F;
+ work[6]=AH2[6]+G;
+ work[7]=AH2[7]+H;
  work[8]=0x80000000;
  work[9]=0x00000000;
  work[10]=0x00000000;
@@ -258,13 +265,15 @@
 
  if((H==0))// && (G<=bestg))
  {
- bestnonce=nonce+it;
+ //bestnonce=nonce+it;
+ atomicExch(&(out[0].m_bestnonce),nonce+it); /*we only need one solution*/
+
  //bestg=G;
  }
 
    }
    
-    out[myid].m_bestnonce=bestnonce;
+    //out[myid].m_bestnonce=bestnonce;
     //out[myid].m_bestg=bestg;
 
 }