求问关于常量存储器变量printf打印问题


#include <cuda_runtime.h>
#include <helper_cuda.h>
#include <stdio.h>
#include <stdlib.h>

__constant__ int a_dev[5];
__global__ void test(int *a_dev)
{
for(int i=0;i<5;i++)
{
   printf("%d ",a_dev[i]);
}
}

int main()
{
int a[5]={1,2,3,4,5};

checkCudaErrors(cudaMemcpyToSymbol(a_dev,a,sizeof(int)*5));
printf("\n");
for(int i=0;i<5;i++)
{
printf("%d ",a[i]);
}
printf("\n");

for(int i=0;i<5;i++)
{
printf("%d ",a_dev[i]);
}
printf("\n");
test<<<1,1>>>(a_dev);
cudaDeviceSynchronize();


return 0;
}

结果如下:
1 2 3 4 5
0 0 0 0 0
请按任意键继续. . .

问题1
为什么我把a数据拷贝到常量内存中,打印出来全是0。。。?(如果是拷贝到全局存储器里,打印就没问题。。。)
问题2
为什么在test()函数里面打印的常量存储器a_dev不显示呢?我已经把计算能力参数改了,显卡计算能力为2.1
之前试过全局存储器,可以打印出来的

坐等啊:2_34:

LZ您好:

您上述代码中,对__constant__定义的数组用法有误,所以会造成您的结果。

__constant__定义的数组/变量隐含了__device__的属性,它相当于一个device端的全局的数组/变量,您无需给kernel传递参数,便可以在kernel内直接使用这个数组/变量。

以及,如果您打算通过参数传递的方法使用该数组,那么请在host端定义一个指针int * p_a_dev,并使用cudaGetSymbolAddress()来得到device端a_dev的实际地址,并保存于p_a_dev中,再将p_a_dev作为kernel的参数。

第三,如果您在host端直接使用a_dev,那么此处是特殊处理的,并非是device端的那个a_dev。所以您直接在host端打印是全零,以及此时kernel的参数应该是不正确的,kernel并无执行(虽然您后面添加了同步),您可以定义cudaError_t err,用err收集cudaDeviceSynchronize()的返回值,并使用cudaGetErrorString()来获得err对应的错误名称,看看kernel是否顺利执行了。

此外,您使用cudaMemcpyToSymbol()这个是正确的。

大致如此,祝您编码顺利~

LZ您好:

已在3#详细回复,刚刚在写回复的时候被您插楼了,请您火速研究一下3#。

祝您好运~

太感谢版主了~~学到了很多:2_31:

不客气的,我和横扫斑竹都十分乐意为您服务。

欢迎您常来论坛。

祝您编码顺利~