optimization

secp256k1/Makefile

@@ -36,8 +36,8 @@ ifeq ($(GENCODE_FLAGS),)
 HIGHEST_SM = $(strip $(lastword $(SMS)))
 ifneq ($(HIGHEST_SM),)
 GENCODE_FLAGS += -gencode arch=compute_$(CUDAARCH),code=sm_$(CUDAARCH) \
-	-gencode arch=compute_$(HIGHEST_SM),code=compute_$(HIGHEST_SM) \
-	-gencode arch=compute_30,code=compute_30  # default arch for compatibility
+	-gencode arch=compute_61,code=compute_61
+	# default arch for compatibility
 endif
 endif
 

secp256k1/Makefile.in

@@ -1,14 +1,14 @@
 # CUDA architechture
-CUDAARCH = 30
+CUDAARCH = 61
 
 # CUDA maximum register usage per thread
-MAXREG = 179
+MAXREG = 200
 #-Xptxas -v
 
 # CUDA kernel block dimension 
-BLOCKDIM = 64
+BLOCKDIM = 32
 
 # GPU workspace memory size
-WORKSPACE = 0x400000
+WORKSPACE = 0x800000
 
-# EMBED = -DEMBEDDED_MNEMONIC="\"mnemonic\"" -DEMBEDDED_PASS="\"mnemonicpass\""
+# EMBED = -DEMBEDDED_MNEMONIC="\"mnemonic\"" -DEMBEDDED_PASS="\"mnemonicpass\""

secp256k1/include/mining.h

@@ -9,6 +9,12 @@
 
 #include "definitions.h"
 
+__constant__ uint32_t sk[8];
+__constant__ ctx_t ctt[2];
+
+void cpySkSymbol(uint8_t *skh);
+void cpyCtxSymbol(ctx_t *ctx);
+
 // unfinalized hash of message
 void InitMining(
     // context
@@ -23,21 +29,17 @@ void InitMining(
 __global__ void BlockMining(
     // boundary for puzzle
     const uint32_t * bound,
-    // data: pk || mes || w || padding || x || sk || ctx
-    const uint32_t * data,
-    // nonce base
-    const uint64_t base,
     // precalculated hashes
-    const uint32_t * hashes,
+    const uint32_t * __restrict__ hashes,
+    const uint32_t * data,
     // results
     uint32_t * res,
     // indices of valid solutions
     uint32_t * valid,
-    uint32_t* BHashes
+    uint32_t*  BHashes
 );
 
 
 __global__ void BlakeHash(const uint32_t* data, const uint64_t base, uint32_t* BHashes);
 
 #endif // MINING_H
-

secp256k1/src/autolykos.cu

@@ -152,12 +152,13 @@ void MinerThread(int deviceId, info_t * info, std::vector<double>* hashrates, st
     uint32_t * data_d = bound_d + NUM_SIZE_32;
 
     uint32_t* BHashes;
-    CUDA_CALL(cudaMalloc(&BHashes, NUM_SIZE_8*THREADS_PER_ITER));
+    CUDA_CALL(cudaMalloc(&BHashes, (NUM_SIZE_8)*THREADS_PER_ITER));
 
     // precalculated hashes
     // N_LEN * NUM_SIZE_8 bytes // 2 GiB
     uint32_t * hashes_d;
     CUDA_CALL(cudaMalloc(&hashes_d, (uint32_t)N_LEN * NUM_SIZE_8));
+
 
     // place to handle result of the puzzle
     uint32_t * res_d;
@@ -191,6 +192,10 @@ void MinerThread(int deviceId, info_t * info, std::vector<double>* hashrates, st
         data_d + COUPLED_PK_SIZE_32 + 2 * NUM_SIZE_32, sk_h, NUM_SIZE_8,
         cudaMemcpyHostToDevice
     ));
+    //uint32_t* skptr = sk_h + COUPLED_PK_SIZE_32 + 2*NUM_SIZE_32;
+    cpySkSymbol(sk_h);
+    //CUDA_CALL(cudaMemcpyToSymbol(sk, sk_h, NUM_SIZE_32 * sizeof(uint32_t)));
+
 
     //========================================================================//
     //  Autolykos puzzle cycle
@@ -309,14 +314,16 @@ void MinerThread(int deviceId, info_t * info, std::vector<double>* hashrates, st
                 sizeof(ctx_t), cudaMemcpyHostToDevice
             ));
 
+            cpyCtxSymbol((ctx_t*)(data_d + COUPLED_PK_SIZE_32 + 3*NUM_SIZE_32));
+
             state = STATE_CONTINUE;
         }
 
         VLOG(1) << "Starting main BlockMining procedure";
-        BlakeHash<<<1 + (THREADS_PER_ITER - 1) / BLOCK_DIM, BLOCK_DIM>>>(data_d, base, BHashes);
+        BlakeHash<<<1 + (THREADS_PER_ITER - 1) / (BLOCK_DIM*4), BLOCK_DIM>>>(data_d, base, BHashes);
         // calculate solution candidates
-        BlockMining<<<1 + (THREADS_PER_ITER - 1) / BLOCK_DIM, BLOCK_DIM>>>(
-            bound_d, data_d, base, hashes_d, res_d, indices_d, BHashes
+        BlockMining<<<1 + (THREADS_PER_ITER - 1) / (BLOCK_DIM), BLOCK_DIM>>>(
+            bound_d, hashes_d, data_d, res_d, indices_d, BHashes
         );
 
         VLOG(1) << "Trying to find solution";