OpenCL vô hướng vs vector

Question

Tôi có hạt nhân đơn giản:

__kernel vecadd(__global const float *A, 
       __global const float *B, 
       __global float *C) 
{ 
    int idx = get_global_id(0); 
    C[idx] = A[idx] + B[idx]; 
} 

Tại sao khi tôi thay đổi phao để float4, kernel chạy chậm hơn 30%?

Tất cả các hướng dẫn nói, rằng việc sử dụng các loại vector tăng tốc độ tính toán ...

Về phía chủ nhà, bộ nhớ alocated cho các đối số float4 là 16 byte liên kết và global_work_size cho clEnqueueNDRangeKernel nhỏ 4 lần.

Hạt nhân chạy trên GPU AMD HD5770, AMD-APP-SDK-v2.6.

thông tin thiết bị cho CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT trả 4.

EDIT:
global_work_size = 1024 * 1024 (và hơn)
local_work_size = 256
Thời gian đo bằng CL_PROFILING_COMMAND_START và CL_PROFILING_COMMAND_END.

Đối với global_work_size nhỏ hơn (8196 cho float/2048 cho float4), phiên bản vectơ nhanh hơn, nhưng tôi muốn biết, tại sao?

Answer 1

Tôi không biết các hướng dẫn bạn đề cập đến là gì, nhưng chúng phải cũ. Cả ATI và NVIDIA đều sử dụng kiến ​​trúc gpu vô hướng trong ít nhất nửa thập kỷ nay. Ngày nay, sử dụng các vectơ trong mã của bạn chỉ để thuận tiện cho cú pháp, nó không mang lại lợi ích hiệu suất nào so với mã vô hướng đơn giản. Nó chỉ ra kiến ​​trúc vô hướng là tốt hơn cho GPU hơn vectơ - nó là tốt hơn lúc sử dụng các tài nguyên phần cứng.

Answer 2

Tôi không chắc chắn lý do tại sao các vectơ sẽ chậm hơn nhiều đối với bạn, mà không cần biết thêm về nhóm làm việc và kích thước toàn cầu. Tôi hy vọng nó sẽ có ít nhất hiệu suất tương tự.

Nếu nó phù hợp với hạt nhân của bạn, bạn có thể bắt đầu với C có các giá trị trong A không? Điều này sẽ cắt giảm truy cập bộ nhớ 33%. Có thể điều này áp dụng cho tình huống của bạn?

__kernel vecadd(__global const float4 *B, 
       __global float4 *C) 
{ 
    int idx = get_global_id(0); 
    C[idx] += B[idx]; 
} 

Ngoài ra, bạn có cảm thấy mệt mỏi khi đọc các giá trị cho vectơ riêng tư, sau đó thêm? Hoặc có thể cả hai chiến lược.

__kernel vecadd(__global const float4 *A, 
       __global const float4 *B, 
       __global float4 *C) 
{ 
    int idx = get_global_id(0); 
    float4 tmp = A[idx] + B[idx]; 
    C[idx] = tmp; 
}