//
// Arquiteturas de Alto Desempenho 2025/2026
//
// DETI Coin Miner - CUDA kernel (optimized for mining)
//

#include "aad_sha1.h"
#include "aad_data_types.h"

//
// Optimized CUDA kernel for DETI coin mining
// Each thread generates coins using the same approach as CPU/SIMD miners
//

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)
{
  u32_t coin[14];
  u32_t hash[5];
  u32_t n;
  u08_t *bytes = (u08_t *)coin;

  // Get thread index (used as offset from base counter)
  n = (u32_t)threadIdx.x + (u32_t)blockDim.x * (u32_t)blockIdx.x;

  // Initialize coin template: "DETI coin 2 " + variable + "\n\x80"
  // Use byte-swapped format to match host expectations (idx ^ 3)
  coin[0] = ('D' << 24) + ('E' << 16) + ('T' << 8) + 'I';
  coin[1] = (' ' << 24) + ('c' << 16) + ('o' << 8) + 'i';
  coin[2] = ('n' << 24) + (' ' << 16) + ('2' << 8) + ' ';

  // Fill the variable part of the coin with a pattern
  for(int i = 3; i < 14; i++)
    coin[i] = 0x41414141; // 'AAAA'

  // End with newline and padding
  bytes[0x36 ^ 3] = '\n';   // Position 54
  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
    // This creates a unique starting point for each thread
    u64_t offset = base_nonce + n + (u64_t)i * gridDim.x * blockDim.x;

    // 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
# 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 this is a valid DETI coin
  if(hash[0] == 0xAAD20250u)
  {
    // Found a coin! Store it atomically
    u32_t idx = atomicAdd(coins_storage_area, 14u);

      // Make sure we don't write outside buffer
      if(idx < 1024u - 14u)
      {
        // Store the complete coin data
        for(int k = 0; k < 14; k++)
          coins_storage_area[idx + k] = coin[k];
      }
    }
  }
}