Compare commits
No commits in common. "ecf0ac85cd26e32626c023bcf62e73fba127d737" and "5febf93e6523ca32b2afd335f4a41a08854b32f3" have entirely different histories.
ecf0ac85cd
...
5febf93e65
|
|
@ -65,9 +65,3 @@ CMakeUserPresets.json
|
||||||
coin_miner_cpu
|
coin_miner_cpu
|
||||||
coin_miner_simd
|
coin_miner_simd
|
||||||
coin_miner_cuda
|
coin_miner_cuda
|
||||||
coin_miner_ocl
|
|
||||||
coin_miner_wasm.js
|
|
||||||
coin_miner_wasm.wasm
|
|
||||||
|
|
||||||
# Vault
|
|
||||||
deti_coins*_vault.txt
|
|
||||||
|
|
|
||||||
|
|
@ -30,14 +30,6 @@ static int is_valid_coin(u32_t *hash)
|
||||||
return hash[0] == 0xAAD20250u;
|
return hash[0] == 0xAAD20250u;
|
||||||
}
|
}
|
||||||
|
|
||||||
// Get current wall time in seconds
|
|
||||||
static double get_wall_time(void)
|
|
||||||
{
|
|
||||||
struct timespec ts;
|
|
||||||
clock_gettime(CLOCK_MONOTONIC, &ts);
|
|
||||||
return (double)ts.tv_sec + (double)ts.tv_nsec * 1.0e-9;
|
|
||||||
}
|
|
||||||
|
|
||||||
//
|
//
|
||||||
// mine DETI coins using the CPU (no SIMD)
|
// mine DETI coins using the CPU (no SIMD)
|
||||||
//
|
//
|
||||||
|
|
@ -77,7 +69,9 @@ static void mine_coins_cpu(u64_t max_attempts)
|
||||||
printf("Mining DETI coins using CPU (no SIMD)...\n");
|
printf("Mining DETI coins using CPU (no SIMD)...\n");
|
||||||
printf("Press Ctrl+C to stop\n\n");
|
printf("Press Ctrl+C to stop\n\n");
|
||||||
|
|
||||||
double start_time = get_wall_time();
|
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))
|
||||||
{
|
{
|
||||||
|
|
@ -123,14 +117,16 @@ static void mine_coins_cpu(u64_t max_attempts)
|
||||||
// Print progress every 1M attempts
|
// Print progress every 1M attempts
|
||||||
if(attempts % 1000000 == 0)
|
if(attempts % 1000000 == 0)
|
||||||
{
|
{
|
||||||
double elapsed = get_wall_time() - start_time;
|
time_measurement();
|
||||||
double rate = attempts / elapsed;
|
double current_time = wall_time_delta() - start_time;
|
||||||
printf("Attempts: %llu, Rate: %.2f MH/s, Coins: %u, Elapsed: %.2fs\n",
|
double rate = attempts / current_time;
|
||||||
(unsigned long long)attempts, rate / 1e6, coins_found, elapsed);
|
printf("Attempts: %llu, Rate: %.2f MH/s, Coins: %u\n",
|
||||||
|
(unsigned long long)attempts, rate / 1e6, coins_found);
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
double total_time = get_wall_time() - start_time;
|
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);
|
||||||
|
|
@ -164,3 +160,4 @@ int main(int argc, char *argv[])
|
||||||
|
|
||||||
return 0;
|
return 0;
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
|
||||||
|
|
@ -38,7 +38,7 @@ static void reconstruct_coin(u32_t *stored_data, u32_t coin[14])
|
||||||
//
|
//
|
||||||
// Mine DETI coins using CUDA
|
// Mine DETI coins using CUDA
|
||||||
//
|
//
|
||||||
static void mine_coins_cuda(u64_t max_attempts)
|
static void mine_coins_cuda(u64_t max_attempts, int use_scan_kernel)
|
||||||
{
|
{
|
||||||
cuda_data_t cd;
|
cuda_data_t cd;
|
||||||
u32_t *host_storage;
|
u32_t *host_storage;
|
||||||
|
|
@ -49,7 +49,7 @@ static void mine_coins_cuda(u64_t max_attempts)
|
||||||
// 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";
|
||||||
cd.kernel_name = "mine_deti_coins_kernel";
|
cd.kernel_name = use_scan_kernel ? "mine_deti_coins_scan_kernel" : "mine_deti_coins_kernel";
|
||||||
cd.data_size[0] = COINS_STORAGE_SIZE * sizeof(u32_t);
|
cd.data_size[0] = COINS_STORAGE_SIZE * sizeof(u32_t);
|
||||||
cd.data_size[1] = 0;
|
cd.data_size[1] = 0;
|
||||||
|
|
||||||
|
|
@ -69,10 +69,12 @@ static void mine_coins_cuda(u64_t max_attempts)
|
||||||
printf("Kernel: %s\n", cd.kernel_name);
|
printf("Kernel: %s\n", cd.kernel_name);
|
||||||
printf("Press Ctrl+C to stop\n\n");
|
printf("Press Ctrl+C to stop\n\n");
|
||||||
|
|
||||||
u64_t base_nonce = 0;
|
u32_t param1 = (u32_t)time(NULL);
|
||||||
u32_t attempts_per_thread = 1024 * 8; // Increased attempts per thread
|
u32_t param2 = 0x12345678u;
|
||||||
|
int scan_pos = 12;
|
||||||
|
|
||||||
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))
|
||||||
{
|
{
|
||||||
|
|
@ -83,10 +85,12 @@ static void mine_coins_cuda(u64_t max_attempts)
|
||||||
host_to_device_copy(&cd, 0);
|
host_to_device_copy(&cd, 0);
|
||||||
|
|
||||||
// Set kernel arguments
|
// Set kernel arguments
|
||||||
cd.n_kernel_arguments = 2;
|
cd.n_kernel_arguments = use_scan_kernel ? 4 : 3;
|
||||||
cd.arg[0] = &cd.device_data[0];
|
cd.arg[0] = &cd.device_data[0];
|
||||||
cd.arg[1] = &base_nonce;
|
cd.arg[1] = ¶m1;
|
||||||
cd.arg[2] = &attempts_per_thread;
|
cd.arg[2] = ¶m2;
|
||||||
|
if(use_scan_kernel)
|
||||||
|
cd.arg[3] = &scan_pos;
|
||||||
|
|
||||||
// Launch the CUDA kernel
|
// Launch the CUDA kernel
|
||||||
launch_kernel(&cd);
|
launch_kernel(&cd);
|
||||||
|
|
@ -100,24 +104,58 @@ static void mine_coins_cuda(u64_t max_attempts)
|
||||||
|
|
||||||
if(n_stored > 0 && host_storage[0] < COINS_STORAGE_SIZE)
|
if(n_stored > 0 && host_storage[0] < COINS_STORAGE_SIZE)
|
||||||
{
|
{
|
||||||
|
printf("DEBUG: host_storage[0] = %u, n_stored = %u\n", host_storage[0], n_stored);
|
||||||
|
|
||||||
for(u32_t i = 0; i < n_stored; i++)
|
for(u32_t i = 0; i < n_stored; i++)
|
||||||
{
|
{
|
||||||
u32_t coin[14];
|
u32_t coin[14];
|
||||||
reconstruct_coin(&host_storage[1 + i * 14], coin);
|
reconstruct_coin(&host_storage[1 + i * 14], coin);
|
||||||
|
|
||||||
|
// Verify it's actually a valid coin
|
||||||
|
u32_t hash[5];
|
||||||
|
sha1(coin, hash);
|
||||||
|
|
||||||
|
printf("DEBUG: Coin %u - hash[0] = 0x%08X (expected 0xAAD20250)\n", i, hash[0]);
|
||||||
|
|
||||||
|
// Print the coin as string
|
||||||
|
if(i == 0) {
|
||||||
|
printf("DEBUG: First coin content: ");
|
||||||
|
u08_t *bytes = (u08_t *)coin;
|
||||||
|
for(int j = 0; j < 55; j++) {
|
||||||
|
char c = bytes[j ^ 3];
|
||||||
|
if(c >= 32 && c <= 126)
|
||||||
|
printf("%c", c);
|
||||||
|
else
|
||||||
|
printf("[0x%02X]", (u08_t)c);
|
||||||
|
}
|
||||||
|
printf("\n");
|
||||||
|
}
|
||||||
|
|
||||||
|
if(hash[0] == 0xAAD20250u)
|
||||||
|
{
|
||||||
coins_found++;
|
coins_found++;
|
||||||
n_coins_this_kernel++;
|
n_coins_this_kernel++;
|
||||||
printf("COIN FOUND! (kernel %u, coin %u in this kernel). Total coins:%u\n",
|
printf("COIN FOUND! (kernel %u, coin %u in this kernel)\n",
|
||||||
kernel_runs, n_coins_this_kernel, coins_found);
|
kernel_runs, n_coins_this_kernel);
|
||||||
save_coin(coin);
|
save_coin(coin);
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
}
|
||||||
|
|
||||||
// Update counters
|
// Update counters
|
||||||
kernel_runs++;
|
kernel_runs++;
|
||||||
u64_t attempts_this_launch = (u64_t)n_threads * attempts_per_thread;
|
if(use_scan_kernel)
|
||||||
attempts += attempts_this_launch;
|
attempts += n_threads * 256; // Each thread tries 256 values
|
||||||
base_nonce += attempts_this_launch;
|
else
|
||||||
|
attempts += n_threads;
|
||||||
|
|
||||||
|
// Update parameters for next iteration
|
||||||
|
param1++;
|
||||||
|
param2 = param2 ^ 0x9E3779B9u;
|
||||||
|
if(use_scan_kernel)
|
||||||
|
scan_pos = (scan_pos + 1) % 42 + 12; // Cycle through positions 12-53
|
||||||
|
|
||||||
|
|
||||||
}
|
}
|
||||||
|
|
||||||
time_measurement();
|
time_measurement();
|
||||||
|
|
@ -140,13 +178,20 @@ static void mine_coins_cuda(u64_t max_attempts)
|
||||||
int main(int argc, char *argv[])
|
int main(int argc, char *argv[])
|
||||||
{
|
{
|
||||||
u64_t max_attempts = 0;
|
u64_t max_attempts = 0;
|
||||||
|
int use_scan_kernel = 0;
|
||||||
|
|
||||||
signal(SIGINT, signal_handler);
|
signal(SIGINT, signal_handler);
|
||||||
|
|
||||||
if(argc > 1)
|
if(argc > 1)
|
||||||
max_attempts = strtoull(argv[1], NULL, 10);
|
max_attempts = strtoull(argv[1], NULL, 10);
|
||||||
|
|
||||||
mine_coins_cuda(max_attempts);
|
if(argc > 2 && strcmp(argv[2], "scan") == 0)
|
||||||
|
{
|
||||||
|
use_scan_kernel = 1;
|
||||||
|
printf("Using scan kernel (tries 256 values per thread)\n");
|
||||||
|
}
|
||||||
|
|
||||||
|
mine_coins_cuda(max_attempts, use_scan_kernel);
|
||||||
|
|
||||||
return 0;
|
return 0;
|
||||||
}
|
}
|
||||||
|
|
|
||||||
|
|
@ -13,7 +13,7 @@
|
||||||
//
|
//
|
||||||
|
|
||||||
extern "C" __global__ __launch_bounds__(RECOMMENDED_CUDA_BLOCK_SIZE,1)
|
extern "C" __global__ __launch_bounds__(RECOMMENDED_CUDA_BLOCK_SIZE,1)
|
||||||
void mine_deti_coins_kernel(u32_t *coins_storage_area, u64_t base_nonce, u32_t attempts_per_thread)
|
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];
|
||||||
|
|
@ -25,27 +25,23 @@ void mine_deti_coins_kernel(u32_t *coins_storage_area, u64_t base_nonce, u32_t a
|
||||||
|
|
||||||
// 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)
|
||||||
coin[0] = ('D' << 24) + ('E' << 16) + ('T' << 8) + 'I';
|
coin[0] = 0x44455449u; // "DETI" with byte swap
|
||||||
coin[1] = (' ' << 24) + ('c' << 16) + ('o' << 8) + 'i';
|
coin[1] = 0x20636F69u; // " coi" with byte swap
|
||||||
coin[2] = ('n' << 24) + (' ' << 16) + ('2' << 8) + ' ';
|
coin[2] = 0x6E203220u; // "n 2 " with byte swap
|
||||||
|
|
||||||
// Fill the variable part of the coin with a pattern
|
// Initialize variable part (positions 12-53, 42 bytes)
|
||||||
for(int i = 3; i < 14; i++)
|
// Start with A-Z pattern like CPU/SIMD miners
|
||||||
coin[i] = 0x41414141; // 'AAAA'
|
for(int i = 12; i < 54; i++)
|
||||||
|
bytes[i ^ 3] = 'A' + ((i - 12) % 26);
|
||||||
|
|
||||||
// End with newline and padding
|
// End with newline and padding
|
||||||
bytes[0x36 ^ 3] = '\n'; // Position 54
|
bytes[0x36 ^ 3] = '\n'; // Position 54
|
||||||
bytes[0x37 ^ 3] = 0x80; // Position 55
|
bytes[0x37 ^ 3] = 0x80; // Position 55
|
||||||
|
|
||||||
for(u32_t i = 0; i < attempts_per_thread; ++i) {
|
|
||||||
// Initialize variable part (positions 12-53, 42 bytes)
|
|
||||||
// Start with A-Z pattern like CPU/SIMD miners
|
|
||||||
for(int j = 12; j < 54; j++)
|
|
||||||
bytes[j ^ 3] = 'A' + ((j - 12) % 26);
|
|
||||||
|
|
||||||
// Calculate offset based on thread index and parameters
|
// Calculate offset based on thread index and parameters
|
||||||
// This creates a unique starting point for each thread
|
// This creates a unique starting point for each thread
|
||||||
u64_t offset = base_nonce + n + (u64_t)i * gridDim.x * blockDim.x;
|
u64_t offset = ((u64_t)param1 << 32) | param2;
|
||||||
|
offset += (u64_t)n;
|
||||||
|
|
||||||
// Apply offset to variable part (increment the coin counter)
|
// Apply offset to variable part (increment the coin counter)
|
||||||
for(int pos = 53; pos >= 12 && offset > 0; pos--)
|
for(int pos = 53; pos >= 12 && offset > 0; pos--)
|
||||||
|
|
@ -88,9 +84,100 @@ void mine_deti_coins_kernel(u32_t *coins_storage_area, u64_t base_nonce, u32_t a
|
||||||
if(idx < 1024u - 14u)
|
if(idx < 1024u - 14u)
|
||||||
{
|
{
|
||||||
// Store the complete coin data
|
// Store the complete coin data
|
||||||
for(int k = 0; k < 14; k++)
|
for(int i = 0; i < 14; i++)
|
||||||
coins_storage_area[idx + k] = coin[k];
|
coins_storage_area[idx + i] = coin[i];
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
//
|
||||||
|
// Kernel that tries all possible values for one character position
|
||||||
|
//
|
||||||
|
|
||||||
|
extern "C" __global__ __launch_bounds__(RECOMMENDED_CUDA_BLOCK_SIZE,1)
|
||||||
|
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 hash[5];
|
||||||
|
u32_t n;
|
||||||
|
u08_t *bytes = (u08_t *)coin;
|
||||||
|
|
||||||
|
n = (u32_t)threadIdx.x + (u32_t)blockDim.x * (u32_t)blockIdx.x;
|
||||||
|
|
||||||
|
// Initialize coin template (with byte swap for idx ^ 3 convention)
|
||||||
|
coin[0] = 0x44455449u; // "DETI" with byte swap
|
||||||
|
coin[1] = 0x20636F69u; // " coi" with byte swap
|
||||||
|
coin[2] = 0x6E203220u; // "n 2 " with byte swap
|
||||||
|
|
||||||
|
// Initialize variable part with A-Z pattern
|
||||||
|
for(int i = 12; i < 54; i++)
|
||||||
|
bytes[i ^ 3] = 'A' + ((i - 12) % 26);
|
||||||
|
|
||||||
|
// End with newline and padding
|
||||||
|
bytes[0x36 ^ 3] = '\n'; // Position 54
|
||||||
|
bytes[0x37 ^ 3] = 0x80; // Position 55
|
||||||
|
|
||||||
|
// 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)
|
||||||
|
continue; // Skip the scan position
|
||||||
|
|
||||||
|
u08_t *byte = &bytes[pos ^ 3];
|
||||||
|
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;
|
||||||
|
|
||||||
|
// Skip newline in the middle (it's only valid at position 54)
|
||||||
|
if(scan_position != 54 && val == '\n')
|
||||||
|
continue;
|
||||||
|
}
|
||||||
|
|
||||||
|
// Compute SHA1 hash
|
||||||
|
# define T u32_t
|
||||||
|
# define C(c) (c)
|
||||||
|
# define ROTATE(x,n) (((x) << (n)) | ((x) >> (32 - (n))))
|
||||||
|
# define DATA(idx) coin[idx]
|
||||||
|
# define HASH(idx) hash[idx]
|
||||||
|
CUSTOM_SHA1_CODE();
|
||||||
|
# undef T
|
||||||
|
# undef C
|
||||||
|
# undef ROTATE
|
||||||
|
# undef DATA
|
||||||
|
# undef HASH
|
||||||
|
|
||||||
|
// Check if valid coin
|
||||||
|
if(hash[0] == 0xAAD20250u)
|
||||||
|
{
|
||||||
|
u32_t idx = atomicAdd(coins_storage_area, 14u);
|
||||||
|
if(idx < 1024u - 14u)
|
||||||
|
{
|
||||||
|
for(int i = 0; i < 14; i++)
|
||||||
|
coins_storage_area[idx + i] = coin[i];
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
|
||||||
|
|
@ -32,14 +32,6 @@ static void reconstruct_coin(u32_t *stored_data, u32_t coin[14])
|
||||||
coin[i] = stored_data[i];
|
coin[i] = stored_data[i];
|
||||||
}
|
}
|
||||||
|
|
||||||
// Get current wall time in seconds
|
|
||||||
static double get_wall_time(void)
|
|
||||||
{
|
|
||||||
struct timespec ts;
|
|
||||||
clock_gettime(CLOCK_MONOTONIC, &ts);
|
|
||||||
return (double)ts.tv_sec + (double)ts.tv_nsec * 1.0e-9;
|
|
||||||
}
|
|
||||||
|
|
||||||
//
|
//
|
||||||
// Mine DETI coins using OpenCL
|
// Mine DETI coins using OpenCL
|
||||||
//
|
//
|
||||||
|
|
@ -65,7 +57,7 @@ static void mine_coins_ocl(u64_t max_attempts, int use_scan_kernel)
|
||||||
|
|
||||||
// Kernel configuration
|
// Kernel configuration
|
||||||
od.local_work_size = RECOMMENDED_OCL_WORK_GROUP_SIZE;
|
od.local_work_size = RECOMMENDED_OCL_WORK_GROUP_SIZE;
|
||||||
od.global_work_size = 4096 * od.local_work_size;
|
od.global_work_size = 4096 * od.local_work_size; // Large grid for maximum GPU utilization
|
||||||
|
|
||||||
u32_t n_threads = od.global_work_size;
|
u32_t n_threads = od.global_work_size;
|
||||||
|
|
||||||
|
|
@ -76,16 +68,58 @@ static void mine_coins_ocl(u64_t max_attempts, int use_scan_kernel)
|
||||||
printf("Kernel: %s\n", od.kernel_name);
|
printf("Kernel: %s\n", od.kernel_name);
|
||||||
printf("Press Ctrl+C to stop\n\n");
|
printf("Press Ctrl+C to stop\n\n");
|
||||||
|
|
||||||
|
// Test SHA1 on host to verify it matches
|
||||||
|
printf("Testing SHA1 implementation on host...\n");
|
||||||
|
u32_t test_coin[14];
|
||||||
|
memset(test_coin, 0, sizeof(test_coin));
|
||||||
|
((u08_t *)test_coin)[0x0 ^ 3] = 'D';
|
||||||
|
((u08_t *)test_coin)[0x1 ^ 3] = 'E';
|
||||||
|
((u08_t *)test_coin)[0x2 ^ 3] = 'T';
|
||||||
|
((u08_t *)test_coin)[0x3 ^ 3] = 'I';
|
||||||
|
((u08_t *)test_coin)[0x4 ^ 3] = ' ';
|
||||||
|
((u08_t *)test_coin)[0x5 ^ 3] = 'c';
|
||||||
|
((u08_t *)test_coin)[0x6 ^ 3] = 'o';
|
||||||
|
((u08_t *)test_coin)[0x7 ^ 3] = 'i';
|
||||||
|
((u08_t *)test_coin)[0x8 ^ 3] = 'n';
|
||||||
|
((u08_t *)test_coin)[0x9 ^ 3] = ' ';
|
||||||
|
((u08_t *)test_coin)[0xa ^ 3] = '2';
|
||||||
|
((u08_t *)test_coin)[0xb ^ 3] = ' ';
|
||||||
|
((u08_t *)test_coin)[0x36 ^ 3] = '\n';
|
||||||
|
((u08_t *)test_coin)[0x37 ^ 3] = 0x80;
|
||||||
|
for(int i = 12; i < 54; i++)
|
||||||
|
((u08_t *)test_coin)[i ^ 3] = 'A' + (i - 12) % 26;
|
||||||
|
|
||||||
|
u32_t test_hash[5];
|
||||||
|
sha1(test_coin, test_hash);
|
||||||
|
printf("Host test hash: 0x%08X 0x%08X 0x%08X 0x%08X 0x%08X\n",
|
||||||
|
test_hash[0], test_hash[1], test_hash[2], test_hash[3], test_hash[4]);
|
||||||
|
|
||||||
|
// Now test on device
|
||||||
|
printf("Testing SHA1 implementation on device...\n");
|
||||||
|
host_storage[0] = 1u;
|
||||||
|
|
||||||
|
// Put the test coin in storage starting at index 1
|
||||||
|
for(int i = 0; i < 14; i++)
|
||||||
|
host_storage[1 + i] = test_coin[i];
|
||||||
|
|
||||||
|
// Copy to device
|
||||||
|
host_to_device_copy(&od, 0);
|
||||||
|
|
||||||
|
// We'll add a test kernel - for now just verify basic kernel launch works
|
||||||
|
printf("Starting mining...\n\n");
|
||||||
|
|
||||||
u32_t param1 = (u32_t)time(NULL);
|
u32_t param1 = (u32_t)time(NULL);
|
||||||
u32_t param2 = 0x12345678u;
|
u32_t param2 = 0x12345678u;
|
||||||
int scan_pos = 12;
|
int scan_pos = 12;
|
||||||
|
|
||||||
double start_time = get_wall_time();
|
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))
|
||||||
{
|
{
|
||||||
// Initialize storage area
|
// Initialize storage area
|
||||||
host_storage[0] = 1u;
|
host_storage[0] = 1u; // First unused index
|
||||||
|
|
||||||
// Copy to device
|
// Copy to device
|
||||||
host_to_device_copy(&od, 0);
|
host_to_device_copy(&od, 0);
|
||||||
|
|
@ -105,6 +139,7 @@ static void mine_coins_ocl(u64_t max_attempts, int use_scan_kernel)
|
||||||
device_to_host_copy(&od, 0);
|
device_to_host_copy(&od, 0);
|
||||||
|
|
||||||
// Process found coins
|
// Process found coins
|
||||||
|
u32_t n_coins_this_kernel = 0;
|
||||||
u32_t n_stored = (host_storage[0] - 1) / 14;
|
u32_t n_stored = (host_storage[0] - 1) / 14;
|
||||||
|
|
||||||
if(n_stored > 0 && host_storage[0] < COINS_STORAGE_SIZE)
|
if(n_stored > 0 && host_storage[0] < COINS_STORAGE_SIZE)
|
||||||
|
|
@ -121,7 +156,9 @@ static void mine_coins_ocl(u64_t max_attempts, int use_scan_kernel)
|
||||||
if(hash[0] == 0xAAD20250u)
|
if(hash[0] == 0xAAD20250u)
|
||||||
{
|
{
|
||||||
coins_found++;
|
coins_found++;
|
||||||
printf("COIN FOUND! (kernel %u)\n", kernel_runs);
|
n_coins_this_kernel++;
|
||||||
|
printf("COIN FOUND! (kernel %u, coin %u in this kernel)\n",
|
||||||
|
kernel_runs, n_coins_this_kernel);
|
||||||
save_coin(coin);
|
save_coin(coin);
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
@ -130,7 +167,7 @@ static void mine_coins_ocl(u64_t max_attempts, int use_scan_kernel)
|
||||||
// Update counters
|
// Update counters
|
||||||
kernel_runs++;
|
kernel_runs++;
|
||||||
if(use_scan_kernel)
|
if(use_scan_kernel)
|
||||||
attempts += n_threads * 95; // Each thread tries 95 printable ASCII values
|
attempts += n_threads * 256; // Each thread tries 256 values
|
||||||
else
|
else
|
||||||
attempts += n_threads;
|
attempts += n_threads;
|
||||||
|
|
||||||
|
|
@ -138,19 +175,21 @@ static void mine_coins_ocl(u64_t max_attempts, int use_scan_kernel)
|
||||||
param1++;
|
param1++;
|
||||||
param2 = param2 ^ 0x9E3779B9u;
|
param2 = param2 ^ 0x9E3779B9u;
|
||||||
if(use_scan_kernel)
|
if(use_scan_kernel)
|
||||||
scan_pos = (scan_pos + 1) % 42 + 12;
|
scan_pos = (scan_pos + 1) % 42 + 12; // Cycle through positions 12-53
|
||||||
|
|
||||||
// Print progress every 10 kernel launches
|
// Print progress every 10 kernel launches
|
||||||
if(kernel_runs % 10 == 0)
|
if(kernel_runs % 10 == 0)
|
||||||
{
|
{
|
||||||
double elapsed = get_wall_time() - start_time;
|
time_measurement();
|
||||||
double rate = attempts / elapsed;
|
double current_time = wall_time_delta() - start_time;
|
||||||
printf("Attempts: %llu, Rate: %.2f MH/s, Coins: %u, Kernels: %u, Elapsed: %.2fs\n",
|
double rate = attempts / current_time;
|
||||||
(unsigned long long)attempts, rate / 1e6, coins_found, kernel_runs, elapsed);
|
printf("Attempts: %llu, Rate: %.2f MH/s, Coins: %u, Kernels: %u\n",
|
||||||
|
(unsigned long long)attempts, rate / 1e6, coins_found, kernel_runs);
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
double total_time = get_wall_time() - start_time;
|
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);
|
||||||
|
|
@ -178,7 +217,7 @@ int main(int argc, char *argv[])
|
||||||
if(argc > 2 && strcmp(argv[2], "scan") == 0)
|
if(argc > 2 && strcmp(argv[2], "scan") == 0)
|
||||||
{
|
{
|
||||||
use_scan_kernel = 1;
|
use_scan_kernel = 1;
|
||||||
printf("Using scan kernel (tries 95 values per thread)\n");
|
printf("Using scan kernel (tries 256 values per thread)\n");
|
||||||
}
|
}
|
||||||
|
|
||||||
mine_coins_ocl(max_attempts, use_scan_kernel);
|
mine_coins_ocl(max_attempts, use_scan_kernel);
|
||||||
|
|
|
||||||
|
|
@ -30,14 +30,6 @@ static int is_valid_coin(u32_t *hash)
|
||||||
return hash[0] == 0xAAD20250u;
|
return hash[0] == 0xAAD20250u;
|
||||||
}
|
}
|
||||||
|
|
||||||
// Get current wall time in seconds
|
|
||||||
static double get_wall_time(void)
|
|
||||||
{
|
|
||||||
struct timespec ts;
|
|
||||||
clock_gettime(CLOCK_MONOTONIC, &ts);
|
|
||||||
return (double)ts.tv_sec + (double)ts.tv_nsec * 1.0e-9;
|
|
||||||
}
|
|
||||||
|
|
||||||
//
|
//
|
||||||
// increment coin variable part using the same logic as CPU miner
|
// increment coin variable part using the same logic as CPU miner
|
||||||
// returns 0 if overflow (all positions wrapped around), 1 otherwise
|
// returns 0 if overflow (all positions wrapped around), 1 otherwise
|
||||||
|
|
@ -162,7 +154,9 @@ static void mine_coins_avx(u64_t max_attempts)
|
||||||
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");
|
||||||
|
|
||||||
double start_time = get_wall_time();
|
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))
|
||||||
{
|
{
|
||||||
|
|
@ -193,14 +187,16 @@ 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)
|
||||||
{
|
{
|
||||||
double elapsed = get_wall_time() - start_time;
|
time_measurement();
|
||||||
double rate = attempts / elapsed;
|
double current_time = wall_time_delta() - start_time;
|
||||||
printf("Attempts: %llu, Rate: %.2f MH/s, Coins: %u, Elapsed: %.2fs\n",
|
double rate = attempts / current_time;
|
||||||
(unsigned long long)attempts, rate / 1e6, coins_found, elapsed);
|
printf("Attempts: %llu, Rate: %.2f MH/s, Coins: %u\n",
|
||||||
|
(unsigned long long)attempts, rate / 1e6, coins_found);
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
double total_time = get_wall_time() - start_time;
|
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);
|
||||||
|
|
@ -250,7 +246,9 @@ static void mine_coins_avx2(u64_t max_attempts)
|
||||||
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");
|
||||||
|
|
||||||
double start_time = get_wall_time();
|
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))
|
||||||
{
|
{
|
||||||
|
|
@ -276,14 +274,16 @@ static void mine_coins_avx2(u64_t max_attempts)
|
||||||
|
|
||||||
if(attempts % 1000000 < SIMD_WIDTH)
|
if(attempts % 1000000 < SIMD_WIDTH)
|
||||||
{
|
{
|
||||||
double elapsed = get_wall_time() - start_time;
|
time_measurement();
|
||||||
double rate = attempts / elapsed;
|
double current_time = wall_time_delta() - start_time;
|
||||||
printf("Attempts: %llu, Rate: %.2f MH/s, Coins: %u, Elapsed: %.2fs\n",
|
double rate = attempts / current_time;
|
||||||
(unsigned long long)attempts, rate / 1e6, coins_found, elapsed);
|
printf("Attempts: %llu, Rate: %.2f MH/s, Coins: %u\n",
|
||||||
|
(unsigned long long)attempts, rate / 1e6, coins_found);
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
double total_time = get_wall_time() - start_time;
|
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);
|
||||||
|
|
@ -294,7 +294,6 @@ static void mine_coins_avx2(u64_t max_attempts)
|
||||||
save_coin(NULL);
|
save_coin(NULL);
|
||||||
}
|
}
|
||||||
#endif
|
#endif
|
||||||
|
|
||||||
#if defined(__AVX2__)
|
#if defined(__AVX2__)
|
||||||
#include <omp.h>
|
#include <omp.h>
|
||||||
//
|
//
|
||||||
|
|
@ -311,13 +310,16 @@ static void mine_coins_avx2_omp(u64_t max_attempts)
|
||||||
printf("Mining DETI coins using AVX2 (8-way SIMD) + OpenMP (%d threads)...\n", num_threads);
|
printf("Mining DETI coins using AVX2 (8-way SIMD) + OpenMP (%d threads)...\n", num_threads);
|
||||||
printf("Press Ctrl+C to stop\n\n");
|
printf("Press Ctrl+C to stop\n\n");
|
||||||
|
|
||||||
double start_time = get_wall_time();
|
time_measurement();
|
||||||
|
time_measurement();
|
||||||
|
double start_time = wall_time_delta();
|
||||||
|
|
||||||
#pragma omp parallel
|
#pragma omp parallel
|
||||||
{
|
{
|
||||||
u32_t base_coin[14];
|
u32_t base_coin[14];
|
||||||
u32_t interleaved_data[14 * SIMD_WIDTH] __attribute__((aligned(32)));
|
u32_t interleaved_data[14 * SIMD_WIDTH] __attribute__((aligned(32)));
|
||||||
u32_t interleaved_hash[5 * 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;
|
u64_t thread_attempts = 0;
|
||||||
|
|
||||||
// Initialize base coin template (unique per thread)
|
// Initialize base coin template (unique per thread)
|
||||||
|
|
@ -373,10 +375,11 @@ static void mine_coins_avx2_omp(u64_t max_attempts)
|
||||||
{
|
{
|
||||||
#pragma omp critical
|
#pragma omp critical
|
||||||
{
|
{
|
||||||
double elapsed = get_wall_time() - start_time;
|
time_measurement();
|
||||||
double rate = (attempts + thread_attempts) / elapsed;
|
double current_time = wall_time_delta() - start_time;
|
||||||
printf("Attempts: %llu, Rate: %.2f MH/s, Coins: %u, Elapsed: %.2fs\n",
|
double rate = (attempts + thread_attempts) / current_time;
|
||||||
(unsigned long long)(attempts + thread_attempts), rate / 1e6, coins_found, elapsed);
|
printf("Attempts: %llu, Rate: %.2f MH/s, Coins: %u\n",
|
||||||
|
(unsigned long long)(attempts + thread_attempts), rate / 1e6, coins_found);
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
@ -385,7 +388,8 @@ static void mine_coins_avx2_omp(u64_t max_attempts)
|
||||||
attempts += thread_attempts;
|
attempts += thread_attempts;
|
||||||
}
|
}
|
||||||
|
|
||||||
double total_time = get_wall_time() - start_time;
|
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);
|
||||||
|
|
|
||||||
Loading…
Reference in New Issue