的OpenCL - 複製由全局結構到本地內存

我有一個結構原始，它的定義如下：的OpenCL - 複製由全局結構到本地內存

typedef struct Primitive { 
    float m[12]; 
    float invm[12]; 
    enum PrimitiveType type; 
    int rayDensity; 
    float util1;    
    float util2;     
} Primitive;

我通過這些結構的陣列，以我的內核在一個恆定的內存緩衝區：

__constant Primitive *objects;

作爲一種優化練習，我想看看加載結構到本地內存，所以我的內核代碼的這個喜歡的一部分：

__kernel void test(int n_objects, __constant Primitives *objects) { 
    local Primitive pFrom, pTo; 

    for(int i = 0; i < n_objects; i++) { 
     pFrom = objects[i]; 
    } 

}

當我運行此我得到一個編譯錯誤說：

ptxas application ptx input, line 42; error: State space mismatch between instruction and address in instruction 'ld'

作爲一個實驗我已經試過先複製結構爲私有變量，然後局部變量如下：

__kernel void test(int n_objects, __constant Primitives *objects) { 
    Primitive pF, Pt; 
    local Primitive pFrom, pTo; 

    for(int i = 0; i < n_objects; i++) { 
     pF = objects[i] 
     pFrom = pF; 
    } 

}

現在編譯並運行，但似乎該對象沒有深入複製到局部變量pFrom中。

請注意，我的代碼示例純粹是示例，爲了簡潔起見，我已刪除所有內容。當我直接從常量全局內存中使用原始結構時，我的代碼也能正常工作。

有沒有人知道我在這裏錯過了什麼，它肯定是深層複製或OpenCL地址空間的一些基本基礎。

來源

2012-06-18 cubiclewar

你需要的是async_work_group_copy function。您可以等待此異步操作完成使用wait_group_events函數。

希望這會有所幫助。

來源

2012-06-18 14:06:15 Ani

帶有wait_group_events的async_work_group_copy函數是答案，我很驚訝這並沒有出現在我的任何搜索中。奇怪的是，但是將結構加載到本地內存而不是使用常量會增加運行時間。也許來自恆定內存的廣播速度比我想象的要快 – cubiclewar

您是否交錯了下一個項目的副本和處理？這是一個巨大的差異。 – Ani

我已經重新調整了操作順序，以便先前調用async_work_group_copy，然後在所有元素上執行一些工作，然後再執行等待屏障並使用對象。這給了一個小的（大約4％）運行時間減少。然而，這種方法不能很好地擴展數量的結構，因爲它會用完本地內存。 – cubiclewar

的OpenCL - 複製由全局結構到本地內存

回答

相關問題