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?
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 –
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 –