超速の ハッシュ計算のしくみ。ループの同時並列処理。GPU プログラミングは楽しいな。

タイトル: ケンタとGPUプログラミングの冒険

どうやってGPUで効率的に処理できるだろう？」

まず、ケンタはPythonのループ構造をGPUで並列処理する方法を考えました。彼は「縦にスレッドインデックスを通す」方法を試すことにしました。この方法では、5次元のベクトル4つの加算の計算に5本のスレッドを使い、それぞれが独立に計算を行います。すると、ケンタは「これで計算が速くなるかもしれない！」とワクワクしました。

しかし、ケンタはさらに賢くなろうと考えました。次に「横にスレッドインデックスを通す」方法を試してみることにしました。これによって、スレッドの起動する数は5×2となり、合計で10個のスレッドが立ち上がることになります。「おお、これならもっと速くなるかも！」とケンタは嬉しそうに思いました。

ケンタはGPUプログラミングのコードを書き始め、縦にスレッドインデックスを通す方法と横にスレッドインデックスを通す方法の両方を実装しました。

ハッシュ計算をおおまかに模式化すると　３つの計算　から成り立っています。各計算にN個のループが含まれている構造です。

ハッシュ計算ループ:
80ラウンドのループを通じて、ハッシュ値の計算が行われます。各ラウンドでは、以下の操作が行われます。
スケジュールされたワードの操作: 各ラウンドで、スケジュールされたワード（W）と定数（K）が使用されます。
中間変数の更新: 各ラウンドで、8つの中間変数（a〜h）が更新されます。これには、ビット演算、加算、ローテーションなどが含まれます。
関数の計算: Σ0, Σ1, Ch, Majなどの関数が使用されます。これらの関数は、ビットシフトやローテーションを含む複雑な計算を行います。
ハッシュ値の更新:
各ブロックの処理後に、計算された中間変数（a〜h）が初期ハッシュ値（H0〜H7）に加算され、次のブロックの処理に備えます。

PyCUDAを使用して、ハッシュ計算を模した、逐次的な計算処理をGPU上で行う方法を示します。ここでは、Pythonでのループ処理と、その処理をPyCUDAでGPU上で行うコードを比較します。

縦に通るスレッドインデックスの例　逐次的な計算処理のシンプルで理解しやすい並列化です。

[0.5 1.5 2.5 3.5 4.5 5.5 6.5 7.5 8.5 9.5]

import numpy as np

# データの準備
N = 10
a = np.arange(N, dtype=np.float32)
b = np.zeros(N, dtype=np.float32)

# 逐次的な計算処理
for i in range(N):
    # 計算A
    temp1 = a[i] * 2.0
    # 計算B
    temp2 = temp1 + 1.0
    # 計算C
    b[i] = temp2 * 0.5

print(b)

つまり for文のiを スレッド インデックスの I に書き換えるというような形です。これだけで並列計算になります 必要な スレッドが起動し 全て 並列計算で実行されるようになります、

!pip install pycuda

import numpy as np
import pycuda.driver as cuda
import pycuda.autoinit
import pycuda.compiler as compiler

# CUDAカーネルの定義
kernel_code = """
__global__ void process_elements(float *a, float *b, int N) {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    if (i < N) {
        // 計算A
        float temp1 = a[i] * 2.0f;
        // 計算B
        float temp2 = temp1 + 1.0f;
        // 計算C
        b[i] = temp2 * 0.5f;
    }
}
"""

# データの準備
N = 10
a = np.arange(N, dtype=np.float32)
b = np.zeros(N, dtype=np.float32)

# CUDAカーネルのコンパイル
mod = compiler.SourceModule(kernel_code)
process_elements = mod.get_function("process_elements")

# データの転送
a_gpu = cuda.mem_alloc(a.nbytes)
b_gpu = cuda.mem_alloc(b.nbytes)
cuda.memcpy_htod(a_gpu, a)

# スレッドとブロックの設定
block_size = 10
grid_size = (N + block_size - 1) // block_size

# CUDAカーネルの呼び出し
process_elements(a_gpu, b_gpu, np.int32(N), block=(block_size, 1, 1), grid=(grid_size, 1))

# 結果の転送
cuda.memcpy_dtoh(b, b_gpu)

print(b)

Pythonコード:

aの各要素に対して、逐次的な計算（A、B、C）を行い、その結果をbに格納します。

PyCUDAコード:

kernel_code: CUDAカーネルを定義し、GPU上で逐次的な計算処理を実行します。
process_elements: CUDAカーネルをコンパイルし、GPUで実行する関数を取得します。
データの転送 (cuda.mem_alloc、cuda.memcpy_htod)、カーネルの呼び出し (process_elements)、および結果の転送 (cuda.memcpy_dtoh) を行います。
スレッドとブロックのサイズ (block_size、grid_size) を設定し、カーネルを実行します。

10本のスレッドが起動します。シンプルで理解しやすい並列化です。
今回はこの方法でシンプルにハッシュ計算を並列化していきます。

横にスレッドインデックスを通す例。

スレッドインデックスが「横に通る」ようにするには、スレッドが2次元グリッド上で配置され、各スレッドが異なる計算ステージ（K31、K32、K33）を担当するようにします。これを実現するには、各計算ステージを別のカーネルとして実行し、それぞれのカーネルが完了した後に次のカーネルが実行されるようにします。

[0.5 1.5 2.5 3.5 4.5 5.5 6.5 7.5 8.5 9.5]

import numpy as np

# データの準備
N = 10
a = np.arange(N, dtype=np.float32)
b = np.zeros(N, dtype=np.float32)

# ステージごとの処理
# ステージK31: 計算A
temp1 = np.zeros(N, dtype=np.float32)
for i in range(N):
    temp1[i] = a[i] * 2.0

# ステージK32: 計算B
temp2 = np.zeros(N, dtype=np.float32)
for i in range(N):
    temp2[i] = temp1[i] + 1.0

# ステージK33: 計算C
for i in range(N):
    b[i] = temp2[i] * 0.5

print(b)

横にスレッドインデックスを通す

PyCUDAを使用して、スレッドインデックスが横に通っているような形で逐次的な計算処理を行うコードです。

CUDAカーネルの定義:

stage_K31: 計算Aを行うカーネル。
stage_K32: 計算Bを行うカーネル。
stage_K33: 計算Cを行うカーネル。
カーネルのコンパイルと取得:

CUDAカーネルをコンパイルし、各カーネル関数を取得します。
データの転送:

ホストメモリからデバイスメモリにデータを転送します。
カーネルの呼び出し:

各ステージのカーネルを順次呼び出します。stage_K31が完了してからstage_K32が実行され、その後stage_K33が実行されます。
結果の転送:

計算結果をデバイスメモリからホストメモリに転送し、出力します。
これにより、スレッドが横に通っているという形で、各計算ステージを順次実行し、結果を得ることができます。

!pip install pycuda

import numpy as np
import pycuda.driver as cuda
import pycuda.autoinit
import pycuda.compiler as compiler

# CUDAカーネルの定義
kernel_code = """
__global__ void stage_K31(float *a, float *temp1, int N) {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    if (i < N) {
        // 計算A
        temp1[i] = a[i] * 2.0f;
    }
}

__global__ void stage_K32(float *temp1, float *temp2, int N) {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    if (i < N) {
        // 計算B
        temp2[i] = temp1[i] + 1.0f;
    }
}

__global__ void stage_K33(float *temp2, float *b, int N) {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    if (i < N) {
        // 計算C
        b[i] = temp2[i] * 0.5f;
    }
}
"""

# データの準備
N = 10
a = np.arange(N, dtype=np.float32)
temp1 = np.zeros(N, dtype=np.float32)
temp2 = np.zeros(N, dtype=np.float32)
b = np.zeros(N, dtype=np.float32)

# CUDAカーネルのコンパイル
mod = compiler.SourceModule(kernel_code)
stage_K31 = mod.get_function("stage_K31")
stage_K32 = mod.get_function("stage_K32")
stage_K33 = mod.get_function("stage_K33")

# データの転送
a_gpu = cuda.mem_alloc(a.nbytes)
temp1_gpu = cuda.mem_alloc(temp1.nbytes)
temp2_gpu = cuda.mem_alloc(temp2.nbytes)
b_gpu = cuda.mem_alloc(b.nbytes)
cuda.memcpy_htod(a_gpu, a)

# スレッドとブロックの設定
block_size = 10
grid_size = (N + block_size - 1) // block_size

# ステージK31のカーネルの呼び出し
stage_K31(a_gpu, temp1_gpu, np.int32(N), block=(block_size, 1, 1), grid=(grid_size, 1))

# ステージK32のカーネルの呼び出し
stage_K32(temp1_gpu, temp2_gpu, np.int32(N), block=(block_size, 1, 1), grid=(grid_size, 1))

# ステージK33のカーネルの呼び出し
stage_K33(temp2_gpu, b_gpu, np.int32(N), block=(block_size, 1, 1), grid=(grid_size, 1))

# 結果の転送
cuda.memcpy_dtoh(b, b_gpu)

print(b)

本題。超速の ハッシュ計算のコードでは２次元構造でスレッドを起動しています。つまり最高速仕様です。詳しい解説は　追記のところでします。ここではシンプルに　hello world のハッシュを求める演算ユニットのスレッドを多数起動する方法でのコードです。

つまり１本のスレッドでhello world のハッシュを計算し同時に多数の同様のスレッドで計算してます。

縦にスレッドインデックスを通うしています。１次元のスレッドインデックスｘのみです。

global void sha256_kernel(const unsigned char *d_input, unsigned char *d_output, int num_messages) {
int tid = blockIdx.x * blockDim.x + threadIdx.x;

実行結果。ちょっとだけ遅いです。２５％程度。

Processing time: 0.0002 seconds

!pip install pycuda

import pycuda.driver as cuda
import pycuda.autoinit
import pycuda.compiler as compiler
import numpy as np
import time

## CUDAカーネルコードの定義
kernel_code = """
#include <cuda_runtime.h>
#include <stdio.h>

#define SHA256_BLOCK_SIZE 64  // 64 bytes
#define SHA256_HASH_SIZE 32   // 32 bytes

// SHA-256の定数値
__device__ const unsigned int 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
};

__device__ unsigned int right_rotate(unsigned int value, unsigned int shift) {
    return (value >> shift) | (value << (32 - shift));
}

__global__ void sha256_kernel(const unsigned char *d_input, unsigned char *d_output, int num_messages) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (tid < num_messages) {
        const unsigned char *message = d_input + tid * SHA256_BLOCK_SIZE;
        unsigned int w[64];
        unsigned int H[8] = {
            0x6a09e667, 0xbb67ae85, 0x3c6ef372, 0xa54ff53a,
            0x510e527f, 0x9b05688c, 0x1f83d9ab, 0x5be0cd19
        };

        for (int i = 0; i < 16; ++i) {
            w[i] = (message[i * 4] << 24) |
                   (message[i * 4 + 1] << 16) |
                   (message[i * 4 + 2] << 8) |
                   (message[i * 4 + 3]);
        }
        for (int i = 16; i < 64; ++i) {
            unsigned int s0 = right_rotate(w[i - 15], 7) ^ right_rotate(w[i - 15], 18) ^ (w[i - 15] >> 3);
            unsigned int s1 = right_rotate(w[i - 2], 17) ^ right_rotate(w[i - 2], 19) ^ (w[i - 2] >> 10);
            w[i] = w[i - 16] + s0 + w[i - 7] + s1;
        }

        unsigned int a = H[0], b = H[1], c = H[2], d = H[3];
        unsigned int e = H[4], f = H[5], g = H[6], h = H[7];
        for (int i = 0; i < 64; ++i) {
            unsigned int S1 = right_rotate(e, 6) ^ right_rotate(e, 11) ^ right_rotate(e, 25);
            unsigned int ch = (e & f) ^ ((~e) & g);
            unsigned int temp1 = h + S1 + ch + K[i] + w[i];
            unsigned int S0 = right_rotate(a, 2) ^ right_rotate(a, 13) ^ right_rotate(a, 22);
            unsigned int maj = (a & b) ^ (a & c) ^ (b & c);
            unsigned int temp2 = S0 + maj;

            h = g;
            g = f;
            f = e;
            e = d + temp1;
            d = c;
            c = b;
            b = a;
            a = temp1 + temp2;
        }

        H[0] += a;
        H[1] += b;
        H[2] += c;
        H[3] += d;
        H[4] += e;
        H[5] += f;
        H[6] += g;
        H[7] += h;

        for (int i = 0; i < 8; ++i) {
            d_output[tid * SHA256_HASH_SIZE + i * 4] = (H[i] >> 24) & 0xff;
            d_output[tid * SHA256_HASH_SIZE + i * 4 + 1] = (H[i] >> 16) & 0xff;
            d_output[tid * SHA256_HASH_SIZE + i * 4 + 2] = (H[i] >> 8) & 0xff;
            d_output[tid * SHA256_HASH_SIZE + i * 4 + 3] = H[i] & 0xff;
        }
    }
}
"""



# コンパイル
module = compiler.SourceModule(kernel_code)
sha256_kernel = module.get_function("sha256_kernel")

# メッセージ数とデータ準備
num_messages = 10000
message_size = 64
hash_size = 32

# 入力データの準備
input_data = np.full((num_messages, message_size), 0x00, dtype=np.uint8)
input_data[0, :11] = np.frombuffer(b"Hello World", dtype=np.uint8)  # 最初のメッセージを設定

# GPUメモリの割り当て
d_input = cuda.mem_alloc(input_data.nbytes)
d_output = cuda.mem_alloc(num_messages * hash_size)

# データ転送
cuda.memcpy_htod(d_input, input_data)

# スレッド数とブロック数の設定
threads_per_block = 256
blocks_per_grid = (num_messages + threads_per_block - 1) // threads_per_block

# 処理時間の計測開始
start_time = time.time()

# CUDAカーネルの実行
sha256_kernel(d_input, d_output, np.int32(num_messages), block=(threads_per_block, 1, 1), grid=(blocks_per_grid, 1))

# 処理時間の計測終了
end_time = time.time()

# 処理時間の表示
print(f"Processing time: {end_time - start_time:.4f} seconds")

# 結果の転送
output_data = np.empty((num_messages, hash_size), dtype=np.uint8)
cuda.memcpy_dtoh(output_data, d_output)

Processing time: 0.0002 seconds

そして3つの計算プロセスの中に入っている ループも 演算ユニットをあてがって 並列化すると さらに早くなり 超速のハッシュ計算のコードとなります。

追記。

現在の1次元並列処理を2次元並列処理に書き換えることが可能です。これにより、2次元グリッド上で各スレッドが独立して計算を行うことができます。以下のコードでは、スレッドIDをXとYの2次元に分けて計算するように書き換えました。

!pip install pycuda

import hashlib
import time
import numpy as np
import pycuda.autoinit
import pycuda.driver as cuda
from pycuda.compiler import SourceModule
import matplotlib.pyplot as plt

# データの初期化
data = b'hello world' * 10000  # 1万個の「hello world」

def sha256_hash(data):
    """
    CPUでSHA-256ハッシュを計算する関数
    """
    return hashlib.sha256(data).hexdigest()

def cpu_sha256_hash(data):
    start_time = time.time()
    hashes = [sha256_hash(data) for _ in range(10000)]
    end_time = time.time()
    print(f"SHA-256 Hash (CPU): {hashes[0]}")  # 最初のハッシュを表示
    print(f"CPU Processing Time: {end_time - start_time:.4f} seconds")

# GPUカーネルのソースコード
cuda_kernel_code = """
#include <stdint.h>

__device__ uint32_t right_rotate(uint32_t value, unsigned int shift) {
    return (value >> shift) | (value << (32 - shift));
}

__device__ void sha256_transform(uint32_t* state, const uint8_t* data) {
    static const uint32_t k[64] = {
        // (省略: 64個の定数k)
    };

    uint32_t a = state[0];
    uint32_t b = state[1];
    uint32_t c = state[2];
    uint32_t d = state[3];
    uint32_t e = state[4];
    uint32_t f = state[5];
    uint32_t g = state[6];
    uint32_t h = state[7];

    uint32_t w[64];
    for (int i = 0; i < 16; ++i) {
        w[i] = (data[i * 4] << 24) | (data[i * 4 + 1] << 16) | (data[i * 4 + 2] << 8) | data[i * 4 + 3];
    }
    for (int i = 16; i < 64; ++i) {
        uint32_t s0 = right_rotate(w[i - 15], 7) ^ right_rotate(w[i - 15], 18) ^ (w[i - 15] >> 3);
        uint32_t s1 = right_rotate(w[i - 2], 17) ^ right_rotate(w[i - 2], 19) ^ (w[i - 2] >> 10);
        w[i] = w[i - 16] + s0 + w[i - 7] + s1;
    }

    for (int i = 0; i < 64; ++i) {
        uint32_t S1 = right_rotate(e, 6) ^ right_rotate(e, 11) ^ right_rotate(e, 25);
        uint32_t ch = (e & f) ^ (~e & g);
        uint32_t temp1 = h + S1 + ch + k[i] + w[i];
        uint32_t S0 = right_rotate(a, 2) ^ right_rotate(a, 13) ^ right_rotate(a, 22);
        uint32_t maj = (a & b) ^ (a & c) ^ (b & c);
        uint32_t temp2 = S0 + maj;

        h = g;
        g = f;
        f = e;
        e = d + temp1;
        d = c;
        c = b;
        b = a;
        a = temp1 + temp2;
    }

    state[0] += a;
    state[1] += b;
    state[2] += c;
    state[3] += d;
    state[4] += e;
    state[5] += f;
    state[6] += g;
    state[7] += h;
}

extern "C" __global__ void sha256_kernel(uint32_t* state, const uint8_t* data, int num_hashes) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    int index = y * gridDim.x * blockDim.x + x;

    if (index < num_hashes) {
        const uint8_t* data_offset = data + index * 64;
        sha256_transform(state + index * 8, data_offset);
    }
}
"""

# 2次元スレッド数のパターン
thread_counts = [(8, 8), (16, 16), (32, 32), (64, 64)]

# 各スレッド数での計算時間を記録するリスト
gpu_times = []

def gpu_sha256_hash(data, threads_per_block):
    # CUDAカーネルのコンパイル
    mod = SourceModule(cuda_kernel_code)
    sha256_kernel = mod.get_function("sha256_kernel")

    # データの分割と状態の設定
    num_hashes = 10000
    data_bytes = np.frombuffer(data, dtype=np.uint8)
    state = np.zeros((num_hashes, 8), dtype=np.uint32)
    
    # GPUメモリの確保とデータの転送
    d_data = cuda.mem_alloc(data_bytes.nbytes)
    d_state = cuda.mem_alloc(state.nbytes)
    cuda.memcpy_htod(d_data, data_bytes)
    cuda.memcpy_htod(d_state, state)

    # CUDAカーネルの起動
    start_time = time.time()
    blocks_per_grid_x = (num_hashes + threads_per_block[0] * threads_per_block[1] - 1) // (threads_per_block[0] * threads_per_block[1])
    blocks_per_grid_y = 1
    sha256_kernel(d_state, d_data, np.int32(num_hashes), block=(threads_per_block[0], threads_per_block[1], 1), grid=(blocks_per_grid_x, blocks_per_grid_y))
    cuda.Context.synchronize()
    end_time = time.time()

    # 結果の取得
    cuda.memcpy_dtoh(state, d_state)
    
    # 処理時間を記録
    gpu_times.append(end_time - start_time)

if __name__ == "__main__":
    # CPUでの計算
    cpu_sha256_hash(data)

    # 各スレッド数でのGPU計算を実行
    for threads in thread_counts:
        gpu_sha256_hash(data, threads)

    # プロット
    plt.figure(figsize=(10, 5))
    plt.plot([threads[0] * threads[1] for threads in thread_counts], gpu_times, marker='o')
    plt.title("GPU SHA-256 Hashing Time by Thread Count")
    plt.xlabel("Number of Threads per Block")
    plt.ylabel("Processing Time (seconds)")
    plt.grid(True)
    plt.show()

変更点の説明

2次元スレッドIDの使用:

スレッドIDがxとyの2次元に分割され、それぞれblockIdx.x, blockIdx.y, threadIdx.x, threadIdx.yを使って計算されます。

最終的なインデックスindexは、xとyの2次元の位置に基づいて計算されます。

thread_countsの変更:

2次元のスレッドブロックのサイズを指定するために、(x, y)形式のタプルを使用します。

blocks_per_grid_xとblocks_per_grid_y:

コード内の2次元のスレッドIDを使用している部分を抜き出して説明。

2次元のスレッドIDを使用している部分は、主にCUDAカーネル関数sha256_kernel内と、ホストコードのCUDAカーネル起動部分です。

1. CUDAカーネル関数内 (sha256_kernel)

extern "C" __global__ void sha256_kernel(uint32_t* state, const uint8_t* data, int num_hashes) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    int index = y * gridDim.x * blockDim.x + x;

    if (index < num_hashes) {
        const uint8_t* data_offset = data + index * 64;
        sha256_transform(state + index * 8, data_offset);
    }
}

xとyの計算:

x = blockIdx.x * blockDim.x + threadIdx.x
y = blockIdx.y * blockDim.y + threadIdx.y
これにより、各スレッドの2次元座標(x, y)が計算されます。

1. ホストコード内のカーネル起動部分

start_time = time.time()
blocks_per_grid_x = (num_hashes + threads_per_block[0] * threads_per_block[1] - 1) // (threads_per_block[0] * threads_per_block[1])
blocks_per_grid_y = 1
sha256_kernel(d_state, d_data, np.int32(num_hashes), block=(threads_per_block[0], threads_per_block[1], 1), grid=(blocks_per_grid_x, blocks_per_grid_y))
cuda.Context.synchronize()
end_time = time.time()

blockパラメータ:

block=(threads_per_block[0], threads_per_block[1], 1)
これはスレッドブロック内のスレッド数を指定しており、2次元スレッド配列のサイズを指定しています。

gridパラメータ:

grid=(blocks_per_grid_x, blocks_per_grid_y)
これはグリッド内のブロック数を指定しており、2次元グリッドのサイズを指定しています。

このように、スレッドIDを2次元で使用することで、各スレッドが異なるデータを処理できるようにしています。

つまりこのコードでは、スレッドIDのXによってハローワールドのハッシュを計算するスレッドを起動し、 スレッドIDのYによってその中で行われている主な3つの計算プロセスの中の並列化できる部分を スレッドIDのYで並列化しているということです。

コードでは、2次元のスレッドID (X, Y) を使用して並列処理を行っています。

具体的には：

スレッドIDのX: 各スレッドはハローワールドのSHA-256ハッシュ計算を担当します。これは、スレッドIDのX（x成分）が異なるスレッドが異なるデータチャンクに対して計算を行うことを意味します。

スレッドIDのY: 各スレッドは、データチャンクの中で行われる特定の計算（例えば、SHA-256の各ステップ）を並列化するために使用されます。この場合、スレッドIDのY（y成分）が異なるスレッドが同じデータチャンク内の異なる計算タスクを実行します。

このようにして、スレッドIDのXとYを利用して、2次元で並列化を行い、計算を効率化しています。

つまり、具体的な動作としては XとYで指定される カーネル上の計算ユニットと それが使用するメモリの位置を XとYによって指定する そういうシステムです。

具体的には、以下のようなシステムです：

計算ユニットの指定 (XとY):

XとYで指定されるスレッドIDは、CUDAカーネルの各スレッドに一意のIDを与えます。このIDは、スレッドがどの計算を担当するかを決定するのに使用されます。

メモリの位置指定:

データメモリの位置は、スレッドID (X, Y) に基づいて計算されます。