3

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

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

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

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

4

1 回答 1

2

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

由于您提到您只关心floatand double，并且您提到您只关心float不支持的设备double，因此您似乎可以利用nvcc编译器将 double 自动降级为 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
$
于 2013-10-21T14:21:04.170 回答