Chapter 6: Performance Considerations¶
This chapter concludes Part I of the book by providing a systematic approach to identifying and fixing performance bottlenecks. It focuses on how to interact with the physical DRAM and how to balance parallel work.
6.1 Memory Coalescing¶
DRAM Bursts: Global memory (DRAM) does not just provide a single byte when requested. Instead, it delivers a "burst" of consecutive locations (e.g., 32, 64, or 128 bytes).
The Technique: If all threads in a warp access a single burst of data in one instruction, the hardware coalesces these many requests into a single memory transaction.
Row-Major Impact: Because C/CUDA use row-major order, threads mapping to adjacent columns (col = ... + threadIdx.x) will naturally access adjacent memory locations.
The Penalty: If threads access data in a "stride" (e.g., each thread accessing the same column but in different rows), the hardware must perform many separate, slow memory transactions. This is "uncoalesced" access.
6.2 Hiding Memory Latency¶
Memory Parallelism: DRAM is organized into Channels and Banks. Multiple requests can be processed simultaneously if they are spread across different banks.
The Symbiosis: High performance depends on a relationship between the threads and the DRAM. To reach peak bandwidth, you need enough active warps (high occupancy) to ensure the memory system is always busy processing requests while the cores are busy doing math.
6.3 Thread Coarsening¶
The "Price of Parallelism": Sometimes, having every thread do the smallest possible unit of work leads to redundant data loading or too much synchronization.
Definition: Thread coarsening is an optimization where the programmer makes one thread responsible for multiple output elements.
Benefit: In matrix multiplication, two adjacent output tiles often share the same input row. By using one coarsened thread to calculate two output points, that row is loaded from global memory only once instead of twice, saving bandwidth.
Risks: Too much coarsening reduces the total number of threads, which can lower occupancy and prevent the GPU from hiding latency.
6.4 A Checklist of Optimizations¶
The authors consolidate the book's foundational lessons into a "Universal Checklist":
- Maximizing Occupancy: Tuning SM resources (registers/shared memory) to keep more warps active.
- Enabling Coalescing: Rearranging thread-to-data mapping so global memory accesses are contiguous.
- Minimizing Control Divergence: Ensuring threads in a warp take the same execution path.
- Tiling of Reused Data: Using shared memory to reduce global memory traffic.
- Privatization: (Introduced later) Giving threads/blocks their own private copies of data to avoid "Atomic" operation contention.
- Thread Coarsening: Reducing redundant work by grouping tasks into a single thread.
6.5 Knowing Your Computation's Bottleneck¶
Resource Limits: Every application has a primary limiting factor (Memory-bound vs. Compute-bound).
Profilers: The authors recommend using tools like the NVIDIA Profiler to identify the "performance cliff" or "bottleneck."
Strategy: Optimization is only useful if it targets the actual bottleneck. For example, adding shared memory tiling won't help if your kernel is already limited by its math operations (compute-bound).
Key Takeaway: Chapter 6 teaches the "economics" of GPU programming. It's not just about writing parallel code; it's about writing code that respects the way hardware moves data (Coalescing) and finding the right balance between doing work in parallel and doing it serially (Coarsening) to maximize efficiency.
Chương này kết thúc Phần I của cuốn sách bằng cách cung cấp phương pháp có hệ thống để xác định và khắc phục các nút thắt cổ chai hiệu năng. Nó tập trung vào cách tương tác với DRAM vật lý và cách cân bằng công việc song song.
6.1 Gộp bộ nhớ (Memory Coalescing)¶
DRAM Burst: Bộ nhớ global (DRAM) không chỉ cung cấp một byte khi được yêu cầu. Thay vào đó, nó cung cấp một "burst" các vị trí liên tiếp (ví dụ: 32, 64 hoặc 128 byte).
Kỹ thuật: Nếu tất cả các luồng trong warp truy cập một burst dữ liệu duy nhất trong một lệnh, phần cứng gộp nhiều yêu cầu này thành một giao dịch bộ nhớ duy nhất.
Tác động thứ tự hàng-chính: Vì C/CUDA sử dụng thứ tự hàng-chính, các luồng ánh xạ đến các cột liền kề (col = ... + threadIdx.x) sẽ tự nhiên truy cập các vị trí bộ nhớ liền kề.
Hình phạt: Nếu các luồng truy cập dữ liệu theo "bước nhảy" (ví dụ: mỗi luồng truy cập cùng một cột nhưng ở các hàng khác nhau), phần cứng phải thực hiện nhiều giao dịch bộ nhớ riêng biệt, chậm. Đây là truy cập "không gộp".
6.2 Che giấu độ trễ bộ nhớ¶
Tính song song bộ nhớ: DRAM được tổ chức thành Kênh và Bank. Nhiều yêu cầu có thể được xử lý đồng thời nếu chúng được phân tán trên các bank khác nhau.
Sự cộng sinh: Hiệu năng cao phụ thuộc vào mối quan hệ giữa các luồng và DRAM. Để đạt băng thông đỉnh, bạn cần đủ warp hoạt động (occupancy cao) để đảm bảo hệ thống bộ nhớ luôn bận xử lý yêu cầu trong khi các lõi bận làm toán.
6.3 Làm thô luồng (Thread Coarsening)¶
"Giá của tính song song": Đôi khi, việc để mỗi luồng thực hiện đơn vị công việc nhỏ nhất có thể dẫn đến tải dữ liệu dư thừa hoặc quá nhiều đồng bộ hóa.
Định nghĩa: Thread coarsening là tối ưu hóa trong đó lập trình viên làm cho một luồng chịu trách nhiệm cho nhiều phần tử đầu ra.
Lợi ích: Trong nhân ma trận, hai tile đầu ra liền kề thường chia sẻ cùng một hàng đầu vào. Bằng cách sử dụng một luồng thô để tính hai điểm đầu ra, hàng đó được tải từ bộ nhớ global chỉ một lần thay vì hai lần, tiết kiệm băng thông.
Rủi ro: Làm thô quá nhiều giảm tổng số luồng, có thể làm giảm occupancy và ngăn GPU che giấu độ trễ.
6.4 Danh sách kiểm tra tối ưu hóa¶
Các tác giả tổng hợp các bài học nền tảng của cuốn sách thành "Danh sách kiểm tra phổ quát":
- Tối đa hóa Occupancy: Điều chỉnh tài nguyên SM (register/shared memory) để giữ nhiều warp hoạt động hơn.
- Kích hoạt Coalescing: Sắp xếp lại ánh xạ luồng-dữ liệu để truy cập bộ nhớ global liền kề.
- Giảm thiểu phân kỳ điều khiển: Đảm bảo các luồng trong warp đi theo cùng một đường thực thi.
- Tiling dữ liệu tái sử dụng: Sử dụng shared memory để giảm lưu lượng bộ nhớ global.
- Privatization: (Giới thiệu sau) Cung cấp cho luồng/block bản sao riêng của dữ liệu để tránh tranh chấp thao tác "Atomic".
- Thread Coarsening: Giảm công việc dư thừa bằng cách nhóm tác vụ vào một luồng duy nhất.
6.5 Biết nút thắt cổ chai của tính toán¶
Giới hạn tài nguyên: Mỗi ứng dụng có yếu tố giới hạn chính (Memory-bound vs. Compute-bound).
Profiler: Các tác giả khuyến nghị sử dụng các công cụ như NVIDIA Profiler để xác định "vách đá hiệu năng" hoặc "nút thắt cổ chai".
Chiến lược: Tối ưu hóa chỉ hữu ích nếu nó nhắm vào nút thắt cổ chai thực tế. Ví dụ: thêm tiling shared memory sẽ không giúp ích nếu kernel của bạn đã bị giới hạn bởi các phép toán (compute-bound).
Điểm chính: Chương 6 dạy "kinh tế học" của lập trình GPU. Không chỉ về viết code song song; mà về viết code tôn trọng cách phần cứng di chuyển dữ liệu (Coalescing) và tìm sự cân bằng đúng giữa làm công việc song song và làm nó tuần tự (Coarsening) để tối đa hóa hiệu quả.