CUDA常量内存最佳实践

Question

CUDA常量内存最佳实践

25

我在这里呈现一些代码

__constant__ int array[1024];

__global__ void kernel1(int *d_dst) {
   int tId = threadIdx.x + blockIdx.x * blockDim.x;
   d_dst[tId] = array[tId];
}

__global__ void kernel2(int *d_dst, int *d_src) {
   int tId = threadIdx.x + blockIdx.x * blockDim.x;
   d_dst[tId] = d_src[tId];
}

int main(int argc, char **argv) {
   int *d_array;
   int *d_src;
   cudaMalloc((void**)&d_array, sizeof(int) * 1024);
   cudaMalloc((void**)&d_src, sizeof(int) * 1024);

   int *test = new int[1024];
   memset(test, 0, sizeof(int) * 1024);

   for (int i = 0; i < 1024; i++) {
     test[i] = 100;
   }

   cudaMemcpyToSymbol(array, test, sizeof(int) * 1024);
   kernel1<<< 1, 1024 >>>(d_array);

   cudaMemcpy(d_src, test, sizeof(int) * 1024, cudaMemcpyHostToDevice);
   kernel2<<<1, 32 >>>(d_array, d_src),

   free(test);
   cudaFree(d_array);
   cudaFree(d_src);

   return 0;
}

这只是显示了常量内存和全局内存的使用情况。在执行过程中，“kernel2”比“kernel1”快大约4倍（根据时间计算）。

根据Cuda C编程指南，我理解这是因为对常量内存的访问被序列化。这让我想到，如果一个warp访问单个常量值（如整数、浮点数、双精度等），那么可以最好地利用常量内存，但是访问数组则没有任何好处。换句话说，为了从常量内存访问中获得任何有益的优化/加速收益，一个warp必须访问单个地址。这样说正确吗？

我还想知道，如果我在常量内存中保留一个结构而不是一个简单类型。在warp内部的线程访问结构的任何位置是否也被视为单个内存访问或多个内存访问？我的意思是，结构可能包含多个简单类型和数组；当访问这些简单类型时，这些访问是否也被序列化？

最后一个问题是，如果我有一个包含常量值的数组，需要由warp内的不同线程访问；为了更快地访问，它应该放在全局内存中而不是常量内存中。这正确吗？

有人能给我推荐一些演示有效使用常量内存的示例代码吗？

问候，

- Psypher

1个回答

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- Robert Crovella · Accepted Answer

我可以说，为了从常量内存访问中获得任何有益的优化/加速收益，一个warp必须访问单个地址。这是正确的吗？

是的，这通常是正确的，并且是使用常量内存/常量缓存的主要目的。常量缓存可以每次为SM提供一定数量的数据。精确措辞如下:

常量内存空间驻留在设备内存中，并被缓存在常量缓存中。

然后将请求分成与初始请求中不同的内存地址一样多的单独请求，将吞吐量降低一个等于单独请求数的因素。

然后，在缓存命中情况下以常量缓存的吞吐量或者在设备内存的吞吐量下服务于结果请求。

上文中的重要内容是希望实现对warp的统一访问，以达到最佳性能。如果一个warp向__constant__内存发出请求，其中不同线程在warp中访问不同位置，这些请求将会被串行化。因此，如果每个线程在warp中访问相同的值：

int i = array[20];

如果每个warp线程都在访问不同的量，那么您将有机会从常数缓存/内存中获得良好的收益。

int i = array[threadIdx.x];

如果访问被序列化，那么常量数据的使用效率会令人失望。

我还想知道，如果我在常量内存中保留一个结构而不是简单类型。线程在warp内对结构的任何访问; 也被视为单个内存访问还是更多？

您可以将结构放入常量内存中。相同的规则适用：

int i = constant_struct_ptr->array[20];

有机会受益，但是

int i = constant_struct_ptr->array[threadIdx.x];

如果您在不同的线程中访问相同的简单类型结构元素，则这对于常量缓存的使用是理想的。

最后一个问题是，如果我有一个包含常量值的数组，需要通过不同的线程在warp内访问；为了更快的访问，它应该保留在全局内存中而不是常量内存中。这正确吗？

是的，如果您知道通常情况下您的访问会打破每个周期一个32位数量的常量内存规则，那么最好将数据保留在普通的全局内存中。

有各种cuda示例代码演示了__constant__数据的用法。以下是其中几个：

还有其他的。

编辑：回答评论中的一个问题，如果我们在常量内存中有这样的结构：

struct Simple { int a, int b, int c} s;

我们可以像这样访问它：

int p = s.a + s.b + s.c;
          ^     ^     ^
          |     |     |
cycle:    1     2     3

我们将充分利用常量内存/缓存。当C代码被编译时，在底层会生成与上图中的1、2、3对应的机器码访问。假设访问1首先发生。由于访问1是对同一内存位置的访问，不管warp中的哪个线程，在第一个周期中，所有线程都将接收s.a中的值，并且它将充分利用缓存以获得最佳效益。对于访问2和3也是如此。另一方面，如果我们有：

struct Simple { int a[32], int b[32], int c[32]} s;
...
int idx = threadIdx.x + blockDim.x * blockIdx.x;
int p = s.a[idx] + s.b[idx] + s.c[idx];

这样做不利于常量内存/缓存的使用。相反，如果我们对s的访问是典型的，那么在普通全局内存中定位s可能会有更好的性能表现。