常数存储器遇到个问题

我在最外面声明:
#define NUM_STENCIL 3
constant float stencil[NUM_STENCIL];

在主机端:
coeff是一个大小为3的float型数组
checkCudaErrors(cudaMemcpyToSymbol(stencil, coeff,
NUM_STENCIL*sizeof(float)));

设备端:
float value = stencil[0] * input[inputIndex]

  • stencil[1] * input[inputIndex-1]
  • stencil[2] * input[inputIndex - stride_y];

从运行结果来看,发现设备端stencil值好像都为0,如果不使用常数存储器就没问题,这是怎么回事?

补充:
我的文件结构:

a.h
里面是constant声明

b.cu
#include <a.h>
主机端代码

c.cu
#include <a.h>
设备端代码
我在网上查到,好像有人说是文件命名空间的问题

现在如果改成:
a.h
去掉constant声明
b.cu
#include <a.h>
#include <c.cu>
c.cu
把constant声明放在这里

这样子就没问题了

能不能详细解释一下为什么?如果我想采用前一种文件安排方式,应该怎么处理?

请在同一个文件中使用__device__和__constant__。

你这样在2个文件中使用,可能会实际上导致你的a.cu用的stencil和b.cu用的stencil是2个独立的__constant__变量。谢谢。

以及,如果你需要在分开的2个文件中使用,您可以尝试这样:

(1)
将你的头文件中的:
constant …;
改成:
extern constant …;

(2)你的2个.cu都#include这个头文件。

(3)在任意的一个.cu里写:constant …;
(注意一个头文件写了,另外一个就不要写了)

然后重新编译。

按照你给的写法:
编译不通过
fdtdKernel.cu:3: error: redefinition of ‘float stencil [3]’
FDTD.h:50: error: ‘float stencil [3]’ previously declared here

显然你写错了(根据你的错误提示得出此结论)。

你不按我说的写,能通过么?

请重复看上文!

你的意思是这样子?

a.h
extern constant float stencil[3];

b.cu --主机端代码
include <a.h>
不写

c.cu --设备端代码
include <a.h>
constant float stencil[3];

另外,我看到cuda-c-programming-guide里面:

When compiling in the whole program compilation mode (see the nvcc user manual for a description of this mode), device, shared, and __constant__variables cannot be defined as external using the extern keyword. The only exception is for dynamically allocated shared variables as described in shared.

When compiling in the separate compilation mode (see the nvcc user manual for a description of this mode), device, shared, and constant variables can be defined as external using the extern keyword. nvlink will generate an error when it cannot find a definition for an external variable (unless it is a dynamically allocated shared variable).

会不会是这两种模式的问题,导致extern失效?第二种模式是cuda 5.0刚添加的,要特别设置编译选项,但我看不懂。。。

您最后的写法是正确的,以及,您需要右键点击您的.cu选中Generate Relocatable Device Code的,这个是对的。

您之前的失败不是您的模式选择的问题,而且您根本就没写对!
您都重复定义了。显然不是我告诉你的写法。

以及,此问题在本论坛讨论了很多很多次了。
您可以参考我之前的若干帖子:
http://cudazone.nvidia.cn/forum/forum.php?mod=viewthread&tid=6344