Giải mã CUDA P1: Khám phá bí ẩn Kernel, Cấu trúc Luồng và Cách GPU Hoạt Động!

Lê Lân

08/07/2025

CUDA Deep Dive: Giải Mã Kernels, Hệ Thống Luồng và Mô Hình Thực Thi GPU

Mở Đầu

CUDA – công nghệ huyền thoại của NVIDIA – đã mở ra một kỷ nguyên mới cho lập trình song song với sức mạnh đồ sộ của GPU.

Bạn đã từng tìm hiểu sơ qua về CUDA và biết rằng nó mở rộng ngôn ngữ C để khai thác sức mạnh của các bộ xử lý đồ họa. Nhưng để thực sự tận dụng được sức mạnh đó, bạn cần hiểu sâu về mô hình thực thi, cách tổ chức luồng, cũng như quản lý bộ nhớ trong GPU. Bài viết này sẽ hướng dẫn bạn qua những kiến thức cốt lõi đó dựa trên các khái niệm trong chương 3 của cuốn Programming Massively Parallel Processors, cùng với nhiều ví dụ thực tế về cách CUDA quản lý và điều phối các luồng song song trên GPU.

Chúng ta sẽ cùng khám phá từ kiến trúc đa dạng của CPU và GPU đến cách CUDA xác định không gian thực thi cho các hàm, chi tiết mô hình thực thi từ lưới (grid), khối (block), đến nhóm luồng (warp) và từng luồng (thread). Cuối cùng, bạn sẽ hiểu rõ hơn về cách thiết kế kernel hiệu quả, xử lý bộ nhớ và nâng cao hiệu suất GPU.

The Dichotomy: Host (CPU) vs. Device (GPU) Architectures

Hai miền tính toán tách biệt

CUDA sử dụng một mô hình tính toán kết hợp giữa:

Host (CPU): Đơn vị xử lý trung tâm của hệ thống, quản lý tài nguyên, điều phối chương trình tổng thể. CPU có bộ nhớ DRAM riêng.

Device (GPU): Bộ xử lý đồ họa đa luồng, với hàng trăm, hàng ngàn nhân CUDA chạy song song, cùng bộ nhớ riêng (thường là GDDR).

Quản lý dữ liệu giữa Host và Device

Dữ liệu phải được sao chép rõ ràng từ bộ nhớ hệ thống sang bộ nhớ GPU để xử lý, và kết quả phải trả về CPU. Quá trình này chạy qua PCIe bus và có thể trở thành nút thắt cổ chai ảnh hưởng nghiêm trọng hiệu năng.

Việc quản lý chuyển đổi dữ liệu giữa CPU và GPU là điểm then chốt để viết chương trình CUDA hiệu quả.

CUDA C Function Specifiers: Định Nghĩa Không Gian Thực Thi

CUDA bổ sung các từ khóa đặc biệt cho C để xác định nơi các hàm được biên dịch và thực thi.

Từ Khóa	Mô Tả	Nơi Gọi Hàm	Ghi Chú
`__host__`	Hàm chạy trên CPU, được biên dịch cho host	Từ hàm `__host__` hoặc toàn cục	Mặc định nếu không có từ khóa CUDA
`__device__`	Hàm chạy trên GPU, chỉ gọi được từ `__global__` hoặc `__device__`	Từ kernel hoặc hàm device khác	Không hỗ trợ đệ quy và con trỏ hàm cũ
`__global__`	Hàm kernel, điểm vào thực thi GPU, void trả về, chạy bất đồng bộ từ host	Từ host	Gọi bất đồng bộ, cần cấu hình grid/block
`__host__ __device__`	Hàm được biên dịch cho cả host và device, hỗ trợ tái sử dụng mã	Từ host và device	Tiện lợi cho các hàm tiện ích chung

Ví dụ hàm dùng `host device`

__host__ __device__
int clamp_value(int val, int min_val, int max_val) {
  if (val < min_val) return min_val;
  if (val > max_val) return max_val;
  return val;
}

Lưu ý: Hàm __device__ không hỗ trợ đệ quy và có những hạn chế về biến tĩnh cũng như con trỏ hàm.

Mô Hình Thực Thi GPU: Grid, Block, Warp và Thread

Cấu trúc phân cấp của CUDA

Khi một kernel __global__ được khởi chạy, nó tạo ra một lưới (grid) gồm nhiều khối (block), mỗi block chứa nhiều luồng (thread).

Threads – Đơn vị thực thi cơ bản

Mỗi luồng chạy đoạn mã kernel

Rất nhẹ, được đánh chỉ số trong block qua biến threadIdx (3 chiều: x, y, z)

Warps – Nhóm luồng

Nhóm 32 luồng chạy đồng thời theo mô hình SIMT (Single Instruction, Multiple Thread)

Divergence (phân nhánh khác nhau bên trong warp) làm giảm hiệu năng do phải xử lý tuần tự

Blocks – Nhóm luồng hợp tác

Tập hợp luồng tối đa 1024/thread block, có thể 1D, 2D hoặc 3D

Có thể chia sẻ dữ liệu trong bộ nhớ __shared__ tốc độ cao nội bộ block

Đồng bộ hóa luồng trong block bằng hàm __syncthreads()

Mỗi block chạy trên một Streaming Multiprocessor (SM)

Grid – Tập hợp các block

Chứa toàn bộ block được chạy bởi kernel gọi tới

Các block hoạt động độc lập, không đồng bộ trực tiếp mà qua bộ nhớ toàn cục hoặc bằng cách khởi chạy kernel mới

Cách tính ID toàn cục của luồng

Để mỗi luồng xử lý dữ liệu riêng, index toàn cục được tính như sau:

// 1D grid và 1D block
int globalThreadId_x = blockIdx.x * blockDim.x + threadIdx.x;

// 2D grid và 2D block
int gx = blockIdx.x * blockDim.x + threadIdx.x;
int gy = blockIdx.y * blockDim.y + threadIdx.y;

Tính toán chính xác global thread ID là bước cần thiết để các luồng truy cập và cập nhật dữ liệu chính xác.

Ví Dụ Thực Tế: Nhân Ma Trận Với Nhiều Blocks

Để nhân hai ma trận M (chiều cao: height_M, chiều rộng: width_M) và N (chiều cao: width_M, chiều rộng: width_N), ta tính phần tử P[i,j] của kết quả:

Kernel CUDA minh họa

__global__
void matrixMulKernel(float *P_d, const float *M_d, const float *N_d,
                     int P_height, int P_width, int M_width) {
  int row = blockIdx.y * blockDim.y + threadIdx.y;
  int col = blockIdx.x * blockDim.x + threadIdx.x;

  if (row < P_height && col < P_width) {
    float p_value = 0.0f;
    for (int k = 0; k < M_width; ++k) {
      p_value += M_d[row * M_width + k] * N_d[k * P_width + col];
    }
    P_d[row * P_width + col] = p_value;
  }
}

Mỗi luồng tính một phần tử P[row][col]. Việc chia nhỏ tác vụ qua block và luồng giúp tận dụng tối đa sức mạnh tính toán song song trên GPU.

Kernel Launch Configuration: Cấu Hình Khởi Chạy Kernel

Cấu trúc gọi kernel như sau:

kernelName<<<dimGrid, dimBlock, sharedMemSize, stream>>>(args);

dimGrid : Kích thước grid – số block theo chiều x, y, z (dim3)

dimBlock : Kích thước block – số luồng/block theo chiều x, y, z (dim3)

sharedMemSize : (Tùy chọn) Kích thước bộ nhớ shared động

stream : (Tùy chọn) CUDA stream để điều phối đa nhiệm

Ví dụ:

dim3 dimGrid(1, 1);
dim3 dimBlock(16, 16);
matrixMulKernel<<<dimGrid, dimBlock>>>(d_Pd, d_Md, d_Nd, WIDTH);

dimGrid và dimBlock tương ứng được truy cập trong kernel qua gridDim và blockDim .

Các Hàm Runtime API Thường Dùng

Hàm	Mục Đích
`cudaMalloc(void **devPtr, size_t size)`	Cấp phát bộ nhớ toàn cục trên thiết bị GPU
`cudaMemcpy(void dst, const void src, size_t count, cudaMemcpyKind kind)`	Sao chép bộ nhớ giữa host và device
`cudaFree(void *devPtr)`	Giải phóng bộ nhớ trên device
`cudaDeviceSynchronize()`	Đồng bộ host đợi GPU hoàn thành các lệnh

Sử dụng cudaDeviceSynchronize() rất cần thiết để đảm bảo rằng các kernel và sao chép bộ nhớ đã hoàn thành khi đo thời gian hay kiểm tra kết quả.

Tính Toán Hiệu Năng và Một Số Lưu Ý

Occupancy

Tỷ lệ warps hoạt động trên SM so với tối đa có thể. Occupancy cao giúp GPU ẩn độ trễ truy cập bộ nhớ bằng cách chuyển đổi luồng sẵn sàng.

Memory Coalescing

Hoạt động truy cập bộ nhớ global của warp cần được lập trình sao cho các truy cập là liên tục và căn chỉnh để GPU có thể gom các giao dịch bộ nhớ lại thành một hoặc một vài truy cập hiệu quả.

Shared Memory Banking

Bộ nhớ chia nhỏ thành "bank". Nếu các luồng truy cập cùng bank đồng thời sẽ gây leo thang xung đột (bank conflicts), làm giảm băng thông thực tế. Lập trình viên cần thiết kế tránh xung đột này.

Thiết kế truy cập bộ nhớ và cấu hình block/thread hợp lý quyết định thành bại của một kernel CUDA hiệu quả.

Các Biến Tổng Hợp Nâng Cao (Đề Cập Nhanh)

Ngoài threadIdx , blockIdx , blockDim , gridDim , CUDA còn cung cấp:

warpSize : Kích thước một warp, thường là 32 threads, hữu ích để viết các idioms lập trình warp-level.

Kết Luận

Qua bài viết này, bạn đã có một cái nhìn sâu sắc về:

Sự khác biệt và tương tác giữa CPU (host) và GPU (device) trong CUDA

Các loại hàm CUDA và từ khóa đặc biệt để xác định không gian thực thi

Mô hình phân cấp thread từ grid, block, warp đến thread, cùng cách hoạt động và cấp phát tài nguyên

Cách cấu hình kernel và các hàm runtime API thiết yếu để quản lý bộ nhớ và đồng bộ

Các nguyên tắc tối ưu hiệu suất về occupancy, memory coalescing và tránh shared memory bank conflicts

Để thành thạo CUDA, bạn cần thực hành viết kernel, đo đạc hiệu năng và tinh chỉnh cấu hình luồng khối, cũng như thiết kế truy cập bộ nhớ hợp lý.

Hãy bắt đầu bằng việc triển khai các ví dụ đơn giản, sau đó tăng dần độ phức tạp với các kỹ thuật chia khối, ghi nhớ tối ưu hóa truy cập bộ nhớ, và đồng bộ hóa luồng!

Tham Khảo

J. Sanders, E. Kandrot. Programming Massively Parallel Processors: A Hands-on Approach, 3rd Edition, Morgan Kaufmann, 2019. (Link PDF)

NVIDIA CUDA Toolkit Documentation. https://docs.nvidia.com/cuda/

CUDA C++ Programming Guide, NVIDIA Developer Zone.

D. Kirk, W. Hwu. Programming Massively Parallel Processors, 2nd Edition, Elsevier, 2016.

June 15, 2024