Có cách nào để bỏ vòng lặp trong hạt nhân AMD OpenCL bằng trình biên dịch không?

Question

Tôi đang cố gắng đánh giá sự khác biệt về hiệu năng giữa OpenCL cho AMD và GPU Nvidia. Tôi có một hạt nhân thực hiện phép nhân vectơ-vector. Tôi đang chạy hạt nhân trên hai hệ thống khác nhau vào thời điểm đó, máy tính xách tay của tôi có GT525m NVidia với Ubuntu 12.04 và CUDA 4.0 (chứa thư viện OpenCL và phần đầu) và một là một máy tính để bàn với AMD Radeon HD7970 một lần nữa với Ubuntu 12.04 và các trình điều khiển Catalyst mới nhất.

Trong hạt nhân, tôi có hai câu lệnh #pragma unroll tạo ra tốc độ lớn cho việc triển khai Nvidia OpenCL (~ 6x). Tuy nhiên, phiên bản AMD OpenCL không tạo ra bất kỳ sự tăng tốc nào. Nhìn vào hạt nhân với bộ phân tích hạt nhân AMD APP cho lỗi mà unroll không được sử dụng vì số lượng chuyến đi không được biết đến. Vì vậy, câu hỏi của tôi là, làm #pragma unroll làm việc với AMD OpenCL hoặc là có một thay thế (có lẽ là một lá cờ biên dịch mà tôi không biết). Tôi đã bao gồm hạt nhân bên dưới

__kernel void mvKernel(__global float* a, const __global float* x, __global float* y, int m, int n) 
{ 
    float sum = 0.0f; 
    __global float* A; 
    int i; 
    int j = 0; 
    int indx = get_global_id(0); 
    __local float xs[12000]; 
#pragma unroll 
    for(i = get_local_id(0); i < n; i+= get_local_size(0)) { 
     xs[i] = x[i]; 
    } 
    barrier(CLK_LOCAL_MEM_FENCE); 
    A = &a[indx]; 
#pragma unroll 256 
    for(i = 0; i < n; i++) { 
     sum += xs[i] * A[j]; 
     j += m; 
    } 
    y[indx] = sum; 
} 

Cùng một hạt nhân này tạo ra kết quả chính xác trong cả hai triển khai nhưng lệnh #pragma unroll không làm bất kỳ điều gì cho AMD (được kiểm tra bằng cách nhận xét).

Answer 1

Nó không được ghi lại, nhưng nó thực sự nên hoạt động với #pragma unroll. Bạn có thể kiểm tra nhật ký trình biên dịch để xem liệu cuộn có được áp dụng không? Tôi không chắc chắn nếu máy phân tích hạt nhân sử dụng trình biên dịch giống như thời gian chạy OpenCL, bạn có thể muốn kiểm tra.

Nếu không, nếu bạn biết rằng n có khối 256, bạn có thể hủy đăng ký theo cách thủ công bằng cách lặp lại khối 256 thành phần và một khối khác bên trong với kích thước cố định là 256, có thể dễ dàng hơn để hủy đăng ký. Điều này chắc chắn sẽ giải quyết được vấn đề rằng số lượng chuyến đi không được biết đến tĩnh.

Tuy nhiên, lưu ý rằng việc bỏ vòng lặp thường không phải là thắng lợi lớn vì bạn không có nhiều sổ đăng ký để lưu trữ tính toán của mình. Việc tăng áp lực đăng ký từ các vòng lặp unrolling có thể dẫn đến đăng ký tràn, mà thậm chí còn chậm hơn. Bạn nên kiểm tra tốc độ hạt nhân thực sự nằm trên thẻ AMD. Một trình biên dịch NVIDIA OpenCL mới hơn cũng có thể không còn lợi ích gì nữa từ pragma chưa đăng ký.