OpenCL AMD vs NVIDIA hiệu suất

Question

Tôi đã triển khai một hạt nhân đơn giản, đó là một loại sắp xếp của một convolution. Tôi đo nó trên NVIDIA GT 240. Mất 70 ms khi viết trên CUDA và 100 ms khi được viết trên OpenCL. Ok, tôi nghĩ, trình biên dịch NVIDIA được tối ưu hóa tốt hơn cho CUDA (hoặc tôi đang làm gì đó sai). Tôi cần chạy trên GPU AMD, vì vậy tôi đã di chuyển sang AMD APP SDK. Chính xác cùng một mã hạt nhân.

Tôi đã thực hiện hai thử nghiệm và kết quả của họ không khuyến khích tôi: 200 ms ở HD 6670 và 70 ms ở HD 5850 (cùng thời điểm với GT 240 + CUDA). Và tôi rất quan tâm đến những lý do của hành vi kỳ lạ như vậy.

Tất cả các dự án được xây dựng trên VS2010 sử dụng cài đặt từ các dự án mẫu của NVIDIA và AMD tương ứng.

Vui lòng không xem bài đăng của tôi là quảng cáo NVIDIA. Tôi khá hiểu rằng HD 5850 mạnh hơn GT 240. Điều duy nhất tôi muốn biết là tại sao sự khác biệt đó là và cách khắc phục vấn đề.

Cập nhật. Dưới đây là mã hạt nhân tìm kiếm 6 hình ảnh mẫu có kích thước bằng nhau trong hình cơ sở. Mỗi điểm ảnh của hình ảnh cơ sở được coi là nguồn gốc có thể của một trong các mẫu và được xử lý bằng một chuỗi riêng biệt. Hạt nhân so sánh các giá trị R, G, B của mỗi pixel của hình ảnh cơ sở và của mẫu, và nếu có ít nhất một khác biệt vượt quá thông số diff, thì pixel tương ứng sẽ được tính không khớp. Nếu số lượng pixel không khớp nhỏ hơn maxNonmatchQt thì mẫu tương ứng sẽ được nhấn.

__constant int tOffset = 8196; // one template size in memory (in bytes) 
__kernel void matchImage6(__global unsigned char* image, // pointer to the base image 
      int imgWidth, // base image width 
      int imgHeight, // base image height 
      int imgPitch, // base image pitch (in bytes) 
      int imgBpp, // base image bytes (!) per pixel 
      __constant unsigned char* templates, // pointer to the array of templates 
      int tWidth, // templates width (the same for all) 
      int tHeight, // templates height (the same for all) 
      int tPitch, // templates pitch (in bytes, the same for all) 
      int tBpp, // templates bytes (!) per pixel (the same for all) 
      int diff, // max allowed difference of intensity 
      int maxNonmatchQt, // max number of nonmatched pixels 
      __global int* result, // results 
          ) { 
int x0 = (int)get_global_id(0); 
int y0 = (int)get_global_id(1); 
if(x0 + tWidth > imgWidth || y0 + tHeight > imgHeight) 
    return; 
int nonmatchQt[] = {0, 0, 0, 0, 0, 0}; 
for(int y = 0; y < tHeight; y++) { 
    int ind = y * tPitch; 
    int baseImgInd = (y0 + y) * imgPitch + x0 * imgBpp; 
    for(int x = 0; x < tWidth; x++) { 
     unsigned char c0 = image[baseImgInd]; 
     unsigned char c1 = image[baseImgInd + 1]; 
     unsigned char c2 = image[baseImgInd + 2]; 
     for(int i = 0; i < 6; i++) 
      if(abs(c0 - templates[i * tOffset + ind]) > diff || 
          abs(c1 - templates[i * tOffset + ind + 1]) > diff || 
          abs(c2 - templates[i * tOffset + ind + 2]) > diff) 
       nonmatchQt[i]++; 
     ind += tBpp; 
     baseImgInd += imgBpp; 
    } 
    if(nonmatchQt[0] > maxNonmatchQt && nonmatchQt[1] > maxNonmatchQt && nonmatchQt[2] > maxNonmatchQt && nonmatchQt[3] > maxNonmatchQt && nonmatchQt[4] > maxNonmatchQt && nonmatchQt[5] > maxNonmatchQt) 
     return; 
} 
for(int i = 0; i < 6; i++) 
    if(nonmatchQt[i] < maxNonmatchQt) { 
     unsigned int pos = atom_inc(&result[0]) * 3; 
     result[pos + 1] = i; 
     result[pos + 2] = x0; 
     result[pos + 3] = y0; 
    } 
} 

Kernel cấu hình chạy: kích thước toàn cầu công việc = (1900, 1200) kích thước làm việc địa phương = (32, 8) cho AMD và (32, 16) cho NVIDIA.

Thời gian thực hiện: HD 5850 - 69 ms, HD 6670 - 200 ms, GT 240 - 100 ms.

Mọi nhận xét về mã của tôi cũng được đánh giá cao.

Answer 1

Có thể không có câu trả lời chính xác hoàn hảo cho việc này. Hiệu suất OpenCL phụ thuộc vào nhiều thông số. Số lượng truy cập vào bộ nhớ toàn cầu, hiệu quả của mã vv. Ngoài ra, nó rất khó so sánh giữa hai thiết bị vì chúng có thể có các ký ức cục bộ, toàn cục, liên tục khác nhau. Số lượng lõi, tần số, băng thông bộ nhớ, quan trọng hơn là kiến ​​trúc phần cứng, v.v.

Mỗi phần cứng cung cấp tăng hiệu suất riêng, ví dụ như native_ từ NVIDIA. Vì vậy, bạn cần phải khám phá thêm về phần cứng mà bạn đang làm việc, điều đó thực sự có thể hiệu quả. Nhưng những gì tôi muốn giới thiệu cá nhân là không sử dụng tối ưu hóa phần cứng cụ thể như vậy nó có thể ảnh hưởng đến sự linh hoạt của mã của bạn.

Bạn cũng có thể tìm thấy một số bài báo được công bố cho thấy, hiệu suất CUDA tốt hơn nhiều so với hiệu năng OpenCL trên cùng một phần cứng NVIDIA.

Vì vậy, luôn luôn tốt hơn để viết mã cung cấp tính linh hoạt tốt, thay vì tối ưu hóa thiết bị cụ thể.

Answer 2

Sự khác biệt về thời gian thực hiện là do trình biên dịch gây ra. Mã của bạn có thể dễ dàng được vector hóa. Xem xét hình ảnh và mẫu dưới dạng mảng của kiểu vector char4 (tọa độ của mỗi véc tơ char4 luôn là 0).Thay vì 3 bộ nhớ đọc:

unsigned char c0 = image[baseImgInd]; 
unsigned char c1 = image[baseImgInd + 1]; 
unsigned char c2 = image[baseImgInd + 2]; 

sử dụng chỉ có một:

unsigned char4 c = image[baseImgInd]; 

Thay vì cồng kềnh nếu:

if(abs(c0 - templates[i * tOffset + ind]) > diff || 
       abs(c1 - templates[i * tOffset + ind + 1]) > diff || 
       abs(c2 - templates[i * tOffset + ind + 2]) > diff) 
     nonmatchQt[i]++; 

sử dụng nhanh:

unsigned char4 t = templates[i * tOffset + ind]; 
    nonmatchQt[i] += any(abs_diff(c,t)>diff); 

Vì vậy bạn tăng hiệu suất mã của bạn tối đa 3 lần (nếu compi ler không tự vector hóa chính nó). Tôi cho rằng trình biên dịch AMD OpenCL không vectorization và tối ưu hóa khác. Từ kinh nghiệm của tôi OpenCL trên GPU NVIDIA thường có thể được thực hiện nhanh hơn so với CUDA, bởi vì nó là cấp thấp hơn.