CUDA/openCL; viết lại chi nhánh là không phân nhánh biểu

Question

Phần lớn thời gian một chi nhánh là cần thiết trong một chương trình CUDA hoặc OpenCL, như:

for (int i=0; i<width; i++) 
{ 
    if(i % threadIdx.x == 0) 
    quantity += i*i; 
} 

mã luôn có thể (hoặc ít nhất, hầu hết thời gian) được viết lại trong không phân nhánh phong cách:

for (int i=0; i<width; i++) 
{ 
    quantity += i*i* (i % threadIdx.x != 0); 
} 

sự cân bằng dường như được hoặc chạy trong một khe dọc duy nhất so với làm nhiều phép tính trên tất cả các chủ đề (trong trường hợp thứ hai, số tiền được thực hiện mọi khi, chỉ là đôi khi giá trị là số không)

Giả sử rằng các hoạt động phân nhánh sẽ lấy nhiều vị trí dọc cho mỗi nhánh có thể, người ta sẽ mong đợi thứ hai liên tục tốt hơn lần đầu tiên, Bây giờ câu hỏi của tôi là; Tôi có thể dựa vào trình biên dịch để tối ưu hóa 1) thành 2) bất cứ khi nào nó có ý nghĩa, hoặc không có tiêu chí áp dụng rộng rãi, ngụ ý rằng nó không thể được quyết định nói chung nào tốt hơn mà không cố gắng và lược tả?

Answer 1

Tôi không có nhiều kỷ niệm về CUDA, nhưng tại sao bạn không song song vòng lặp của mình? Bạn nên sử dụng các phép toán nguyên tử [1] để thêm tính toán của bạn. Tôi hy vọng điều này sẽ giúp bạn! Xin lỗi nếu nó không phải là trường hợp.

1. Atomic Operations: http://supercomputingblog.com/cuda/cuda-tutorial-4-atomic-operations/

Answer 2

Theo kinh nghiệm của tôi - đó là hoàn toàn lên đến trình biên dịch-nhà văn để tối ưu hóa các loại trường hợp cạnh.

Vì vậy, tôi có thể nghĩ đến bất kỳ trường hợp nào trong đó 1) không thể được chuyển thành 2)? Dưới đây là một: Tôi đã viết hạt nhân ở đó hiệu quả hơn để chạy các phần tính toán nhất định trong 10 chủ đề hoặc thứ gì đó tương tự như vậy trong trường hợp đó tối ưu hóa không thể suy ra được ngay cả khi có phép toán (chia trừ trừ) mang lại kết quả tương tự không phụ thuộc vào điều kiện so với "chạy trên tất cả nhưng mang lại kết quả bằng không".

Tuy nhiên, ngay cả khi kiểm tra cho threadId == 0 là một kịch bản đủ phổ biến, tôi không có ý tưởng nếu nó thực sự được tối ưu hóa cho. Tôi sẽ cược rằng nó phụ thuộc vào việc thực hiện và thậm chí cả thiết bị chính nó (CPU vs GPU). Bạn sẽ phải thử nó để thực sự tìm ra những gì hoạt động tốt nhất, không chỉ vì lý do ở trên mà còn vì công việc lên lịch có thể hoạt động khác nhau dựa trên mức độ đắt tiền để lên lịch/bắt đầu/dừng một tập hợp các chủ đề trái ngược với việc tất cả chúng đều chạy (và hầu hết cung cấp một kết quả không/nhận dạng).

Hy vọng điều này sẽ hữu ích!

Answer 3

Hoạt động Modulo là khá đắt: tôi chắc chắn việc thêm vào modulo sẽ sử dụng nhiều thời gian hơn là chỉ có một lệnh duy nhất chỉ thực hiện 1 chuỗi. Tuyên bố phân nhánh duy nhất của bạn, một if không có else, sẽ chỉ treo các luồng khác trong khi đó nếu statment đang được thực hiện. Bởi vì gpus được tối ưu hóa cho việc chuyển đổi bối cảnh rất nhanh, nên có rất ít chi phí cho việc đó.

Bạn nên tránh sử dụng câu lệnh phân nhánh dài: quá nhiều tính toán nối tiếp trên GPU (nghĩa là một chuỗi làm tất cả công việc) phủ nhận lợi thế của tính song song.