Tại sao phải biết về CUDA Warps?

Question

Tôi có GeForce GTX460 SE, vì vậy nó là: 6 SM x 48 CUDA = 288 CUDA lõi. Được biết, trong một Warp có 32 luồng, và trong một khối cùng một lúc (tại một thời điểm) có thể được thực hiện chỉ có một Warp. Tức là, trong một bộ xử lý đa (SM) chỉ có thể thực hiện đồng thời một khối, một Warp và chỉ 32 luồng, ngay cả khi có 48 lõi có sẵn?

Và ngoài ra, một ví dụ để phân phối Thread và Block cụ thể có thể được sử dụng threadIdx.x và blockIdx.x. Để phân bổ chúng, hãy sử dụng kernel < < < Blocks, Threads >>>(). Nhưng làm thế nào để phân bổ một số cụ thể của Warp-s và phân phối chúng, và nếu nó là không thể thì tại sao bận tâm để biết về Warps?

Answer 1

Tình hình hơi phức tạp hơn những gì bạn mô tả.

ALUs (lõi), đơn vị tải/lưu trữ (LD/ST) và Đơn vị chức năng đặc biệt (SFU) (màu xanh lá cây trong hình ảnh) là đơn vị đường ống. Họ giữ kết quả của nhiều tính toán hoặc hoạt động cùng một lúc, trong các giai đoạn hoàn thành khác nhau. Vì vậy, trong một chu kỳ, họ có thể chấp nhận một hoạt động mới và cung cấp kết quả của một hoạt động khác đã được bắt đầu từ lâu (khoảng 20 chu kỳ cho ALU, nếu tôi nhớ chính xác). Vì vậy, một SM duy nhất trong lý thuyết có nguồn lực để xử lý 48 * 20 chu kỳ = 960 ALU hoạt động cùng một lúc, đó là 960/32 chủ đề cho mỗi warp = 30 warps. Ngoài ra, nó có thể xử lý các hoạt động LD/ST và các hoạt động SFU ở bất kỳ độ trễ và thông lượng nào của chúng.

Bộ lập lịch dọc (màu vàng trong hình ảnh) có thể lên lịch 2 * 32 luồng mỗi sợi dọc = 64 chủ đề cho đường ống mỗi chu kỳ. Vì vậy, đó là số lượng kết quả có thể thu được trên mỗi đồng hồ. Vì vậy, cho rằng có một sự kết hợp của tài nguyên máy tính, 48 lõi, 16 LD/ST, 8 SFU, mỗi có độ trễ khác nhau, một hỗn hợp của warps đang được xử lý cùng một lúc. Tại bất kỳ chu kỳ nhất định nào, các bộ lập lịch dọc sẽ cố gắng "ghép nối" hai xung đột để lên lịch, để tối đa hóa việc sử dụng SM.

Bộ lập lịch warp có thể phát hành cảnh báo từ các khối khác nhau hoặc từ các vị trí khác nhau trong cùng một khối, nếu các lệnh độc lập. Vì vậy, warps từ nhiều khối có thể được xử lý cùng một lúc.

Thêm vào sự phức tạp, warps đang thực hiện các hướng dẫn có ít hơn 32 tài nguyên, phải được phát hành nhiều lần cho tất cả các chuỗi được phục vụ. Ví dụ, có 8 SFUs, do đó có nghĩa là một sợi dọc chứa một lệnh yêu cầu các SFU phải được lên lịch 4 lần.

Mô tả này được đơn giản hóa. Có những hạn chế khác cũng được đưa ra để xác định cách GPU lên lịch công việc. Bạn có thể tìm thêm thông tin bằng cách tìm kiếm trên web cho "kiến trúc fermi".

Vì vậy, sắp tới cho câu hỏi thực tế của bạn,

tại sao bận tâm để biết về warps?

Biết số lượng chủ đề trong sợi dọc và xem xét chủ đề trở nên quan trọng khi bạn cố gắng tối đa hóa hiệu suất của thuật toán.Nếu bạn không làm theo các quy tắc, bạn sẽ mất hiệu suất:

• Trong invocation hạt nhân, <<<Blocks, Threads>>>, hãy cố gắng chọn một số chủ đề mà chia đồng đều với số lượng bài trong một warp. Nếu không, bạn kết thúc bằng việc khởi chạy một khối có chứa các chuỗi không hoạt động.

• Trong hạt nhân, hãy cố gắng để mỗi sợi trong một sợi dọc theo cùng một đường dẫn mã. Nếu bạn không, bạn sẽ có được những gì được gọi là phân kỳ dọc. Điều này xảy ra vì GPU phải chạy toàn bộ dọc qua từng đường dẫn mã khác nhau.

• Trong hạt nhân của bạn, hãy cố gắng có từng luồng trong tải dọc và lưu trữ dữ liệu trong các mẫu cụ thể. Ví dụ, có các luồng trong một truy vấn dọc tiếp cận từ 32 bit trong bộ nhớ toàn cục.

Answer 2

Các chủ đề được nhóm thành Warps nhất thiết theo thứ tự, 1 - 32, 33 - 64 ...?

Có, mô hình lập trình đảm bảo rằng các chuỗi được nhóm thành warps theo thứ tự cụ thể đó.

Ví dụ đơn giản về tối ưu hóa các đường dẫn mã khác nhau có thể được sử dụng để tách tất cả các chuỗi trong khối theo nhóm 32 chủ đề? Ví dụ: switch (threadIdx.s/32) {case 0:/* 1 warp */break; trường hợp 1:/* 2 warp */break;/* Etc * /}

Chính xác :)

bao nhiêu byte phải được đọc cùng một lúc cho Warp duy nhất: 4 byte * 32 Chuyên, 8 byte * 32 Chủ đề hoặc 16 byte * 32 Chủ đề? Theo như tôi biết, một giao dịch vào bộ nhớ toàn cục tại một thời điểm nhận được 128 byte.

Có, giao dịch với bộ nhớ toàn cục là 128 byte. Vì vậy, nếu mỗi luồng đọc một từ 32 bit từ các địa chỉ liên tiếp (chúng có thể cần phải được căn chỉnh 128 byte), tất cả các luồng trong dọc có thể được phục vụ với một giao dịch đơn lẻ (4 byte * 32 chủ đề = 128 byte). Nếu mỗi luồng đọc nhiều byte hơn, hoặc nếu các địa chỉ không liên tiếp, cần phải có nhiều giao dịch hơn (với các giao dịch riêng biệt cho mỗi dòng 128 byte riêng biệt được chạm vào).

Điều này được mô tả trong Hướng dẫn lập trình CUDA 4.2, mục F.4.2, "Bộ nhớ chung". Ngoài ra còn có một blurb trong đó nói rằng tình hình là khác nhau với dữ liệu được lưu trữ chỉ trong L2, như bộ nhớ cache L2 có 32-byte dòng bộ nhớ cache. Tôi không biết làm thế nào để sắp xếp cho dữ liệu được lưu trữ chỉ trong L2 hoặc có bao nhiêu giao dịch một kết thúc với.