Kiến trúc GPU và CUDA

GPU đã biến đổi AI bằng cách cung cấp hàng ngàn lõi cho tính song song quy mô lớn. File này bao gồm triết lý thiết kế GPU vs CPU, hệ thống phân cấp bộ nhớ GPU, lập trình CUDA trong C++, mô hình thực thi SIMT, các mẫu truy cập bộ nhớ, đồng bộ hóa, streams, profiling, và các thế hệ GPU NVIDIA — kiến thức cần thiết để viết và hiểu các kernel GPU.

• Để có các hướng dẫn CUDA thực hành với các ví dụ hoạt động đầy đủ, hãy xem kho lưu trữ đồng hành: github.com/HenryNdubuaku/cuda-tutorials.

• Một GPU NVIDIA hiện đại có hơn 10,000 nhân CUDA. Một CPU có 4-128 lõi. Lợi thế 100-1000x lõi chính là lý do GPU thống trị ML: huấn luyện một transformer yêu cầu hàng nghìn tỷ phép nhân-cộng, và GPU xử lý chúng song song ở quy mô CPU không thể sánh kịp.

• Ngay cả khi bạn không bao giờ tự viết kernel CUDA, hiểu kiến trúc GPU giải thích: tại sao kích thước batch lại quan trọng (cần đủ công việc để làm đầy GPU), tại sao bộ nhớ thường là nút thắt (không phải tính toán), và tại sao một số phép toán (scatter, rẽ nhánh có điều kiện) lại chậm trên GPU.

GPU vs CPU: Các Thiết kế Khác nhau Cơ bản

• CPU được thiết kế cho độ trễ: giảm thiểu thời gian hoàn thành một tác vụ. Nó dành phần lớn ngân sách bóng bán dẫn cho bộ nhớ cache, bộ dự đoán rẽ nhánh, và thực thi ngoài thứ tự — tất cả các mẹo để làm cho một luồng nhanh.

• GPU được thiết kế cho thông lượng: tối đa hóa số lượng tác vụ hoàn thành mỗi giây. Nó dành phần lớn bóng bán dẫn cho các đơn vị thực thi (ALU). Các luồng riêng lẻ chậm, nhưng có hàng ngàn luồng.

	CPU	GPU
Lõi	4-128 (phức tạp, nhanh)	1,000-20,000 (đơn giản, chậm)
Tốc độ xung nhịp	3-5 GHz	1-2.5 GHz
Bộ nhớ cache	Lớn (32 MB+ L3)	Nhỏ (shared memory mỗi SM)
Dự đoán rẽ nhánh	Tinh vi	Không có (mọi luồng theo cùng đường)
Tốt nhất cho	Độ trễ thấp, luồng điều khiển phức tạp	Thông lượng cao, công việc song song hóa dữ liệu
FLOPS điển hình (FP32)	1-5 TFLOPS	30-80 TFLOPS
Băng thông bộ nhớ	50-100 GB/s	1-3 TB/s

• Lợi thế băng thông bộ nhớ của GPU (10-30x) thường quan trọng hơn lợi thế tính toán. Nhiều phép toán ML bị giới hạn bởi bộ nhớ (các phép toán theo phần tử, chuẩn hóa, attention), và băng thông của GPU cho phép nó cấp dữ liệu cho các lõi đủ nhanh.

Hệ thống Phân cấp Bộ nhớ GPU

• Hiểu bộ nhớ GPU là rất quan trọng vì truy cập bộ nhớ là nút thắt chính, không phải tính toán.

Bộ nhớ	Kích thước	Độ trễ	Băng thông	Phạm vi
Thanh ghi	~256 KB mỗi SM	0 chu kỳ	Cao nhất	Mỗi luồng
Shared memory	48-228 KB mỗi SM	~5 chu kỳ	~20 TB/s	Mỗi block luồng
L1 cache	128-256 KB mỗi SM	~30 chu kỳ		Mỗi SM
L2 cache	4-96 MB	~200 chu kỳ	~6 TB/s	Toàn cục
Global memory (HBM)	24-192 GB	~400 chu kỳ	1-3.3 TB/s	Toàn cục

• Thanh ghi là nhanh nhất nhưng hạn chế nhất. Mỗi luồng có một tập thanh ghi riêng (thường tối đa 255). Dùng quá nhiều thanh ghi mỗi luồng làm giảm occupancy (ít luồng có thể chạy đồng thời hơn).

• Shared memory là bộ nhớ cache do lập trình viên quản lý, được chia sẻ bởi tất cả luồng trong một block. Nó là chìa khóa để viết kernel CUDA nhanh: nạp một tile dữ liệu từ global memory chậm vào shared memory nhanh, sau đó tính toán trên đó. Đây là mẫu tiling thống trị lập trình GPU.

• Global memory (HBM): bộ nhớ GPU chính (VRAM). Lớn nhưng chậm (400 chu kỳ độ trễ). Mọi dữ liệu bắt đầu và kết thúc ở đây. Mục tiêu của tối ưu hóa kernel là giảm thiểu truy cập global memory.

Mô hình Lập trình CUDA

• CUDA (Compute Unified Device Architecture) là mô hình lập trình GPU của NVIDIA. Bạn viết kernel: các hàm chạy trên GPU, được thực thi bởi hàng ngàn luồng đồng thời.

Hệ thống Phân cấp: Grid, Block, Thread

Grid (toàn bộ lần khởi chạy)
├── Block (0,0)
│   ├── Thread (0,0)
│   ├── Thread (1,0)
│   ├── Thread (2,0)
│   └── ... (tối đa 1024 luồng mỗi block)
├── Block (1,0)
│   ├── Thread (0,0)
│   └── ...
└── ... (hàng triệu block có thể)

• Thread: đơn vị nhỏ nhất. Mỗi luồng có một ID duy nhất (threadIdx.x) trong block của nó.
• Block: một nhóm các luồng có thể chia sẻ bộ nhớ và đồng bộ hóa. Block ID: blockIdx.x. Kích thước block: blockDim.x (tối đa 1024 luồng).

• Grid: tất cả các block được khởi chạy bởi một kernel. Có thể là 1D, 2D, hoặc 3D.

• Mỗi luồng tính chỉ số toàn cục của nó: int idx = blockIdx.x * blockDim.x + threadIdx.x;

Kernel CUDA Đầu tiên của Bạn

// vector_add.cu — Tệp nguồn CUDA (phần mở rộng .cu)

#include <stdio.h>

// __global__ đánh dấu đây là kernel GPU (được gọi từ CPU, chạy trên GPU)
__global__ void vector_add(const float* a, const float* b, float* c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {           // kiểm tra ranh giới (grid có thể lớn hơn dữ liệu)
        c[idx] = a[idx] + b[idx];
    }
}

int main() {
    int n = 1 << 20;  // ~1 triệu phần tử
    size_t bytes = n * sizeof(float);

    // Cấp phát bộ nhớ host (CPU)
    float *h_a = new float[n];
    float *h_b = new float[n];
    float *h_c = new float[n];

    // Khởi tạo
    for (int i = 0; i < n; i++) {
        h_a[i] = 1.0f;
        h_b[i] = 2.0f;
    }

    // Cấp phát bộ nhớ device (GPU)
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);

    // Sao chép dữ liệu từ CPU sang GPU
    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);

    // Khởi chạy kernel: 256 luồng mỗi block, đủ block để phủ n phần tử
    int block_size = 256;
    int grid_size = (n + block_size - 1) / block_size;  // chia làm tròn lên
    vector_add<<<grid_size, block_size>>>(d_a, d_b, d_c, n);

    // Sao chép kết quả từ GPU về CPU
    cudaMemcpy(h_c, d_a, bytes, cudaMemcpyDeviceToHost);

    // Xác minh
    printf("c[0] = %f (expected 3.0)\n", h_c[0]);

    // Giải phóng bộ nhớ
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    delete[] h_a; delete[] h_b; delete[] h_c;

    return 0;
}

# Biên dịch với trình biên dịch NVIDIA
nvcc -O3 -o vector_add vector_add.cu
./vector_add

• Các khái niệm C++ quan trọng trong CUDA:
  • __global__: một từ khóa CUDA đánh dấu hàm kernel. Gọi từ CPU (host), chạy trên GPU (device).
  • <<<grid_size, block_size>>>: cú pháp khởi chạy kernel. Chỉ định bao nhiêu block và luồng để dùng.
  • cudaMalloc / cudaFree: cấp phát/giải phóng bộ nhớ GPU (như new/delete nhưng cho GPU).
  • cudaMemcpy: sao chép dữ liệu giữa CPU và GPU. Đây thường là nút thắt lớn nhất (băng thông PCIe ~32 GB/s, trong khi băng thông bộ nhớ GPU ~3 TB/s).

Warp và SIMT

• GPU thực thi các luồng theo nhóm 32 gọi là warp. Tất cả 32 luồng trong một warp thực thi cùng một chỉ thị tại cùng một thời điểm (Single Instruction, Multiple Threads — SIMT). Đây là tương đương của SIMD trên GPU, nhưng ở cấp độ luồng.

• Warp divergence (phân kỳ warp) xảy ra khi các luồng trong cùng một warp đi theo các nhánh khác nhau của câu lệnh if. GPU không thể thực thi hai chỉ thị khác nhau đồng thời trong một warp, nên nó thực thi cả hai nhánh tuần tự, mặt nạ các luồng không nên tham gia. Điều này giảm một nửa hiệu năng (hoặc tệ hơn).

// XẤU: warp divergence (các luồng trong cùng warp đi theo đường khác nhau)
if (threadIdx.x % 2 == 0) {
    c[idx] = a[idx] + b[idx];    // luồng chẵn làm cái này
} else {
    c[idx] = a[idx] - b[idx];    // luồng lẻ làm cái này (cùng warp, tuần tự hóa)
}

// TỐT HƠN: không rẽ nhánh (không divergence)
float sign = (threadIdx.x % 2 == 0) ? 1.0f : -1.0f;
c[idx] = a[idx] + sign * b[idx];  // mọi luồng thực thi cùng chỉ thị

Kết hợp Bộ nhớ (Memory Coalescing)

• Truy cập kết hợp: khi các luồng liên tiếp truy cập các địa chỉ bộ nhớ liên tiếp, GPU kết hợp chúng thành một giao dịch bộ nhớ duy nhất. Điều này rất quan trọng cho hiệu năng.

// TỐT: coalesced — luồng 0 đọc a[0], luồng 1 đọc a[1], ...
c[idx] = a[idx] + b[idx];

// XẤU: strided — luồng 0 đọc a[0], luồng 1 đọc a[stride], ...
c[idx] = a[idx * stride] + b[idx * stride];  // stride > 1 lãng phí băng thông

• Với một warp 32 luồng, truy cập kết hợp tải 128 byte (32 × 4 byte cho float32) trong một giao dịch. Truy cập stride yêu cầu nhiều giao dịch, mỗi giao dịch tải 128 byte nhưng chỉ dùng một phần. Stride 32 là trường hợp tệ nhất: mỗi giao dịch tải 128 byte nhưng chỉ một luồng dùng 4 byte (sử dụng 3%).

Shared Memory và Tiling

• Mẫu tiling là kỹ thuật tối ưu GPU quan trọng nhất. Ý tưởng: nạp một khối dữ liệu từ global memory chậm vào shared memory nhanh, tính toán trên nó, sau đó ghi kết quả trở lại.

// Nhân ma trận với shared memory tiling (đơn giản hóa)
__global__ void matmul_tiled(const float* A, const float* B, float* C,
                              int M, int N, int K) {
    // Shared memory cho một tile của A và một tile của B
    __shared__ float tile_A[TILE_SIZE][TILE_SIZE];
    __shared__ float tile_B[TILE_SIZE][TILE_SIZE];

    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    float sum = 0.0f;

    // Lặp qua các tile
    for (int t = 0; t < (K + TILE_SIZE - 1) / TILE_SIZE; t++) {
        // Nạp một tile của A và B vào shared memory
        if (row < M && t * TILE_SIZE + threadIdx.x < K)
            tile_A[threadIdx.y][threadIdx.x] = A[row * K + t * TILE_SIZE + threadIdx.x];
        else
            tile_A[threadIdx.y][threadIdx.x] = 0.0f;

        if (col < N && t * TILE_SIZE + threadIdx.y < K)
            tile_B[threadIdx.y][threadIdx.x] = B[(t * TILE_SIZE + threadIdx.y) * N + col];
        else
            tile_B[threadIdx.y][threadIdx.x] = 0.0f;

        __syncthreads();  // chờ mọi luồng nạp xong

        // Tính tích vô hướng riêng phần từ tile này
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += tile_A[threadIdx.y][k] * tile_B[k][threadIdx.x];
        }

        __syncthreads();  // chờ trước khi nạp tile tiếp theo
    }

    if (row < M && col < N)
        C[row * N + col] = sum;
}

• __shared__: khai báo bộ nhớ được chia sẻ bởi tất cả luồng trong block (nhanh, trên chip).
• __syncthreads(): một barrier chờ cho đến khi mọi luồng trong block đạt đến điểm này. Cần thiết giữa việc ghi vào shared memory và đọc từ nó (nếu không một số luồng đọc dữ liệu cũ).
• Tại sao tiling hiệu quả: không có nó, mỗi luồng tải từ global memory cho mỗi phép nhân. Với tiling, một khối dữ liệu TILE_SIZE × TILE_SIZE được nạp một lần vào shared memory và được tái sử dụng bởi tất cả luồng trong block. Hệ số tái sử dụng là TILE_SIZE, giảm lưu lượng global memory theo hệ số đó.

Streams và Tính đồng thời (Concurrency)

• Theo mặc định, các thao tác CUDA là tuần tự: CPU khởi chạy một kernel, chờ nó hoàn thành, sau đó khởi chạy cái tiếp theo. Streams cho phép chồng lấp:

cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

// Các thao tác này có thể chồng lấp: các stream khác nhau thực thi đồng thời
cudaMemcpyAsync(d_a, h_a, bytes, cudaMemcpyHostToDevice, stream1);
cudaMemcpyAsync(d_b, h_b, bytes, cudaMemcpyHostToDevice, stream2);

kernel1<<<grid, block, 0, stream1>>>(d_a, d_c);
kernel2<<<grid, block, 0, stream2>>>(d_b, d_d);

• Streams chồng lấp truyền dữ liệu với tính toán: trong khi kernel của một stream chạy, stream khác sao chép dữ liệu. Điều này ẩn độ trễ truyền PCIe và giữ cho GPU bận rộn.

Profiling Mã CUDA

# NVIDIA Nsight Compute: profiling cấp kernel
ncu --set full ./my_program

# NVIDIA Nsight Systems: dòng thời gian cấp hệ thống
nsys profile ./my_program

# Số liệu nhanh
ncu --metrics sm__throughput,dram__throughput ./my_program

• Những gì cần tìm kiếm:
  • Occupancy: tỷ lệ dung lượng của SM được sử dụng. Occupancy thấp (< 50%) nghĩa là quá ít luồng để ẩn độ trễ bộ nhớ. Nguyên nhân: quá nhiều thanh ghi trên mỗi luồng, quá nhiều shared memory trên mỗi block.
  • Memory throughput: so sánh với băng thông đỉnh. Nếu đạt < 50% đỉnh, các mẫu truy cập bộ nhớ không hiệu quả (non-coalesced, xung đột bank).
  • Compute throughput: so sánh với đỉnh FLOPS. Nếu cả memory và compute throughput đều thấp, kernel bị giới hạn bởi độ trễ (không đủ song song).

Các Kỹ thuật Tối ưu Nâng cao

• Ngoài những điều cơ bản về coalescing và shared memory tiling, mã GPU (và CPU) hiệu năng cao dùng một số kỹ thuật nâng cao:

Bố cục Dữ liệu: AoS vs SoA

• Array of Structures (AoS): mỗi phần tử lưu tất cả các trường của nó cùng nhau. [{x,y,z}, {x,y,z}, {x,y,z}].
• Structure of Arrays (SoA): mỗi trường được lưu trong một mảng liên tiếp riêng. {[x,x,x], [y,y,y], [z,z,z]}.

// AoS: XẤU cho SIMD/GPU (truy cập tất cả giá trị x chạm vào bộ nhớ không liên tiếp)
struct Particle { float x, y, z, mass; };
Particle particles[N];
// particles[0].x, particles[1].x cách nhau 16 byte

// SoA: TỐT cho SIMD/GPU (tất cả giá trị x là liên tiếp)
struct Particles {
    float x[N], y[N], z[N], mass[N];
};
// x[0], x[1] cách nhau 4 byte — hoàn hảo cho coalesced access và SIMD

• SoA hầu như luôn nhanh hơn cho các khối lượng công việc song song hóa dữ liệu (SIMD, GPU). AoS tốt hơn khi bạn luôn truy cập tất cả các trường của một phần tử cùng nhau (hiếm trong mã số học). Tensor PyTorch vốn dĩ là SoA: mỗi đặc trưng là một chiều liên tiếp.

Software Prefetching

• CPU có thể được yêu cầu bắt đầu tải dữ liệu trước khi nó cần, ẩn độ trễ bộ nhớ:

#include <xmmintrin.h>  // cho _mm_prefetch

for (int i = 0; i < n; i += 4) {
    _mm_prefetch((char*)(a + i + 64), _MM_HINT_T0);  // prefetch 64 phần tử phía trước
    // xử lý a[i:i+4] với SIMD
    __m128 va = _mm_load_ps(a + i);
    // ...
}

• Chỉ thị prefetch là một gợi ý: nếu dữ liệu đã có trong cache, nó là no-op. Nếu không, CPU bắt đầu tải nó trong nền khi thực thi các chỉ thị khác. Khoảng cách prefetch (64 phần tử phía trước trong ví dụ này) nên được tùy chỉnh để khớp với độ trễ bộ nhớ và thời gian vòng lặp.

Kernel Fusion

• Kernel fusion kết hợp nhiều phép toán thành một kernel duy nhất để tránh ghi kết quả trung gian vào bộ nhớ. Đây là tối ưu hóa GPU tác động nhất duy nhất cho ML:

// KHÔNG FUSE: 3 lần khởi chạy kernel, 3 chuyến khứ hồi global memory
y = matmul(x, W)     // ghi y vào global memory
z = y + bias          // đọc y, ghi z
out = relu(z)         // đọc z, ghi out

// ĐÃ FUSE: 1 lần khởi chạy kernel, 1 lần ghi global memory
out = fused_matmul_bias_relu(x, W, bias)  // y và z không bao giờ rời SRAM

• Với các phép toán bị giới hạn bởi bộ nhớ (add bias, ReLU, layer norm), lưu lượng bộ nhớ thống trị thời gian thực thi. Fusion loại bỏ hoàn toàn lưu lượng. torch.compile của PyTorch và Triton cho phép fusion tự động hoặc với nỗ lực tối thiểu.

Kernel Mixed-Precision

• Dùng độ chính xác thấp hơn (FP16, BF16, INT8) cho tính toán và độ chính xác cao hơn (FP32) cho tích lũy cho ta điều tốt nhất của cả hai thế giới:

// Tensor Core: nhân ma trận FP16, tích lũy trong FP32
// Mỗi chỉ thị Tensor Core: D (FP32) = A (FP16) × B (FP16) + C (FP32)
nvcuda::wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);

• FP16 nhỏ hơn 2x so với FP32, nên nó nhân đôi băng thông bộ nhớ (nút thắt thường gặp) và đặt vừa 2x dữ liệu hơn trong cache. Tensor Core xử lý FP16 với tốc độ 8-16x tốc độ của các nhân CUDA FP32. Đây là lý do huấn luyện mixed-precision (chương 6) mang lại tăng tốc 2-3x với mất mát độ chính xác tối thiểu.

Bộ Cấp phát Memory Pool

• cudaMalloc chậm (~1 ms mỗi lần gọi) vì nó đồng bộ với GPU. Trong một vòng huấn luyện cấp phát buffer tạm thời mỗi lần lặp, điều này tích lũy lên.

• Memory pools (bộ cấp phát cache của PyTorch, memory pool CUDA) cấp phát trước một khối lớn bộ nhớ GPU và cấp phát con từ nó mà không cần gọi hệ thống:

# PyTorch làm điều này tự động — nhưng hiểu tại sao lại quan trọng
# Mỗi torch.empty() tái sử dụng bộ nhớ từ pool, không cần cudaMalloc
temp = torch.empty(1024, 1024, device='cuda')  # micro-giây, không phải mili-giây

• Đây là lý do torch.cuda.memory_allocated() và torch.cuda.max_memory_allocated() của PyTorch khác nhau: allocated là lượng đang dùng, max là đỉnh (pool có thể giữ nhiều hơn lượng đang dùng).

Profile-Guided Optimisation

• Đừng tối ưu mù quáng. Profile trước, xác định nút thắt, tối ưu nó, và profile lại. Mô hình roofline (file 01) cho bạn biết nút thắt là bộ nhớ hay tính toán:

  • Memory-bound (cường độ tính toán thấp): tối ưu bố cục dữ liệu (SoA), fuse kernel, dùng độ chính xác thấp hơn, prefetch.
  • Compute-bound (cường độ tính toán cao): dùng Tensor Core, tăng tính song song, dùng các chỉ thị nhanh hơn (FMA).
  • Latency-bound (không đủ song song): tăng occupancy, giảm sử dụng thanh ghi, khởi chạy nhiều luồng hơn.

• Hầu hết khối lượng công việc ML đều memory-bound. Hệ quả đáng ngạc nhiên: một GPU nhanh hơn (nhiều FLOPS hơn) thường không giúp ích. Bộ nhớ nhanh hơn (HBM3 vs HBM2e) giúp ích nhiều hơn. Đây là lý do việc nâng cấp A100→H100 không chỉ về FLOPS — H100 cũng có 2x băng thông bộ nhớ.

Các Thế hệ GPU NVIDIA

Thế hệ	Năm	Đổi mới Chính	Liên quan đến AI
Pascal (P100)	2016	HBM2, NVLink	GPU học sâu nghiêm túc đầu tiên
Volta (V100)	2017	Tensor Cores (matmul mixed-precision)	Cho phép huấn luyện FP16, 125 TFLOPS TF32
Ampere (A100)	2020	TF32, Sparsity, Tensor Cores thế hệ 3	312 TFLOPS TF32, thưa thớt cấu trúc 2:4
Hopper (H100)	2022	Transformer Engine (FP8), HBM3	989 TFLOPS FP8, chuyển đổi chính xác động
Blackwell (B200)	2024	Transformer Engine thế hệ 2, NVLink 5	2.5 PFLOPS FP4, thiết kế đa chip

• Tensor Cores là các đơn vị nhân ma trận chuyên dụng. Một chỉ thị Tensor Core duy nhất tính phép nhân ma trận 4×4 (D = A×B + C) trong một chu kỳ. Các nhân CUDA thông thường cần 64 phép FMA. Tensor Cores là lý do huấn luyện mixed-precision (tính float16, tích lũy float32) nhanh.

• Transformer Engine (Hopper+) chuyển đổi động giữa độ chính xác FP8 và FP16 trong một lớp duy nhất, chọn độ chính xác cao hơn chỉ khi cần. Điều này tối đa hóa thông lượng mà không hy sinh chất lượng mô hình. Nó được thiết kế riêng cho kiến trúc transformer (attention + MLP), thứ thống trị AI hiện đại.

Bài tập Lập trình (biên dịch với nvcc)

1. Viết một kernel CUDA áp dụng ReLU lên một mảng. Đo thời gian bao gồm cả truyền bộ nhớ. Bài này dạy viết kernel, cudaMalloc/cudaMemcpy, và nút thắt truyền host↔device.

   // task1_relu.cu
   // Biên dịch: nvcc -O3 -o task1_relu task1_relu.cu
   
   #include <stdio.h>
   #include <stdlib.h>
   #include <cuda_runtime.h>
   
   __global__ void relu_kernel(const float* input, float* output, int n) {
       int idx = blockIdx.x * blockDim.x + threadIdx.x;
       if (idx < n) {
           output[idx] = input[idx] > 0.0f ? input[idx] : 0.0f;
       }
   }
   
   int main() {
       const int N = 1 << 24;  // ~16M phần tử
       size_t bytes = N * sizeof(float);
   
       // Cấp phát bộ nhớ host
       float* h_input  = (float*)malloc(bytes);
       float* h_output = (float*)malloc(bytes);
       for (int i = 0; i < N; i++) {
           h_input[i] = (float)(i % 100) - 50.0f;  // hỗn hợp dương và âm
       }
   
       // Cấp phát bộ nhớ device
       float *d_input, *d_output;
       cudaMalloc(&d_input, bytes);
       cudaMalloc(&d_output, bytes);
   
       // Đo thời gian pipeline đầy đủ: copy lên GPU, tính toán, copy về
       cudaEvent_t start, stop;
       cudaEventCreate(&start);
       cudaEventCreate(&stop);
   
       cudaEventRecord(start);
       cudaMemcpy(d_input, h_input, bytes, cudaMemcpyHostToDevice);
   
       int block_size = 256;
       int grid_size = (N + block_size - 1) / block_size;
       relu_kernel<<<grid_size, block_size>>>(d_input, d_output, N);
   
       cudaMemcpy(h_output, d_output, bytes, cudaMemcpyDeviceToHost);
       cudaEventRecord(stop);
       cudaEventSynchronize(stop);
   
       float ms = 0;
       cudaEventElapsedTime(&ms, start, stop);
   
       // Xác minh
       int errors = 0;
       for (int i = 0; i < N; i++) {
           float expected = h_input[i] > 0.0f ? h_input[i] : 0.0f;
           if (h_output[i] != expected) errors++;
       }
   
       printf("Time (including transfers): %.2f ms\n", ms);
       printf("Bandwidth: %.1f GB/s\n", 2.0 * bytes / ms / 1e6);  // đọc + ghi
       printf("Errors: %d / %d\n", errors, N);
   
       cudaFree(d_input); cudaFree(d_output);
       free(h_input); free(h_output);
       return 0;
   }
   

2. Viết một phép nhân ma trận dạng tiled trong CUDA dùng shared memory. So sánh hiệu năng với phiên bản naive (không tiled). Bài này dạy shared memory, __syncthreads, và tại sao tiling quan trọng.

   // task2_matmul.cu
   // Biên dịch: nvcc -O3 -o task2_matmul task2_matmul.cu
   
   #include <stdio.h>
   #include <cuda_runtime.h>
   
   #define TILE 16
   
   // Naive matmul: mỗi luồng tính một phần tử của C
   __global__ void matmul_naive(const float* A, const float* B, float* C, int N) {
       int row = blockIdx.y * blockDim.y + threadIdx.y;
       int col = blockIdx.x * blockDim.x + threadIdx.x;
       if (row < N && col < N) {
           float sum = 0.0f;
           for (int k = 0; k < N; k++) {
               sum += A[row * N + k] * B[k * N + col];
           }
           C[row * N + col] = sum;
       }
   }
   
   // Tiled matmul: dùng shared memory để giảm truy cập global memory
   __global__ void matmul_tiled(const float* A, const float* B, float* C, int N) {
       __shared__ float sA[TILE][TILE];
       __shared__ float sB[TILE][TILE];
   
       int row = blockIdx.y * TILE + threadIdx.y;
       int col = blockIdx.x * TILE + threadIdx.x;
       float sum = 0.0f;
   
       for (int t = 0; t < (N + TILE - 1) / TILE; t++) {
           sA[threadIdx.y][threadIdx.x] = (row < N && t*TILE+threadIdx.x < N)
               ? A[row * N + t*TILE + threadIdx.x] : 0.0f;
           sB[threadIdx.y][threadIdx.x] = (t*TILE+threadIdx.y < N && col < N)
               ? B[(t*TILE + threadIdx.y) * N + col] : 0.0f;
   
           __syncthreads();
           for (int k = 0; k < TILE; k++)
               sum += sA[threadIdx.y][k] * sB[k][threadIdx.x];
           __syncthreads();
       }
   
       if (row < N && col < N)
           C[row * N + col] = sum;
   }
   
   int main() {
       const int N = 1024;
       size_t bytes = N * N * sizeof(float);
   
       float *d_A, *d_B, *d_C;
       cudaMalloc(&d_A, bytes); cudaMalloc(&d_B, bytes); cudaMalloc(&d_C, bytes);
   
       // Khởi tạo với các số 1 (dễ xác minh: C phải toàn N)
       float* h_A = new float[N*N];
       for (int i = 0; i < N*N; i++) h_A[i] = 1.0f;
       cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
       cudaMemcpy(d_B, h_A, bytes, cudaMemcpyHostToDevice);
   
       dim3 block(TILE, TILE);
       dim3 grid((N+TILE-1)/TILE, (N+TILE-1)/TILE);