Ngày nay, nhiều máy chủ chứa 8 GPU hoặc nhiều hơn. Trên nguyên tắc đó, chạy một ứng dụng sử dụng một đến nhiều GPU đúng ra sẽ mang lại một hiệu suất gấp bội. Nhưng thực tế thì khá khó để tận dụng được lợi thế này. Có hai nguyên nhân cơ bản phía sau sự thiếu tính khả mở của hệ thống đa-GPU. Đầu tiên ứng không có đủ luồng xử lý song song để dùng hết các vi xử lý trên GPU. Lý do thứ hai gây khó khăn khi tận dụng mở rộng ở chỗ các bộ vi xử lý sẽ phải dành thời gian để trao đổi dữ liệu với nhau quá nhiều thay vì tập trung chạy thuật toán. Để giảm tình trạng nút cổ chai ấy, điều chính yếu cần thiết phải tạo được băng thông lớn nhất sẵn sàng cho các kết nối liên GPU. NCCL là một đáp án.
NCCL (phát âm là “Nickel“) là một thư viện tập lệnh gốc đa-GPU có hướng cấu trúc liên kết và có thể dễ dàng tích hợp vào các ứng dụng. Khởi thủy từ một dự án nghiên cứu mã nguồn mở, NCCL được thiết kế để có dung lượng thật nhỏ dựa trên C ++ thông thường và các thư viện CUDA. NCCL có thể được ứng dụng vào trong các phần mềm đơn luồng hoặc đa luồng, để quản lý truyền nhận tin giữa các luồng xử lý hiệu quả. Sau nữa, NCCL đã bao gồm sẵn các API vốn rất quen thuộc với bất cứ ai có kinh nghiệm sử dụng các tập lệnh MPI.
Tập lệnh giao thức
Các chương trình giao thức tổ hợp là mẫu hình chung của truyền dữ liệu giữa nhiều bộ vi xử lý. Nếu đã có kinh nghiệm với MPI thì có lẽ cũng đã quen thuộc với một số phép toán tổ hợp lệnh. Ví dụ, All-Reduce (Giảm toàn phần) bắt đầu với mảng độc lập $latex V_k$ của N giá trị trên K bộ vi xử lý và kết thúc với mảng giống hệt S của N giá trị, trong đó $S[k] = V_0[k] + V_1[k] + … + V_k[k]$, cho mỗi bộ xử lý k . Xem hình 1.
Một giao thức tổ hợp thông thường khác là All-Gather (Nhóm toàn phần), trong đó mỗi bộ xử lý trong tập hợp K bắt đầu với một mảng độc lập của N giá trị và thu thập dữ liệu từ tất cả các bộ xử lý khác để tạo thành một kết quả mảng N*K , như minh họa ở hình 2.
Broadcast (Tán xạ) là một ví dụ thứ ba. Dưới đây là một thành phần đệm gồm N phẩn tử trên một bộ xử lý được sao chép vào tất cả các bộ xử lý khác, như hình 3 cho thấy.
Tất cả các tổ hợp trên có thể được thực hiện ở vùng “ngoại vi”, có bộ đệm vào và ra riêng biệt, hoặc “nội hàm” có chồng lấn vào và ra.
Thuật toán nhánh cây hai bước là một lựa chọn phổ biến ở đây. Tại bước đầu tiên, dữ liệu được gửi từ GPU0 tới một GPU khác, và trong bước thứ hai cả hai sẽ gửi dữ liệu đến các bộ vi xử lý còn lại. Dù thế thì vẫn có nhiều cách thực thi khác nhau. Dữ liệu có thể được gửi từ GPU0 đến GPU1 trong bước đầu tiên và sau đó từ GPU0 đến GPU2 và từ GPU1 đến GPU3 ở bước hai, hoặc có thể thực hiện dưới hình thức bản sao ban đầu từ GPU0 đến GPU2, kế đó ở bước hai thì từ GPU0 đến GPU1 và GPU2 đến GPU3. Kiểm tra cấu trúc liên kết, cách thứ hai tốt hơn rõ rệt, vì nếu gửi dữ liệu đồng thời từ GPU0 đến GPU2 và GPU1 đến GPU3 sẽ gây tranh chấp tài nguyên trên các liên kết PCIe, hiệu suất băng thông ở bước này bị giảm một nửa. Nói chung, để đạt được hiệu quả cao đối với tổ hợp lệnh đòi hỏi sự quan tâm cẩn trọng đến cấu trúc liên kết.
Để tối ưu hóa băng thông Broadcast, còn một cách tiếp cận tốt hơn là tạo liên kết cấu trúc vòng PCIe.
Tán xạ được thực hiện bằng việc sắp xếp các tập nhỏ dữ liệu đầu vào xung quanh vòng từ GPU0 để GPU3. Điều thú vị là, các thuật toán vòng cung cấp băng thông cận tối ưu cho hầu hết tất cả các thực thi tổ hợp lệnh tiêu chuẩn, thậm chí khi áp dụng cho cấu trúc PCIe kiểu “nhánh cây”. Tất nhiên phải lưu ý rằng việc lựa chọn đúng chiều tuần hoàn của vòng vẫn rất quan trọng.
Tập lệnh GPU với NCCL
Hình 5: Vòng tuần hoàn của GPU trong cây PCIe.
Nhiều tổ hợp lệnh đòi hỏi một bộ đệm chứa kết quả trung gian. Để giảm thiểu dung lượng chiếm bộ nhớ xuống còn một vài MB trên mỗi GPU, NCCL chia tách tổ hợp lệnh lớn thành nhiều phần nhỏ. Hiệu quả cực thấp nếu khởi tạo lõi xử lý và thực thi gọi hàm cudaMemcpy riêng biệt cho mỗi công đoạn hay từng phần nhỏ của cùng một thuật toán tổ hợp. Thay vào đó, mỗi tổ hợp được thực hiện trọn vẹn trên một nhân CUDA. NCCL tận dụng tối đa truy cập trực tiếp của giao thức GPUDirect Peer-to-Peer để truyền dữ liệu giữa các bộ xử lý. Trường hợp không lấy được truy cập trực tiếp peer-to-peer (ví dụ, khi phải truyền qua một kết nối QPI), thì các dữ liệu đã truyền đi được tổ chức thông qua một bộ đệm trong một vùng xác định trên bộ nhớ hệ thống. Tương tự như vậy, đồng bộ dũ liệu được thực hiện bằng cách kiểm tra các biến số biến đổi trong vùng nhớ của card (hoặc vùng nhớ xác định của hệ thống).
Bên trong, NCCL quy việc thực hiện mỗi tổ hợp lệnh về 3 nguyên hàm: Copy, Reduce, và ReduceAndCopy. Mỗi nguyên hàm được tối ưu hóa để truyền đi một lượng dữ liệu tinh giản (4-16KB) giữa các GPU. Các lõi xử lý cũng đã được tối ưu hóa đặc biệt để đạt được băng thông tối đa trong khi lại chiếm dụng thấp. Kết quả là, NCCL có thể sử dụng tối đa tải một kết nối PCIe 3.0 x16 mà chỉ bằng một đơn vị tổ hợp luồng xử lý CUDA. Điều này khiến giải phóng phần lớn GPU dùng vào chạy các tác vụ tính toán đồng thời với truyền tải thông tin.
NCCL hiện hỗ trợ các tập lệnh all-gather, all-reduce, broadcast, reduce, và reduce-scatter, không hạn chế số lượng GPU được sử dụng, miễn là nằm trên cùng một máy.
Sử dụng NCCL
Mã nguồn mở của NCCL có sẵn trên Github . Các thư viện có thể chạy trên bất kỳ phiên bản Linux phổ biến và cũng tương thích với CUDA 7.0 hay mới hơn. Hiện giờ, CUDA chỉ hỗ trợ truy cập trực tiếp giữa các GPU cùng kiến trúc và phải nằm trên cùng một máy (PCIe root hub). GPU không phù hợp các tiêu chí này vẫn được hỗ trợ bởi NCCL, mặc dù hiệu suất sẽ được giảm do việc chuyển dữ liệu sẽ đi vòng thông qua vùng nhớ xác định trong bộ nhớ hệ thống.
Các NCCL API đều tuân thủ chặt theo MPI. Trước khi thực thi các tập lệnh, một đối tượng truyền tin phải được tạo ra trên mỗi GPU. Các đối tượng này xác định các thiết lập của các GPU tham gia vào tổ hợp và lập bản đồ các luồng thông tin liên lạc giữa chúng. Tập hợp đối tượng giao tiếp liên quan với nhau được gọi là một clique (phe). Có hai cách để khởi tạo đối tượng truyền tin trong NCCL. Phương pháp phổ thông nhất là gọi hàm ncclCommInitRank() tương ứng với mỗi GPU.
Chức năng này giả định rằng các GPU thuộc các hàng xác định đã được thiết lập thông qua cudaSetDevice(). nGPUs là số lượng GPU trong một clique. cliqueId cho phép các hàng trong cùng clique nhận diện nhau. Các hàng thì sẽ có cùng một cliqueId. Để làm vậy, hàm ncclGetUniqueId() được gọi từ trong một hàng và sẽ phát tán kết quả trả về uniqueId đến các hàng khác còn lại trong clique thông qua mô hình truyền thông tùy chọn (ví dụ, MPI_Bcast).
Các biến số cuối cùng trong ncclCommInitRank() xác định các chỉ số của GPU hiện tại trong clique. Nó phải là duy nhất cho mỗi cấp trong clique và trong tập hợp [0, nGPUs]. Chỉ số này được sử dụng, ví dụ, để xác định các GPU nguồn cho lệnh broadcast, hoặc để truy xuất đến theo lệnh all-gather.
Sau khi khởi tạo thành công, ncclCommInitRank() trả về ncclSuccess và *comm được thiết lập chứa đối tượng truyền tin mới của NCCL.
Bên trong, ncclInitRank() thực hiện đồng bộ giữa tất cả các đối tượng truyền tin trong cùng một clique. Như vậy, nó phải được gọi từ các luồng khởi tạo khác nhau cho từng GPU, hoặc từ các xử lý riêng biệt (ví dụ, MPI ranks). Cụ thể, nếu gọi ncclInitRank() trong cùng một vòng lặp xử lý cho các GPUs sẽ dẫn đến tình trạng xung đột khóa chết (deadlocks)
Cách thứ hai để khởi tạo các đối tượng truyền tin là sử dụng ncclCommInitAll(). Đây thực chất là một thói quen tiện dụng để tiết kiệm việc phải sinh thêm các luồng khởi tạo mới để chạy NCCL trong trình ứng dụng đơn nhiệm.
Biến comms đang trỏ đến một mảng các đối tượng ncclComm_t, tương ứng cho mỗi hàng chứa nGPUs trong cùng 1 clique. devList thì quy định cụ thể thiết bị CUDA nào tương ứng với từng hàng. Khi khởi tạo các đối tượng truyền tin, ta có thể gọi các tập lệnh thông qua các hàm gốc của chúng, chẳng hạn như ncclAllReduce().
Tài liệu chi tiết cho mỗi tập lệnh được cung cấp trong file nccl.h. Hầu hết các biến đều tương tự như trong tập lệnh API của MPI. Trường hợp ngoại lệ nổi bật là các biến dòng (stream). Tương tự như rất nhiều chương trình CUDA “Async”, lệnh NCCL sẽ lên lịch thực thi trong một dòng, nhưng có thể trả lại trước khi tập lệnh được hoàn tất. Bằng cách xếp các tập lệnh vào hàng đợi còn nhân tính toán riêng khác ở các dòng độc lập, GPU có thể chồng lấn nhiều hơn công việc đòi hỏi tính toán tập trung cao vào cùng thời điểm truyền tập lệnh. Nhằm tối đa hóa sự chồng lấn này, các tập lệnh NCCL được lập lịch chạy trên dòng có độ ưu tiên cao để cho phép chúng trượt qua lại trong các lưới tính toán tập trung.
Ứng với mỗi đối tượng truyền tin trong một clique đều phải gọi đến ncclAllReduce(), mỗi lần gọi được thực hiện qua bộ đệm gửi và nhận riêng biệt, v.v. Tập lệnh NCCL sẽ giả định ngữ cảnh CUDA tương ứng là mới nhất. Và bởi vì NCCL chạy không đồng bộ, một vòng lặp đơn giản có thể được sử dụng để khởi tạo một tập lệnh từ một ứng dụng đơn luồng.
Hiệu suất
Hiệu suất mà tập lệnh NCCL đạt được phụ thuộc vào chính cấu trúc liên kết của máy tính. Trong trường hợp tốt nhất, tất cả các GPU đều chiếm phần truy cập như nhau. Đây là mô hình phổ biến nhất ở các máy trạm được trang bị vài GPU. Các máy chủ lớn hơn thường gắn với CPU kép, và một số lượng các GPU tách biệt trên các hub IO khác nhau. Hình 6 cho thấy băng thông NCCL tương ứng đối với các tập hợp lệnh khác nhau đo trên máy NVIDIA Digits DevBox trang bị 4 GPU GeForce GTX Titan X được ráp theo mô hình ở hình 4 trên đây.
Đường màu đỏ 10.4 GB/ s đại diện cho băng thông sử dụng bở một lệnh cudaMemcpy dự liệu lớn giữa hai trong số các GPU (cụ thể là GPU 0 và 2 trong hình 4). NCCL vẫn có thể duy trì tỷ lệ cao khả năng chiếm dụng băng thông cao điểm đó khi thực hiện giao tiếp giữa tất cả bốn GPU.
Hướng phát triển trong tương lai và một số điểm lưu tâm
Mục tiêu của NCCL là để cung cấp tập lệnh tổ hợp tự nhận thức được cấu trúc liên kết nhằm có thể cải thiện độ khả mở của các ứng dụng đa GPU. Sử dụng NCCL sẽ có thể giúp đạt được hiệu suất tuyệt vời mà không phải chú ý về các chi tiết phần cứng bậc thấp.
Nhữn mẩu mã nguồn ví dụ đơn giản cho cả hai loại ứng dụng trình đơn và đa luồng MPI cũng được kèm sẵn trong bộ thư viện NCCL. Xuất phát từ một dự án nghiên cứu, mọi phản hồi cuả các nhà phát triển đểu được hoanh nghênh vì dự án này vẫn còn tiếp tục. Hãy dùng thử và chia sẻ bạn nghĩ gì nhé! Để cập nhật về những thông tin mới nhất, hãy đến dự buổi nói chuyện NCCL của tôi ở GTC.