c++ - 使用常量内存的 CUDA 应用程序模板

Question

我正在尝试编写一个针对 float 和 double 模板化的 CUDA 应用程序，因为我希望能够在单精度和 double 卡上运行。应用程序使用动态分配的全局内存、动态分配的共享内存以及常量内存和静态全局内存。

我见过模板化动态分配的全局和共享内存变量的例子。而且我意识到常量内存是静态的，因此通常不可能进行模板化，如这篇文章所述:Defining templated constant variables in cuda .

我一直无法找到解决此持续内存问题的任何解决方法，这让我感到惊讶，因为我确定我不是第一个遇到此问题的人。目前，如果我想使用常量内存，我似乎面临着必须编写同一应用程序的两份拷贝，一份用于 double ，一份用于 float 。我希望情况并非如此。

作为一种解决方法，我正在考虑编写一个(虚拟的？)基类，它是模板化的并实现除常量内存变量声明之外的所有内容。然后我想编写两个从基类继承的类(一个用于 float ，一个用于 double )，主要只处理常量变量声明。我的问题是这个策略是否有效，或者是否存在明显的缺陷？我只是想在实现设计之前问一下，结果发现它不起作用。如果此策略不起作用，是否还有其他经过验证的策略至少可以缓解问题？或者我只需要编写两份应用程序拷贝，一份用于 float，一份用于 double？

Answer 1

请注意，此答案仅具有历史意义，或者适用于使用 CUDA 工具包 6.5 或更早版本的用户。从 CUDA 7.0 开始，没有支持的 CUDA 设备，只支持 float , 所以 nvcc CUDA 编译器不再保留下面描述的自动降级的能力 double至 float .

既然你提到你只关心 float和 double ，你提到你只关心 float在不支持 double 的设备上, 看起来你可以利用 nvcc编译器 automatic demotion of double to float为了处理这个。

这是一个使用 __constant__ 的例子内存:

$ cat t264.cu
#include <stdio.h>

#define DSIZE 64

__constant__ double my_const_data[DSIZE];

__global__ void my_kernel(double *data){
  data[1] = my_const_data[0];
  data[0] = sqrt(my_const_data[0]);
}

int main(){
  double my_data[DSIZE], h_data[DSIZE], *d_data;
  my_data[0] = 256.0;
  cudaMemcpyToSymbol(my_const_data, my_data, sizeof(double)*DSIZE);
  printf("hello\n");
  cudaMalloc((void **)&d_data, sizeof(double)*DSIZE);
  my_kernel<<<1,1>>>(d_data);
  cudaMemcpy(&h_data, d_data, sizeof(double)*DSIZE, cudaMemcpyDeviceToHost);
  printf("data = %lf\n", h_data[1]);
  printf("sqrt = %lf\n", h_data[0]);
  return 0;
}

$ nvcc -o t264 t264.cu
ptxas /tmp/tmpxft_00003228_00000000-5_t264.ptx, line 62; warning : Double is not supported. Demoting to float
$ ./t264
hello
data = 256.000000
sqrt = 16.000000
$