OpenCL的代價get_local_id（）

我有一個簡單的掃描內核，它可以計算循環中幾個塊的掃描。我注意到，當get_local_id（）存儲在局部變量中而不是在循環內調用時，性能有所提高。因此，爲了與代碼總結，這樣的：OpenCL的代價get_local_id（）

__kernel void LocalScan_v0(__global const int *p_array, int n_array_size, __global int *p_scan) 
{ 
    const int n_group_offset = get_group_id(0) * SCAN_BLOCK_SIZE; 
    p_array += n_group_offset; 
    p_scan += n_group_offset; 
    // calculate group offset 

    const int li = get_local_id(0); // *** local id cached *** 
    const int gn = get_num_groups(0); 
    __local int p_workspace[SCAN_BLOCK_SIZE]; 
    for(int i = n_group_offset; i < n_array_size; i += SCAN_BLOCK_SIZE * gn) { 
     LocalScan_SingleBlock(p_array, p_scan, p_workspace, li); 

     p_array += SCAN_BLOCK_SIZE * gn; 
     p_scan += SCAN_BLOCK_SIZE * gn; 
    } 
    // process all the blocks in the array (each block size SCAN_BLOCK_SIZE) 
}

有吞吐量74 GB/s的GTX-780，而這一點：

__kernel void LocalScan_v0(__global const int *p_array, int n_array_size, __global int *p_scan) 
{ 
    const int n_group_offset = get_group_id(0) * SCAN_BLOCK_SIZE; 
    p_array += n_group_offset; 
    p_scan += n_group_offset; 
    // calculate group offset 

    const int gn = get_num_groups(0); 
    __local int p_workspace[SCAN_BLOCK_SIZE]; 
    for(int i = n_group_offset; i < n_array_size; i += SCAN_BLOCK_SIZE * gn) { 
     LocalScan_SingleBlock(p_array, p_scan, p_workspace, get_local_id(0)); 
     // *** local id polled inside the loop *** 

     p_array += SCAN_BLOCK_SIZE * gn; 
     p_scan += SCAN_BLOCK_SIZE * gn; 
    } 
    // process all the blocks in the array (each block size SCAN_BLOCK_SIZE) 
}

在相同的硬件上只有70 GB /秒。唯一的區別是對get_local_id（）的調用是在循環內部還是外部。 LocalScan_SingleBlock（）中的代碼在this GPU Gems article中有很多描述。

現在，這帶來了一些問題。我一直認爲線程標識存儲在某個寄存器中，並且對任何線程局部變量的訪問速度都很快。這似乎並非如此。我總是習慣於把一個本地ID緩存在一個變量中，而這個變量不願意讓一個老的「C」程序員不願意在一個循環中調用一個函數，如果他希望每次都返回相同的值，不要認爲這會有所作爲。

任何想法，爲什麼這可能是？我沒有做任何檢查編譯的二進制代碼。有沒有人有相同的經歷？ CUDA中的threadIdx.x與此相同嗎？ ATI平臺如何？這種行爲是在什麼地方描述的？我通過CUDA最佳實踐快速掃描，但沒有發現任何內容。

來源

2013-09-26 the swine

請不要刪除CUDA標記。雖然代碼本身不在CUDA中，但問題體現在NVIDIA硬件上，與CUDA的threadIdx的實現方式密切相關，以及它如何影響程序的runitme。 –

這只是一種猜測，但按照Khronos的頁面

http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/get_local_id.html

get_local_id（）沒有定義爲返回一個恆定值（僅僅爲size_t）。這可能意味着，就編譯器所知，與常量local_id相比，可能不允許執行某些優化，因爲函數值的返回可能會在編譯器的眼中發生變化（即使它不會按每個線程）

來源

2013-09-27 16:17:26

這對NVIDA來說真的很愚蠢，特別是因爲在CUDA中，threadIdx是一個變量而不是函數。通過將get_local_id（）聲明爲宏可以很容易地解決這個問題。此外，人們希望在某處閱讀它。儘管如此，一個很好的猜測。 –

那麼，它不是nvidia什麼opencl規範說，如果問題是編譯器優化與它是一個非const函數，它不能被優化，那麼它可能獨立於threadidx如何在硬件中表示。另外，宏不是恆定的而是不恆定的？根據鏈接中對規範的實際引用，特別是在「內置函數」一節和「與工作項相關的函數」一節中，這意味着它可能另外在技術上不適合作爲宏實現。只是更多的猜測 –

NVIDIA正在編寫這個編譯器。你會發現當涉及到供應商的實現時，規範不是法律:)。我的意思是說，OpenCL編譯器將只是#define get_local_id（coord）（threadIdx.x *（〜（coord | coord >> 1）＆1）+ threadIdx.y * ...），它看起來像一個函數並進行評估到編譯時常量。並不是說他們需要這樣做，但是對於圖像來說可能更簡單。 –

OpenCL的代價get_local_id（）

回答

相關問題