diff --git a/src/main.cu b/src/main.cu index bc4d3b1..d8b33d8 100644 --- a/src/main.cu +++ b/src/main.cu @@ -17,6 +17,9 @@ #include "mining.h" #include "getopt.h" + +std::atomic found_solutions{0}; + typedef std::chrono::high_resolution_clock Time; typedef std::chrono::duration duration_t; typedef std::chrono::time_point 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++) diff --git a/src/mining.h b/src/mining.h index 568d447..679ae4d 100644 --- a/src/mining.h +++ b/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<<grid_size, worker->block_size, 0, worker->stream>>>(worker->device_hasher); MINER_IMPL(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 diff --git a/src/template.h b/src/template.h index c0495f3..810bd4e 100644 --- a/src/template.h +++ b/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); }