我在使用这两个函数时发现，不管对应的参数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.yblockDim.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, xsizeof(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分钟就给你瞅出来了。是吧。
版主冤枉我了，“char *s”俨然列在main函数代码的第5行， 估计我们主要都从cuda函数调用的角度找错了。