Merge pull request #41 from biospb/master

memory leaks fix and other improvements

src/main.cu

@@ -17,6 +17,9 @@
 #include "mining.h"
 #include "getopt.h"
 
+
+std::atomic<uint32_t> found_solutions{0};
+
 typedef std::chrono::high_resolution_clock Time;
 typedef std::chrono::duration<double> duration_t;
 typedef std::chrono::time_point<std::chrono::high_resolution_clock> time_point_t;
@@ -68,6 +71,7 @@ void submit_new_block(mining_worker_t *worker)
     uint32_t buf_count = 1;
 
     uv_write(write_req, tcp, &buf, buf_count, on_write_end);
+    found_solutions.fetch_add(1, std::memory_order_relaxed);
 }
 
 void mine_with_timer(uv_timer_t *timer);
@@ -192,7 +196,7 @@ void log_hashrate(uv_timer_t *timer)
         {
             printf("gpu%d: %.0f MH/s ", i, device_mining_count[i].load() / eplased.count() / 1000000);
         }
-        printf("\n");
+        printf("solutions: %u\n", found_solutions.load(std::memory_order_relaxed));
     }
 }
 
@@ -332,7 +336,7 @@ int main(int argc, char **argv)
 
     printf("Running gpu-miner version : %s\n", MINER_VERSION);
 
-    int gpu_count;
+    int gpu_count = 0;
     cudaGetDeviceCount(&gpu_count);
     printf("GPU count: %d\n", gpu_count);
     for (int i = 0; i < gpu_count; i++)

src/mining.h

@@ -3,6 +3,8 @@
 
 #include "blake3.cu"
 
+//#define SHOW_MINING_TIME  1
+
 void worker_stream_callback(cudaStream_t stream, cudaError_t status, void *data);
 
 void start_worker_mining(mining_worker_t *worker)
@@ -12,23 +14,31 @@ void start_worker_mining(mining_worker_t *worker)
     reset_worker(worker);
     TRY( cudaMemcpyAsync(hasher(worker, false), hasher(worker, true), hasher_len(worker), cudaMemcpyHostToDevice, worker->stream) );
 
+#ifdef SHOW_MINING_TIME
     cudaEvent_t startEvent, stopEvent;
     TRY( cudaEventCreate(&startEvent) );
     TRY( cudaEventCreate(&stopEvent) );
-
     TRY( cudaEventRecord(startEvent, worker->stream) );
+#endif
+
     // blake3_hasher_mine<<<worker->grid_size, worker->block_size, 0, worker->stream>>>(worker->device_hasher);
     MINER_IMPL(worker)<<<worker->grid_size, worker->block_size, 0, worker->stream>>>(worker->device_hasher.inline_hasher);
 
+#ifdef SHOW_MINING_TIME
     TRY( cudaEventRecord(stopEvent, worker->stream) );
+#endif
 
     TRY(cudaMemcpyAsync(hasher(worker, true), hasher(worker, false), hasher_len(worker), cudaMemcpyDeviceToHost, worker->stream));
 
     TRY( cudaStreamAddCallback(worker->stream, worker_stream_callback, worker, 0) );
 
+#ifdef SHOW_MINING_TIME
     float time;
     TRY( cudaEventElapsedTime(&time, startEvent, stopEvent) );
-    // printf(" === mining time: %f\n", time);
+    TRY( cudaEventDestroy(&startEvent) );
+    TRY( cudaEventDestroy(&stopEvent) );
+    printf(" === mining time: %f\n", time);
+#endif
 }
 
 #endif // ALEPHIUM_MINING_H

src/template.h

@@ -38,7 +38,7 @@ uint32_t sub_template__ref_count(mining_template_t *template_ptr, uint32_t value
 void free_template(mining_template_t *template_ptr)
 {
     uint32_t old_count = sub_template__ref_count(template_ptr, 1);
-    if (old_count == 0) {
+    if (old_count == 1) { // fetch_sub returns original value
         free_job(template_ptr->job);
         free(template_ptr);
     }