CUDA恒定记忆最佳实践

Question

我在这里介绍一些代码

__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中的不同线程访问; 为了更快地访问它应该保存在全局内存而不是常量内存中.那是对的吗？

任何人都可以向我推荐一些示例代码,其中显示了有效的常量内存使用.

问候,

Answer 1

我可以说warp必须访问单个地址才能从常量内存访问中获得任何有益的优化/加速增益.它是否正确？

是的,这通常是正确的,是使用常量内存/常量缓存的主要目的.常量高速缓存可以为每个SM每个周期提供一个32位数量.因此,如果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];

才不是.如果跨线程访问相同的简单类型结构元素,那么这对于常量缓存使用是理想的.

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

是的,如果你知道一般来说你的访问会打破每个循环规则的一个32位数量的常量内存,那么你可能最好将数据留在普通的全局内存中.

有各种cuda示例代码可以演示__constant__数据的使用.以下是一些:

1. 图形volumeRender
2. 成像双边过滤器
3. 成像卷积纹理
4. 财务MonteCarloGPU

还有其他人.

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

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在普通的全局内存中.