Cuda - OpenCL CPU 4x nhanh hơn so với OpenCL hoặc CUDA phiên bản GPU

Question

Một mô phỏng sóng Tôi đã làm việc trên với C# + Cudafy (C# -> CUDA hoặc OpenCL dịch) hoạt động tốt, ngoại trừ một thực tế rằng chạy OpenCL Phiên bản CPU (trình điều khiển Intel, 15 "MacBook Pro Retina i7 2.7GHz, GeForce 650M (Kepler, 384 lõi)) nhanh gấp bốn lần so với phiên bản GPU.

(Điều này xảy ra cho dù tôi sử dụng GPU CL hoặc CUDA Phiên bản OpenCL GPU và CUDA hoạt động gần như giống hệt nhau.)

Để làm rõ, đối với một sự cố mẫu:

• CPU OpenCL 1200 Hz
• OpenCL GPU 320 Hz
• CUDA GPU - ~ 330 Hz

Tôi đang ở một mất mát để giải thích lý do tại sao phiên bản CPU sẽ nhanh hơn GPU. Trong trường hợp này, mã hạt nhân đang thực thi (trong trường hợp CL) trên CPU và GPU giống hệt nhau. Tôi chọn thiết bị CPU hoặc GPU trong quá trình khởi tạo, nhưng ngoài ra, mọi thứ đều giống nhau.

Sửa

Dưới đây là mã C# mà ra mắt một trong những hạt nhân. (Những người khác là rất giống nhau.)

public override void UpdateEz(Source source, float Time, float ca, float cb) 
    { 
     var blockSize = new dim3(1); 
     var gridSize = new dim3(_gpuEz.Field.GetLength(0),_gpuEz.Field.GetLength(1)); 

     Gpu.Launch(gridSize, blockSize) 
      .CudaUpdateEz(
       Time 
       , ca 
       , cb 
       , source.Position.X 
       , source.Position.Y 
       , source.Value 
       , _gpuHx.Field 
       , _gpuHy.Field 
       , _gpuEz.Field 
      ); 

    } 

Và, đây là chức năng hạt nhân CUDA có liên quan được tạo ra bởi Cudafy:

extern "C" __global__ void CudaUpdateEz(float time, float ca, float cb, int sourceX, int sourceY, float sourceValue, float* hx, int hxLen0, int hxLen1, float* hy, int hyLen0, int hyLen1, float* ez, int ezLen0, int ezLen1) 
{ 
    int x = blockIdx.x; 
    int y = blockIdx.y; 
    if (x > 0 && x < ezLen0 - 1 && y > 0 && y < ezLen1 - 1) 
    { 
     ez[(x) * ezLen1 + (y)] = ca * ez[(x) * ezLen1 + (y)] + cb * (hy[(x) * hyLen1 + (y)] - hy[(x - 1) * hyLen1 + (y)]) - cb * (hx[(x) * hxLen1 + (y)] - hx[(x) * hxLen1 + (y - 1)]); 
    } 
    if (x == sourceX && y == sourceY) 
    { 
     ez[(x) * ezLen1 + (y)] += sourceValue; 
    } 
} 

Chỉ cần cho đầy đủ, đây là C# được sử dụng để tạo ra các CUDA:

[Cudafy] 
    public static void CudaUpdateEz(
     GThread thread 
     , float time 
     , float ca 
     , float cb 
     , int sourceX 
     , int sourceY 
     , float sourceValue 
     , float[,] hx 
     , float[,] hy 
     , float[,] ez 
     ) 
    { 
     var i = thread.blockIdx.x; 
     var j = thread.blockIdx.y; 

     if (i > 0 && i < ez.GetLength(0) - 1 && j > 0 && j < ez.GetLength(1) - 1) 
      ez[i, j] = 
       ca * ez[i, j] 
       + 
       cb * (hy[i, j] - hy[i - 1, j]) 
       - 
       cb * (hx[i, j] - hx[i, j - 1]) 
       ; 

     if (i == sourceX && j == sourceY) 
      ez[i, j] += sourceValue; 
    } 

Rõ ràng, các if trong hạt nhân này là xấu, nhưng ngay cả các gian hàng đường ống dẫn đến không nên gây ra một đồng bằng hiệu suất cực đoan.

Điều duy nhất khác nhảy vào tôi là tôi đang sử dụng lược đồ phân bổ lưới/khối què - nghĩa là lưới là kích thước của mảng được cập nhật và mỗi khối là một chuỗi. Tôi chắc chắn điều này có một số tác động đến hiệu suất, nhưng tôi không thể nhìn thấy nó gây ra nó là 1/4 tốc độ của mã CL chạy trên CPU. ARGH!

Answer 1

Trả lời câu hỏi này để thoát khỏi danh sách chưa được trả lời.

Mã được đăng cho biết rằng khởi chạy hạt nhân chỉ định một chuỗi chặn 1 (chủ động). Đây không phải là cách để viết mã GPU nhanh, vì nó sẽ khiến hầu hết khả năng của GPU không hoạt động.

Kích thước chuỗi quảng cáo điển hình phải có ít nhất 128 luồng cho mỗi khối và cao hơn thường là tốt hơn, trong bội số của 32, tối đa giới hạn 512 hoặc 1024 cho mỗi khối, tùy thuộc vào GPU.

GPU "thích" để ẩn độ trễ bằng cách có nhiều công việc song song "có sẵn". Chỉ định thêm chủ đề cho mỗi khối hỗ trợ với mục tiêu này. (Có số lượng luồng hợp lý lớn trong lưới cũng có thể hữu ích.)

Ngoài ra, GPU thực hiện chuỗi trong nhóm 32.Chỉ định 1 thread cho mỗi block hoặc một non-multiple của 32 sẽ để lại một số slot thực thi nhàn rỗi, trong mỗi threadblock được thực hiện. 1 chuỗi cho mỗi khối đặc biệt xấu.