的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 - 复制由全局结构到本地内存

回答

相关问题