Nsight查看constant memory的值

问题有三:
1、定义全局二维常量数组:constant int de[9][2]={{0,0},{1,0},{0,1},{-1,0},{0,-1},{1,1},{-1,1},{-1,-1},{1,-1}};
在kernel中设置一处断点,启动Nsight调试程序,检测de的值,发现跟上面初始化的不一样
[attach]3272[/attach]
然后我在kernel中将de数组取一元素(假设为de[0][0]),返回给kernel实参,并拷贝给host端进行输出,结果证实跟初始化的值一样。那么Nsight检测的值为何显示错误?!
global void kernel(int *a)
{
*a=de[0][0];
}
2、kernel函数中使用printf时,有时可正确显示,有时却不行,我已把vs2010项目属性相关项改为compute_20,sm_20
3、定义constant二维数组,是否能够像cudaMemcpy一样进行初始化。

@ice 及各位大神,谢谢。

LZ您好:

先说下后两个问题吧:

2:如果您是有时候kernel里面的printf没有显示出来,那么请您在您的程序结束前加上cudaDeviceReset(),这将保证您的所有被缓冲的kernel的printf信息都输出后,程序才结束。

如果是其他问题,欢迎跟帖说明。

3:对于__device__ 声明的device端的全局变量/数组(__constant__自动隐含__device__属性 ),可以使用cudaMemcpyToSymbol系列函数赋值,用cudaMemcpyFromSymbol()回拷数据。

以及可以使用cudaGetSymbolAddress()获取该缓冲区的地址,使用保存在host端的指针指向该地址以后,就可以用这个指针作为普通的cudaMemcpy的参数了。

因为我很少用__device__或者__constant__,所以1:的问题需要稍微上机验证下,随后再回答您。

祝您编码顺利~

我来说下第一个问题:

这个真心奇葩了,按理说不应该出现这种情况的。您确定您将鼠标移动到de上,看到的仍然是这种情况?(请尝试下)。

以及,一种情况可能造成这种,就是您在kernel中写入显存的时候,不小心写多了(例如越界),写到了de所在的那段内存,从而改变了它的值。

我不认为还有其他情况可能会这样。。
以及,请尝试下刚才的建议。

谢ice版主。
2、在main函数结束处加入cudaDeviceReset()后,打印都正常,反之不打印。
3、若我申请一个全局__constant__ int de[9][2],但并不给初始值,然后在host端 int e[9][2]并赋值。
我通过cudaMemcpyToSymbol(de, e, sizeof(int)92) 拷贝后没有问题,de中的值是正确的
但是,若拷贝如下,则不成功:
for (int i=0; i<9; i++)
{
cudaMemcpyToSymbol(de[i], e[i], sizeof(int)*2);
}

1、Nsight对于一维和二维constant数组的监测都有问题,这才是最苦恼的。

LZ您好:

说一下3:

你可以参阅一下 CUDA Toolkit Reference Manual里面关于cudaMemcpyToSymbol()的说明,在copy的时候使用参数设定offset的,而不是直接对指针加上偏移量。

因为这里是特殊处理的,所以不能按照C语言的习惯写。

祝您编码顺利~

。。。。
请勿乱下结论,真心无问题的。我长期使用他无事的。

以及,您的那个:
for (int i=0; i<9; i++)
{
cudaMemcpyToSymbol(de, e, sizeof(int)*2);
}
这个是啥意思?

失误,"[i]"误以为是斜体,应该插入代码。

for (int i=0; i<9; i++)
{
cudaMemcpyToSymbol(de[i], e[i], sizeof(int)*2);
}

你在host上看到的device上的de变量是无法直接加上偏移的,实际上这个是特殊处理的,要么用de, 要么不用,不能直接加偏移的,
你可以这样修改:
for (int i=0; i<9; i++)
{
cudaMemcpyToSymbol(de, e[ i ], sizeof(int)*2, 8 * i );
}

LZ您好:

请您说一下您软件环境情况:CUDA Toolkit版本,VS版本,nsight版本。

我和横扫斑竹都上次测试了一下,他是nsight 3.0+VS2010,确实可以正常显示无任何问题的。
我是nsight 2.2+VS2008 SP1,此时在kernel中看__constant__地址显示0x 00000000值是未知的,但是将这个数组地址复制给一个合适的指针,直接查看这个指针是正确的。

以及,我这里的情况和您截图中的情况也不太一致。

所以,建议您使用更新版本的软件环境,或许能解决这个问题。

祝您编码顺利~

非常感谢@横扫千军、@ice 两位版主对问题3的详细解释。

关于问题1,还是直接上两个测试程序吧。

代码一:

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
__constant__ int de[9][2]={{0,0},{1,0},{0,1},{-1,0},{0,-1},{1,1},{-1,1},{-1,-1},{1,-1}};

__global__ void kernel()
{
	printf("hello\n");
	printf("%d\n",de[8][1]);
}
int main()
{
	kernel<<<1,1>>>();
	cudaDeviceReset();
	return 0;
}

Nsight调试正常。kernel中打印无问题。
[attach]3275[/attach][indent]
代码二:

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>

__constant__ int de[9][2]={{0,0},{1,0},{0,1},{-1,0},{0,-1},{1,1},{-1,1},{-1,-1},{1,-1}};
__constant__ int dw[9];

__global__ void kernel(int *a)
{
	printf("hello\n");
	printf("%d\n",de[8][1]);
	*a=de[8][0];
}
int main()
{
	int w[9];
	for (int i=0;i<9;i++)
	{
		w[i]=i;
	}
	cudaMemcpyToSymbol(dw, w, sizeof(int)*9);

	int *a,b;
	cudaMalloc((void **) &a, sizeof(int));
	kernel<<<1,1>>>(a);
	cudaMemcpy(&b,a,sizeof(int), cudaMemcpyDeviceToHost);
	printf("%d\n",b);
	cudaDeviceReset(); //保证kernel中的printf能正常打印
	return 0;
}

此时,de的值显示错误,但dw的值显示正确,kernel中和main函数中的打印输出都正确。
[attach]3276[/attach]

[attach]3277[/attach]
[/indent]

ice版主你好,我使用的是Nsight 3.0+vs2010旗舰版,我上午找了另一个网友用Nsight测试了一下,同样显示错误。

LZ您好:

看了您的代码和截图,真心不知道为何会这样了。

祝您好运~

目前真心奇葩了,不知道是否是已知或者新BUG:

您连续定义了2个:
constant int de

device int dw
在您cudaMemcpyToSymbol复制改写了dw后,用nsight查看会发现de的值是你复制过去的值??

这个真心不知道原因。难道是nsight的BUG?

版主你好,您能否将程序二的代码在你的机器上测试一下?看看de和dw是否都显示正确。

楼主您好,您的每段代码都被我们在自己机器上试验过,

在我这里是无问题的,2个数组的显示都非常正确,但真心不知道为何你会出现这种奇葩现象。

真心无法解释。

我是nsight 3.0 + VS 2010。其中VS 2010是英文正版。