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

Question

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

3

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

我看过一些模板化动态分配的全局和共享内存变量的示例。我意识到常量内存是静态的，因此通常无法进行模板化，正如在这篇文章中所述：Defining templated constant variables in cuda。

我一直找不到任何解决这个常量内存问题的方法，这让我感到惊讶，因为我确定我不是第一个遇到这个问题的人。目前，如果我想使用常量内存，似乎我面临着必须编写两份相同应用程序的问题，一份针对双精度，一份针对浮点数。我希望这不是真的。

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

- Michael Puglia

1个回答

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- Robert Crovella · Accepted Answer

请注意，本答案仅具有历史参考价值，或供使用CUDA Toolkit版本6.5或更早版本的用户使用。自CUDA 7.0以来，不再支持仅支持float的CUDA设备，因此nvcc CUDA编译器不再保留以下所述自动降级double为float的功能。

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