Làm thế nào để tối ưu hóa trò chơi cuộc sống của Conway cho CUDA?

Question

Tôi đã viết kernel CUDA này cho trò chơi của cuộc sống Conway:

__global__ void gameOfLife(float* returnBuffer, int width, int height) { 
    unsigned int x = blockIdx.x*blockDim.x + threadIdx.x; 
    unsigned int y = blockIdx.y*blockDim.y + threadIdx.y; 
    float p = tex2D(inputTex, x, y); 
    float neighbors = 0; 
    neighbors += tex2D(inputTex, x+1, y); 
    neighbors += tex2D(inputTex, x-1, y); 
    neighbors += tex2D(inputTex, x, y+1); 
    neighbors += tex2D(inputTex, x, y-1); 
    neighbors += tex2D(inputTex, x+1, y+1); 
    neighbors += tex2D(inputTex, x-1, y-1); 
    neighbors += tex2D(inputTex, x-1, y+1); 
    neighbors += tex2D(inputTex, x+1, y-1); 
    __syncthreads(); 
    float final = 0; 
    if(neighbors < 2) final = 0; 
    else if(neighbors > 3) final = 0; 
    else if(p != 0) final = 1; 
    else if(neighbors == 3) final = 1; 
    __syncthreads(); 
    returnBuffer[x + y*width] = final; 
} 

Tôi đang tìm kiếm các lỗi/tối ưu hóa. Lập trình song song khá mới đối với tôi và tôi không chắc liệu tôi có làm cách nào để làm điều đó đúng không.

Phần còn lại là memcpy từ mảng đầu vào cho đầu vào kết cấu 2D được liên kết với mảng CUDA. Đầu ra là memcpy-ed từ bộ nhớ toàn cục đến máy chủ và sau đó được xử lý.

Như bạn có thể thấy một chuỗi giao dịch với một pixel. Tôi không chắc liệu đó có phải là cách nhanh nhất vì một số nguồn đề xuất thực hiện một hàng hay nhiều hơn cho mỗi luồng. Nếu tôi hiểu chính xác thì NVidia nói rằng càng nhiều chủ đề càng tốt. Tôi rất thích lời khuyên này từ một người có kinh nghiệm thực tế.

Answer 1

Hai xu của tôi.

Toàn bộ điều có thể bị ràng buộc bởi độ trễ của giao tiếp giữa các bộ xử lý đa và bộ nhớ GPU. Bạn có mã phải mất một cái gì đó giống như 30-50 đồng hồ ve để thực hiện trên riêng của mình, và nó tạo ra ít nhất 3 bộ nhớ truy cập trong đó có 200 + đồng hồ ve mỗi nếu dữ liệu cần thiết không có trong bộ nhớ cache.

Sử dụng bộ nhớ kết cấu là một cách hay để giải quyết vấn đề đó, nhưng không nhất thiết phải là cách tối ưu.

Ít nhất, hãy thử làm 4 pixel mỗi lần (theo chiều ngang) cho mỗi chuỗi. Bộ nhớ toàn cầu có thể được truy cập 128 byte tại một thời điểm (miễn là bạn có một warp cố gắng truy cập bất kỳ byte nào trong khoảng thời gian 128 byte, bạn cũng có thể kéo toàn bộ dòng bộ nhớ cache với hầu như không mất thêm chi phí). Vì một sợi dọc là 32 luồng, nên mỗi thread hoạt động trên 4 pixel sẽ hiệu quả.

Hơn nữa, bạn muốn có các điểm ảnh liền kề theo chiều dọc được thực hiện bởi cùng một bộ xử lý đa. Lý do là các hàng lân cận chia sẻ cùng một dữ liệu đầu vào. Nếu bạn có pixel (x = 0, y = 0) làm việc trên một MP và pixel (x = 0, y = 1) được làm việc bởi một MP khác, cả hai MP phải đưa ra ba yêu cầu bộ nhớ toàn cục. Nếu cả hai đều làm việc trên cùng một MP và kết quả được lưu trữ đúng cách (ngầm hoặc rõ ràng), bạn chỉ cần tổng cộng bốn. Điều này có thể được thực hiện bằng cách mỗi thread hoạt động trên một số pixel dọc hoặc bằng cách có blockDim.y> 1.

Nói chung, bạn có thể muốn có mỗi tải dọc 32 luồng càng nhiều bộ nhớ khi bạn có sẵn trên MP (16-48 kb hoặc ít nhất là khối 128x128) và sau đó xử lý tất cả các pixel trong cửa sổ đó.

Trên các thiết bị tương thích tính toán trước 2.0, bạn sẽ muốn sử dụng bộ nhớ dùng chung. Trên các thiết bị tương thích tính toán 2.0 và 2.1, khả năng lưu bộ nhớ đệm được cải thiện nhiều, vì vậy bộ nhớ toàn cầu có thể ổn.

Một số khoản tiết kiệm không có khả năng có thể có bằng cách đảm bảo rằng mỗi sợi chỉ truy cập hai dòng bộ nhớ cache trong mỗi hàng ngang pixel đầu vào, thay vì ba, như sẽ xảy ra trong triển khai ngây thơ hoạt động trên 4 pixel cho mỗi chuỗi, 32 luồng mỗi sợi dọc.

Không có lý do chính đáng để sử dụng phao làm loại bộ đệm. Bạn không chỉ kết thúc với băng thông bộ nhớ gấp bốn lần, mà mã trở nên không đáng tin cậy và dễ bị lỗi. (Ví dụ, bạn có chắc chắn rằng if(neighbors == 3) hoạt động chính xác, vì bạn đang so sánh một phao và một số nguyên?) Sử dụng unsigned char. Tốt hơn, sử dụng uint8_t và typedef nó có nghĩa là unsigned char nếu nó không được định nghĩa.

Cuối cùng, đừng đánh giá thấp giá trị của thử nghiệm.Khá thường xuyên thực hiện mã CUDA không thể được giải thích một cách dễ dàng bởi logic và bạn phải sử dụng để tinh chỉnh các tham số và xem điều gì xảy ra.

Answer 2

có một cái nhìn tại thread này, chúng tôi đã phân phối của cải trên đó ...

http://forums.nvidia.com/index.php?showtopic=152757&st=60

Answer 3

TL; DR: xem: http://golly.sourceforge.net

Vấn đề là hầu hết các trường CUDA theo não chết ý tưởng của hướng dẫn đếm của hàng xóm. Điều này rất chậm chạp đến nỗi bất kỳ việc triển khai CPU nối tiếp thông minh nào cũng sẽ hoạt động tốt hơn.

Cách duy nhất hợp lý để thực hiện tính toán GoL là sử dụng bảng tra cứu.
Việc triển khai nhanh nhất hiện nay trên một CPU sử dụng tra cứu một khối vuông 4x4 = 16 bit để xem có được các ô 2x2 trong tương lai bên trong.

trong thiết lập này các tế bào được đặt ra như sau:

0xxxxxxxx //byte0 
1xxxxxxxx //byte1 
2 etc 
3 
4 
5 
6 
7 

Một số bit chuyển được sử dụng để có được một khối 4x4 để phù hợp với một từ và từ đó được nhìn bằng cách sử dụng một bảng tra cứu. Các bảng tra cứu cũng chứa các từ, theo cách này có thể lưu trữ 4 phiên bản khác nhau của kết quả trong bảng tra cứu, do đó bạn có thể giảm thiểu số bit cần thực hiện trên đầu vào và/hoặc đầu ra.

Bên cạnh đó các thế hệ khác nhau đang loạng choạng, do đó bạn chỉ cần nhìn vào 4 tấm láng giềng, thay vì 9. Giống như vậy:

AAAAAAAA 
AAAAAAAA BBBBBBBB 
AAAAAAAA BBBBBBBB 
AAAAAAAA BBBBBBBB 
AAAAAAAA BBBBBBBB 
AAAAAAAA BBBBBBBB 
AAAAAAAA BBBBBBBB 
AAAAAAAA BBBBBBBB 
      BBBBBBBB 
//odd generations (A) are 1 pixel above and to the right of B, 
//even generations (B) are 1 pixels below and to the left of A. 

này mình kết quả trong một 1000x + tốc độ lên so với triển khai đếm ngớ ngẩn.

Sau đó là tối ưu hóa không tính toán tấm mà là tĩnh hoặc có một chu kỳ của 2.

Và sau đó là HashLife, nhưng đó là một con thú khác nhau hoàn toàn.
HashLife có thể tạo ra các mẫu Cuộc sống trong thời gian O (log n), thay vì thời gian thực hiện bình thường O (n) có thể. Điều này cho phép bạn tính toán thế hệ: 6,366,548,773,467,669,985,195,496,000 (6 octillion) chỉ trong vài giây.
Thật không may Hashlife yêu cầu đệ quy, và do đó rất khó trên CUDA.