NCA-AIIO 2: CUDA và lập trình song song trên GPU

Loạt NCA-AIIO (9 phần): 1. Tổng quan & GPU cơ bản · 2. CUDA & lập trình song song · 3. Container & inference serving · 4. H100, H200 & DGX · 5. Networking data center · 6. GPUDirect & NCCL · 7. Blackwell & NVL72 · 8. HGX, DGX & SuperPOD · 9. Vận hành, điện & sizing

Phần 2/9 trong loạt NCA-AIIO: đọc sau Phần 1. Phần này đi xuống tầng lập trình: GPU tổ chức hàng nghìn luồng thực thi ra sao, và các định dạng số quyết định hiệu năng thế nào.

Cập nhật: 2026-10-23

Mục lục

1. Vì sao GPU cần một mô hình lập trình riêng?
2. CUDA là gì?
3. Bộ thư viện quanh CUDA
4. Kiến trúc thực thi: thread, warp, block, grid, SM
5. Host và device
6. Các tầng bộ nhớ trong CUDA
7. Quy trình xử lý của một chương trình CUDA
8. Độ chính xác số: từ FP64 xuống FP4
9. Tensor core khác CUDA core thế nào?
10. Compute capability, CUDA toolkit và driver
11. Quản lý nhiều phiên bản CUDA
12. Thực hành: chương trình CUDA đầu tiên
13. ResNet-50, model tham chiếu hay gặp trong đề thi
Nguồn tham khảo

1. Vì sao GPU cần một mô hình lập trình riêng?

Một vòng lặp cộng hai mảng một triệu phần tử viết bằng C chạy trên CPU mất vài mili giây. Cũng phép tính đó trên GPU nhanh hơn hàng chục lần, nhưng không thể viết bằng vòng lặp quen thuộc.

Nguyên nhân nằm ở chỗ CPU và GPU giải bài toán song song theo hai cách khác hẳn nhau. CPU có 4 tới vài chục nhân mạnh, mỗi nhân tự xử lý một luồng lệnh riêng, có dự đoán rẽ nhánh và cache lớn để chạy nhanh các đoạn code phức tạp. GPU có hàng nghìn nhân đơn giản, và chúng không chạy độc lập: các nhân được gom thành nhóm cùng thực hiện chung một lệnh trên các dữ liệu khác nhau. Mô hình này gọi là SIMT.

Hệ quả cho người lập trình: thay vì viết "lặp qua một triệu phần tử", phải viết "đây là việc cần làm với một phần tử" rồi nói với GPU hãy chạy đoạn đó một triệu lần song song. CUDA tồn tại để diễn đạt được cách tư duy ngược lại đó.

Ba câu hỏi phần này trả lời:

CUDA tổ chức hàng nghìn luồng thực thi theo thứ bậc nào.
Dữ liệu nằm ở đâu trong GPU, và vì sao chuyện đó quyết định hiệu năng.
Các định dạng số FP32, FP16, FP8, FP4 khác nhau ra sao và dùng khi nào.

2. CUDA là gì?

CUDA là viết tắt của Compute Unified Device Architecture, một nền tảng tính toán song song kèm API do NVIDIA phát triển, cho phép dùng GPU cho tính toán tổng quát chứ không chỉ cho đồ hoạ. Khái niệm này thường được gọi là GPGPU (general-purpose computing on GPUs). Hầu hết card NVIDIA hiện nay đều hỗ trợ CUDA.

NVIDIA phát triển CUDA từ khoảng năm 2007, sớm hơn nhiều so với làn sóng deep learning.

Cài đặt qua CUDA Toolkit, trong đó có trình biên dịch NVCC. CUDA thường viết bằng C++, nhưng có wrapper cho các ngôn ngữ cấp cao hơn như Python và .NET, trong đó có cả thư viện mang phong cách NumPy chạy trên nền CUDA.

3. Bộ thư viện quanh CUDA

Phần khiến CUDA khó bị thay thế không phải bản thân API, mà là lớp thư viện dựng sẵn bên trên nó. Mọi framework deep learning lớn đều gọi xuống các thư viện này thay vì tự viết kernel:

Thư viện	Làm gì
cuDNN	Các phép của mạng neural đã tối ưu sâu: convolution chiều tiến và chiều lùi, pooling, normalization, hàm kích hoạt
cuBLAS	Đại số tuyến tính, trọng tâm là nhân ma trận
NCCL	Giao tiếp tập thể giữa nhiều GPU và nhiều node
cuFFT	Biến đổi Fourier nhanh
cuDF	Dataframe chạy trên GPU, thuộc bộ RAPIDS

Một nhóm nghiên cứu muốn chuyển sang phần cứng của hãng khác không chỉ phải viết lại code của mình, mà còn chờ hãng đó có đủ bộ thư viện tương đương với chất lượng tương đương. Đó là ý nghĩa thực tế của cụm "hệ sinh thái CUDA".

4. Kiến trúc thực thi: thread, warp, block, grid, SM

Đây là phần dễ nhớ sai nhất, nên cần đi từ dưới lên và phân biệt rạch ròi đâu là khái niệm phần mềm và đâu là phần cứng.

Phía phần mềm, lập trình viên tổ chức công việc theo ba tầng:

Thread là đơn vị nhỏ nhất, một luồng thực thi chạy kernel trên một phần dữ liệu.
Block (thread block) là một nhóm thread, tối đa 1.024 thread mỗi block. Các thread trong cùng block dùng chung được shared memory và đồng bộ được với nhau.
Grid là tập hợp mọi block của một lần gọi kernel. Grid và block đều khai được theo 1, 2 hoặc 3 chiều cho tiện ánh xạ với hình dạng dữ liệu.

Phía phần cứng:

CUDA core là đơn vị số học đơn giản, thực hiện phép số nguyên và dấu phẩy động. Nhỏ và đơn giản hơn nhiều so với một nhân CPU.
Streaming Multiprocessor (SM) là đơn vị thực thi chính của GPU, chứa CUDA core, tensor core, đơn vị load/store, thanh ghi và shared memory. H100 bản SXM có 132 SM.
Warp là nhóm 32 thread mà phần cứng thực thi đồng loạt cùng một lệnh.

Chỗ nối giữa hai phía: khi một block được giao cho một SM, phần cứng tự cắt block đó thành các warp 32 thread liên tiếp. Một block 256 thread sinh ra 8 warp. Một block 100 thread sinh ra 4 warp, trong đó warp cuối chỉ có 4 thread hoạt động và 28 thread chạy không, lãng phí. Vì vậy kích thước block nên là bội số của 32.

Block không phải là warp. Block là khái niệm do lập trình viên khai báo và có thể lên tới 1.024 thread; warp là đơn vị cố định 32 thread do phần cứng tạo ra. Nhầm hai thứ này dẫn tới hiểu sai toàn bộ phần còn lại.

💡 Hình dung như thế này: GPU giống một nhà máy lớn. Nhà máy có nhiều phân xưởng, đó là các SM. Quản đốc chia công việc thành các tổ, mỗi tổ tối đa 1.024 người, đó là block. Nhưng dây chuyền trong xưởng chỉ vận hành theo kíp đúng 32 người thao tác đồng loạt cùng một động tác, đó là warp, và kíp là thứ do dây chuyền quy định chứ quản đốc không đổi được. Giao một tổ 100 người thì dây chuyền vẫn phải chạy 4 kíp, kíp cuối chỉ có 4 người làm còn 28 chỗ bỏ trống.

Khi một chương trình CUDA đã biên dịch chạy lên, CUDA runtime quyết định phân bổ các block cho các SM, rải đều theo số SM mà GPU đang có. Lập trình viên không chỉ định block nào chạy trên SM nào, và cũng không nên phụ thuộc vào thứ tự thực thi giữa các block.

5. Host và device

Hai thuật ngữ xuất hiện khắp tài liệu CUDA:

Host là CPU của hệ thống, và host memory là RAM thường.
Device là GPU, và device memory là bộ nhớ HBM hoặc GDDR trên card.

Phần lớn lỗi của người mới học CUDA đến từ việc nhầm con trỏ của hai bên: truyền một con trỏ host vào kernel chạy trên device, hoặc quên chép dữ liệu qua lại.

Công cụ gỡ lỗi đáng biết khi làm việc với CUDA gồm Nsight, cuda-gdb và compute-sanitizer. Các công cụ này không vào đề thi nhưng nên có trong bộ đồ nghề.

6. Các tầng bộ nhớ trong CUDA

GPU có một hệ phân cấp bộ nhớ riêng, và biết dữ liệu đang nằm ở tầng nào quyết định hiệu năng nhiều hơn bất cứ tối ưu nào khác.

Tầng	Phạm vi	Tốc độ và dung lượng
Register	Riêng từng thread	Nhanh nhất, rất nhỏ
Shared memory / L1	Dùng chung trong một block, nằm trên chip	Rất nhanh, nhỏ
Read-only (constant, texture)	Chỉ đọc trong kernel, mỗi SM một bản cache	Rất nhanh, rất nhỏ
L2 cache	Dùng chung giữa các SM	Nhanh vừa
Global memory	Toàn GPU, chính là HBM	Chậm nhất trong nhóm, lớn nhất
Local memory	Riêng từng thread nhưng thực chất nằm trong global memory	Chậm, xuất hiện khi thanh ghi không đủ chỗ

Đặc điểm của local memory hay gây bất ngờ: tên gọi nghe như bộ nhớ cục bộ nhanh, nhưng khi một thread cần nhiều biến hơn số thanh ghi khả dụng, phần dư bị đẩy xuống global memory. Hiện tượng này gọi là register spilling và làm hiệu năng tụt mạnh.

Hệ phân cấp này song song với hệ phân cấp L1, L2, L3 và DRAM của CPU đã nhắc ở phần đầu loạt bài. Khác biệt lớn nhất: shared memory của GPU do lập trình viên quản lý tường minh, còn cache của CPU do phần cứng tự quyết.

💡 Hình dung như thế này: register giống đồ nghề cầm sẵn trên tay người thợ, lấy ra dùng tức thì nhưng chỉ cầm được vài món. Shared memory là bàn làm việc chung của cả tổ, ai trong tổ cũng với tới được. L2 là kệ dụng cụ chung của cả xưởng. Global memory là kho ở cuối nhà máy: chứa được mọi thứ nhưng mỗi lần cần gì phải cử người đi lấy, và trong lúc chờ thì cả tổ ngồi không.

7. Quy trình xử lý của một chương trình CUDA

Bốn bước kinh điển:

Chép dữ liệu từ host memory sang device memory.
CPU khởi động kernel trên GPU.
Các CUDA core thực thi kernel song song, đọc và ghi trên device memory.
Chép kết quả từ device memory về host memory.

Bước 1 và bước 4 chính là chỗ mà GPUDirect nhắm tới, chủ đề của phần sau trong loạt bài. Với dữ liệu lớn, thời gian chép qua lại có thể vượt cả thời gian tính toán, khiến tăng tốc trên giấy không thành hiện thực.

Chứng chỉ này không đòi hỏi viết được CUDA ở mức chi tiết, nhưng các khái niệm trên xuất hiện dày đặc trong hệ sinh thái NVIDIA nên vẫn cần nắm.

8. Độ chính xác số: từ FP64 xuống FP4

Phần data sheet của mọi GPU data center đều liệt kê hiệu năng theo từng định dạng số, và hiểu bảng đó là điều kiện để đọc được data sheet.

Định dạng	Số bit	Byte mỗi giá trị	Dùng cho
FP64	64	8	Tính toán khoa học, mô phỏng cần độ chính xác cao
FP32	32	4	Độ chính xác chuẩn, mặc định cũ của deep learning
TF32	19 bit hữu dụng, vỏ 32 bit	4	Định dạng riêng của Tensor Core, thay thế FP32 gần như trong suốt
FP16	16	2	Training và inference, dải giá trị hẹp
BF16	16	2	Như FP16 nhưng giữ dải giá trị rộng của FP32, ít tràn số hơn
FP8	8	1	Hopper trở đi, cần Transformer Engine
FP4 / NVFP4	4	0,5	Blackwell trở đi, cần kỹ thuật chia tỉ lệ theo khối

Khác biệt giữa FP16 và BF16 đáng nắm vì hay vào đề: cả hai cùng 16 bit, nhưng FP16 dành nhiều bit hơn cho phần định trị nên chính xác hơn trong dải hẹp, còn BF16 dành nhiều bit hơn cho phần mũ nên bao được dải giá trị rộng như FP32. Trong huấn luyện, dải giá trị quan trọng hơn độ mịn, nên BF16 thường ổn định hơn và không cần kỹ thuật loss scaling.

Mỗi lần giảm một nửa số bit mang lại hai lợi ích cùng lúc: Tensor Core thực hiện được gấp đôi số phép tính, và lượng dữ liệu phải chảy qua bộ nhớ giảm một nửa. Đó là lý do chuỗi tiến hoá FP32, FP16, FP8, FP4 lặp lại qua từng đời kiến trúc.

Quy đổi nhanh khi tính bộ nhớ cho model: nhân số tham số với số byte ở cột thứ ba. Model 70 tỉ tham số ở FP16 cần 140 GB, ở FP8 cần 70 GB, ở FP4 cần 35 GB.

9. Tensor core khác CUDA core thế nào?

Hai loại nhân này cùng nằm trong một SM nhưng làm việc khác nhau.

CUDA core thực hiện các phép vô hướng: cộng hai số, nhân hai số, so sánh. Mỗi CUDA core xử lý một thread.

Tensor core thực hiện trọn một phép nhân ma trận cộng dồn trong một lệnh: lấy hai ma trận nhỏ nhân với nhau rồi cộng vào ma trận thứ ba. Phép toán này chiếm phần lớn khối lượng tính toán của mạng neural, nên gom nó vào một lệnh phần cứng cho tốc độ cao hơn nhiều lần so với ghép từ các phép vô hướng.

	CUDA core	Tensor core
Phép toán	Vô hướng, từng số một	Nhân ma trận cộng dồn
Độ chính xác hỗ trợ	FP64, FP32, INT	TF32, FP16, BF16, FP8, FP4 tuỳ thế hệ
Xuất hiện từ	Mọi đời	Volta trở đi

Chi tiết này giải thích con số trong data sheet: H100 SXM cho 67 TFLOPS ở FP32 bằng CUDA core, nhưng 989 TFLOPS ở FP16 Tensor Core khi tính dense (1.979 TFLOPS nếu lấy con số có sparsity như trang sản phẩm công bố), chênh khoảng mười lăm lần. Hai con số đo hai thứ khác nhau, không so trực tiếp được.

10. Compute capability, CUDA toolkit và driver

Ba con số phiên bản dễ nhầm lẫn, và nhầm chúng là nguyên nhân rất phổ biến của lỗi khi triển khai.

Compute capability là số hiệu kiến trúc của GPU, ví dụ 8.0 cho Ampere, 9.0 cho Hopper. Nó cố định theo phần cứng, không nâng cấp được. Con số này quyết định GPU chạy được những tính năng CUDA nào, ví dụ FP8 của Transformer Engine cần từ 9.0 trở lên.

CUDA Toolkit là bộ công cụ phát triển, gồm NVCC và các thư viện. Phiên bản toolkit quyết định code viết được dùng những API nào.

Driver là phần chạy trên hệ điều hành. Mỗi phiên bản driver hỗ trợ tới một phiên bản CUDA nhất định.

Quy tắc tương thích: driver phải mới hơn hoặc bằng mức mà CUDA toolkit yêu cầu. Driver cũ kèm toolkit mới sẽ báo lỗi, trong khi driver mới chạy được code biên dịch bằng toolkit cũ.

# Phiên bản CUDA mà driver hỗ trợ, xem ở góc trên bên phải
nvidia-smi
+-----------------------------------------------------------------------------+
| NVIDIA-SMI 560.35.03    Driver Version: 560.35.03    CUDA Version: 12.6      |
+-----------------------------------------------------------------------------+

# Phiên bản của trình biên dịch trong CUDA Toolkit đang cài
nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Cuda compilation tools, release 12.4, V12.4.131

Hai con số trên không bắt buộc phải trùng nhau. Ở ví dụ này driver hỗ trợ tới 12.6 còn toolkit là 12.4, hoàn toàn hợp lệ. Thấy chúng khác nhau rồi kết luận có vấn đề là hiểu sai.

11. Quản lý nhiều phiên bản CUDA

Làm việc với nhiều phiên bản CUDA cùng lúc là chuyện bình thường, vì các dự án khác nhau khoá vào các phiên bản khác nhau. Một số tính năng bị loại bỏ hoặc đổi tên giữa các bản, nên có chương trình chỉ chạy đúng với một phiên bản rất cụ thể.

CUDA không có trình quản lý phiên bản riêng kiểu như nvm của Node hay pyenv của Python. Cách làm phổ biến là cô lập môi trường:

Conda hoặc Mamba. Mamba là bản viết lại conda bằng C++, giữ nguyên cách dùng nhưng giải phụ thuộc nhanh hơn nhiều nhờ thư viện libsolv và tải song song. Mỗi environment khai một phiên bản cudatoolkit riêng, chuyển qua lại bằng lệnh kích hoạt environment. Bản rút gọn liên kết tĩnh của nó tên là micromamba.
Container. Mỗi image mang sẵn đúng phiên bản CUDA cần dùng, chỉ driver nằm ở host. Đây là cách sạch nhất và là lý do các image từ NGC phổ biến tới vậy.

CUDA SDK cung cấp hai tầng API: CUDA driver API ở mức thấp, kiểm soát chi tiết hơn, và CUDA runtime API ở mức cao, tiện hơn cho phần lớn trường hợp.

12. Thực hành: chương trình CUDA đầu tiên

Cài CUDA Toolkit bản mới nhất, trên Windows có thể dùng qua WSL. File mã nguồn CUDA có phần mở rộng .cu.

Một chương trình cộng hai vector, bài hello world chuẩn của CUDA:

// vector_add.cu
#include <stdio.h>

// Kernel chạy trên device, mỗi thread lo đúng một phần tử
__global__ void vector_add(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) {
    int n = 1 << 20;                 // hơn một triệu phần tử
    size_t bytes = n * sizeof(float);
    float *h_a = (float *)malloc(bytes);   // bộ nhớ phía host
    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; }

    float *d_a, *d_b, *d_c;          // con trỏ phía device
    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);
    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);   // bước 1
    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);

    int threads = 256;               // bội số của 32
    int blocks  = (n + threads - 1) / threads;
    vector_add<<<blocks, threads>>>(d_a, d_b, d_c, n);     // bước 2 và 3

    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);   // bước 4
    printf("c[0] = %f, c[%d] = %f\n", h_c[0], n - 1, h_c[n - 1]);
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    free(h_a); free(h_b); free(h_c);
    return 0;
}

Biên dịch và chạy:

nvcc vector_add.cu -o vector_add
./vector_add
c[0] = 3.000000, c[1048575] = 3.000000

Bốn chi tiết đáng chú ý trong đoạn code trên. Thứ nhất, __global__ đánh dấu hàm chạy trên device nhưng gọi từ host. Thứ hai, công thức blockIdx.x * blockDim.x + threadIdx.x là cách mỗi thread tự tính xem mình phụ trách phần tử nào, mẫu hình lặp lại trong mọi chương trình CUDA. Thứ ba, câu lệnh if (i < n) cần thiết vì số thread khởi chạy luôn là bội số của 256 nên thường nhiều hơn số phần tử. Riêng ví dụ này n chia hết cho 256 nên không thread nào bị chặn, nhưng bỏ if đi là code sai với mọi n khác. Thứ tư, cú pháp ba dấu ngoặc nhọn <<<blocks, threads>>> là cách khai cấu hình grid và block khi gọi kernel.

Bốn bước của quy trình xử lý ở mục 7 xuất hiện đúng theo thứ tự trong hàm main.

13. ResNet-50, model tham chiếu hay gặp trong đề thi

ResNet-50 là một backbone vision model phổ biến, xuất hiện nhiều trong đề thi và trong các ví dụ của hệ sinh thái NVIDIA.

Đây là một convolutional neural network dùng cho phân loại ảnh, thuộc họ ResNet (residual network) cùng với ResNet-18 và ResNet-34. Con số 50 là số lớp có trọng số. Model phát hành năm 2015, dung lượng chỉ khoảng 100 MB.

Vì nhỏ, ResNet-50 chạy được trên phần cứng khá cũ, từ CPU Intel thế hệ 4 tới điện thoại đời 2017 hay GPU GTX 700 series. Chính tính phổ quát và nhẹ nhàng đó khiến nó được dùng làm chuẩn so sánh suốt nhiều năm, và một model cũ không đồng nghĩa với một model tệ.

Điểm liên quan tới hạ tầng: 100 MB nằm gọn trong bộ nhớ của bất kỳ GPU nào, nên inference với ResNet-50 không bao giờ là bài toán bộ nhớ. Đối chiếu với một LLM 70 tỉ tham số cần 140 GB ở FP16, chênh lệch hơn một nghìn lần. Đó là lý do hạ tầng cho vision model và hạ tầng cho LLM là hai bài toán khác hẳn nhau.

ResNet-50 có sẵn bản ở định dạng ONNX trong ONNX model zoo, định dạng model trung lập framework được trình bày ở phần sau. Muốn xem model này xử lý qua từng lớp thế nào, công cụ trực quan hoá Netron nạp được file ONNX và hiển thị toàn bộ đồ thị ngay trên trình duyệt.

Nguồn tham khảo

Nguồn gốc: NVIDIA-Certified Associate AI Infrastructure and Operations (NCA AIIO) Free Study Course, kênh freeCodeCamp.org.

Đã fact-check (truy cập 2026-10-23):

Thứ bậc grid, block, thread và giới hạn 1.024 thread mỗi block; warp gồm 32 thread do phần cứng cắt ra từ block theo các dải thread ID liên tiếp, nên block T thread sinh ra ceil(T/32) warp; grid và block khai được 1, 2 hoặc 3 chiều: CUDA Programming Guide, Compute Capabilities, CUDA Programming Guide, Writing SIMT Kernels
Đính chính so với bản trước của bài: block và warp là hai khái niệm khác nhau. Bản cũ mô tả block là "một tập hợp thread, thường gồm 32 thread" rồi đồng nhất nó với warp. Đúng ra block do lập trình viên khai, tối đa 1.024 thread, còn warp là đơn vị cố định 32 thread do phần cứng tạo.
Đính chính so với bản trước của bài: Mamba là bản viết lại conda bằng C++ dùng libsolv để giải phụ thuộc nhanh hơn và tải song song, không phải "một phiên bản gọn nhẹ hơn của Docker". Bản liên kết tĩnh của nó tên micromamba. Nguồn: mamba-org/mamba
Số SM của H100 SXM (132 SM): NVIDIA Hopper Architecture In-Depth
Hiệu năng FP32 (67 TFLOPS) và FP16 Tensor Core của H100 SXM (989 TFLOPS bản dense, 1.979 TFLOPS là con số có sparsity mà trang sản phẩm công bố): NVIDIA H100 Tensor Core GPU
Transformer Engine cùng định dạng FP8 xuất hiện từ Hopper, tức compute capability 9.0 trở lên: NVIDIA Hopper Architecture In-Depth
Output của nvidia-smi và nvcc --version trong bài là minh hoạ định dạng, số liệu thay đổi theo từng máy.

Mục lục​

1. Vì sao GPU cần một mô hình lập trình riêng?​

2. CUDA là gì?​

3. Bộ thư viện quanh CUDA​

4. Kiến trúc thực thi: thread, warp, block, grid, SM​

5. Host và device​

6. Các tầng bộ nhớ trong CUDA​

7. Quy trình xử lý của một chương trình CUDA​

8. Độ chính xác số: từ FP64 xuống FP4​

9. Tensor core khác CUDA core thế nào?​

10. Compute capability, CUDA toolkit và driver​

11. Quản lý nhiều phiên bản CUDA​

12. Thực hành: chương trình CUDA đầu tiên​

13. ResNet-50, model tham chiếu hay gặp trong đề thi​

Nguồn tham khảo​