2012-08-03 17 views
5

Tôi đang học OpenACC (với trình biên dịch PGI) và cố gắng tối ưu hóa ví dụ nhân ma trận. Triển khai nhanh nhất mà tôi đã đưa ra cho đến thời điểm này là:cách tối ưu hóa phép nhân bằng cách sử dụng OpenACC?

void matrix_mul(float *restrict r, float *a, float *b, int N, int accelerate){ 

#pragma acc data copyin (a[0: N * N ], b[0: N * N]) copyout (r [0: N * N ]) if(accelerate) 
{ 
# pragma acc region if(accelerate) 
{ 
# pragma acc loop independent vector(32) 
for (int j = 0; j < N; j ++) 
{  
    # pragma acc loop independent vector(32) 
    for (int i = 0; i < N ; i ++) 
    { 
     float sum = 0; 
     for (int k = 0; k < N ; k ++) { 
     sum += a [ i + k*N ] * b [ k + j * N ]; 
     } 
     r[i + j * N ] = sum ; 
    } 
} 
} 
} 

Kết quả này trong các chuỗi chuỗi có kích thước 32x32 và mang lại hiệu suất tốt nhất cho đến nay. Sau đây là các tiêu chuẩn:

Matrix multiplication (1500x1500): 
GPU: Geforce GT650 M, 64-bit Linux 

Data sz    : 1500  
Unaccelerated: 
    matrix_mul() time : 5873.255333 msec 
Accelerated: 
    matrix_mul() time : 420.414700 msec 

Data size    : 1750 x 1750  
    matrix_mul() time : 876.271200 msec 
Data size    : 2000 x 2000  
    matrix_mul() time : 1147.783400 msec 
Data size    : 2250 x 2250  
    matrix_mul() time : 1863.458100 msec 
Data size    : 2500 x 2500  
    matrix_mul() time : 2516.493200 msec 

Đáng tiếc là tôi nhận ra rằng mã CUDA tạo ra là khá thô sơ (ví dụ nó thậm chí không sử dụng bộ nhớ chia sẻ) và do đó không thể cạnh tranh với chương trình CUDA tay được tối ưu hóa. Là một tài liệu tham khảo thực hiện tôi đã Arrayfire lib với kết quả như sau:

Arrayfire 1500 x 1500 matrix mul 
CUDA toolkit 4.2, driver 295.59 
GPU0 GeForce GT 650M, 2048 MB, Compute 3.0 (single,double) 
Memory Usage: 1932 MB free (2048 MB total) 
af: 0.03166 seconds 

Arrayfire 1750 x 1750 matrix mul 
af: 0.05042 seconds 
Arrayfire 2000 x 2000 matrix mul 
af: 0.07493 seconds 
Arrayfire 2250 x 2250 matrix mul 
af: 0.10786 seconds 
Arrayfire 2500 x 2500 matrix mul 
af: 0.14795 seconds 

Tôi tự hỏi nếu có bất cứ đề nghị làm thế nào để có được hiệu suất tốt hơn từ OpenACC? Có lẽ lựa chọn chỉ thị của tôi không đúng?

+1

Vấn đề này minh họa cách tiếp cận khác nhau của Trình biên dịch chỉ thị so với CUDA/OpenCL. CUDA/OpenCL gần gũi hơn với H/W; nơi bạn có thể tối ưu hóa và tinh chỉnh cho nền tảng H/W. Bạn có thể unroll các vòng lặp bên trong tính toán 2,4, hoặc 8, ... Tổng số như vậy làm giảm số lượng các vòng bên trong –

+1

huh ý tưởng tốt, cảm ơn .. Vâng tôi biết, CUDA/OpenCL có thể được coi là "cấp thấp" API, Bản thân tôi là từ trường CUDA cũ. Mặt khác, OpenACC có nhiều tiềm năng hơn trong tương lai vì nó không chỉ giới hạn ở GPU và tất nhiên là chi phí phát triển. Tuy nhiên, nó sẽ là tốt đẹp nếu các trình biên dịch OpenACC có thể khai thác bộ nhớ chia sẻ của GPU để tính toán: Tôi biết có OpenACC 'cache' chỉ thị nhưng tôi không thể quản lý để làm cho nó hoạt động –

Trả lời

4

Bạn đang nhận được ngay với tốc độ 14x, điều này khá tốt cho trình biên dịch PGI theo kinh nghiệm của tôi.

Trước hết, bạn có đang biên dịch bằng -Minfo không? Điều đó sẽ cung cấp cho bạn rất nhiều phản hồi từ trình biên dịch về các lựa chọn tối ưu hóa.

Bạn đang sử dụng khối chuỗi 32x32, nhưng trong các khối chuỗi 16x16 kinh nghiệm của tôi có xu hướng có hiệu suất tốt hơn. Nếu bạn bỏ qua mệnh đề vectơ (32), thì trình biên dịch sẽ chọn gì?

Khai báo a và b bằng các giới hạn có thể cho phép trình biên dịch tạo mã tốt hơn.

Chỉ bằng cách xem mã của bạn, tôi không chắc chắn rằng bộ nhớ dùng chung sẽ giúp hiệu suất. Bộ nhớ chia sẻ chỉ giúp cải thiện hiệu suất nếu mã của bạn có thể lưu trữ và sử dụng lại các giá trị ở đó thay vì chuyển sang bộ nhớ chung. Trong trường hợp này, bạn không sử dụng lại bất kỳ phần nào của a hoặc b sau khi đọc nó.

Điều đáng lưu ý là tôi đã có trải nghiệm xấu với trình biên dịch PGI khi nói đến việc sử dụng bộ nhớ dùng chung. Nó đôi khi sẽ làm công cụ vui và bộ nhớ cache các giá trị sai (dường như chủ yếu xảy ra nếu bạn lặp lại một vòng lặp lạc hậu), tạo ra kết quả sai. Tôi thực sự phải biên dịch ứng dụng hiện tại của tôi bằng cách sử dụng tùy chọn không có giấy phép -ta = nvidia, nocache để làm cho nó hoạt động chính xác, bằng cách bỏ qua việc sử dụng bộ nhớ chia sẻ hoàn toàn.

+0

vâng tôi đã thử trường hợp 16x16 nhưng nó thực sự chạy chậm hơn. Tôi cho rằng điều này là chính xác vì không có bộ nhớ dùng chung nào được sử dụng. Do đó, các chủ đề chúng tôi nhận được cho mỗi khối càng lớn thì hiệu quả của các kết quả trung gian "caching" trong thanh ghi. Có thực sự là một cách làm thế nào bộ nhớ chia sẻ có thể giúp hiệu suất nếu bạn có một cái nhìn tại ví dụ nhân ma trận trong CUDA SDK. Nếu tôi loại bỏ các mệnh đề vectơ (32), trình biên dịch chỉ vector hóa theo các hàng của ma trận (không phải bằng các ô 2D) và hiệu năng sẽ giảm xuống. Dù sao cảm ơn cho một lời khuyên tốt –

Các vấn đề liên quan