cuda常量记忆参考

2019年8月4日 164次阅读

我在常量内存中有一个数组(它是一个全局变量),并通过函数调用cudaGetSymbolAddress获得对它的引用.当我使用此引用来获取常量数据而不是使用全局变量时,我的内核运行缓慢.这是什么原因？

__constant__ int g[2] = {1,2};
// __device__ int g[2] = {1,2};

// kernel: use by reference
__global__ void add_1( int *a, int *b, int *c, int *f )
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    c[tid] = f[0] * a[tid] + f[1] * b[tid];
}

// kernel: use global variable
__global__ void add_2( int *a, int *b, int *c, int *f )
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    c[tid] = g[0] * a[tid] + f[1] * b[tid];
}

int main()
{
    ......
    // a,b,c are large arrays in device memory of size 40960.

    int *f;
    cudaGetSymbolAddress( (void **)&f, (char *)&g);

    add_1 <<< 160, 256 >>> ( a, b, c, f );

    ......
}

这是示例代码,warp中的所有线程同时加载相同的位置.注释代码是通过直接访问常量内存

解释为什么不使用常量内存缓存(由talonmies)

原因是缺乏恒定的缓存.
缓存访问仅在编译器在标记为处于常量状态空间的显式变量上发出特定PTX指令(ld.const)时发生.编译器知道这样做的方式是当一个变量被声明为__constant__时 – 它是一个静态的编译时属性,它会影响代码生成.在运行时不能发生相同的过程.

如果在全局内存中传递指针并且编译器无法确定指针处于常量状态空间,则它将不会生成正确的PTX以通过常量缓存访问该内存.结果访问速度会变慢.

没有回答的问题

为什么即使将数组g声明为__device__变量,代码在使用时也会变慢.通过查看PTX代码,将全局内存加载到寄存器：

>使用ld.global.s32的2条指令,它将4个字节加载到寄存器. (在代码中使用参考)
>使用ld.global.v2.s32的1条指令,它将8个字节加载到2个寄存器,(在使用全局变量的代码中)

有什么区别,任何文档参考将不胜感激？

最佳答案 与全局内存不同,对常量内存的访问将被序列化(拆分为多个事务),如果它们不一致(a(计算能力1.x的一半)的所有线程都会转换为相同的地址.

因此,如果访问可能是一致的,则只使用常量内存.

点赞