请注意,实际上有两种类型。在分配缓冲区时,您有
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;
}
编译后的结果为:
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
return *p_src;
#endif
}
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
*p_dest = n_value;
#endif
}
即使在使用sm_35
设备进行简单的memcpy内核操作时,也可以提供约15个额外的GB / sec(已在GTX 780和K40上进行了测试)。在sm_30
上没有看到明显的加速(不确定甚至是否支持该版本-虽然指令未从ptx中删除)。请注意,您需要自己定义NVIDIA
(或参见在内核代码中检测OpenCL设备供应商)。
error: access qualifier can only be used for pipe and image type
- undefined