CUDA প্রোগ্রামিং মডেল — হার্ডওয়্যারকে কোড দিয়ে address করা

Launch করার পর আসলে কী ঘটে

Kernel launch লেখার পর CPU একটা asynchronous request পাঠায় GPU driver-কে — CPU সাথে সাথেই পরের লাইনে চলে যায় (kernel শেষ হওয়ার জন্য অপেক্ষা করে না, যদি না explicitly cudaDeviceSynchronize() ডাকা হয়)।

GPU-র পাশে:

1. Block scheduler প্রতিটা block-কে একটা উপলব্ধ SM-এ বরাদ্দ করে (গত লেসনের occupancy সীমার সাপেক্ষে — কতগুলো block একটা SM-এ একসাথে আঁটবে)
2. প্রতিটা SM তার বরাদ্দকৃত block-এর thread-গুলোকে ৩২-এর দলে (warp) ভাগ করে
3. Warp scheduler প্রতি cycle-এ একটা প্রস্তুত warp বেছে চালায়; memory-র জন্য অপেক্ষারত warp সরিয়ে অন্য warp চালানো হয় (latency hiding, গত লেসন)
4. একটা block শেষ হলে সেই SM-এর জায়গায় নতুন block আসতে পারে (যদি grid-এ আরও block অপেক্ষা করে থাকে)

সম্পূর্ণ প্রোগ্রাম — host থেকে device, আবার host-এ

একটা kernel একা কিছু করতে পারে না — data-কে CPU-র memory (host) থেকে GPU-র memory (device)-এ পাঠাতে হয়, কাজ শেষে ফেরত আনতে হয়। এখানে সম্পূর্ণ, বাস্তব vector-add প্রোগ্রাম:

#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>

#define N 1048576              /* ~১০ লক্ষ element (2^20) */
#define THREADS_PER_BLOCK 256

__global__ void vecAdd(const float *A, const float *B, float *C, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i \< n) {
        C[i] = A[i] + B[i];
    }
}

int main(void) {
    size_t bytes = N * sizeof(float);

    /* ১. Host memory allocate ও ভরাট করা */
    float *h_A = (float *)malloc(bytes);
    float *h_B = (float *)malloc(bytes);
    float *h_C = (float *)malloc(bytes);
    for (int i = 0; i \< N; i++) {
        h_A[i] = 1.0f;
        h_B[i] = 2.0f;
    }

    /* ২. Device memory allocate করা — GPU-র নিজস্ব DRAM-এ */
    float *d_A, *d_B, *d_C;
    cudaMalloc((void **)&d_A, bytes);
    cudaMalloc((void **)&d_B, bytes);
    cudaMalloc((void **)&d_C, bytes);

    /* ৩. Host থেকে device-এ copy — PCIe/NVLink-এর মধ্য দিয়ে */
    cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice);

    /* ৪. Kernel launch */
    int numBlocks = (N + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;
    vecAdd<<<numBlocks, THREADS_PER_BLOCK>>>(d_A, d_B, d_C, N);

    /* ৫. ফলাফল device থেকে host-এ ফেরত */
    cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);

    printf("h_C[0] = %f (আশা করি 3.0)\n", h_C[0]);

    /* ৬. পরিষ্কার — device ও host memory উভয়ই */
    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
    free(h_A); free(h_B); free(h_C);
    return 0;
}

প্রতিটা মেমরি ধরনের ভূমিকা:

Shared memory ও __syncthreads() — block-এর ভেতরে সহযোগিতা

গত লেসনে দেখেছি shared memory একটা block-এর সব thread ভাগাভাগি করে, কিন্তু programmer-managed — hardware automatically সিদ্ধান্ত নেয় না। __shared__ keyword দিয়ে ঘোষণা করা হয়:

__global__ void blockSum(const float *input, float *blockSums, int n) {
    __shared__ float cache[THREADS_PER_BLOCK];   /* এই block-এর সব thread ভাগ করে */

    int tid = threadIdx.x;
    int i = blockIdx.x * blockDim.x + tid;

    /* ধাপ ১ — global memory থেকে shared memory-তে লোড */
    cache[tid] = (i \< n) ? input[i] : 0.0f;
    __syncthreads();   /* বাধ্যতামূলক — নিচে ব্যাখ্যা */

    /* ধাপ ২ — cache-এর ভেতরে tree-আকৃতির যোগফল (বিস্তারিত পরের লেসনে) */
    for (int stride = blockDim.x / 2; stride > 0; stride /= 2) {
        if (tid \< stride) {
            cache[tid] += cache[tid + stride];
        }
        __syncthreads();   /* প্রতিটা ধাপের পরে — পরের ধাপ শুরুর আগে */
    }

    /* ধাপ ৩ — শুধু thread ০ চূড়ান্ত ফলাফল লেখে */
    if (tid == 0) {
        blockSums[blockIdx.x] = cache[0];
    }
}

কেন __syncthreads() ছাড়া চলবে না: একটা block-এর সব thread একসাথে, lockstep-এ চলে না (বিভিন্ন warp ভিন্ন গতিতে এগোতে পারে, বিশেষ করে memory access-এর সময়)। cache[tid] = ... লেখার পরপরই cache[tid + stride] পড়া শুরু করলে — যদি অন্য warp তখনও তার cache[...] না লিখে থাকে — আপনি stale বা undefined data পড়বেন। __syncthreads() একটা barrier: block-এর প্রতিটা thread এই লাইনে না পৌঁছানো পর্যন্ত কেউ এগোবে না। এটা lesson ৪-এ (atomics/memory ordering) দেখা memory barrier ধারণারই একটা রূপ, শুধু এখানে scope একটা GPU thread block-এ সীমাবদ্ধ, পুরো system-এ নয়।