Chương 2: Mô hình Lập trình GPU
Các Đơn vị Xử lý Đồ họa (GPUs) đã phát triển từ các bộ gia tốc đồ họa cố định sang các động cơ máy tính song song, có thể lập trình được, có khả năng gia tốc một loạt các ứng dụng rộng rãi. Để cho phép các lập trình viên khai thác hiệu quả sự song song hóa khổng lồ trong các GPU, một số mô hình lập trình song song và API đã được phát triển, như NVIDIA CUDA, OpenCL và DirectCompute. Những mô hình lập trình này cung cấp các trừu tượng cho phép lập trình viên diễn đạt sự song song hóa trong các ứng dụng của họ trong khi che giấu các chi tiết cấp thấp của phần cứng GPU.
Trong chương này, chúng ta sẽ khám phá các khái niệm và nguyên lý chính đằng sau các mô hình lập trình song song cho GPU, tập trung vào mô hình thực thi, kiến trúc tập lệnh GPU (ISAs), ISA của GPU NVIDIA và ISA Graphics Core Next (GCN) của AMD. Chúng tôi cũng sẽ cung cấp các ví dụ để minh họa cách các khái niệm này được áp dụng trong thực tế.
Mô hình Thực thi
Mô hình thực thi của các mô hình lập trình GPU hiện đại dựa trên khái niệm về lõi, là các chức năng được thực thi song song bởi một số lượng lớn các luồng trên GPU. Khi khởi chạy một lõi, lập trình viên chỉ định số lượng luồng sẽ được tạo và cách chúng được tổ chức thành một cấu trúc phân cấp của lưới, khối (hoặc mảng luồng hợp tác - CTA) và các luồng riêng lẻ.
- Một lưới biểu diễn toàn bộ không gian vấn đề và bao gồm một hoặc nhiều khối.
- Một khối là một nhóm các luồng có thể hợp tác và đồng bộ hóa với nhau thông qua bộ nhớ chia sẻ và các rào cản. Các luồng trong một khối được thực thi trên cùng một lõi GPU (gọi là đa bộ xử lý kết hợp hoặc đơn vị tính toán).
- Mỗi luồng có một ID duy nhất trong khối và lưới của nó, có thể được sử dụng để tính toán các địa chỉ bộ nhớ và đưa ra các quyết định luồng điều khiển.
Cấu trúc phân cấp này cho phép lập trình viên diễn đạt cả song song dữ liệu (nơi cùng một thao tác được áp dụng cho nhiều yếu tố dữ liệu) và song song nhiệm vụ (nơi các nhiệm vụ khác nhau được thực hiện song song).
Hình...Here is the Vietnamese translation of the provided Markdown file, with the code comments translated:
Hình 2.1 minh họa cấu trúc chuỗi luồng trong mô hình thực thi GPU.
Lưới
________________
/ / / / /
/ / / / /
/ / / / /
/ / / / /
/__/__/__/__/__/
| | | |
| | Khối |
| | | |
Luồng Luồng ...Hình 2.1: Cấu trúc chuỗi luồng trong mô hình thực thi GPU.
Thực thi SIMT
Các mô hình lập trình GPU như CUDA và OpenCL tuân theo mô hình thực thi Lệnh Đơn, Luồng Đa (SIMT). Trong mô hình SIMT, các luồng được thực thi theo nhóm gọi là warp (thuật ngữ của NVIDIA) hoặc wavefront (thuật ngữ của AMD). Tất cả các luồng trong một warp thực thi cùng một lệnh tại cùng một thời điểm, nhưng mỗi luồng hoạt động trên dữ liệu khác nhau.
Tuy nhiên, khác với mô hình Lệnh Đơn, Dữ liệu Đa (SIMD) truyền thống, nơi tất cả các phần tử xử lý thực thi đồng bộ, SIMT cho phép các luồng có các đường dẫn thực thi độc lập và phân kỳ tại các hướng dẫn phân nhánh. Khi một warp gặp một hướng dẫn phân nhánh, phần cứng GPU sẽ đánh giá điều kiện phân nhánh cho từng luồng trong warp. Nếu tất cả các luồng lấy cùng một đường (hội tụ), warp sẽ tiếp tục thực thi bình thường. Nếu một số luồng lấy các đường khác nhau (phân kỳ), warp sẽ được chia thành hai hoặc nhiều sub-warp, mỗi sub-warp sẽ theo một đường khác nhau. Phần cứng GPU sẽ tuần tự thực thi các đường phân kỳ này, che lấp các luồng không hoạt động trong mỗi sub-warp. Khi tất cả các đường hoàn thành, các sub-warp sẽ hội tụ lại và tiếp tục thực thi đồng bộ.
Hình 2.2 minh họa việc thực thi SIMT với các luồng phân kỳ.
Warp
________________
/ / / / /
/ / / / /
/ / / / /
| | |
| Nhánh |
| | |
/ \ / \ / \
/ X \ \
/ / \ \ \
/ \ \
/ \ \
/ \ \
/ \ \
\
\
\
Hội tụ lạiHình 2.2: Thực thi SIMT với các luồng phân kỳ.
Cơ chế xử lý phân kỳ này cho phép SIMT hỗ trợ các luồng có luồng điều khiển linh hoạt hĐây là bản dịch tiếng Việt của tệp Markdown:
Phân cấp bộ nhớ
GPU có một phân cấp bộ nhớ phức tạp để hỗ trợ các yêu cầu băng thông cao và độ trễ thấp của các tải công việc song song. Phân cấp bộ nhớ thường bao gồm:
- Bộ nhớ toàn cục: Không gian bộ nhớ lớn nhất nhưng chậm nhất, có thể truy cập được bởi tất cả các luồng trong một hạt nhân. Bộ nhớ toàn cục thường được thực hiện bằng bộ nhớ GDDR hoặc HBM có băng thông cao.
- Bộ nhớ chia sẻ: Một không gian bộ nhớ nhanh, được tích hợp trên chip, được chia sẻ bởi tất cả các luồng trong một khối. Bộ nhớ chia sẻ được sử dụng để giao tiếp giữa các luồng và chia sẻ dữ liệu trong một khối.
- Bộ nhớ hằng số: Một không gian bộ nhớ chỉ đọc được sử dụng để phát sóng dữ liệu chỉ đọc đến tất cả các luồng.
- Bộ nhớ kết cấu: Một không gian bộ nhớ chỉ đọc được tối ưu hóa cho tính định vị không gian và được truy cập thông qua bộ nhớ cache kết cấu. Bộ nhớ kết cấu thường được sử dụng nhiều hơn trong các tải công việc đồ họa.
- Bộ nhớ cục bộ: Một không gian bộ nhớ riêng tư cho mỗi luồng, được sử dụng để tràn đăng ký và các cấu trúc dữ liệu lớn. Bộ nhớ cục bộ thường được ánh xạ đến bộ nhớ toàn cục.
Sử dụng hiệu quả phân cấp bộ nhớ là rất quan trọng để đạt hiệu suất cao trên GPU. Lập trình viên nên cố gắng tối đa hóa việc sử dụng bộ nhớ chia sẻ và giảm thiểu truy cập vào bộ nhớ toàn cục để giảm độ trễ bộ nhớ và các điểm nghẽn băng thông.
Hình 2.3 minh họa phân cấp bộ nhớ của GPU.
____________
| |
| Global |
| Memory |
____________
|
____________
| |
| Constant |
| Memory |
____________
|
____________
| |
| Texture |
| Memory |
____________
|
|
____________
| |
| Shared |
| Memory |
____________
|
____________
| |
| Local |
| Memory |
____________HìnhĐây là bản dịch tiếng Việt của tệp Markdown:
Kiến trúc Tập lệnh GPU
Các kiến trúc tập lệnh GPU (Instruction Set Architecture - ISA) xác định giao diện cấp thấp giữa phần mềm và phần cứng. Chúng chỉ định các lệnh, thanh ghi và chế độ địa chỉ bộ nhớ được hỗ trợ bởi GPU. Hiểu biết về các ISA GPU là rất cần thiết để phát triển mã GPU hiệu quả và tối ưu hóa hiệu suất.
Trong phần này, chúng ta sẽ khám phá các ISA của hai nhà cung cấp GPU chính: NVIDIA và AMD. Chúng ta sẽ tập trung vào ISA Parallel Thread Execution (PTX) và SASS của NVIDIA, cũng như ISA Graphics Core Next (GCN) của AMD.
Các ISA GPU của NVIDIA
Các GPU NVIDIA hỗ trợ hai cấp độ ISA: PTX (Parallel Thread Execution) và SASS (Streaming ASSembler). PTX là một ISA ảo cung cấp một mục tiêu ổn định cho trình biên dịch CUDA, trong khi SASS là ISA gốc của các GPU NVIDIA.
PTX (Parallel Thread Execution)
PTX là một ISA cấp thấp, ảo thiết kế cho các GPU NVIDIA. Nó giống như LLVM IR hoặc Java bytecode trong việc cung cấp một mục tiêu ổn định và độc lập với kiến trúc cho các trình biên dịch. Các chương trình CUDA thường được biên dịch thành mã PTX, sau đó được dịch sang các hướng dẫn SASS gốc bởi trình điều khiển GPU NVIDIA.
PTX hỗ trợ một loạt các lệnh toán học, bộ nhớ và điều khiển luồng. Nó có một số lượng thanh ghi ảo vô hạn và hỗ trợ điều khiển luồng với cơ chế "predication". PTX cũng cung cấp các lệnh đặc biệt để đồng bộ hóa luồng, thực hiện các thao tác nguyên tử và lấy mẫu texture.
Dưới đây là một ví dụ về mã PTX cho một kernel vector addition đơn giản:
.version 7.0
.target sm_70
.address_size 64
.visible .entry vecAdd(
.param .u64 vecAdd_param_0,
.param .u64 vecAdd_param_1,
.param .u64 vecAdd_param_2,
.param .u32 vecAdd_param_3
)
{
.reg .b32 %r<4>;
.reg .b64 %rd<6>;
ld.param.u64 %rd1, [vecAdd_param_0];
ld.param.u64 %rd2, [vecAdd_param_1];
ld.param.u64 %rd3, [vecAdd_param_2];
ld.param.u32 %r1, [vecAdd_param_3];
cvta.to.global.u64 %rd4, %rd1;
cvta
}
```Đây là bản dịch tiếng Việt của tệp Markdown được cung cấp:
.to.global.u64 %rd5, %rd2; mov.u32 %r2, %tid.x; mul.wide.u32 %rd6, %r2, 4; add.s64 %rd7, %rd4, %rd6; add.s64 %rd8, %rd5, %rd6;
ld.global.f32 %f1, [%rd7]; ld.global.f32 %f2, [%rd8]; add.f32 %f3, %f1, %f2;
cvta.to.global.u64 %rd9, %rd3; add.s64 %rd10, %rd9, %rd6; st.global.f32 [%rd10], %f3;
ret; }
Mã PTX này định nghĩa một hàm kernel `vecAdd` có bốn tham số: các con trỏ đến các vector đầu vào và đầu ra, và kích thước của các vector. Hàm kernel tính toán ID của luồng toàn cục, tải các phần tử tương ứng từ các vector đầu vào, thực hiện phép cộng, và lưu kết quả vào vector đầu ra.
#### SASS (Streaming ASSembler)
SASS là ISA bản gốc của GPU NVIDIA. Đây là một ISA cấp thấp, máy-cụ thể, trực tiếp ánh xạ đến phần cứng GPU. Các lệnh SASS được GPU driver của NVIDIA tạo ra từ mã PTX và thông thường không được hiển thị cho lập trình viên.
Các lệnh SASS được mã hóa ở một định dạng gọn gàng để giảm băng thông bộ nhớ và lượng bộ nhớ cache chỉ thị. Chúng hỗ trợ một loạt các kiểu toán hạng, bao gồm các thanh ghi, các giá trị trực tiếp, và các chế độ địa chỉ khác nhau để truy cập bộ nhớ.
Dưới đây là ví dụ về mã SASS cho kernel cộng vector:
```sass
code_version_number 90
@P0 LDG.E R2, [R8];
@P0 LDG.E R0, [R10];
@P0 FADD R0, R2, R0;
@P0 STG.E [R12], R0;
EXIT;Mã SASS này tương ứng với mã PTX được hiển thị trước đó. Nó tải các phần tử vector đầu vào từ bộ nhớ toàn cục (LDG.E), thực hiện phép cộng (FADD), lưu kết quả trở lại bộ nhớ toàn cục (STG.E), và thoát khỏi kernel.
AMD Graphics Core Next ISA
GPU AMD sử dụng kiến trúc và ISA Graphics Core Next (GCN). GCN là một ISA dựa trên RISC hỗ trợ cả các tác vụ đồ họa và tính toán. Nó được thiết kế để có hiệu suất cao, có thể mở rộng và tiết kiệm năng lượng.
GCN giới thiệu một số tính năng chính như:
- Một ....Here is the Vietnamese translation of the given markdown file, with the code comments translated:
Bộ ALU cánh (alar) để thực thi hiệu quả các phép toán cô lập và điều khiển luồng.
- Một bộ ALU vector để thực thi song song các thao tác song song dữ liệu.
- Một hệ thống bộ nhớ băng thông cao, với hỗ trợ các phép toán nguyên tử và truy cập độ trễ thấp đến bộ nhớ chung.
- Chế độ địa chỉ linh hoạt cho các thao tác bộ nhớ, hỗ trợ địa chỉ cơ sở+offset và cô lập+vector.
Đây là một ví dụ về mã ISA GCN cho một hạt nhân cộng vector:
.text
.globl vecAdd
.p2align 2
.type vecAdd,@function
vecAdd:
.set DPTR, 0
# Tải các đối số của hạt nhân từ bộ nhớ
s_load_dwordx4 s[0:3], s[4:5], 0x0
s_load_dword s4, s[4:5], 0x10
s_waitcnt lgkmcnt(0)
# Chuẩn bị các địa chỉ để tải các phần tử của vector input
v_lshlrev_b32 v0, 2, v0
v_add_u32 v1, vcc, s1, v0
v_mov_b32 v3, s3
v_addc_u32 v2, vcc, s2, v3, vcc
flat_load_dword v1, v[1:2]
v_add_u32 v3, vcc, s0, v0
v_mov_b32 v5, s3
v_addc_u32 v4, vcc, s2, v5, vcc
flat_load_dword v0, v[3:4]
# Thực hiện phép cộng vector
v_add_f32 v0, v0, v1
flat_store_dword v[3:4], v0
s_endpgmĐoạn mã GCN này tải các phần tử vector đầu vào bằng flat_load_dword, thực hiện phép cộng bằng v_add_f32, và lưu kết quả trở lại bộ nhớ bằng flat_store_dword. Các chỉ thị s_load_dwordx4 và s_load_dword được sử dụng để tải các đối số của hạt nhân từ bộ nhớ.
Ánh xạ Thuật toán lên Kiến trúc GPU
Ánh xạ hiệu quả các thuật toán lên kiến trúc GPU là rất quan trọng để đạt được hiệu suất cao. Các yếu tố chính cần xem xét bao gồm:
Phơi bày đủ sự song song
Thuật toán nên được phân rã thành nhiều luồng nhỏ có thể thực hiện đồng thời để tận dụng tối đa khả năng xử lý song song của GPU. Điều này thường liên quan đến việc xác định các phần dữ liệu song song của thuật toán có thể được thực hiện độc lập trên các phần tử dữ liệu khác nhau.
Giảm thiểu phân nhánh chệch
Luồng điều khiển chệch trong một warp/wavefront có thể dẫn đến tuần tự hóa và hiệu suất SIMD bị giảm. Các thuật toán nên được cấu trúc để giảm thiểu phân nhánh chệch ở mức có thể. Điều này có thể đạt được bằng cách giảm việc sử dụng các nhánh phụ thuộc vào dữ liệuĐây là bản dịch tiếng Việt của tập tin markdown được cung cấp:
Khai thác Phân cấp Bộ nhớ
Truy cập bộ nhớ toàn cục là tốn kém. Các thuật toán nên tối đa hóa việc sử dụng bộ nhớ chung và thanh ghi để giảm số lần truy cập bộ nhớ toàn cục. Dữ liệu cũng nên được sắp xếp trong bộ nhớ để cho phép các truy cập bộ nhớ liên kết, nơi các luồng trong một warp truy cập các vị trí bộ nhớ liên tiếp. Sử dụng hiệu quả phân cấp bộ nhớ có thể giảm đáng kể độ trễ bộ nhớ và các điểm nghẽn băng thông.
Cân bằng Tính toán và Truy cập Bộ nhớ
Các thuật toán nên có tỷ lệ cao các phép tính số học so với các thao tác bộ nhớ để che giấu hiệu quả độ trễ bộ nhớ và đạt được thông lượng tính toán cao. Điều này có thể đạt được bằng cách tối đa hóa việc tái sử dụng dữ liệu, trích xuất dữ liệu trước và chồng các tính toán với các truy cập bộ nhớ.
Giảm thiểu Chuyển Dữ liệu giữa Host-Device
Chuyển dữ liệu giữa bộ nhớ của host (CPU) và device (GPU) là chậm. Các thuật toán nên giảm thiểu những chuyển đổi này bằng cách thực hiện càng nhiều tính toán trên GPU càng tốt. Dữ liệu nên được chuyển sang GPU theo từng lô lớn và được giữ trên thiết bị trong thời gian cần thiết để phân tán chi phí chuyển đổi.
Một số mẫu thiết kế thuật toán song song thường được sử dụng khi phát triển các hạt nhân GPU:
- Map: Mỗi luồng thực hiện cùng một thao tác trên một phần tử dữ liệu khác nhau, cho phép xử lý song song đơn giản các tập dữ liệu lớn.
- Reduce: Giảm song song được sử dụng để tính hiệu quả một giá trị đơn (ví dụ: tổng, giá trị lớn nhất) từ một tập dữ liệu đầu vào lớn. Các luồng thực hiện các phép giảm cục bộ, sau đó được kết hợp để tạo ra kết quả cuối cùng.
- Scan: Còn được gọi là tiền tố tổng, scan được sử dụng để tính tổng lũy tiến của các phần tử trong một mảng. Các thuật toán scan song song hiệu quả là các khối xây dựng quan trọng cho nhiều ứng dụng được gia tốc bởi GPU.
- Stencil: Mỗi luồng tính toán một giá trị dựa trên các phần tử dữ liệu lân cận. Các tính toán stencil rất phổ biến trong các mô phỏng khoa học và các ứng dụng xử lý hình ảnh.Dưới đây là bản dịch tiếng Việt của tệp markdown:
ications.
- Gather/Scatter: Các luồng đọc từ (gather) hoặc ghi vào (scatter) các vị trí tùy ý trong bộ nhớ toàn cục. Bố cục dữ liệu và mẫu truy cập cẩn thận là cần thiết để đạt hiệu quả.
Hình 3.20 minh họa một ví dụ về mẫu map, trong đó mỗi luồng áp dụng một hàm (ví dụ: căn bậc hai) cho một phần tử khác nhau của mảng đầu vào.
Mảng đầu vào:
| | | | | | | |
v v v v v v v v
______________________________
Các luồng: | | | | | | | |
|____|____|____|____|____|____|____|
| | | | | | |
v v v v v v v
Mảng đầu ra:Hình 3.20: Ví dụ về mẫu map trong lập trình GPU.
Kết luận
Các mô hình lập trình GPU như CUDA và OpenCL phơi bày khả năng xử lý song song của các GPU hiện đại cho các nhà phát triển, cho phép họ tăng tốc một loạt các ứng dụng. Các mô hình lập trình này cung cấp các trừu tượng cho phép các tải công việc song song cấp độ chi tiết được áu tự hiệu quả lên phần cứng GPU.
Hiểu mô hình thực thi, phân cấp bộ nhớ và các nguyên tố đồng bộ hóa được cung cấp bởi các mô hình lập trình này là thiết yếu để viết mã GPU có hiệu suất cao. Các nhà phát triển phải xem xét cẩn thận các yếu tố như tổ chức luồng, phân nhánh phân kỳ, mẫu truy cập bộ nhớ và thiết kế thuật toán để khai thác trọn vẹn sức mạnh tính toán của GPU.
Khi kiến trúc GPU tiếp tục phát triển, các mô hình lập trình và công cụ cũng phải tiến bộ để cho phép các nhà phát triển khai thác hiệu quả các tính năng và khả năng phần cứng mới. Nghiên cứu liên tục trong các lĩnh vực như thiết kế ngôn ngữ lập trình, tối ưu hóa trình biên dịch và auto-tuning sẽ rất quan trọng để cải thiện năng suất lập trình và khả năng mang tính di động về hiệu suất trong kỷ nguyên máy tính dị hợm.