// OpenCLカーネルのコンパイル時にSHA-256実装のヘッダーは利用できないため、
// 必要な定数と最小限のSHA-256関数をカーネル内で定義します。
// 以下のSHA-256実装は、OpenCL Cで動作するように調整されたものです。

// SHA-256の定数 (K)
// OpenCLではプリプロセッサマクロで配列を定義できないため、constグローバル変数として定義
constant uint K[64] = {
    0x428a2f98, 0x71374491, 0xb5c0fbcf, 0xe9b5dba5, 0x3956c25b, 0x59f111f1, 0x923f82a4, 0xab1c5ed5,
    0xd807aa98, 0x12835b01, 0x243185be, 0x550c7dc3, 0x72be5d74, 0x80deb1fe, 0x9bdc06a7, 0xc19bf174,
    0xe49b69c1, 0xefbe4786, 0x0fc19dc6, 0x240ca1cc, 0x2de92c6f, 0x4a7484aa, 0x5cb0a9dc, 0x76f988da,
    0x983e5152, 0xa831c66d, 0xb00327c8, 0xbf597fc7, 0xc6e00bf3, 0xd5a79147, 0x06ca6351, 0x14292967,
    0x27b70a85, 0x2e1b2138, 0x4d2c6dfc, 0x53380d13, 0x650a7354, 0x766a0abb, 0x81c2c92e, 0x92722c85,
    0xa2bfe8a1, 0xa81a664b, 0xc24b8b70, 0xc76c51a3, 0xd192e819, 0xd6990624, 0xf40e3585, 0x106aa070,
    0x19a4c116, 0x1e376c08, 0x2748774c, 0x34b0bcb5, 0x391c0cb3, 0x4ed8aa4a, 0x5b9cca4f, 0x682e6ff3,
    0x748f82ee, 0x78a5636f, 0x84c87814, 0x8cc70208, 0x90befffa, 0xa4506ceb, 0xbef9a3f7, 0xc67178f2
};

// 初期ハッシュ値 (H)
constant uint H_INIT[8] = {
    0x6a09e667, 0xbb67ae85, 0x3c6ef372, 0xa54ff53a, 
    0x510e527f, 0x9b05688c, 0x1f83d9ab, 0x5be0cd19
};

// ビット操作マクロ
#define ROR(x, n) (((x) >> (n)) | ((x) << (32 - (n))))
#define SHR(x, n) ((x) >> (n))
#define CH(x, y, z) (((x) & (y)) ^ ((~(x)) & (z)) )
#define MAJ(x, y, z) (((x) & (y)) ^ ((x) & (z)) ^ ((y) & (z)))
#define SIG0(x) (ROR((x), 2) ^ ROR((x), 13) ^ ROR((x), 22))
#define SIG1(x) (ROR((x), 6) ^ ROR((x), 11) ^ ROR((x), 25))
#define sig0(x) (ROR((x), 7) ^ ROR((x), 18) ^ SHR((x), 3))
#define sig1(x) (ROR((x), 17) ^ ROR((x), 19) ^ SHR((x), 10))

/**
 * @brief ターゲット比較のためのユーティリティ関数
 * @param a 最初のハッシュ (プライベートメモリ)
 * @param b ターゲットハッシュ (グローバルメモリ)
 * @return a < b なら 1, その他は 0
 */
int big_int_compare_cl(const __private uint *a, const __global uint *b) {
    for (int i = 0; i < 8; i++) { // 256 bits = 8 * 32-bit words
        if (a[i] < b[i]) return 1; // a < b
        if (a[i] > b[i]) return 0; // a > b
    }
    return 0; // a == b
}

/**
 * @brief SHA-256 コア関数
 */
void sha256_core(const __private uchar *input, uint input_len, __private uint *H) {
    uint a, b, c, d, e, f, g, h;
    uint T1, T2;
    uint W[64];
    
    // 1. Initial Hash Values (H_INIT)
    a = H_INIT[0]; b = H_INIT[1]; c = H_INIT[2]; d = H_INIT[3];
    e = H_INIT[4]; f = H_INIT[5]; g = H_INIT[6]; h = H_INIT[7];

    // 2. Padding and Parsing (メッセージブロックの準備)
    
    // W[0..15]の準備 (ビッグエンディアンのバイトをuintに変換)
    // input は private メモリ
    for (int i = 0; i < 16; i++) {
        // inputは uchar 配列 (64バイト)
        W[i] = (input[i * 4] << 24) | (input[i * 4 + 1] << 16) | 
               (input[i * 4 + 2] << 8) | (input[i * 4 + 3]);
    }
    
    // 3. メッセージスケジュールの拡張 W[16..63]
    for (int i = 16; i < 64; i++) {
        W[i] = sig1(W[i - 2]) + W[i - 7] + sig0(W[i - 15]) + W[i - 16];
    }

    // 4. 圧縮ループ
    for (int i = 0; i < 64; i++) {
        T1 = h + SIG1(e) + CH(e, f, g) + K[i] + W[i];
        T2 = SIG0(a) + MAJ(a, b, c);
        
        h = g;
        g = f;
        f = e;
        e = d + T1;
        d = c;
        c = b;
        b = a;
        a = T1 + T2;
    }

    // 5. 最終ハッシュ値の計算
    // H は private メモリ
    H[0] = a + H_INIT[0]; H[1] = b + H_INIT[1]; H[2] = c + H_INIT[2]; H[3] = d + H_INIT[3];
    H[4] = e + H_INIT[4]; H[5] = f + H_INIT[5]; H[6] = g + H_INIT[6]; H[7] = h + H_INIT[7];
}


/**
 * @brief PoWハッシュカーネル
 * @param target_words ターゲットハッシュ (uint[8]) 
 * @param base_block SHA-256の最初のブロック (uchar[64]) のベース部分
 * @param nonce_start 探索開始Nonce
 * @param nonce_increment Nonceの増分 (グローバルワークサイズ)
 * @param base_len Nonce文字列の開始位置 (バイト単位)
 * @param max_nonce_len Nonce文字列の最大長
 * @param result_buffer 見つかったNonceを格納するバッファ (volatile uint[1])
 * @param final_hash_buffer 見つかったハッシュを格納するバッファ (uint[8])
 * * 修正点: 64ビット原子操作が利用できないため、Nonceの型を long (64bit) から uint (32bit) に変更しました。
 */
__kernel void pow_hasher(
    __global const uint *target_words,
    __global uchar *base_block,
    uint nonce_start, // 修正: long -> uint
    uint nonce_increment, // 修正: long -> uint
    uint base_len, // Nonceが始まるブロック内のバイト位置
    uint max_nonce_len, // Nonce文字列の最大長
    __global volatile uint *result_buffer, // 修正: long -> uint
    __global uint *final_hash_buffer
) {
    // グローバルID
    size_t gid = get_global_id(0);
    
    // このワークアイテムが検証するNonce
    uint nonce = nonce_start + (uint)gid; // 修正: long -> uint
    
    // result_buffer の初期値チェック (uintで-1は0xFFFFFFFF)
    // 32ビット整数の「見つかっていない」状態を示す値
    const uint NOT_FOUND = (uint)-1;

    // 結果が既に見つかっている場合は終了
    if (result_buffer[0] != NOT_FOUND) return; 
    
    // ローカルコピー (current_block は __private なので OK)
    uchar current_block[64]; 
    for(int i = 0; i < 64; i++) {
        current_block[i] = base_block[i];
    }
    
    uint H[8]; // ハッシュ結果 (private)
    
    const int BATCH_SIZE = 100000; 

    for (int j = 0; j < BATCH_SIZE; j++) {
        
        // --- 1. Nonceを10進数文字列に変換し、ブロックに埋め込む ---
        // Nonceがuint (32bit) になったため、最大約10桁の文字列
        uint temp_nonce = nonce; // 修正: long -> uint
        int k = 0;
        
        // Nonce文字列を逆順に生成
        // uint は最大10桁なので、バッファサイズは11 (10桁 + NULL) で十分
        uchar nonce_str_rev[11];
        if (temp_nonce == 0) {
            nonce_str_rev[k++] = '0';
        } else {
            // Nonceが uint になっても、文字列変換ロジックは変わらない
            while (temp_nonce > 0) {
                nonce_str_rev[k++] = (uchar)('0' + (temp_nonce % 10));
                temp_nonce /= 10;
            }
        }
        
        // Nonce文字列長
        int nonce_len = k;
        
        // Nonce文字列を正しい順序で current_block にコピー
        for (int i = 0; i < nonce_len; i++) {
            current_block[base_len + i] = nonce_str_rev[nonce_len - 1 - i];
        }
        
        // --- 2. SHA-256パディングの処理 ---
        uint total_bits = (base_len + nonce_len) * 8; // 全入力のビット長
        
        // a. Nonce文字列の直後に '0x80' を追加
        current_block[base_len + nonce_len] = 0x80;
        
        // b. '0x00' を追加して、ブロック末尾8バイトを除く部分を埋める
        // 埋め込み開始位置
        int pad_start = base_len + nonce_len + 1;
        // 埋め込み終了位置 (56)
        for (int i = pad_start; i < 56; i++) {
            current_block[i] = 0x00;
        }
        
        // c. 最終の8バイトにメッセージのビット長 (ビッグエンディアン) を追加
        // total_bits は uint (32bit) なので、上位4バイトは0
        current_block[63] = (uchar)(total_bits);
        current_block[62] = (uchar)(total_bits >> 8);
        current_block[61] = (uchar)(total_bits >> 16);
        current_block[60] = (uchar)(total_bits >> 24);
        current_block[59] = 0x00; // 上位32ビット
        current_block[58] = 0x00;
        current_block[57] = 0x00;
        current_block[56] = 0x00; // 修正: total_bitsがuintのため、上位バイトを0に

        // SHA-256コアの実行
        sha256_core(current_block, 64, H); 

        // --- 3. 検証 ---
        if (big_int_compare_cl(H, target_words) == 1) {
            
            // 解決策が見つかった！
            // atomic_cmpxchg は uint 型に対応
            if (atomic_cmpxchg(result_buffer, NOT_FOUND, nonce) == NOT_FOUND) {
                // 最初に発見したワークアイテムのみが更新
                for (int i = 0; i < 8; i++) {
                    final_hash_buffer[i] = H[i];
                }
            }
            // ループを終了
            return; 
        }

        // 次のNonceへ (グローバルインクリメント)
        nonce += nonce_increment; 
        
        // Nonceのオーバーフローチェック（uintの範囲を超えないことを保証）
        // uintで加算した結果が元の値より小さくなった場合、オーバーフロー
        if (nonce < nonce_increment) return; 

        // 結果が既に見つかっている場合は終了
        if (result_buffer[0] != NOT_FOUND) return;
    }
}
