Compare commits
No commits in common. "a906816cd45acff3311940081667113a56890a47" and "d3076198337fd05ec92c849a6d4ab8d45708a7be" have entirely different histories.
a906816cd4
...
d307619833
|
|
@ -60,8 +60,3 @@ CMakeUserPresets.json
|
||||||
*.ptx
|
*.ptx
|
||||||
*.cubin
|
*.cubin
|
||||||
*.fatbin
|
*.fatbin
|
||||||
|
|
||||||
# Coin miner executables
|
|
||||||
coin_miner_cpu
|
|
||||||
coin_miner_simd
|
|
||||||
coin_miner_cuda
|
|
||||||
|
|
|
||||||
|
|
@ -26,13 +26,59 @@ void signal_handler(int signum)
|
||||||
keep_running = 0;
|
keep_running = 0;
|
||||||
}
|
}
|
||||||
|
|
||||||
|
// Histogram data structures
|
||||||
|
typedef struct {
|
||||||
|
u32_t bins[MAX_HISTOGRAM_BINS];
|
||||||
|
u32_t count;
|
||||||
|
double min_value;
|
||||||
|
double max_value;
|
||||||
|
} histogram_t;
|
||||||
|
|
||||||
|
static void histogram_init(histogram_t *h)
|
||||||
|
{
|
||||||
|
memset(h->bins, 0, sizeof(h->bins));
|
||||||
|
h->count = 0;
|
||||||
|
h->min_value = 1e99;
|
||||||
|
h->max_value = 0.0;
|
||||||
|
}
|
||||||
|
|
||||||
|
static void histogram_add(histogram_t *h, double value)
|
||||||
|
{
|
||||||
|
if(value < h->min_value)
|
||||||
|
h->min_value = value;
|
||||||
|
if(value > h->max_value)
|
||||||
|
h->max_value = value;
|
||||||
|
h->count++;
|
||||||
|
|
||||||
|
// For now, just count - we'll bin them later when printing
|
||||||
|
}
|
||||||
|
|
||||||
|
static void histogram_print(histogram_t *h, const char *title, int n_bins)
|
||||||
|
{
|
||||||
|
if(h->count == 0)
|
||||||
|
{
|
||||||
|
printf("%s: No data\n", title);
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
printf("\n%s:\n", title);
|
||||||
|
printf(" Count: %u\n", h->count);
|
||||||
|
printf(" Min: %.6f\n", h->min_value);
|
||||||
|
printf(" Max: %.6f\n", h->max_value);
|
||||||
|
printf(" Avg: %.6f\n", (h->min_value + h->max_value) / 2.0);
|
||||||
|
}
|
||||||
|
|
||||||
// Coin reconstruction from stored data
|
// Coin reconstruction from stored data
|
||||||
static void reconstruct_coin(u32_t *stored_data, u32_t coin[14])
|
static void reconstruct_coin(u32_t *stored_data, u32_t coin[14])
|
||||||
{
|
{
|
||||||
// Simply copy the complete coin data from storage
|
// Fixed parts (must match kernel byte order)
|
||||||
for(int i = 0; i < 14; i++)
|
coin[0] = 0x44455449u; // "DETI" with byte swap (idx ^ 3)
|
||||||
coin[i] = stored_data[i];
|
coin[1] = 0x20636F69u; // " coi" with byte swap (idx ^ 3)
|
||||||
|
coin[2] = 0x6E203220u; // "n 2 " with byte swap (idx ^ 3)
|
||||||
|
|
||||||
|
// Variable parts (restore from storage)
|
||||||
|
for(int i = 0; i < 11; i++)
|
||||||
|
coin[3 + i] = stored_data[i];
|
||||||
}
|
}
|
||||||
|
|
||||||
//
|
//
|
||||||
|
|
@ -46,6 +92,19 @@ static void mine_coins_cuda(u64_t max_attempts, int use_scan_kernel)
|
||||||
u32_t coins_found = 0;
|
u32_t coins_found = 0;
|
||||||
u32_t kernel_runs = 0;
|
u32_t kernel_runs = 0;
|
||||||
|
|
||||||
|
// Histograms
|
||||||
|
histogram_t time_histogram;
|
||||||
|
histogram_t coins_histogram;
|
||||||
|
double *kernel_times = NULL;
|
||||||
|
u32_t *kernel_coin_counts = NULL;
|
||||||
|
u32_t histogram_capacity = 10000;
|
||||||
|
|
||||||
|
histogram_init(&time_histogram);
|
||||||
|
histogram_init(&coins_histogram);
|
||||||
|
|
||||||
|
kernel_times = (double *)malloc(histogram_capacity * sizeof(double));
|
||||||
|
kernel_coin_counts = (u32_t *)malloc(histogram_capacity * sizeof(u32_t));
|
||||||
|
|
||||||
// Initialize CUDA
|
// Initialize CUDA
|
||||||
cd.device_number = 0;
|
cd.device_number = 0;
|
||||||
cd.cubin_file_name = "coin_miner_cuda_kernel.cubin";
|
cd.cubin_file_name = "coin_miner_cuda_kernel.cubin";
|
||||||
|
|
@ -58,7 +117,7 @@ static void mine_coins_cuda(u64_t max_attempts, int use_scan_kernel)
|
||||||
host_storage = (u32_t *)cd.host_data[0];
|
host_storage = (u32_t *)cd.host_data[0];
|
||||||
|
|
||||||
// Kernel configuration
|
// Kernel configuration
|
||||||
cd.block_dim_x = RECOMMENDED_CUDA_BLOCK_SIZE;
|
cd.block_dim_x = RECOMENDED_CUDA_BLOCK_SIZE;
|
||||||
cd.grid_dim_x = 4096; // Large grid for maximum GPU utilization
|
cd.grid_dim_x = 4096; // Large grid for maximum GPU utilization
|
||||||
|
|
||||||
u32_t n_threads = cd.grid_dim_x * cd.block_dim_x;
|
u32_t n_threads = cd.grid_dim_x * cd.block_dim_x;
|
||||||
|
|
@ -74,7 +133,9 @@ static void mine_coins_cuda(u64_t max_attempts, int use_scan_kernel)
|
||||||
int scan_pos = 12;
|
int scan_pos = 12;
|
||||||
|
|
||||||
time_measurement();
|
time_measurement();
|
||||||
// double start_time = wall_time_delta();
|
time_measurement();
|
||||||
|
double start_time = wall_time_delta();
|
||||||
|
double last_report_time = 0.0;
|
||||||
|
|
||||||
while(keep_running && (max_attempts == 0 || attempts < max_attempts))
|
while(keep_running && (max_attempts == 0 || attempts < max_attempts))
|
||||||
{
|
{
|
||||||
|
|
@ -92,8 +153,13 @@ static void mine_coins_cuda(u64_t max_attempts, int use_scan_kernel)
|
||||||
if(use_scan_kernel)
|
if(use_scan_kernel)
|
||||||
cd.arg[3] = &scan_pos;
|
cd.arg[3] = &scan_pos;
|
||||||
|
|
||||||
// Launch the CUDA kernel
|
// Launch kernel and measure time
|
||||||
launch_kernel(&cd);
|
time_measurement();
|
||||||
|
double kernel_start = cpu_time_delta();
|
||||||
|
lauch_kernel(&cd);
|
||||||
|
time_measurement();
|
||||||
|
double kernel_end = cpu_time_delta();
|
||||||
|
double kernel_time = kernel_end - kernel_start;
|
||||||
|
|
||||||
// Copy results back
|
// Copy results back
|
||||||
device_to_host_copy(&cd, 0);
|
device_to_host_copy(&cd, 0);
|
||||||
|
|
@ -142,6 +208,16 @@ static void mine_coins_cuda(u64_t max_attempts, int use_scan_kernel)
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
|
// Update histograms
|
||||||
|
if(kernel_runs < histogram_capacity)
|
||||||
|
{
|
||||||
|
kernel_times[kernel_runs] = kernel_time;
|
||||||
|
kernel_coin_counts[kernel_runs] = n_coins_this_kernel;
|
||||||
|
}
|
||||||
|
|
||||||
|
histogram_add(&time_histogram, kernel_time);
|
||||||
|
histogram_add(&coins_histogram, (double)n_coins_this_kernel);
|
||||||
|
|
||||||
// Update counters
|
// Update counters
|
||||||
kernel_runs++;
|
kernel_runs++;
|
||||||
if(use_scan_kernel)
|
if(use_scan_kernel)
|
||||||
|
|
@ -155,11 +231,21 @@ static void mine_coins_cuda(u64_t max_attempts, int use_scan_kernel)
|
||||||
if(use_scan_kernel)
|
if(use_scan_kernel)
|
||||||
scan_pos = (scan_pos + 1) % 42 + 12; // Cycle through positions 12-53
|
scan_pos = (scan_pos + 1) % 42 + 12; // Cycle through positions 12-53
|
||||||
|
|
||||||
|
// Print progress every second
|
||||||
|
time_measurement();
|
||||||
|
double current_time = wall_time_delta() - start_time;
|
||||||
|
if(current_time - last_report_time >= 1.0)
|
||||||
|
{
|
||||||
|
double rate = attempts / current_time;
|
||||||
|
printf("Attempts: %llu, Rate: %.2f MH/s, Coins: %u, Kernels: %u, Avg time: %.6f s\n",
|
||||||
|
(unsigned long long)attempts, rate / 1e6, coins_found, kernel_runs,
|
||||||
|
current_time / kernel_runs);
|
||||||
|
last_report_time = current_time;
|
||||||
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
time_measurement();
|
time_measurement();
|
||||||
double total_time = cpu_time_delta();
|
double total_time = wall_time_delta() - start_time;
|
||||||
|
|
||||||
printf("\n=== Mining Statistics ===\n");
|
printf("\n=== Mining Statistics ===\n");
|
||||||
printf("Total attempts: %llu\n", (unsigned long long)attempts);
|
printf("Total attempts: %llu\n", (unsigned long long)attempts);
|
||||||
|
|
@ -168,10 +254,30 @@ static void mine_coins_cuda(u64_t max_attempts, int use_scan_kernel)
|
||||||
printf("Coins found: %u\n", coins_found);
|
printf("Coins found: %u\n", coins_found);
|
||||||
printf("Kernel launches: %u\n", kernel_runs);
|
printf("Kernel launches: %u\n", kernel_runs);
|
||||||
|
|
||||||
|
// Print histograms
|
||||||
|
histogram_print(&time_histogram, "Kernel Execution Time Histogram", 20);
|
||||||
|
histogram_print(&coins_histogram, "Coins Found Per Kernel Histogram", 10);
|
||||||
|
|
||||||
|
// Save detailed histogram data
|
||||||
|
FILE *fp = fopen("cuda_kernel_stats.csv", "w");
|
||||||
|
if(fp != NULL)
|
||||||
|
{
|
||||||
|
fprintf(fp, "kernel_id,time_seconds,coins_found\n");
|
||||||
|
u32_t n_to_save = (kernel_runs < histogram_capacity) ? kernel_runs : histogram_capacity;
|
||||||
|
for(u32_t i = 0; i < n_to_save; i++)
|
||||||
|
{
|
||||||
|
fprintf(fp, "%u,%.9f,%u\n", i, kernel_times[i], kernel_coin_counts[i]);
|
||||||
|
}
|
||||||
|
fclose(fp);
|
||||||
|
printf("\nDetailed statistics saved to cuda_kernel_stats.csv\n");
|
||||||
|
}
|
||||||
|
|
||||||
// Save any remaining coins
|
// Save any remaining coins
|
||||||
save_coin(NULL);
|
save_coin(NULL);
|
||||||
|
|
||||||
|
// Cleanup
|
||||||
|
free(kernel_times);
|
||||||
|
free(kernel_coin_counts);
|
||||||
terminate_cuda(&cd);
|
terminate_cuda(&cd);
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
|
||||||
|
|
@ -5,23 +5,26 @@
|
||||||
//
|
//
|
||||||
|
|
||||||
#include "aad_sha1.h"
|
#include "aad_sha1.h"
|
||||||
#include "aad_data_types.h"
|
|
||||||
|
typedef unsigned int u32_t;
|
||||||
|
typedef unsigned char u08_t;
|
||||||
|
|
||||||
//
|
//
|
||||||
// Optimized CUDA kernel for DETI coin mining
|
// Optimized CUDA kernel for DETI coin mining
|
||||||
// Each thread generates coins using the same approach as CPU/SIMD miners
|
// Each thread generates its own message based on thread coordinates and external parameters
|
||||||
//
|
//
|
||||||
|
|
||||||
extern "C" __global__ __launch_bounds__(RECOMMENDED_CUDA_BLOCK_SIZE,1)
|
extern "C" __global__ __launch_bounds__(RECOMENDED_CUDA_BLOCK_SIZE,1)
|
||||||
void mine_deti_coins_kernel(u32_t *coins_storage_area, u32_t param1, u32_t param2)
|
void mine_deti_coins_kernel(u32_t *coins_storage_area, u32_t param1, u32_t param2)
|
||||||
{
|
{
|
||||||
u32_t coin[14];
|
u32_t coin[14];
|
||||||
u32_t hash[5];
|
u32_t hash[5];
|
||||||
u32_t n;
|
u32_t n, warp_id, lane_id;
|
||||||
u08_t *bytes = (u08_t *)coin;
|
|
||||||
|
|
||||||
// Get thread index (used as offset from base counter)
|
// Get thread coordinates
|
||||||
n = (u32_t)threadIdx.x + (u32_t)blockDim.x * (u32_t)blockIdx.x;
|
n = (u32_t)threadIdx.x + (u32_t)blockDim.x * (u32_t)blockIdx.x;
|
||||||
|
warp_id = n >> 5u;
|
||||||
|
lane_id = n & 31u;
|
||||||
|
|
||||||
// Initialize coin template: "DETI coin 2 " + variable + "\n\x80"
|
// Initialize coin template: "DETI coin 2 " + variable + "\n\x80"
|
||||||
// Use byte-swapped format to match host expectations (idx ^ 3)
|
// Use byte-swapped format to match host expectations (idx ^ 3)
|
||||||
|
|
@ -29,37 +32,25 @@ void mine_deti_coins_kernel(u32_t *coins_storage_area, u32_t param1, u32_t param
|
||||||
coin[1] = 0x20636F69u; // " coi" with byte swap
|
coin[1] = 0x20636F69u; // " coi" with byte swap
|
||||||
coin[2] = 0x6E203220u; // "n 2 " with byte swap
|
coin[2] = 0x6E203220u; // "n 2 " with byte swap
|
||||||
|
|
||||||
// Initialize variable part (positions 12-53, 42 bytes)
|
// Variable part: encode thread ID and parameters
|
||||||
// Start with A-Z pattern like CPU/SIMD miners
|
// This ensures each thread works on a different message
|
||||||
for(int i = 12; i < 54; i++)
|
coin[3] = n; // Global thread ID
|
||||||
bytes[i ^ 3] = 'A' + ((i - 12) % 26);
|
coin[4] = param1; // External parameter 1
|
||||||
|
coin[5] = param2; // External parameter 2
|
||||||
|
coin[6] = blockIdx.x; // Block index
|
||||||
|
coin[7] = threadIdx.x; // Thread index
|
||||||
|
coin[8] = warp_id; // Warp ID
|
||||||
|
coin[9] = lane_id; // Lane ID
|
||||||
|
coin[10] = n ^ param1; // XOR combination
|
||||||
|
coin[11] = n ^ param2; // XOR combination
|
||||||
|
coin[12] = (n * 0x9E3779B9u); // Hash-like mixing
|
||||||
|
|
||||||
// End with newline and padding
|
// Last word: bytes 52-55
|
||||||
bytes[0x36 ^ 3] = '\n'; // Position 54
|
// Memory layout: coin[13]=0xAABBCCDD -> mem[52]=DD, [53]=CC, [54]=BB, [55]=AA
|
||||||
bytes[0x37 ^ 3] = 0x80; // Position 55
|
// With idx^3: bytes[52^3]=bytes[55]=AA, bytes[53^3]=bytes[54]=BB, bytes[54^3]=bytes[53]=CC, bytes[55^3]=bytes[52]=DD
|
||||||
|
// We want: bytes[54^3]=0x0A (newline), bytes[55^3]=0x80 (padding)
|
||||||
// Calculate offset based on thread index and parameters
|
// So: bytes[53]=0x0A, bytes[52]=0x80 -> coin[13]=0x????0A80
|
||||||
// This creates a unique starting point for each thread
|
coin[13] = ((n & 0xFFFFu) << 16) | 0x0A80u; // Top 2 bytes: variable, bottom: 0x80 0x0A
|
||||||
u64_t offset = ((u64_t)param1 << 32) | param2;
|
|
||||||
offset += (u64_t)n;
|
|
||||||
|
|
||||||
// Apply offset to variable part (increment the coin counter)
|
|
||||||
for(int pos = 53; pos >= 12 && offset > 0; pos--)
|
|
||||||
{
|
|
||||||
u08_t *byte = &bytes[pos ^ 3];
|
|
||||||
u64_t add = offset % 127;
|
|
||||||
offset /= 127;
|
|
||||||
|
|
||||||
u32_t val = *byte + add;
|
|
||||||
u08_t new_val = val % 127;
|
|
||||||
|
|
||||||
// Skip newline character (ASCII 10) in the variable part
|
|
||||||
if(new_val == '\n')
|
|
||||||
new_val++;
|
|
||||||
|
|
||||||
*byte = new_val;
|
|
||||||
offset += val / 127; // Carry
|
|
||||||
}
|
|
||||||
|
|
||||||
// Compute SHA1 hash
|
// Compute SHA1 hash
|
||||||
# define T u32_t
|
# define T u32_t
|
||||||
|
|
@ -83,9 +74,22 @@ void mine_deti_coins_kernel(u32_t *coins_storage_area, u32_t param1, u32_t param
|
||||||
// Make sure we don't write outside buffer
|
// Make sure we don't write outside buffer
|
||||||
if(idx < 1024u - 14u)
|
if(idx < 1024u - 14u)
|
||||||
{
|
{
|
||||||
// Store the complete coin data
|
// Store the coin data (only variable parts needed)
|
||||||
for(int i = 0; i < 14; i++)
|
coins_storage_area[idx + 0] = coin[ 3];
|
||||||
coins_storage_area[idx + i] = coin[i];
|
coins_storage_area[idx + 1] = coin[ 4];
|
||||||
|
coins_storage_area[idx + 2] = coin[ 5];
|
||||||
|
coins_storage_area[idx + 3] = coin[ 6];
|
||||||
|
coins_storage_area[idx + 4] = coin[ 7];
|
||||||
|
coins_storage_area[idx + 5] = coin[ 8];
|
||||||
|
coins_storage_area[idx + 6] = coin[ 9];
|
||||||
|
coins_storage_area[idx + 7] = coin[10];
|
||||||
|
coins_storage_area[idx + 8] = coin[11];
|
||||||
|
coins_storage_area[idx + 9] = coin[12];
|
||||||
|
coins_storage_area[idx + 10] = coin[13];
|
||||||
|
// Store hash value for verification
|
||||||
|
coins_storage_area[idx + 11] = hash[1];
|
||||||
|
coins_storage_area[idx + 12] = hash[2];
|
||||||
|
coins_storage_area[idx + 13] = hash[3];
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
@ -94,13 +98,12 @@ void mine_deti_coins_kernel(u32_t *coins_storage_area, u32_t param1, u32_t param
|
||||||
// Kernel that tries all possible values for one character position
|
// Kernel that tries all possible values for one character position
|
||||||
//
|
//
|
||||||
|
|
||||||
extern "C" __global__ __launch_bounds__(RECOMMENDED_CUDA_BLOCK_SIZE,1)
|
extern "C" __global__ __launch_bounds__(RECOMENDED_CUDA_BLOCK_SIZE,1)
|
||||||
void mine_deti_coins_scan_kernel(u32_t *coins_storage_area, u32_t param1, u32_t param2, int scan_position)
|
void mine_deti_coins_scan_kernel(u32_t *coins_storage_area, u32_t param1, u32_t param2, int scan_position)
|
||||||
{
|
{
|
||||||
u32_t coin[14];
|
u32_t coin[14];
|
||||||
u32_t hash[5];
|
u32_t hash[5];
|
||||||
u32_t n;
|
u32_t n;
|
||||||
u08_t *bytes = (u08_t *)coin;
|
|
||||||
|
|
||||||
n = (u32_t)threadIdx.x + (u32_t)blockDim.x * (u32_t)blockIdx.x;
|
n = (u32_t)threadIdx.x + (u32_t)blockDim.x * (u32_t)blockIdx.x;
|
||||||
|
|
||||||
|
|
@ -109,49 +112,36 @@ void mine_deti_coins_scan_kernel(u32_t *coins_storage_area, u32_t param1, u32_t
|
||||||
coin[1] = 0x20636F69u; // " coi" with byte swap
|
coin[1] = 0x20636F69u; // " coi" with byte swap
|
||||||
coin[2] = 0x6E203220u; // "n 2 " with byte swap
|
coin[2] = 0x6E203220u; // "n 2 " with byte swap
|
||||||
|
|
||||||
// Initialize variable part with A-Z pattern
|
// Variable part
|
||||||
for(int i = 12; i < 54; i++)
|
coin[3] = param1;
|
||||||
bytes[i ^ 3] = 'A' + ((i - 12) % 26);
|
coin[4] = param2;
|
||||||
|
coin[5] = n >> 8; // High bits of n
|
||||||
|
coin[6] = blockIdx.x;
|
||||||
|
coin[7] = threadIdx.x;
|
||||||
|
coin[8] = param1 ^ param2;
|
||||||
|
coin[9] = n & 0xFFu; // Low 8 bits of n
|
||||||
|
coin[10] = param1 + n;
|
||||||
|
coin[11] = param2 - n;
|
||||||
|
coin[12] = (n * 0x9E3779B9u);
|
||||||
|
coin[13] = ((n & 0xFFFFu) << 16) | 0x0A80u; // Top 2 bytes: variable, bottom: 0x80 0x0A
|
||||||
|
|
||||||
// End with newline and padding
|
// Try all possible values for the scan position (0-255)
|
||||||
bytes[0x36 ^ 3] = '\n'; // Position 54
|
// This allows exploring a full byte range in a single kernel launch
|
||||||
bytes[0x37 ^ 3] = 0x80; // Position 55
|
for(u32_t val = 0; val < 256u; val++)
|
||||||
|
|
||||||
// Apply base offset from parameters (similar to main kernel)
|
|
||||||
u64_t offset = ((u64_t)param1 << 32) | param2;
|
|
||||||
offset += (u64_t)n;
|
|
||||||
|
|
||||||
// Apply offset to all positions except the scan position
|
|
||||||
for(int pos = 53; pos >= 12 && offset > 0; pos--)
|
|
||||||
{
|
{
|
||||||
if(pos == scan_position)
|
// Insert the test value at the scan position
|
||||||
continue; // Skip the scan position
|
u32_t word_idx = scan_position / 4;
|
||||||
|
u32_t byte_pos = scan_position % 4;
|
||||||
|
u32_t shift = byte_pos * 8;
|
||||||
|
|
||||||
u08_t *byte = &bytes[pos ^ 3];
|
if(word_idx >= 3 && word_idx < 13)
|
||||||
u64_t add = offset % 95;
|
|
||||||
offset /= 95;
|
|
||||||
|
|
||||||
u32_t val = (*byte - 32 + add);
|
|
||||||
u08_t new_val = 32 + (val % 95);
|
|
||||||
|
|
||||||
// Skip newline character (ASCII 10) in the variable part
|
|
||||||
if(new_val == '\n')
|
|
||||||
new_val++;
|
|
||||||
|
|
||||||
*byte = new_val;
|
|
||||||
offset += val / 95;
|
|
||||||
}
|
|
||||||
|
|
||||||
// Try all possible printable ASCII values for the scan position (32-126)
|
|
||||||
for(u32_t val = 32; val < 127; val++)
|
|
||||||
{
|
|
||||||
// Set the test value at the scan position
|
|
||||||
if(scan_position >= 12 && scan_position < 54)
|
|
||||||
{
|
{
|
||||||
bytes[scan_position ^ 3] = (u08_t)val;
|
u32_t mask = ~(0xFFu << shift);
|
||||||
|
coin[word_idx] = (coin[word_idx] & mask) | (val << shift);
|
||||||
|
|
||||||
// Skip newline in the middle (it's only valid at position 54)
|
// Make sure we don't use newline in the middle
|
||||||
if(scan_position != 54 && val == '\n')
|
u08_t *bytes = (u08_t *)coin;
|
||||||
|
if(scan_position < 54 && bytes[scan_position ^ 3] == 0x0A)
|
||||||
continue;
|
continue;
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
@ -174,8 +164,20 @@ void mine_deti_coins_scan_kernel(u32_t *coins_storage_area, u32_t param1, u32_t
|
||||||
u32_t idx = atomicAdd(coins_storage_area, 14u);
|
u32_t idx = atomicAdd(coins_storage_area, 14u);
|
||||||
if(idx < 1024u - 14u)
|
if(idx < 1024u - 14u)
|
||||||
{
|
{
|
||||||
for(int i = 0; i < 14; i++)
|
coins_storage_area[idx + 0] = coin[ 3];
|
||||||
coins_storage_area[idx + i] = coin[i];
|
coins_storage_area[idx + 1] = coin[ 4];
|
||||||
|
coins_storage_area[idx + 2] = coin[ 5];
|
||||||
|
coins_storage_area[idx + 3] = coin[ 6];
|
||||||
|
coins_storage_area[idx + 4] = coin[ 7];
|
||||||
|
coins_storage_area[idx + 5] = coin[ 8];
|
||||||
|
coins_storage_area[idx + 6] = coin[ 9];
|
||||||
|
coins_storage_area[idx + 7] = coin[10];
|
||||||
|
coins_storage_area[idx + 8] = coin[11];
|
||||||
|
coins_storage_area[idx + 9] = coin[12];
|
||||||
|
coins_storage_area[idx + 10] = coin[13];
|
||||||
|
coins_storage_area[idx + 11] = hash[1];
|
||||||
|
coins_storage_area[idx + 12] = hash[2];
|
||||||
|
coins_storage_area[idx + 13] = hash[3];
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
|
||||||
|
|
@ -30,61 +30,39 @@ static int is_valid_coin(u32_t *hash)
|
||||||
return hash[0] == 0xAAD20250u;
|
return hash[0] == 0xAAD20250u;
|
||||||
}
|
}
|
||||||
|
|
||||||
//
|
|
||||||
// increment coin variable part using the same logic as CPU miner
|
|
||||||
// returns 0 if overflow (all positions wrapped around), 1 otherwise
|
|
||||||
//
|
|
||||||
static int increment_coin(u32_t coin[14])
|
|
||||||
{
|
|
||||||
// Increment the variable part using byte-by-byte logic with carry
|
|
||||||
// Increment from the end to beginning (positions 53 down to 12)
|
|
||||||
int pos = 53;
|
|
||||||
while(pos >= 12)
|
|
||||||
{
|
|
||||||
u08_t *byte = &((u08_t *)coin)[pos ^ 3];
|
|
||||||
if(*byte == '\n' || *byte == 0x80)
|
|
||||||
*byte = 32; // Start from space
|
|
||||||
|
|
||||||
(*byte)++;
|
|
||||||
|
|
||||||
// Skip newline character
|
|
||||||
if(*byte == '\n')
|
|
||||||
(*byte)++;
|
|
||||||
|
|
||||||
// Wrap around at 127 (printable ASCII limit)
|
|
||||||
if(*byte >= 127)
|
|
||||||
{
|
|
||||||
*byte = 32; // Reset to space
|
|
||||||
pos--; // Carry to next position
|
|
||||||
}
|
|
||||||
else
|
|
||||||
{
|
|
||||||
break; // No carry needed
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
||||||
// Return 0 if we carried all the way through (overflow), 1 otherwise
|
|
||||||
return (pos >= 12);
|
|
||||||
}
|
|
||||||
|
|
||||||
//
|
//
|
||||||
// prepare interleaved data for SIMD processing
|
// prepare interleaved data for SIMD processing
|
||||||
//
|
//
|
||||||
static void prepare_coins(u32_t base_coin[14], u32_t *interleaved_data, int simd_width)
|
static void prepare_coins(u32_t base_coin[14], u32_t *interleaved_data, int simd_width, u64_t base_counter)
|
||||||
{
|
{
|
||||||
for(int lane = 0; lane < simd_width; lane++)
|
for(int lane = 0; lane < simd_width; lane++)
|
||||||
{
|
{
|
||||||
u32_t coin[14];
|
u32_t coin[14];
|
||||||
memcpy(coin, base_coin, sizeof(coin));
|
memcpy(coin, base_coin, sizeof(coin));
|
||||||
|
|
||||||
|
// Modify the coin for this lane (encode counter in the variable part)
|
||||||
|
u64_t counter = base_counter + lane;
|
||||||
|
for(int i = 12; i < 20 && i < 54; i++)
|
||||||
|
{
|
||||||
|
int shift = (19 - i) * 8;
|
||||||
|
if(shift >= 0 && shift < 64)
|
||||||
|
{
|
||||||
|
u08_t byte = (counter >> shift) & 0xFF;
|
||||||
|
// Map to ASCII printable range (32-126, excluding newline position)
|
||||||
|
if(byte == '\n' || byte >= 0x80)
|
||||||
|
byte = 'X';
|
||||||
|
// Ensure byte is in printable ASCII range (32-126)
|
||||||
|
// Map all values to ASCII printable characters: space (32) to tilde (126)
|
||||||
|
// byte = 32 + (byte % 95);
|
||||||
|
((u08_t *)coin)[i ^ 3] = byte;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
// Interleave the data
|
// Interleave the data
|
||||||
for(int idx = 0; idx < 14; idx++)
|
for(int idx = 0; idx < 14; idx++)
|
||||||
{
|
{
|
||||||
interleaved_data[idx * simd_width + lane] = coin[idx];
|
interleaved_data[idx * simd_width + lane] = coin[idx];
|
||||||
}
|
}
|
||||||
|
|
||||||
// Increment the base coin for the next lane
|
|
||||||
increment_coin(base_coin);
|
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
@ -147,21 +125,19 @@ static void mine_coins_avx(u64_t max_attempts)
|
||||||
((u08_t *)base_coin)[0x36 ^ 3] = '\n';
|
((u08_t *)base_coin)[0x36 ^ 3] = '\n';
|
||||||
((u08_t *)base_coin)[0x37 ^ 3] = 0x80;
|
((u08_t *)base_coin)[0x37 ^ 3] = 0x80;
|
||||||
|
|
||||||
// Initialize variable part with A-Z cycling pattern (same as CPU miner)
|
// Initialize variable part
|
||||||
for(int i = 12; i < 54; i++)
|
for(int i = 12; i < 54; i++)
|
||||||
((u08_t *)base_coin)[i ^ 3] = 'A' + (i - 12) % 26;
|
((u08_t *)base_coin)[i ^ 3] = 'A';
|
||||||
|
|
||||||
printf("Mining DETI coins using AVX (4-way SIMD)...\n");
|
printf("Mining DETI coins using AVX (4-way SIMD)...\n");
|
||||||
printf("Press Ctrl+C to stop\n\n");
|
printf("Press Ctrl+C to stop\n\n");
|
||||||
|
|
||||||
time_measurement();
|
time_measurement();
|
||||||
time_measurement();
|
|
||||||
double start_time = wall_time_delta();
|
|
||||||
|
|
||||||
while(keep_running && (max_attempts == 0 || attempts < max_attempts))
|
while(keep_running && (max_attempts == 0 || attempts < max_attempts))
|
||||||
{
|
{
|
||||||
// Prepare coins for this batch
|
// Prepare coins for this batch
|
||||||
prepare_coins(base_coin, interleaved_data, SIMD_WIDTH);
|
prepare_coins(base_coin, interleaved_data, SIMD_WIDTH, attempts);
|
||||||
|
|
||||||
// Compute SHA1 hashes
|
// Compute SHA1 hashes
|
||||||
sha1_avx((v4si *)interleaved_data, (v4si *)interleaved_hash);
|
sha1_avx((v4si *)interleaved_data, (v4si *)interleaved_hash);
|
||||||
|
|
@ -187,16 +163,13 @@ static void mine_coins_avx(u64_t max_attempts)
|
||||||
// Print progress every 1M attempts
|
// Print progress every 1M attempts
|
||||||
if(attempts % 1000000 < SIMD_WIDTH)
|
if(attempts % 1000000 < SIMD_WIDTH)
|
||||||
{
|
{
|
||||||
time_measurement();
|
printf("Attempts: %llu, Coins: %u\n",
|
||||||
double current_time = wall_time_delta() - start_time;
|
(unsigned long long)attempts, coins_found);
|
||||||
double rate = attempts / current_time;
|
|
||||||
printf("Attempts: %llu, Rate: %.2f MH/s, Coins: %u\n",
|
|
||||||
(unsigned long long)attempts, rate / 1e6, coins_found);
|
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
time_measurement();
|
time_measurement();
|
||||||
double total_time = wall_time_delta() - start_time;
|
double total_time = wall_time_delta();
|
||||||
|
|
||||||
printf("\n=== Mining Statistics ===\n");
|
printf("\n=== Mining Statistics ===\n");
|
||||||
printf("Total attempts: %llu\n", (unsigned long long)attempts);
|
printf("Total attempts: %llu\n", (unsigned long long)attempts);
|
||||||
|
|
@ -239,20 +212,17 @@ static void mine_coins_avx2(u64_t max_attempts)
|
||||||
((u08_t *)base_coin)[0x36 ^ 3] = '\n';
|
((u08_t *)base_coin)[0x36 ^ 3] = '\n';
|
||||||
((u08_t *)base_coin)[0x37 ^ 3] = 0x80;
|
((u08_t *)base_coin)[0x37 ^ 3] = 0x80;
|
||||||
|
|
||||||
// Initialize variable part with A-Z cycling pattern (same as CPU miner)
|
|
||||||
for(int i = 12; i < 54; i++)
|
for(int i = 12; i < 54; i++)
|
||||||
((u08_t *)base_coin)[i ^ 3] = 'A' + (i - 12) % 26;
|
((u08_t *)base_coin)[i ^ 3] = 'A';
|
||||||
|
|
||||||
printf("Mining DETI coins using AVX2 (8-way SIMD)...\n");
|
printf("Mining DETI coins using AVX2 (8-way SIMD)...\n");
|
||||||
printf("Press Ctrl+C to stop\n\n");
|
printf("Press Ctrl+C to stop\n\n");
|
||||||
|
|
||||||
time_measurement();
|
time_measurement();
|
||||||
time_measurement();
|
|
||||||
double start_time = wall_time_delta();
|
|
||||||
|
|
||||||
while(keep_running && (max_attempts == 0 || attempts < max_attempts))
|
while(keep_running && (max_attempts == 0 || attempts < max_attempts))
|
||||||
{
|
{
|
||||||
prepare_coins(base_coin, interleaved_data, SIMD_WIDTH);
|
prepare_coins(base_coin, interleaved_data, SIMD_WIDTH, attempts);
|
||||||
sha1_avx2((v8si *)interleaved_data, (v8si *)interleaved_hash);
|
sha1_avx2((v8si *)interleaved_data, (v8si *)interleaved_hash);
|
||||||
attempts += SIMD_WIDTH;
|
attempts += SIMD_WIDTH;
|
||||||
|
|
||||||
|
|
@ -274,122 +244,13 @@ static void mine_coins_avx2(u64_t max_attempts)
|
||||||
|
|
||||||
if(attempts % 1000000 < SIMD_WIDTH)
|
if(attempts % 1000000 < SIMD_WIDTH)
|
||||||
{
|
{
|
||||||
time_measurement();
|
printf("Attempts: %llu, Coins: %u\n",
|
||||||
double current_time = wall_time_delta() - start_time;
|
(unsigned long long)attempts, coins_found);
|
||||||
double rate = attempts / current_time;
|
|
||||||
printf("Attempts: %llu, Rate: %.2f MH/s, Coins: %u\n",
|
|
||||||
(unsigned long long)attempts, rate / 1e6, coins_found);
|
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
time_measurement();
|
time_measurement();
|
||||||
double total_time = wall_time_delta() - start_time;
|
double total_time = wall_time_delta();
|
||||||
|
|
||||||
printf("\n=== Mining Statistics ===\n");
|
|
||||||
printf("Total attempts: %llu\n", (unsigned long long)attempts);
|
|
||||||
printf("Total time: %.2f seconds\n", total_time);
|
|
||||||
printf("Average rate: %.2f attempts/second\n", attempts / total_time);
|
|
||||||
printf("Coins found: %u\n", coins_found);
|
|
||||||
|
|
||||||
save_coin(NULL);
|
|
||||||
}
|
|
||||||
#endif
|
|
||||||
#if defined(__AVX2__)
|
|
||||||
#include <omp.h>
|
|
||||||
//
|
|
||||||
// mine DETI coins using AVX2 (8-way SIMD) + OpenMP
|
|
||||||
//
|
|
||||||
__attribute__((unused))
|
|
||||||
static void mine_coins_avx2_omp(u64_t max_attempts)
|
|
||||||
{
|
|
||||||
const int SIMD_WIDTH = 8;
|
|
||||||
int num_threads = omp_get_max_threads();
|
|
||||||
u64_t attempts = 0;
|
|
||||||
u32_t coins_found = 0;
|
|
||||||
|
|
||||||
printf("Mining DETI coins using AVX2 (8-way SIMD) + OpenMP (%d threads)...\n", num_threads);
|
|
||||||
printf("Press Ctrl+C to stop\n\n");
|
|
||||||
|
|
||||||
time_measurement();
|
|
||||||
time_measurement();
|
|
||||||
double start_time = wall_time_delta();
|
|
||||||
|
|
||||||
#pragma omp parallel
|
|
||||||
{
|
|
||||||
u32_t base_coin[14];
|
|
||||||
u32_t interleaved_data[14 * SIMD_WIDTH] __attribute__((aligned(32)));
|
|
||||||
u32_t interleaved_hash[5 * SIMD_WIDTH] __attribute__((aligned(32)));
|
|
||||||
// u32_t thread_coins_found = 0;
|
|
||||||
u64_t thread_attempts = 0;
|
|
||||||
|
|
||||||
// Initialize base coin template (unique per thread)
|
|
||||||
memset(base_coin, 0, sizeof(base_coin));
|
|
||||||
((u08_t *)base_coin)[0x0 ^ 3] = 'D';
|
|
||||||
((u08_t *)base_coin)[0x1 ^ 3] = 'E';
|
|
||||||
((u08_t *)base_coin)[0x2 ^ 3] = 'T';
|
|
||||||
((u08_t *)base_coin)[0x3 ^ 3] = 'I';
|
|
||||||
((u08_t *)base_coin)[0x4 ^ 3] = ' ';
|
|
||||||
((u08_t *)base_coin)[0x5 ^ 3] = 'c';
|
|
||||||
((u08_t *)base_coin)[0x6 ^ 3] = 'o';
|
|
||||||
((u08_t *)base_coin)[0x7 ^ 3] = 'i';
|
|
||||||
((u08_t *)base_coin)[0x8 ^ 3] = 'n';
|
|
||||||
((u08_t *)base_coin)[0x9 ^ 3] = ' ';
|
|
||||||
((u08_t *)base_coin)[0xa ^ 3] = '2';
|
|
||||||
((u08_t *)base_coin)[0xb ^ 3] = ' ';
|
|
||||||
((u08_t *)base_coin)[0x36 ^ 3] = '\n';
|
|
||||||
((u08_t *)base_coin)[0x37 ^ 3] = 0x80;
|
|
||||||
|
|
||||||
// Initialize variable part with A-Z cycling pattern (offset per thread)
|
|
||||||
int thread_id = omp_get_thread_num();
|
|
||||||
for(int i = 12; i < 54; i++)
|
|
||||||
((u08_t *)base_coin)[i ^ 3] = 'A' + ((i - 12 + thread_id * SIMD_WIDTH) % 26);
|
|
||||||
|
|
||||||
while(keep_running && (max_attempts == 0 || (attempts + thread_attempts) < max_attempts))
|
|
||||||
{
|
|
||||||
prepare_coins(base_coin, interleaved_data, SIMD_WIDTH);
|
|
||||||
sha1_avx2((v8si *)interleaved_data, (v8si *)interleaved_hash);
|
|
||||||
thread_attempts += SIMD_WIDTH;
|
|
||||||
|
|
||||||
u32_t hashes[SIMD_WIDTH][5];
|
|
||||||
extract_hashes(interleaved_hash, hashes, SIMD_WIDTH);
|
|
||||||
|
|
||||||
for(int lane = 0; lane < SIMD_WIDTH; lane++)
|
|
||||||
{
|
|
||||||
if(is_valid_coin(hashes[lane]))
|
|
||||||
{
|
|
||||||
#pragma omp critical
|
|
||||||
{
|
|
||||||
coins_found++;
|
|
||||||
u32_t coins[SIMD_WIDTH][14];
|
|
||||||
extract_coins(interleaved_data, coins, SIMD_WIDTH);
|
|
||||||
printf("COIN FOUND! (attempt %llu, thread %d, lane %d)\n",
|
|
||||||
(unsigned long long)(attempts + thread_attempts - SIMD_WIDTH + lane),
|
|
||||||
thread_id, lane);
|
|
||||||
save_coin(coins[lane]);
|
|
||||||
}
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
||||||
// Print progress every 1M attempts (per thread)
|
|
||||||
if(thread_attempts % 1000000 < SIMD_WIDTH)
|
|
||||||
{
|
|
||||||
#pragma omp critical
|
|
||||||
{
|
|
||||||
time_measurement();
|
|
||||||
double current_time = wall_time_delta() - start_time;
|
|
||||||
double rate = (attempts + thread_attempts) / current_time;
|
|
||||||
printf("Attempts: %llu, Rate: %.2f MH/s, Coins: %u\n",
|
|
||||||
(unsigned long long)(attempts + thread_attempts), rate / 1e6, coins_found);
|
|
||||||
}
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
||||||
#pragma omp atomic
|
|
||||||
attempts += thread_attempts;
|
|
||||||
}
|
|
||||||
|
|
||||||
time_measurement();
|
|
||||||
double total_time = wall_time_delta() - start_time;
|
|
||||||
|
|
||||||
printf("\n=== Mining Statistics ===\n");
|
printf("\n=== Mining Statistics ===\n");
|
||||||
printf("Total attempts: %llu\n", (unsigned long long)attempts);
|
printf("Total attempts: %llu\n", (unsigned long long)attempts);
|
||||||
|
|
@ -411,13 +272,8 @@ int main(int argc, char *argv[])
|
||||||
max_attempts = strtoull(argv[1], NULL, 10);
|
max_attempts = strtoull(argv[1], NULL, 10);
|
||||||
|
|
||||||
#if defined(__AVX2__)
|
#if defined(__AVX2__)
|
||||||
if(argc > 2 && strcmp(argv[2], "omp") == 0) {
|
printf("Using AVX2 implementation\n");
|
||||||
printf("Using AVX2 + OpenMP implementation\n");
|
mine_coins_avx2(max_attempts);
|
||||||
mine_coins_avx2_omp(max_attempts);
|
|
||||||
} else {
|
|
||||||
printf("Using AVX2 implementation\n");
|
|
||||||
mine_coins_avx2(max_attempts);
|
|
||||||
}
|
|
||||||
#elif defined(__AVX__)
|
#elif defined(__AVX__)
|
||||||
printf("Using AVX implementation\n");
|
printf("Using AVX implementation\n");
|
||||||
mine_coins_avx(max_attempts);
|
mine_coins_avx(max_attempts);
|
||||||
|
|
@ -428,3 +284,4 @@ int main(int argc, char *argv[])
|
||||||
|
|
||||||
return 0;
|
return 0;
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
|
||||||
|
|
@ -1,303 +0,0 @@
|
||||||
//
|
|
||||||
// Arquiteturas de Alto Desempenho 2025/2026
|
|
||||||
//
|
|
||||||
// DETI Coin Miner - WebAssembly implementation
|
|
||||||
//
|
|
||||||
|
|
||||||
#include <time.h>
|
|
||||||
#include <stdio.h>
|
|
||||||
#include <stdlib.h>
|
|
||||||
#include <string.h>
|
|
||||||
#include "aad_data_types.h"
|
|
||||||
|
|
||||||
#ifdef __EMSCRIPTEN__
|
|
||||||
#include <emscripten.h>
|
|
||||||
#include "aad_sha1_wasm.h"
|
|
||||||
#else
|
|
||||||
#include "aad_sha1_cpu.h"
|
|
||||||
#include "aad_utilities.h"
|
|
||||||
#include "aad_vault.h"
|
|
||||||
#endif
|
|
||||||
|
|
||||||
// Global mining state
|
|
||||||
static volatile int keep_running = 1;
|
|
||||||
static u64_t total_attempts = 0;
|
|
||||||
static u32_t coins_found = 0;
|
|
||||||
static double mining_start_time = 0;
|
|
||||||
static double pause_time_offset = 0; // Track paused time
|
|
||||||
static double last_pause_time = 0; // When mining was paused
|
|
||||||
static u32_t found_coins[1024][14]; // Store up to 1024 found coins
|
|
||||||
static u32_t found_coins_count = 0;
|
|
||||||
|
|
||||||
//
|
|
||||||
// Check if a hash starts with aad20250
|
|
||||||
//
|
|
||||||
static int is_valid_coin(u32_t *hash)
|
|
||||||
{
|
|
||||||
return hash[0] == 0xAAD20250u;
|
|
||||||
}
|
|
||||||
|
|
||||||
//
|
|
||||||
// Increment coin variable part (positions 12-53)
|
|
||||||
//
|
|
||||||
static int increment_coin(u32_t coin[14])
|
|
||||||
{
|
|
||||||
int pos = 53;
|
|
||||||
while(pos >= 12)
|
|
||||||
{
|
|
||||||
u08_t *byte = &((u08_t *)coin)[pos ^ 3];
|
|
||||||
if(*byte == '\n' || *byte == 0x80)
|
|
||||||
*byte = 32;
|
|
||||||
|
|
||||||
(*byte)++;
|
|
||||||
|
|
||||||
if(*byte == '\n')
|
|
||||||
(*byte)++;
|
|
||||||
|
|
||||||
if(*byte >= 127)
|
|
||||||
{
|
|
||||||
*byte = 32;
|
|
||||||
pos--;
|
|
||||||
}
|
|
||||||
else
|
|
||||||
{
|
|
||||||
break;
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
||||||
return (pos >= 12);
|
|
||||||
}
|
|
||||||
|
|
||||||
//
|
|
||||||
// Get current time in seconds
|
|
||||||
//
|
|
||||||
static double get_time()
|
|
||||||
{
|
|
||||||
#ifdef __EMSCRIPTEN__
|
|
||||||
return emscripten_get_now() / 1000.0;
|
|
||||||
#else
|
|
||||||
struct timespec ts;
|
|
||||||
clock_gettime(CLOCK_MONOTONIC_RAW, &ts);
|
|
||||||
return (double)ts.tv_sec + (double)ts.tv_nsec / 1e9;
|
|
||||||
#endif
|
|
||||||
}
|
|
||||||
|
|
||||||
//
|
|
||||||
// Main mining iteration (called from JavaScript)
|
|
||||||
//
|
|
||||||
#ifdef __EMSCRIPTEN__
|
|
||||||
EMSCRIPTEN_KEEPALIVE
|
|
||||||
#endif
|
|
||||||
int mine_coins_wasm(u32_t iterations_per_call)
|
|
||||||
{
|
|
||||||
static u32_t coin[14];
|
|
||||||
static int initialized = 0;
|
|
||||||
u32_t hash[5];
|
|
||||||
|
|
||||||
// Initialize coin template on first call
|
|
||||||
if(!initialized)
|
|
||||||
{
|
|
||||||
memset(coin, 0, sizeof(coin));
|
|
||||||
((u08_t *)coin)[0x0 ^ 3] = 'D';
|
|
||||||
((u08_t *)coin)[0x1 ^ 3] = 'E';
|
|
||||||
((u08_t *)coin)[0x2 ^ 3] = 'T';
|
|
||||||
((u08_t *)coin)[0x3 ^ 3] = 'I';
|
|
||||||
((u08_t *)coin)[0x4 ^ 3] = ' ';
|
|
||||||
((u08_t *)coin)[0x5 ^ 3] = 'c';
|
|
||||||
((u08_t *)coin)[0x6 ^ 3] = 'o';
|
|
||||||
((u08_t *)coin)[0x7 ^ 3] = 'i';
|
|
||||||
((u08_t *)coin)[0x8 ^ 3] = 'n';
|
|
||||||
((u08_t *)coin)[0x9 ^ 3] = ' ';
|
|
||||||
((u08_t *)coin)[0xa ^ 3] = '2';
|
|
||||||
((u08_t *)coin)[0xb ^ 3] = ' ';
|
|
||||||
((u08_t *)coin)[0x36 ^ 3] = '\n';
|
|
||||||
((u08_t *)coin)[0x37 ^ 3] = 0x80;
|
|
||||||
|
|
||||||
for(int i = 12; i < 54; i++)
|
|
||||||
((u08_t *)coin)[i ^ 3] = 'A' + (i - 12) % 26;
|
|
||||||
|
|
||||||
mining_start_time = get_time();
|
|
||||||
initialized = 1;
|
|
||||||
}
|
|
||||||
|
|
||||||
if(!keep_running)
|
|
||||||
return 0;
|
|
||||||
|
|
||||||
// Mine for the specified number of iterations
|
|
||||||
for(u32_t i = 0; i < iterations_per_call && keep_running; i++)
|
|
||||||
{
|
|
||||||
sha1(coin, hash);
|
|
||||||
total_attempts++;
|
|
||||||
|
|
||||||
if(is_valid_coin(hash))
|
|
||||||
{
|
|
||||||
if(found_coins_count < 1024)
|
|
||||||
{
|
|
||||||
memcpy(found_coins[found_coins_count], coin, sizeof(coin));
|
|
||||||
found_coins_count++;
|
|
||||||
}
|
|
||||||
coins_found++;
|
|
||||||
|
|
||||||
#ifndef __EMSCRIPTEN__
|
|
||||||
printf("COIN FOUND! (attempt %llu)\n", (unsigned long long)total_attempts);
|
|
||||||
save_coin(coin);
|
|
||||||
#endif
|
|
||||||
}
|
|
||||||
|
|
||||||
increment_coin(coin);
|
|
||||||
}
|
|
||||||
|
|
||||||
return keep_running;
|
|
||||||
}
|
|
||||||
|
|
||||||
//
|
|
||||||
// Get mining statistics
|
|
||||||
//
|
|
||||||
#ifdef __EMSCRIPTEN__
|
|
||||||
EMSCRIPTEN_KEEPALIVE
|
|
||||||
#endif
|
|
||||||
void get_statistics(u64_t *attempts, u32_t *coins, double *hash_rate, double *elapsed_time)
|
|
||||||
{
|
|
||||||
*attempts = total_attempts;
|
|
||||||
*coins = coins_found;
|
|
||||||
|
|
||||||
double current_time = get_time();
|
|
||||||
double actual_elapsed;
|
|
||||||
|
|
||||||
if(!keep_running && last_pause_time > 0) {
|
|
||||||
// If paused, use the paused time
|
|
||||||
actual_elapsed = last_pause_time - mining_start_time - pause_time_offset;
|
|
||||||
} else {
|
|
||||||
// If running, calculate normally
|
|
||||||
actual_elapsed = current_time - mining_start_time - pause_time_offset;
|
|
||||||
}
|
|
||||||
|
|
||||||
*elapsed_time = actual_elapsed;
|
|
||||||
*hash_rate = (actual_elapsed > 0) ? (total_attempts / actual_elapsed) : 0;
|
|
||||||
}
|
|
||||||
|
|
||||||
//
|
|
||||||
// Stop mining
|
|
||||||
//
|
|
||||||
#ifdef __EMSCRIPTEN__
|
|
||||||
EMSCRIPTEN_KEEPALIVE
|
|
||||||
#endif
|
|
||||||
void stop_mining()
|
|
||||||
{
|
|
||||||
if(keep_running) {
|
|
||||||
keep_running = 0;
|
|
||||||
last_pause_time = get_time();
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
||||||
//
|
|
||||||
// Resume mining
|
|
||||||
//
|
|
||||||
#ifdef __EMSCRIPTEN__
|
|
||||||
EMSCRIPTEN_KEEPALIVE
|
|
||||||
#endif
|
|
||||||
void resume_mining()
|
|
||||||
{
|
|
||||||
if(!keep_running && last_pause_time > 0) {
|
|
||||||
double pause_duration = get_time() - last_pause_time;
|
|
||||||
pause_time_offset += pause_duration;
|
|
||||||
keep_running = 1;
|
|
||||||
last_pause_time = 0;
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
||||||
//
|
|
||||||
// Get found coin data (returns pointer to coin array)
|
|
||||||
//
|
|
||||||
#ifdef __EMSCRIPTEN__
|
|
||||||
EMSCRIPTEN_KEEPALIVE
|
|
||||||
#endif
|
|
||||||
u32_t* get_found_coin(u32_t index)
|
|
||||||
{
|
|
||||||
if(index < found_coins_count)
|
|
||||||
return found_coins[index];
|
|
||||||
return NULL;
|
|
||||||
}
|
|
||||||
|
|
||||||
//
|
|
||||||
// Get number of found coins
|
|
||||||
//
|
|
||||||
#ifdef __EMSCRIPTEN__
|
|
||||||
EMSCRIPTEN_KEEPALIVE
|
|
||||||
#endif
|
|
||||||
u32_t get_found_coins_count()
|
|
||||||
{
|
|
||||||
return found_coins_count;
|
|
||||||
}
|
|
||||||
|
|
||||||
//
|
|
||||||
// Reset mining state
|
|
||||||
//
|
|
||||||
#ifdef __EMSCRIPTEN__
|
|
||||||
EMSCRIPTEN_KEEPALIVE
|
|
||||||
#endif
|
|
||||||
void reset_mining()
|
|
||||||
{
|
|
||||||
keep_running = 1;
|
|
||||||
total_attempts = 0;
|
|
||||||
coins_found = 0;
|
|
||||||
found_coins_count = 0;
|
|
||||||
pause_time_offset = 0;
|
|
||||||
last_pause_time = 0;
|
|
||||||
mining_start_time = get_time();
|
|
||||||
}
|
|
||||||
|
|
||||||
//
|
|
||||||
// Main function (for standalone compilation/testing)
|
|
||||||
//
|
|
||||||
#ifndef __EMSCRIPTEN__
|
|
||||||
int main(int argc, char *argv[])
|
|
||||||
{
|
|
||||||
u64_t max_attempts = 0;
|
|
||||||
|
|
||||||
if(argc > 1)
|
|
||||||
max_attempts = strtoull(argv[1], NULL, 10);
|
|
||||||
|
|
||||||
printf("Mining DETI coins using WebAssembly implementation (standalone mode)...\n");
|
|
||||||
printf("Press Ctrl+C to stop\n\n");
|
|
||||||
|
|
||||||
time_measurement();
|
|
||||||
time_measurement();
|
|
||||||
double start_time = wall_time_delta();
|
|
||||||
double last_report = start_time;
|
|
||||||
|
|
||||||
while(keep_running && (max_attempts == 0 || total_attempts < max_attempts))
|
|
||||||
{
|
|
||||||
mine_coins_wasm(100000);
|
|
||||||
|
|
||||||
time_measurement();
|
|
||||||
double current_time = wall_time_delta() - start_time;
|
|
||||||
|
|
||||||
if(current_time - last_report >= 1.0)
|
|
||||||
{
|
|
||||||
u64_t attempts;
|
|
||||||
u32_t coins;
|
|
||||||
double hash_rate, elapsed;
|
|
||||||
get_statistics(&attempts, &coins, &hash_rate, &elapsed);
|
|
||||||
|
|
||||||
printf("Attempts: %llu, Rate: %.2f MH/s, Coins: %u\n",
|
|
||||||
(unsigned long long)attempts, hash_rate / 1e6, coins);
|
|
||||||
|
|
||||||
last_report = current_time;
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
||||||
time_measurement();
|
|
||||||
double total_time = wall_time_delta() - start_time;
|
|
||||||
|
|
||||||
printf("\n=== Mining Statistics ===\n");
|
|
||||||
printf("Total attempts: %llu\n", (unsigned long long)total_attempts);
|
|
||||||
printf("Total time: %.2f seconds\n", total_time);
|
|
||||||
printf("Average rate: %.2f attempts/second\n", total_attempts / total_time);
|
|
||||||
printf("Coins found: %u\n", coins_found);
|
|
||||||
|
|
||||||
save_coin(NULL);
|
|
||||||
|
|
||||||
return 0;
|
|
||||||
}
|
|
||||||
#endif
|
|
||||||
|
|
@ -36,7 +36,7 @@ typedef struct
|
||||||
CUdeviceptr device_data[2]; // the pointers to the device data
|
CUdeviceptr device_data[2]; // the pointers to the device data
|
||||||
// launch kernel data
|
// launch kernel data
|
||||||
unsigned int grid_dim_x; // the number of grid blocks (in the X dimension, the only one we will use here)
|
unsigned int grid_dim_x; // the number of grid blocks (in the X dimension, the only one we will use here)
|
||||||
unsigned int block_dim_x; // the number of threads in a block (in the X dimension, the only one we will use here, should be equal to RECOMMENDED_CUDA_BLOCK_SIZE)
|
unsigned int block_dim_x; // the number of threads in a block (in the X dimension, the only one we will use here, should be equal to RECOMENDED_CUDA_BLOCK_SIZE)
|
||||||
int n_kernel_arguments; // number of kernel arguments
|
int n_kernel_arguments; // number of kernel arguments
|
||||||
void *arg[MAX_N_ARGUMENTS]; // pointers to the kernel argument data
|
void *arg[MAX_N_ARGUMENTS]; // pointers to the kernel argument data
|
||||||
|
|
||||||
|
|
@ -260,10 +260,10 @@ static void device_to_host_copy(cuda_data_t *cd,int idx)
|
||||||
// launch a CUDA kernel (with 0 bytes of shared memory and no extra options)
|
// launch a CUDA kernel (with 0 bytes of shared memory and no extra options)
|
||||||
//
|
//
|
||||||
|
|
||||||
static void launch_kernel(cuda_data_t *cd)
|
static void lauch_kernel(cuda_data_t *cd)
|
||||||
{
|
{
|
||||||
if(cd->block_dim_x != (unsigned int)RECOMMENDED_CUDA_BLOCK_SIZE)
|
if(cd->block_dim_x != (unsigned int)RECOMENDED_CUDA_BLOCK_SIZE)
|
||||||
fprintf(stderr,"launch_kernel(): block_dim_x should be equal to %d\n",RECOMMENDED_CUDA_BLOCK_SIZE);
|
fprintf(stderr,"lauch_kernel(): block_dim_x should be equal to %d\n",RECOMENDED_CUDA_BLOCK_SIZE);
|
||||||
CU_CALL( cuLaunchKernel , (cd->cu_kernel,cd->grid_dim_x,1u,1u,cd->block_dim_x,1u,1u,0u,cd->cu_stream,&cd->arg[0],NULL) );
|
CU_CALL( cuLaunchKernel , (cd->cu_kernel,cd->grid_dim_x,1u,1u,cd->block_dim_x,1u,1u,0u,cd->cu_stream,&cd->arg[0],NULL) );
|
||||||
synchronize_cuda(cd);
|
synchronize_cuda(cd);
|
||||||
}
|
}
|
||||||
|
|
|
||||||
|
|
@ -28,7 +28,7 @@
|
||||||
//
|
//
|
||||||
// we place this here to simplify things (aad_sha1_cuda_kernel.cu includes this file...)
|
// we place this here to simplify things (aad_sha1_cuda_kernel.cu includes this file...)
|
||||||
//
|
//
|
||||||
#define RECOMMENDED_CUDA_BLOCK_SIZE 128
|
#define RECOMENDED_CUDA_BLOCK_SIZE 128
|
||||||
|
|
||||||
|
|
||||||
//
|
//
|
||||||
|
|
|
||||||
|
|
@ -165,7 +165,7 @@ static void test_sha1_avx(int n_tests,int n_measurements)
|
||||||
for(n = 0;n < n_measurements;n++)
|
for(n = 0;n < n_measurements;n++)
|
||||||
{
|
{
|
||||||
interleaved_data[0][0]++;
|
interleaved_data[0][0]++;
|
||||||
sha1_avx((v4si *)&interleaved_data[0],(v4si *)&interleaved_hash[0]);
|
sha1(&data[lane].i[0],&hash[lane].i[0]);
|
||||||
sum += interleaved_hash[4][0];
|
sum += interleaved_hash[4][0];
|
||||||
}
|
}
|
||||||
time_measurement();
|
time_measurement();
|
||||||
|
|
@ -235,7 +235,7 @@ static void test_sha1_avx2(int n_tests,int n_measurements)
|
||||||
for(n = 0;n < n_measurements;n++)
|
for(n = 0;n < n_measurements;n++)
|
||||||
{
|
{
|
||||||
interleaved_data[0][0]++;
|
interleaved_data[0][0]++;
|
||||||
sha1_avx2((v8si *)&interleaved_data[0],(v8si *)&interleaved_hash[0]);
|
sha1(&data[lane].i[0],&hash[lane].i[0]);
|
||||||
sum += interleaved_hash[4][0];
|
sum += interleaved_hash[4][0];
|
||||||
}
|
}
|
||||||
time_measurement();
|
time_measurement();
|
||||||
|
|
@ -305,7 +305,7 @@ static void test_sha1_avx512f(int n_tests,int n_measurements)
|
||||||
for(n = 0;n < n_measurements;n++)
|
for(n = 0;n < n_measurements;n++)
|
||||||
{
|
{
|
||||||
interleaved_data[0][0]++;
|
interleaved_data[0][0]++;
|
||||||
sha1_avx512f((v16si *)&interleaved_data[0],(v16si *)&interleaved_hash[0]);
|
sha1(&data[lane].i[0],&hash[lane].i[0]);
|
||||||
sum += interleaved_hash[4][0];
|
sum += interleaved_hash[4][0];
|
||||||
}
|
}
|
||||||
time_measurement();
|
time_measurement();
|
||||||
|
|
@ -375,7 +375,7 @@ static void test_sha1_neon(int n_tests,int n_measurements)
|
||||||
for(n = 0;n < n_measurements;n++)
|
for(n = 0;n < n_measurements;n++)
|
||||||
{
|
{
|
||||||
interleaved_data[0][0]++;
|
interleaved_data[0][0]++;
|
||||||
sha1_neon((uint32x4_t *)&interleaved_data[0],(uint32x4_t *)&interleaved_hash[0]);
|
sha1(&data[lane].i[0],&hash[lane].i[0]);
|
||||||
sum += interleaved_hash[4][0];
|
sum += interleaved_hash[4][0];
|
||||||
}
|
}
|
||||||
time_measurement();
|
time_measurement();
|
||||||
|
|
|
||||||
|
|
@ -12,7 +12,8 @@
|
||||||
//
|
//
|
||||||
|
|
||||||
#include "aad_sha1.h"
|
#include "aad_sha1.h"
|
||||||
#include "aad_data_types.h"
|
|
||||||
|
typedef unsigned int u32_t;
|
||||||
|
|
||||||
//
|
//
|
||||||
// the nvcc compiler stores w[] in registers (constant indices!)
|
// the nvcc compiler stores w[] in registers (constant indices!)
|
||||||
|
|
@ -22,7 +23,7 @@
|
||||||
// warp thread number: n & 31 -- the lane
|
// warp thread number: n & 31 -- the lane
|
||||||
//
|
//
|
||||||
|
|
||||||
extern "C" __global__ __launch_bounds__(RECOMMENDED_CUDA_BLOCK_SIZE,1)
|
extern "C" __global__ __launch_bounds__(RECOMENDED_CUDA_BLOCK_SIZE,1)
|
||||||
void sha1_cuda_kernel(u32_t *interleaved32_data,u32_t *interleaved32_hash)
|
void sha1_cuda_kernel(u32_t *interleaved32_data,u32_t *interleaved32_hash)
|
||||||
{
|
{
|
||||||
u32_t n;
|
u32_t n;
|
||||||
|
|
|
||||||
|
|
@ -19,7 +19,7 @@ static void test_sha1_cuda(int n_tests)
|
||||||
double host_to_device_time,kernel_time,device_to_host_time,hashes_per_second;
|
double host_to_device_time,kernel_time,device_to_host_time,hashes_per_second;
|
||||||
cuda_data_t cd;
|
cuda_data_t cd;
|
||||||
|
|
||||||
if(n_tests <= 0 || n_tests > (1 << 24) || n_tests % RECOMMENDED_CUDA_BLOCK_SIZE != 0)
|
if(n_tests <= 0 || n_tests > (1 << 24) || n_tests % RECOMENDED_CUDA_BLOCK_SIZE != 0)
|
||||||
{
|
{
|
||||||
fprintf(stderr,"test_sha1_cuda(): bad number of tests\n");
|
fprintf(stderr,"test_sha1_cuda(): bad number of tests\n");
|
||||||
exit(1);
|
exit(1);
|
||||||
|
|
@ -44,13 +44,13 @@ static void test_sha1_cuda(int n_tests)
|
||||||
host_to_device_copy(&cd,0); // idx=0 means that the interleaved32_data is copied to the CUDA device
|
host_to_device_copy(&cd,0); // idx=0 means that the interleaved32_data is copied to the CUDA device
|
||||||
time_measurement();
|
time_measurement();
|
||||||
host_to_device_time = wall_time_delta();
|
host_to_device_time = wall_time_delta();
|
||||||
cd.grid_dim_x = (u32_t)n_tests / (u32_t)RECOMMENDED_CUDA_BLOCK_SIZE;
|
cd.grid_dim_x = (u32_t)n_tests / (u32_t)RECOMENDED_CUDA_BLOCK_SIZE;
|
||||||
cd.block_dim_x = (u32_t)RECOMMENDED_CUDA_BLOCK_SIZE;
|
cd.block_dim_x = (u32_t)RECOMENDED_CUDA_BLOCK_SIZE;
|
||||||
cd.n_kernel_arguments = 2;
|
cd.n_kernel_arguments = 2;
|
||||||
cd.arg[0] = &cd.device_data[0]; // interleaved32_data
|
cd.arg[0] = &cd.device_data[0]; // interleaved32_data
|
||||||
cd.arg[1] = &cd.device_data[1]; // interleaved32_hash
|
cd.arg[1] = &cd.device_data[1]; // interleaved32_hash
|
||||||
time_measurement();
|
time_measurement();
|
||||||
launch_kernel(&cd);
|
lauch_kernel(&cd);
|
||||||
time_measurement();
|
time_measurement();
|
||||||
kernel_time = wall_time_delta();
|
kernel_time = wall_time_delta();
|
||||||
time_measurement();
|
time_measurement();
|
||||||
|
|
|
||||||
|
|
@ -1,35 +0,0 @@
|
||||||
#ifndef AAD_SHA1_WASM
|
|
||||||
#define AAD_SHA1_WASM
|
|
||||||
|
|
||||||
#include "aad_data_types.h"
|
|
||||||
#include "aad_sha1.h"
|
|
||||||
|
|
||||||
//
|
|
||||||
// SHA1 hash computation for WebAssembly (scalar implementation)
|
|
||||||
//
|
|
||||||
|
|
||||||
static inline u32_t rotate_left(u32_t x, int n)
|
|
||||||
{
|
|
||||||
return (x << n) | (x >> (32 - n));
|
|
||||||
}
|
|
||||||
|
|
||||||
static void sha1(u32_t *coin, u32_t *hash)
|
|
||||||
{
|
|
||||||
// Define the macros needed by CUSTOM_SHA1_CODE
|
|
||||||
#define T u32_t
|
|
||||||
#define C(c) (c)
|
|
||||||
#define ROTATE(x,n) rotate_left(x,n)
|
|
||||||
#define DATA(idx) coin[idx]
|
|
||||||
#define HASH(idx) hash[idx]
|
|
||||||
|
|
||||||
// Use the standard SHA1 template from aad_sha1.h
|
|
||||||
CUSTOM_SHA1_CODE();
|
|
||||||
|
|
||||||
#undef T
|
|
||||||
#undef C
|
|
||||||
#undef ROTATE
|
|
||||||
#undef DATA
|
|
||||||
#undef HASH
|
|
||||||
}
|
|
||||||
|
|
||||||
#endif
|
|
||||||
258
index.html
258
index.html
|
|
@ -1,258 +0,0 @@
|
||||||
<!DOCTYPE html>
|
|
||||||
<html>
|
|
||||||
<head>
|
|
||||||
<title>DETI Coin Miner - WebAssembly</title>
|
|
||||||
<style>
|
|
||||||
body {
|
|
||||||
font-family: monospace;
|
|
||||||
padding: 20px;
|
|
||||||
max-width: 1200px;
|
|
||||||
margin: 0 auto;
|
|
||||||
}
|
|
||||||
button {
|
|
||||||
padding: 10px 20px;
|
|
||||||
margin: 5px;
|
|
||||||
font-size: 16px;
|
|
||||||
}
|
|
||||||
#stats {
|
|
||||||
margin-top: 20px;
|
|
||||||
padding: 10px;
|
|
||||||
background: #f0f0f0;
|
|
||||||
border-radius: 5px;
|
|
||||||
box-sizing: border-box;
|
|
||||||
}
|
|
||||||
#coins-container {
|
|
||||||
margin-top: 20px;
|
|
||||||
}
|
|
||||||
#coins {
|
|
||||||
width: 100%;
|
|
||||||
height: 300px;
|
|
||||||
padding: 10px;
|
|
||||||
background: #f9f9f9;
|
|
||||||
border: 2px solid #ccc;
|
|
||||||
border-radius: 5px;
|
|
||||||
font-family: monospace;
|
|
||||||
font-size: 12px;
|
|
||||||
overflow-y: auto;
|
|
||||||
white-space: pre-wrap;
|
|
||||||
word-wrap: break-word;
|
|
||||||
box-sizing: border-box;
|
|
||||||
}
|
|
||||||
.control-group {
|
|
||||||
margin: 10px 0;
|
|
||||||
}
|
|
||||||
label {
|
|
||||||
display: inline-block;
|
|
||||||
width: 200px;
|
|
||||||
}
|
|
||||||
h2 {
|
|
||||||
margin-top: 20px;
|
|
||||||
margin-bottom: 10px;
|
|
||||||
}
|
|
||||||
.coin-entry {
|
|
||||||
color: #006400;
|
|
||||||
font-weight: bold;
|
|
||||||
}
|
|
||||||
.coin-data {
|
|
||||||
color: #000080;
|
|
||||||
}
|
|
||||||
</style>
|
|
||||||
</head>
|
|
||||||
<body>
|
|
||||||
<h1>DETI Coin Miner (WebAssembly)</h1>
|
|
||||||
<div class="control-group">
|
|
||||||
<label>Iterations per batch:</label>
|
|
||||||
<input type="number" id="batchSize" value="1000000" step="100000">
|
|
||||||
</div>
|
|
||||||
<div class="control-group">
|
|
||||||
<label>Update interval (ms):</label>
|
|
||||||
<input type="number" id="updateInterval" value="100" step="50">
|
|
||||||
</div>
|
|
||||||
<button id="start">Start Mining</button>
|
|
||||||
<button id="stop">Stop Mining</button>
|
|
||||||
<button id="reset">Reset</button>
|
|
||||||
<button id="clearCoins">Clear Coins Display</button>
|
|
||||||
|
|
||||||
<div id="stats">
|
|
||||||
Waiting to start...
|
|
||||||
</div>
|
|
||||||
|
|
||||||
<div id="coins-container">
|
|
||||||
<h2>Found Coins (<span id="coin-count">0</span>)</h2>
|
|
||||||
<div id="coins"></div>
|
|
||||||
</div>
|
|
||||||
|
|
||||||
<script src="coin_miner_wasm.js"></script>
|
|
||||||
<script>
|
|
||||||
let mining = false;
|
|
||||||
let Module;
|
|
||||||
let miningInterval;
|
|
||||||
let updateInterval;
|
|
||||||
let lastDisplayedCoinCount = 0;
|
|
||||||
let pausedStats = false;
|
|
||||||
|
|
||||||
CoinMinerModule().then(mod => {
|
|
||||||
Module = mod;
|
|
||||||
console.log('WebAssembly module loaded');
|
|
||||||
|
|
||||||
document.getElementById('start').onclick = () => {
|
|
||||||
if (!mining) {
|
|
||||||
mining = true;
|
|
||||||
Module._resume_mining();
|
|
||||||
pausedStats = false;
|
|
||||||
console.log('Starting mining...');
|
|
||||||
startMining();
|
|
||||||
}
|
|
||||||
};
|
|
||||||
|
|
||||||
document.getElementById('stop').onclick = () => {
|
|
||||||
mining = false;
|
|
||||||
Module._stop_mining();
|
|
||||||
clearInterval(miningInterval);
|
|
||||||
clearInterval(updateInterval);
|
|
||||||
|
|
||||||
updateStats();
|
|
||||||
|
|
||||||
let currentHTML = document.getElementById('stats').innerHTML;
|
|
||||||
document.getElementById('stats').innerHTML = currentHTML.replace('Mining Statistics:', 'Mining Statistics (PAUSED):');
|
|
||||||
|
|
||||||
pausedStats = true;
|
|
||||||
|
|
||||||
console.log('Mining stopped');
|
|
||||||
};
|
|
||||||
|
|
||||||
document.getElementById('reset').onclick = () => {
|
|
||||||
Module._reset_mining();
|
|
||||||
mining = false;
|
|
||||||
pausedStats = false;
|
|
||||||
clearInterval(miningInterval);
|
|
||||||
clearInterval(updateInterval);
|
|
||||||
lastDisplayedCoinCount = 0;
|
|
||||||
document.getElementById('stats').innerHTML = 'Reset complete. Click Start to begin.';
|
|
||||||
document.getElementById('coins').innerHTML = '';
|
|
||||||
document.getElementById('coin-count').textContent = '0';
|
|
||||||
console.log('Mining reset');
|
|
||||||
};
|
|
||||||
|
|
||||||
document.getElementById('clearCoins').onclick = () => {
|
|
||||||
document.getElementById('coins').innerHTML = '';
|
|
||||||
console.log('Coins display cleared');
|
|
||||||
};
|
|
||||||
|
|
||||||
function coinToString(coinPtr) {
|
|
||||||
let coinStr = '';
|
|
||||||
for (let i = 0; i < 55; i++) {
|
|
||||||
const byteIdx = i ^ 3;
|
|
||||||
const wordIdx = Math.floor(byteIdx / 4);
|
|
||||||
const byteInWord = byteIdx % 4;
|
|
||||||
|
|
||||||
const word = Module.getValue(coinPtr + wordIdx * 4, 'i32');
|
|
||||||
const byte = (word >> (byteInWord * 8)) & 0xFF;
|
|
||||||
|
|
||||||
if (byte >= 32 && byte <= 126) {
|
|
||||||
coinStr += String.fromCharCode(byte);
|
|
||||||
} else if (byte === 10) {
|
|
||||||
coinStr += '\\n';
|
|
||||||
} else if (byte === 0x80) {
|
|
||||||
coinStr += '[0x80]';
|
|
||||||
} else {
|
|
||||||
coinStr += `[0x${byte.toString(16).padStart(2, '0')}]`;
|
|
||||||
}
|
|
||||||
}
|
|
||||||
return coinStr;
|
|
||||||
}
|
|
||||||
|
|
||||||
function displayNewCoins() {
|
|
||||||
const totalCoins = Module._get_found_coins_count();
|
|
||||||
|
|
||||||
if (totalCoins > lastDisplayedCoinCount) {
|
|
||||||
const coinsDiv = document.getElementById('coins');
|
|
||||||
|
|
||||||
for (let i = lastDisplayedCoinCount; i < totalCoins; i++) {
|
|
||||||
const coinPtr = Module._get_found_coin(i);
|
|
||||||
if (coinPtr !== 0) {
|
|
||||||
const coinStr = coinToString(coinPtr);
|
|
||||||
const timestamp = new Date().toLocaleTimeString();
|
|
||||||
|
|
||||||
const entry = document.createElement('div');
|
|
||||||
entry.innerHTML = `<span class="coin-entry">[${timestamp}] Coin #${i + 1}:</span> <span class="coin-data">${coinStr}</span>`;
|
|
||||||
coinsDiv.appendChild(entry);
|
|
||||||
|
|
||||||
coinsDiv.scrollTop = coinsDiv.scrollHeight;
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
||||||
lastDisplayedCoinCount = totalCoins;
|
|
||||||
document.getElementById('coin-count').textContent = totalCoins;
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
||||||
function updateStats() {
|
|
||||||
const attemptsPtr = Module._malloc(8);
|
|
||||||
const coinsPtr = Module._malloc(4);
|
|
||||||
const hashRatePtr = Module._malloc(8);
|
|
||||||
const elapsedPtr = Module._malloc(8);
|
|
||||||
|
|
||||||
// Zero-initialize the memory before calling
|
|
||||||
Module.setValue(attemptsPtr, 0, 'i32');
|
|
||||||
Module.setValue(attemptsPtr + 4, 0, 'i32');
|
|
||||||
|
|
||||||
Module._get_statistics(attemptsPtr, coinsPtr, hashRatePtr, elapsedPtr);
|
|
||||||
|
|
||||||
// Read 64-bit unsigned value correctly
|
|
||||||
// On little-endian, low 32 bits come first
|
|
||||||
const attemptsLowUnsigned = Module.getValue(attemptsPtr, 'i32') >>> 0;
|
|
||||||
const attemptsHighUnsigned = Module.getValue(attemptsPtr + 4, 'i32') >>> 0;
|
|
||||||
|
|
||||||
// Combine - for display purposes, if high part is 0, just show low part
|
|
||||||
let attempts;
|
|
||||||
if (attemptsHighUnsigned === 0) {
|
|
||||||
attempts = attemptsLowUnsigned;
|
|
||||||
} else {
|
|
||||||
// Use BigInt for values > 32 bits
|
|
||||||
const low = BigInt(attemptsLowUnsigned);
|
|
||||||
const high = BigInt(attemptsHighUnsigned);
|
|
||||||
attempts = (high * BigInt(4294967296)) + low;
|
|
||||||
}
|
|
||||||
|
|
||||||
const coins = Module.getValue(coinsPtr, 'i32');
|
|
||||||
const hashRate = Module.getValue(hashRatePtr, 'double');
|
|
||||||
const elapsed = Module.getValue(elapsedPtr, 'double');
|
|
||||||
|
|
||||||
Module._free(attemptsPtr);
|
|
||||||
Module._free(coinsPtr);
|
|
||||||
Module._free(hashRatePtr);
|
|
||||||
Module._free(elapsedPtr);
|
|
||||||
|
|
||||||
document.getElementById('stats').innerHTML = `
|
|
||||||
<strong>Mining Statistics:</strong><br>
|
|
||||||
Attempts: ${attempts.toString()}<br>
|
|
||||||
Coins Found: ${coins}<br>
|
|
||||||
Hash Rate: ${(hashRate / 1e6).toFixed(2)} MH/s<br>
|
|
||||||
Elapsed Time: ${elapsed.toFixed(2)} seconds
|
|
||||||
`;
|
|
||||||
|
|
||||||
displayNewCoins();
|
|
||||||
}
|
|
||||||
|
|
||||||
function startMining() {
|
|
||||||
const batchSize = parseInt(document.getElementById('batchSize').value);
|
|
||||||
const updateMs = parseInt(document.getElementById('updateInterval').value);
|
|
||||||
|
|
||||||
miningInterval = setInterval(() => {
|
|
||||||
if (!mining) {
|
|
||||||
clearInterval(miningInterval);
|
|
||||||
return;
|
|
||||||
}
|
|
||||||
Module._mine_coins_wasm(batchSize);
|
|
||||||
}, 0);
|
|
||||||
|
|
||||||
updateInterval = setInterval(updateStats, updateMs);
|
|
||||||
}
|
|
||||||
}).catch(err => {
|
|
||||||
console.error('Failed to load WebAssembly module:', err);
|
|
||||||
document.getElementById('stats').innerHTML = 'Error loading module: ' + err.message;
|
|
||||||
});
|
|
||||||
</script>
|
|
||||||
</body>
|
|
||||||
</html>
|
|
||||||
41
makefile
41
makefile
|
|
@ -32,7 +32,7 @@ OPENCL_DIR = $(CUDA_DIR)
|
||||||
# RTX 4070 -------------- sm_89
|
# RTX 4070 -------------- sm_89
|
||||||
#
|
#
|
||||||
|
|
||||||
CUDA_ARCH = sm_86
|
CUDA_ARCH = sm_75
|
||||||
|
|
||||||
|
|
||||||
#
|
#
|
||||||
|
|
@ -42,9 +42,6 @@ CUDA_ARCH = sm_86
|
||||||
clean:
|
clean:
|
||||||
rm -f sha1_tests
|
rm -f sha1_tests
|
||||||
rm -f sha1_cuda_test sha1_cuda_kernel.cubin
|
rm -f sha1_cuda_test sha1_cuda_kernel.cubin
|
||||||
rm -f coin_miner_cpu coin_miner_simd coin_miner_cuda coin_miner_cuda_kernel.cubin
|
|
||||||
rm -f coin_miner_wasm.js coin_miner_wasm.wasm
|
|
||||||
rm -f benchmark
|
|
||||||
rm -f a.out
|
rm -f a.out
|
||||||
|
|
||||||
|
|
||||||
|
|
@ -67,39 +64,3 @@ sha1_cuda_kernel.cubin: aad_sha1_cuda_kernel.cu aad_sha1.h makefile
|
||||||
nvcc -arch=$(CUDA_ARCH) --compiler-options -O2,-Wall -I$(CUDA_DIR)/include --cubin $< -o $@
|
nvcc -arch=$(CUDA_ARCH) --compiler-options -O2,-Wall -I$(CUDA_DIR)/include --cubin $< -o $@
|
||||||
|
|
||||||
all: sha1_tests sha1_cuda_test sha1_cuda_kernel.cubin
|
all: sha1_tests sha1_cuda_test sha1_cuda_kernel.cubin
|
||||||
|
|
||||||
|
|
||||||
#
|
|
||||||
# DETI coin miners
|
|
||||||
#
|
|
||||||
|
|
||||||
coin_miner_cpu: aad_coin_miner_cpu.c aad_sha1.h aad_sha1_cpu.h aad_data_types.h aad_utilities.h aad_vault.h makefile
|
|
||||||
cc -march=native -Wall -Wshadow -Werror -O3 $< -o $@
|
|
||||||
|
|
||||||
coin_miner_simd: aad_coin_miner_simd.c aad_sha1.h aad_sha1_cpu.h aad_data_types.h aad_utilities.h aad_vault.h makefile
|
|
||||||
cc -march=native -Wall -Wshadow -Werror -fopenmp -mavx2 -O3 $< -o $@
|
|
||||||
|
|
||||||
coin_miner_cuda_kernel.cubin: aad_coin_miner_cuda_kernel.cu aad_sha1.h makefile
|
|
||||||
nvcc -arch=$(CUDA_ARCH) --compiler-options -O2,-Wall -I$(CUDA_DIR)/include --cubin $< -o $@
|
|
||||||
|
|
||||||
coin_miner_cuda: aad_coin_miner_cuda.c coin_miner_cuda_kernel.cubin aad_sha1.h aad_sha1_cpu.h aad_data_types.h aad_utilities.h aad_vault.h aad_cuda_utilities.h makefile
|
|
||||||
cc -march=native -Wall -Wshadow -Werror -O3 -I$(CUDA_DIR)/include $< -o $@ -lcuda
|
|
||||||
|
|
||||||
coin_miner_wasm: aad_coin_miner_wasm.c aad_sha1.h aad_sha1_cpu.h aad_sha1_wasm.h aad_data_types.h aad_utilities.h aad_vault.h makefile
|
|
||||||
emcc -O3 -flto -o coin_miner_wasm.js aad_coin_miner_wasm.c \
|
|
||||||
-s WASM=1 \
|
|
||||||
-s EXPORTED_FUNCTIONS='["_mine_coins_wasm","_get_statistics","_stop_mining","_reset_mining","_get_found_coin","_get_found_coins_count","_malloc","_free"]' \
|
|
||||||
-s EXPORTED_RUNTIME_METHODS='["cwrap","ccall","getValue","setValue"]' \
|
|
||||||
-s ALLOW_MEMORY_GROWTH=1 \
|
|
||||||
-s MODULARIZE=1 \
|
|
||||||
-s EXPORT_NAME='CoinMinerModule' \
|
|
||||||
-s INITIAL_MEMORY=67108864
|
|
||||||
|
|
||||||
benchmark: aad_benchmark.c aad_sha1.h aad_sha1_cpu.h aad_data_types.h aad_utilities.h makefile
|
|
||||||
cc -march=native -Wall -Wshadow -Werror -O3 $< -o $@
|
|
||||||
|
|
||||||
miners: coin_miner_cpu coin_miner_simd coin_miner_wasm coin_miner_cuda benchmark
|
|
||||||
|
|
||||||
all: sha1_tests sha1_cuda_test sha1_cuda_kernel.cubin \
|
|
||||||
coin_miner_cpu coin_miner_simd coin_miner_wasm coin_miner_cuda coin_miner_cuda_kernel.cubin \
|
|
||||||
benchmark
|
|
||||||
|
|
|
||||||
Binary file not shown.
Binary file not shown.
Loading…
Reference in New Issue