Cuda内核中的静态常量数组

Question

我需要在 Cuda 内核中包含以下内容：

static const float PREDEFINED_CONSTS[16] = {...}; // 16 constants.

float c = PREDEFINED_CONSTS[threadId.x % 16];
/// Use c in computations.

提供 PREDEFINED_CONSTS 的最佳方式是什么？

const 内存看起来不太好，因为不同的线程会访问不同的位置。
如果我这样定义它们，PREDEFINED_CONSTS会存储在全局内存中吗？

这个怎么样：

float c;
if      ( threadId.x % 16 == 0 ) c = VAL0;
else if ( threadId.x % 16 == 1 ) c = VAL1;
...
else if ( threadId.x % 16 ==15 ) c = VAL15;

虽然最后一个例子有线程分歧，文字VAL*值是指令操作码的一部分，所以不会从内存中读取。

Answer 1

What's the best way to provide PREDEFINED_CONSTS ?

如果是我，我会简单地将第一个示例中的内容放入您的 CUDA 内核中，然后继续使用它。这很可能是最好的方法。稍后，如果您觉得您的代码存在性能问题，您可以使用分析器引导您朝着需要解决的方向前进。我怀疑会是这样。对于常量，真的只有两种可能：

从某种内存中加载它们
将它们作为指令流的一部分加载。

您已经表示您知道这一点，如果您真的很担心，可以简单地对两者进行基准测试。基准测试需要的不仅仅是您在此处显示的内容，可能没有定论，并且还可能取决于其他因素，例如您加载这些常量的次数和方式。

正如您已经指出的那样，__constant__ 似乎不是一个明智的选择，因为整个经线的负载模式明显不均匀。

If I define them as above, will PREDEFINED_CONSTS be stored in global memory?

是的，您的第一个方法将存储在全局内存中。这可以通过使用 -Xptxas -v 仔细研究和编译来确认。您的第二种方法有可能（至少）通过指令流加载常量。由于从编码的角度来看第二种方法非常丑陋，而且与第一种方法相比也非常不灵活（如果我在代码的不同位置需要每个线程的不同常量怎么办？），这不是我会选择的。

我认为这是过早的优化。从代码灵活性和简洁性的角度来看，第一种方法显然更受欢迎，并且没有真正的理由认为仅仅因为您是从内存中加载它就是一个问题。第二种方法丑陋、不灵活，从性能角度来看可能也好不到哪里去。即使数据是指令流的一部分，它仍然必须从内存中加载。

这是一个示例测试用例，向我建议第一个案例是首选。如果您提出不同类型的测试用例，您可能会得到不同的观察结果：

$ cat t97.cu
#include <cstdio>
const float VAL0 = 1.1;
const float VAL1 = 2.2;
const float VAL2 = 3;
const float VAL3 = 4;
const float VAL4 = 5;
const float VAL5 = 6;
const float VAL6 = 7;
const float VAL7 = 8;
const float VAL8 = 9;
const float VAL9 = 10;
const float VAL10 = 11;
const float VAL11 = 12;
const float VAL12 = 13;
const float VAL13 = 14;
const float VAL14 = 15;
const float VAL15 = 16;


__global__ void k1(int l){
        static const float PREDEFINED_CONSTS[16] = {VAL0, VAL1, VAL2, VAL3, VAL4, VAL5, VAL6, VAL7, VAL8, VAL9, VAL10, VAL11, VAL12, VAL13, VAL14, VAL15};
        float sum = 0.0;
  for (int i = 0; i < l; i++)
    sum += PREDEFINED_CONSTS[(threadIdx.x+i) & 15];
  if (sum == 0.0) printf("%f\n", sum);
}
__device__ float get_const(int i){
  float c = VAL15;
  unsigned t = (threadIdx.x+i) & 15;
  if      (t == 0)  c = VAL0;
  else if (t == 1)  c = VAL1;
  else if (t == 2)  c = VAL2;
  else if (t == 3)  c = VAL3;
  else if (t == 4)  c = VAL4;
  else if (t == 5)  c = VAL5;
  else if (t == 6)  c = VAL6;
  else if (t == 7)  c = VAL7;
  else if (t == 8)  c = VAL8;
  else if (t == 9)  c = VAL9;
  else if (t == 10) c = VAL10;
  else if (t == 11) c = VAL11;
  else if (t == 12) c = VAL12;
  else if (t == 13) c = VAL13;
  else if (t == 14) c = VAL14;
  return c;
}

__global__ void k2(int l){
        float sum = 0.0;
  for (int i = 0; i < l; i++)
    sum += get_const(i);
  if (sum == 0.0) printf("%f\n", sum);
}

int main(){
        int l = 1048576;
  k1<<<1,16>>>(l);
  k2<<<1,16>>>(l);
  cudaDeviceSynchronize();
}
$ nvcc -o t97 t97.cu -Xptxas -v
ptxas info    : 68 bytes gmem
ptxas info    : Compiling entry function '_Z2k2i' for 'sm_52'
ptxas info    : Function properties for _Z2k2i
    8 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 9 registers, 324 bytes cmem[0], 8 bytes cmem[2]
ptxas info    : Compiling entry function '_Z2k1i' for 'sm_52'
ptxas info    : Function properties for _Z2k1i
    8 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 32 registers, 324 bytes cmem[0]
$ nvprof ./t97
==22848== NVPROF is profiling process 22848, command: ./t97
==22848== Profiling application: ./t97
==22848== Profiling result:
            Type  Time(%)      Time     Calls       Avg       Min       Max  Name
 GPU activities:   91.76%  239.39ms         1  239.39ms  239.39ms  239.39ms  k2(int)
                    8.24%  21.508ms         1  21.508ms  21.508ms  21.508ms  k1(int)
      API calls:   62.34%  260.89ms         1  260.89ms  260.89ms  260.89ms  cudaDeviceSynchronize
                   37.48%  156.85ms         2  78.427ms  10.319us  156.84ms  cudaLaunchKernel
                    0.13%  542.39us       202  2.6850us     192ns  117.71us  cuDeviceGetAttribute
                    0.04%  156.19us         2  78.094us  58.411us  97.777us  cuDeviceTotalMem
                    0.01%  59.150us         2  29.575us  26.891us  32.259us  cuDeviceGetName
                    0.00%  10.845us         2  5.4220us  1.7280us  9.1170us  cuDeviceGetPCIBusId
                    0.00%  1.6860us         4     421ns     216ns     957ns  cuDeviceGet
                    0.00%  1.5850us         3     528ns     283ns     904ns  cuDeviceGetCount
                    0.00%     667ns         2     333ns     296ns     371ns  cuDeviceGetUuid
$

Cuda内核中的静态常量数组

Static const array in Cuda kernel

cuda