OpenCL - 爲什麼要使用READ_ONLY或WRITE_ONLY緩衝區

在OpenCL中，標記緩衝區是否有任何性能優勢，如READ_ONLY或WRITE_ONLY？OpenCL - 爲什麼要使用READ_ONLY或WRITE_ONLY緩衝區

這kernel是我經常看到（一個是READ_ONLY和b是WRITE_ONLY）：

__kernel void two_buffer_double(__global float* a, __global float* b) 
{ 
    int i = get_global_id(0); 
    b[i] = a[i] * 2; 
}

這kernel似乎更好，因爲它使用較少的全局內存（一個是READ_WRITE）：

__kernel void one_buffer_double(__global float* a) 
{ 
    int i = get_global_id(0); 
    a[i] = a[i] * 2; 
}

是否存在READ_ONLY和WRITE_ONLY標誌僅用於幫助調試和捕獲錯誤？

來源

2013-07-27 benshope

要回答直截了當地你的問題我會說：不，這些標誌不只是存在，以幫助調試和捕捉錯誤。然而，很難給出任何實施方式如何使用這些標誌以及它們如何影響性能的參考。

我的理解（不幸的是沒有任何文檔備份）是使用這些標誌，當你把更多的約束在緩衝區將如何使用，因此你可以幫助運行/驅動器/編譯器做一些假設可能會改善表演。例如我想，因爲工作項不應該寫入內核，所以在內核使用它時，不應該擔心與只讀緩衝區的內存一致性。因此可以跳過一些檢查...雖然在Opencl中，你應該使用障礙等來自己處理這個問題。

還要注意，由於Opencl 1.2引入了一些其他標誌，這次與主機需要如何訪問緩衝區相關。主要有：

CL_MEM_HOST_NO_ACCESS, 
CL_MEM_HOST_{READ, WRITE}_ONLY, 
CL_MEM_{USE, ALLOC, COPY}_HOST_PTR

我猜一遍它必須幫助實現的OpenCL，以提高性能的人，但我想我們需要一些AMD或NVIDIA專家的輸入。

請注意，到目前爲止我所說的全部是只有我的想法，並沒有根據任何嚴重的文件（我沒有設法找到任何）。

在另一方面，我可以告訴你肯定的是，標準不被迫只讀緩衝器是在恆定的空間@Quonux說。這可能是一些實現爲小緩衝區執行此操作。我們不要忘記，恆定的空間內存很小，所以你可以只讀緩衝區太大而不適合。確保緩衝區在恆定空間內存中的唯一方法是在內核代碼中使用常量關鍵字，如here所述。當然，在主機端，如果你想使用常量緩衝區，你必須使用只讀標誌。

來源

2013-07-27 21:21:36 CaptainObvious

這取決於，

一個READ_ONLY __global存儲器位置存儲在「全局/恆定存儲器數據高速緩存」，這是比在GPU上的正常高速緩存或RAM更快多（見here），一個CPU它上沒關係。

我不知道WRITE_ONLY的任何優勢，也許它也有幫助，因爲GPU知道它可以將數據流出來，而不需要緩存。

剛去，如果你不確定衡量...

來源

2013-07-27 19:31:07 Quonux

請注意，其實有兩種。在分配緩衝區時，您有CL_MEM_READ_ONLY,CL_MEM_WRITE_ONLY和CL_MEM_READ_WRITE，但是您也有__read_only,__write_only和__read_write用於在內核代碼中修飾您的指針。

這些可以用於優化和錯誤檢查。讓我們先看看演出。如果遇到只寫緩衝區，寫入不需要被緩存（如在寫入緩存中一樣），爲讀取保存更多緩存。這很大程度上取決於GPU硬件，至少NVIDIA硬件確實需要實際執行此指令（.cs和.lu修改器）。你可以參考他們的PTX ISA。我還沒有看到實際進行優化編譯器的任何證據，如：

__kernel void Memset4(__global __write_only unsigned int *p_dest, 
    const unsigned int n_dword_num) 
{ 
    unsigned int i = get_global_id(0); 
    if(i < n_dword_num) 
     p_dest[i] = 0; // this 
}

被編譯爲：

st.global.u32 [%r10], %r11; // no cache operation specified

這使得作爲CUDA沒有爲那些預選賽所以編譯器等同意義很可能會默默地忽略這些。但把它們放在那裏並沒有什麼壞處，我們未來可能會更幸運。在CUDA中，使用__ldg函數以及使用編譯器標誌來選擇是否在L1中緩存全局內存傳輸（-Xptxas -dlcm=cg），從而公開了一些此功能。如果發現繞過緩存產生主要優勢，您也可以始終使用asm。

至於錯誤檢查，在內核聲明中使用const說明符很容易避免寫入只讀緩衝區。在純粹的「C」中不允許從只寫緩衝區讀取數據。

將這些緩衝區映射到主機內存時發生另一種可能的優化。映射CL_MEM_READ_ONLY緩衝區時，映射區域可能會保持未初始化狀態，因爲主機只會寫入該內存，因此設備只能讀取該內存。同樣，在取消映射CL_MEM_WRITE_ONLY緩衝區時，驅動程序不需要將主機內存（可能由主機修改）的內容從主機內存複製到設備內存。我沒有衡量這一點。

作爲一個方面說明，我已經嘗試使用：

inline unsigned int n_StreamingLoad(__global __read_only const unsigned int *p_src) 
{ 
#ifdef NVIDIA 
    unsigned int n_result; 
    asm("ld.global.cs.u32 %r0, [%r1];" : "=r" (n_result) : "r" (p_src)); 
    return n_result; 
#else // NVIDIA 
    return *p_src; // generic 
#endif // NVIDIA 
} 

inline void StreamingWrite(__global __write_only unsigned int *p_dest, const unsigned int n_value) 
{ 
#ifdef NVIDIA 
    asm("st.global.cs.u32 [%r0], %r1;" : : "r" (p_dest), "r" (n_value) : "memory"); 
#else // NVIDIA 
    *p_dest = n_value; // generic 
#endif // NVIDIA 
}

它給你約15多GB /秒甚至與sm_35設備的簡單的memcpy內核（對GTX 780和K40測試）。在sm_30上沒有看到明顯的加速（不知道它是否意味着在那裏支持 - 雖然指令不會從ptx中剝離）。請注意，您需要自己定義NVIDIA（或參閱Detect OpenCL device vendor in kernel code）。

來源

2015-12-10 14:30:49

OpenCL - 爲什麼要使用READ_ONLY或WRITE_ONLY緩衝區

回答

相關問題