兩個幾乎相同的OpenCL內核之間的性能差距

Question

我有兩個幾乎相同的OpenCL內核，我想用GFLOPS來計算它們的性能。內核＃1：

__kernel void Test41(__global float *data, __global float *rands, int index, int rand_max){ 

    float16 temp; 
    int gid = get_global_id(0); 

    temp = data[gid]; 
    temp = (float) rands[1] * temp; 
    temp = (float) rands[2] * temp; 
    temp = (float) rands[3] * temp; 
    temp = (float) rands[4] * temp; 
    . 
    . 
    . 
    temp = (float) rands[497] * temp; 
    temp = (float) rands[498] * temp; 
    temp = (float) rands[499] * temp; 
    data[gid] = temp.s0; 

} 

第二核心是：

__kernel void Test42(__global float *data, __global float *rands, int index, int rand_max){ 

    float16 temp[500]; 
    int gid = get_global_id(0); 

    temp[0] = data[gid]; 
    temp[1] = (float) rands[1] * temp[0]; 
    temp[2] = (float) rands[2] * temp[1]; 
    temp[3] = (float) rands[3] * temp[2]; 
    temp[4] = (float) rands[4] * temp[3]; 
    . 
    . 
    . 
    temp[497] = (float) rands[497] * temp[496]; 
    temp[498] = (float) rands[498] * temp[497]; 
    temp[499] = (float) rands[499] * temp[498]; 
    data[gid] = temp[index].s0; 

} 

正如你可以在代碼中看到的，我使用的16流大小每個內核具有500線運營，其中的每一個其中只有一個浮點操作。我還總共部署了大約1048576個內核，因此我將有大約1048576個工作項目並行執行。

爲了計算我做的觸發器：

flops = #numWorkItems(1048576) * (500) * StreamSize(16)/timeTaken; 

不幸的是，第一個內核我避開1.4 TFLOPs的，但對於第二個內核，我得到38個GFLOPS。我無法解釋這個巨大的差距。使用一個溫度矢量而不是單個溫度似乎是一個巨大的交易。也似乎真正的應用程序大多像第二個內核。第一個內核對於真正的應用來說太簡單了。

任何人都可以幫助我理解這裏究竟發生了什麼，以及第二個內核性能如何能夠達到第一個？一般來說，如果我要對我的設備進行基準測試，我希望看到性能接近理論值嗎？

P.S.我知道我需要將rands複製到__local內存中，但現在讓我們跳過。

Answer 1

有兩個可能的問題：

• 你宣佈float16臨時緩衝區爲__private {這是OpenCL的默認}，並最有可能將在global內存空間與相當高的訪問延遲進行分配。如果它適合您的設備本地內存，您可能會嘗試將其聲明爲__local float16。
• 添加temp緩衝創造了編譯器的一些問題...原代碼是在某些GPU架構易於量化的（英特爾爲例），並通過添加store + load

其實我提交添加人工的依賴爲編譯器發佈報告。編譯器應該很容易找出依賴關係，進行優化，甚至擺脫您的temp緩衝區。

Answer 2

由於@pmdj在評論中提出，第二個內核的主要問題是註冊壓力：您正在使用大量硬件寄存器，這會減少同時執行的工作組數量。 一般來說，大型私有數組在OpenCL/CUDA中並不是一個好主意。 在這種情況下，編譯器可以做的很少，以優化性能。 您可以使用本地內存作爲陣列，但是您需要添加適當的同步來訪問它。