用于使用常量内存的 CUDA 应用程序的模板

Question

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

我已经看到了模板化动态分配的全局和共享内存变量的示例。而且我意识到常量内存是静态的，因此模板通常是不可能的，如这篇文章所述：Defining templated constant variables in cuda。

我一直无法找到解决此持续内存问题的任何解决方法，这让我感到惊讶，因为我确信我不是第一个遇到此问题的人。目前，如果我想使用常量内存，我似乎不得不为同一个应用程序编写两份副本，一份用于双打，一份用于浮动。我希望不是这样。

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

Answer 1

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

由于您提到您只关心float 和double，并且您提到您只关心不支持double 的设备上的float，看来您可以采取利用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
$