关于surf3Dread/write的问题

我在使用这两个函数时发现,不管对应的参数z给什么值,读写的结果跟z=0时的情形一样。不知有没有哪位大虾用过这两个函数,有什么注意事项吗?谢过了先!

楼主您好。

在您的帖子中我读到:
“在使用这两个函数时发现,不管对应的参数z给什么值,读写的结果跟z=0时的情形一样”
–请问可否详细的说明下具体是怎么读/写的?以及“和z=0时情形一样“是什么情形?

根据您1楼提供的信息,我猜测一下可能会导致问题的原因:
(1)您将array绑定到您的surface, 但您的array里面的元素的值都是一样的,导致您使用不同的z来读取,值是一样的。
(2)您写入了不同的值。您先对surface进行写入,然后再对surface进行元素读取来检验写入是否成功。此时有注意事项:您的写入在本次kernel里是无效的,可能会导致您读取的值依然是老的值,而此时旧值会导致(1)里的现象。

对此给出的建议:
如果需要对surface进行读写,请注意“写入的结果在下一次kernel启动才会生效”! 所以您进行了写入后,本次不能使用,只能下次用!因为surface使用texture cache, 这个cache是只读的而且只有在下一次kernel启动才会被刷新。

如果根据您1楼的帖子里的信息量进行的猜测不是您想要的情况,建议您进一步跟贴,给出更详细的信息(例如代码)。

祝您编码愉快。

谢谢版主的回复!
我的操作过程是这样的:先在host初始化了一个3D的cudaArray(对应几张图片的像素值), 然后绑定到一个3D的surface,之后在device端读出某元素的值,修改之后再写回原位置。最后在host端读出结果并输出相应的图片。此时的图片应该是已经修改过的吧?而且每次我最后检查的时候,无论我在transformKernel中给f赋什么值,包括超界的值,系统都不报错,而且只有第一张图片被修改了(对应f=0)。

具体代码如下:
global void
transformKernel()
{
// calculate this thread’s data point
int x = blockIdx.xblockDim.x + threadIdx.x;
int y = blockIdx.y
blockDim.y + threadIdx.y;
int4 rgbv;
int f=1;

if(x>=WIDTH)
{printf(“x=%d\n”,x);}
if(y>=HEIGHT)
{printf(“y=%d\n”,y);}

surf3Dread(&rgbv, surfRef, xsizeof(int4), y, f, cudaBoundaryModeZero);
rgbv.x = rgbv.z + rgbv.y + rgbv.x;
rgbv.y = 0;
rgbv.z = 0;
surf3Dwrite(rgbv, surfRef, x
sizeof(int4), y, f, cudaBoundaryModeZero);
}

////////////////////////////////////////////////////////////////////////////////
// Program main
////////////////////////////////////////////////////////////////////////////////
int main()
{

CImg img;
int width;
int height;
int4 imgbuf[FRAME_NUM][WIDTH][HEIGHT];
char* s;
//unsigned int size = FRAME_NUM * WIDTH * HEIGHT * sizeof(int4);
//Create a 2dxint4 device array descriptor, each element composed of 4 ints
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(8sizeof(int), 8sizeof(int), 8sizeof(int), 8sizeof(int), cudaChannelFormatKindSigned);
dim3 dimBlock(8, 8, 1);
dim3 dimGrid(WIDTH / dimBlock.x, HEIGHT / dimBlock.y, 1);
checkCudaErrors( cudaMalloc3DArray( &cu3dArray, &channelDesc, make_cudaExtent(WIDTH, HEIGHT, FRAME_NUM), cudaArraySurfaceLoadStore));

for(int fi = 0; fi < (FRAME_NUM);fi++)
{
sprintf(s, “./dino/dino%04d.png”,fi+1);
img.load(s);
width = img.width();
height = img.height();

//Change each rgb value into int type
for(int i=0;i<width;i++)
{
for(int j=0;j<height;j++)
{
imgbuf[fi][i][j].x = int(img(i,j,0,0));
imgbuf[fi][i][j].y = int(img(i,j,0,1));
imgbuf[fi][i][j].z = int(img(i,j,0,2));
}
}

}
cudaMemcpy3DParms myparms = {0};
myparms.srcPos = make_cudaPos(0,0,0);
myparms.dstPos = make_cudaPos(0,0,0);//fi
myparms.srcPtr = make_cudaPitchedPtr(imgbuf, WIDTH * sizeof(int4), WIDTH, HEIGHT);
myparms.dstArray = cu3dArray;
myparms.extent = make_cudaExtent(WIDTH, HEIGHT, FRAME_NUM);//1
myparms.kind = cudaMemcpyHostToDevice;
checkCudaErrors( cudaMemcpy3D( &myparms));

checkCudaErrors( cudaBindSurfaceToArray( surfRef, cu3dArray, channelDesc));
////////////////////

printf(“Covering 2D data array of %d x %d: Grid size is %d x %d, each block has 8 x 8 threads\n”, WIDTH, HEIGHT, dimGrid.x, dimGrid.y);

transformKernel<<< dimGrid, dimBlock >>>();//WIDTH, HEIGHT); // warmup (for better timing)

// check if kernel execution generated an error
getLastCudaError(“warmup Kernel execution failed”);

checkCudaErrors( cudaDeviceSynchronize() );

// copy result from device to host[0]
//checkCudaErrors( cudaMemcpyFromArray(imgbuf, cu3dArray, 0, 0, size, cudaMemcpyDeviceToHost));//not work for 3d
cudaMemcpy3DParms myparms1 = {0};
myparms1.srcPos = make_cudaPos(0,0,19);
myparms1.dstPos = make_cudaPos(0,0,0);
myparms1.extent = make_cudaExtent(WIDTH, HEIGHT, 1);//FRAME_NUM
myparms1.srcArray = cu3dArray;
myparms1.dstPtr = make_cudaPitchedPtr(imgbuf[0], WIDTH * sizeof(int4), WIDTH, HEIGHT);
myparms1.kind = cudaMemcpyDeviceToHost;
checkCudaErrors( cudaMemcpy3D( &myparms1));

for(int i=0;i<width;i++)
{
for(int j=0;j<height;j++)
{
img(i,j,0,0) = (unsigned char)(imgbuf[0][i][j].x);
img(i,j,0,1) = (unsigned char)(imgbuf[0][i][j].y);
img(i,j,0,2) = (unsigned char)(imgbuf[0][i][j].z);
}
}

img.display(“modified image”);

// cleanup memory
checkCudaErrors(cudaFreeArray(cu3dArray));

cudaDeviceReset();

}

[/i][/i][/i][/i][/i][/i]

楼主您好,看到您3楼的跟贴了:

(1)关于您的“无论我在transformKernel中给f赋什么值,包括超界的值,系统都不报错”的问题,原因如下:
您的代码行surf3Dread(&rgbv, surfRef, x*sizeof(int4), y, f, cudaBoundaryModeZero);
此处指定了cudaBoundaryModeZero, 所以系统不会报错。而只是会忽略错误。

(2)您的host代码里的width, height, depth顺序可能和您的surface中使用的顺序不符合。
您的代码行:int4 imgbuf[FRAME_NUM][WIDTH][HEIGHT];表明了(y,x,z)的顺序(按存储看)
但您的kernel要求的(x,y,z)顺序(按存储看,surface固定的坐标顺序)。

(3)您的代码行dim3 dimGrid(WIDTH / dimBlock.x, HEIGHT / dimBlock.y, 1);
此处没有处理当WIDTH和HEIGHT不是8的整数倍的问题。您这样会导致处理不完全。

(4)您的多行代码
for(int i=0;i<width;i++)
{
for(int j=0;j<height;j++)
{
imgbuf[fi][j].x = int(img(i,j,0,0));
imgbuf[fi][j].y = int(img(i,j,0,1));
imgbuf[fi][j].z = int(img(i,j,0,2));
}
}
此处您是否少粘帖了部分内容呢?您应该写入[fi] [ i ] j.

(5)您没有给出您的cu3dArray的声明的行。

(6)您没有给出您的surfRef的声明的行。导致:
您的代码行: checkCudaErrors( cudaBindSurfaceToArray( surfRef, cu3dArray, channelDesc));
此处可能需要指针类型。

(7)您的上一个代码行的最后一个参数需要指针类型。

(8)您的代码行:myparms1.srcPos = make_cudaPos(0,0,19);
可能导致您复制来的frame(19)和您的代码修改的frame(1)不符合。
可能的您的笔误会导致您观察的不是您的运行结果。所以您可能因此看不见写入的效果。

建议的解决方案:

(1)建议您将cudaBoundaryModeZero改成cudaBoundaryModeTrap. 这样您可以在访问surface越界的时候,得到出错的提示。

(2)建议您将[FRAME_NUM][WIDTH][HEIGHT]改为[FRAME_NUM][HEIGHT][WIDTH]。如果您认为需要的话。并依次修改所有对您的imgbuf访问的语句。

(3)建议您可以简单的将WIDTH / dimBlock.x之类的改为WIDTH/dimBlock.x + 1, 并在kernel加上范围判断。如果您认为您需要,即您的WIDTH和HEIGHT可能存在不是8的整数倍的情况的话。

(4)建议您检查imgbuf[fi][j].x 是否忘记写入[ i ]下标。排版错误?

(5)没有建议。只是提示。

(6)建议您将surfRef改为&surfRef, 如果surfRef已经是一个指针(您上文没有给出代码), 请忽略此建议。

(7)建议您将channelDesc改为&channelDesc, 如果channelDesc已经是一个指针(您上文没有给出代码), 请忽略此建议。

(8)建议您将(0,0,19)改为(0,0,1), 以便观察您在(,,1)上的修改结果。这样您可以不必抱怨结果不变。

祝您调试愉快!

非常感谢版主的细心和耐心!

针对您提的几点建议回复如下:
1)cudaBoundaryModeZero改成cudaBoundaryModeTrap后,系统仍然没有报错。

2)Host端定义这个3维数组的目的是想实现在memory中按帧分段存放(其中WIDTH/HEIGHT分别对应于cuda3DArray的x/y),从调试的结果看(虽然不完全正确),我最终从surface Memory中读回的数验证了我的FrameNum是与3DArray的z相对应的(按我理解, surface Memory也应该是按z分段存放的)。不知您怎么看。

3)这个建议很中肯,但现在的WIDTH/HEIGHT都是8的整数倍,所以系统输出不符合预期暂时还没不是由此问题引起的。

4)这个好像是因为粘贴时本论坛系统将我代码中的“[I]”部分解释成了斜体效果?我的源代码是完整的。

5)抱歉, 声明如下:
device cudaArray* cu3dArray;

6) surface<void, cudaSurfaceType3D> surfRef;

7) 这在Main()段代码的第8行有定义。 我又看了下Nvidia关于此函数的定义,它有C++和C两种版本,前者为引用,后者为指针。现在按您的建议给它们都加上了&,但系统输出仍然没有变。

8)这里的1和19是我调试时临时修改的。经过赋给f不同的值,我才得到了我第一个帖子里说的那个结论:不管我在device端操作哪个f,我在host端只能读到第一帧被修改了,其他帧都没有变。

还望不吝赐教,谢谢!

谢谢版主,前一问题已经解决了。我必须将(width,height)传递给kernal,宏定义貌似不行?

还有另一个小问题:最后cudaFreeArray(cuImgArray)时系统报错:
Segmentation fault (core dumped)

不知是何原因,请赐教,谢谢!

友情提示一下,贴代码使用“代码模式”应该效果好一些。

楼主看到您5楼和6楼的跟贴了:

1)cudaBoundaryModeZero改成cudaBoundaryModeTrap后,系统仍然没有报错。
–这不可能。改为ModeTrap后,当surface访问越界后,是会报错的。这个经过了长时间的验证。
在nsight下会提示access violation并给出位置,在脱离nsight独立运行会返回cudaErrorUnknown. 而后者(未知错误), 是典型的越界导致。楼主您可以检查参数范围看是否不会越界。

2)Host端定义这个3维数组的目的是想实现在memory中按帧分段存放(其中WIDTH/HEIGHT分别对应于cuda3DArray的x/y),从调试的结果看(虽然不完全正确),我最终从surface Memory中读回的数验证了我的FrameNum是与3DArray的z相对应的(按我理解, surface Memory也应该是按z分段存放的)。不知您怎么看。
–您的使用是正确的。此外,需要说明的是,没有surface memory, 您的数据被复制到的是global memory里,但具体位置您和我都不知道(cudaArray是用户透明类型), 然后通过surface访问语句访问的。

3)这个建议很中肯,但现在的WIDTH/HEIGHT都是8的整数倍,所以系统输出不符合预期暂时还没不是由此问题引起的。
4)这个好像是因为粘贴时本论坛系统将我代码中的“[I]”部分解释成了斜体效果?我的源代码是完整的。
5)抱歉, 声明如下:
device cudaArray* cu3dArray;
–请问您确定吗?如果加上__device__是错误的。您的cu3dArray的声明需要在host代码上。

6) surface<void, cudaSurfaceType3D> surfRef;
7) 这在Main()段代码的第8行有定义。 我又看了下Nvidia关于此函数的定义,它有C++和C两种版本,前者为引用,后者为指针。现在按您的建议给它们都加上了&,但系统输出仍然没有变。
–原来还有C++的不同重载。我不知道。

8)这里的1和19是我调试时临时修改的。经过赋给f不同的值,我才得到了我第一个帖子里说的那个结论:不管我在device端操作哪个f,我在host端只能读到第一帧被修改了,其他帧都没有变。

谢谢版主,前一问题已经解决了。我必须将(width,height)传递给kernal,宏定义貌似不行?
–除了楼主您的5上的问题,没有其他问题被发现。此外,宏定义必然是可以的。你将WIDTH和HEIGHT作为参数传递,然后告诉我解决了问题。这不可能。
(1)您的宏定义可以在kernel里正确被使用。因为这个不涉及编译的过程(是贵电脑上的C/C++预处理器负责的)。
(2) 至少__device__ cudaArray *你的array;此处是不对的。

所以建议楼主报告实情。如果不愿意报告(例如涉及隐私), 请不要胡乱找个借口。

最后cudaFreeArray(cuImgArray)时系统报错:
Segmentation fault (core dumped)

–您的代码里没有出现cuImgArray的地方。
–请发送您对cuImgArray的声明和使用的过程。

此外,您使用的是linux. 这个我不懂。不过依然建议您发送相关代码。以便其他版主/会员为您回复。

呵呵,版主说笑了。现汇报最新战果:
此问题终于最后解决: 主要原因在于其中char * s没有初始化,所以导致程序运行出奇怪结果。将这点修改过来之后,程序就完全正确了。

我刚开始通过修改宏定义为参数传递貌似部分解决了问题,但其实质是避免了未初始化的指针对于宏定义的影响,所以kernel部分结果对了。 版主确实厉害,谢谢了!

嗯嗯。恭喜楼主。

下次建议楼主发代码发全乎点。你看你说的什么s没赋值神马的。你的代码里根本就没有体现。结果5天后你才发现。如果你能发代码。这里人多眼多,说不定5分钟就给你瞅出来了。是吧。:slight_smile:

版主冤枉我了,“char *s”俨然列在main函数代码的第5行, 估计我们主要都从cuda函数调用的角度找错了。