22

我在这里展示一些代码

__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 编程指南中了解到，这是因为对常量内存的访问正在被序列化。这让我想到，如果扭曲访问单个常量值（例如整数、浮点数、双精度等），则可以最好地利用常量内存，但访问数组根本没有好处。换句话说，我可以说扭曲必须访问单个地址，以便从持续的内存访问中获得任何有益的优化/加速收益。这个对吗？

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

最后一个问题是，如果我确实有一个具有常量值的数组，需要通过扭曲中的不同线程访问它；为了更快地访问它应该保存在全局内存而不是常量内存中。那是对的吗？

任何人都可以向我推荐一些示例代码，其中显示了有效的恒定内存使用情况。

问候，

4

1 回答 1

32

我可以说，warp 必须访问单个地址才能从持续的内存访问中获得任何有益的优化/加速收益。这个对吗？

是的，这通常是正确的，并且是使用常量内存/常量缓存的主要目的。常量缓存可以“一次”为每个 SM 提供一个数量。准确的措辞如下：

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

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

然后在缓存命中的情况下以常量缓存的吞吐量为结果请求提供服务，否则以设备内存的吞吐量提供服务。

从上面的文本中得到的一个重要结论是希望通过 warp 进行统一访问以实现最佳性能。如果 warp 向__constant__内存中的不同线程访问不同位置的请求发出请求，则这些请求将被序列化。因此，如果 warp 中的每个线程都访问相同的值：

int i = array[20];

那么您将有机会从恒定的缓存/内存中获益。如果 warp 中的每个线程都访问一个唯一的数量：

int i = array[threadIdx.x];

然后访问将被序列化，并且持续的数据使用将令人失望，性能方面。

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

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

int i = constant_struct_ptr->array[20];

有机会受益，但

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

才不是。如果您跨线程访问相同的简单类型结构元素，则非常适合常量缓存使用。

最后一个问题是，如果我确实有一个具有常量值的数组，需要通过扭曲中的不同线程访问它；为了更快地访问它应该保存在全局内存而不是常量内存中。那是对的吗？

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

有多种cuda 示例代码演示了__constant__数据的使用。这里有几个：

  1. 图形体积渲染
  2. 成像双边滤波器
  3. 成像卷积纹理
  4. 金融蒙特卡洛GPU

还有其他人。

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

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 中的哪个线程无关，因此在周期 1 期间，所有线程都将接收到该值，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

于 2013-08-02T16:00:58.303 回答