4

我在常量内存中有一个数组（它是一个全局变量），并通过函数调用 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代码，将全局内存加载到寄存器：

  • 使用了 2 条指令ld.global.s32，将 4 个字节加载到寄存器中。（在使用参考的代码中）
  • 使用1 条指令ld.global.v2.s32，将 8 个字节加载到 2 个寄存器，（在使用全局变量的代码中）

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

4

1 回答 1

2

与全局内存不同，对常量内存的访问将被序列化（分成多个事务），如果它们不统一（一个（半计算能力 1.x）warp 的所有线程访问相同的地址。

因此，如果访问可能是统一的，则仅使用常量内存。

于 2012-10-17T12:10:11.773 回答