我有一个1G左右的图像，显卡也是1G，一次肯定处理不了整幅图像，于是采用分块的方法。但是分得的小块图像最大不超过大约60MB时核函数才能正常执行。
请问有什么别的办法吗？
显存1GB，却只能分配60MB?
请问是在大于60MB的时候分配出错，还是在大约60MB的时候，kernel挂掉。
前者应该不会存在（除了你正好将显存用到了剩余60MB…)
后者应该是你的kernel自身写的问题, 但具体是哪里导致只能处理60MB就不好说了。
请问现在出错是分配还是kernel?
是kernel的话，提示神马错误？（例如cuda error unknown? unspecified launch failure?)
我算了一下，我是把图像分为20块传进kernel，每块大约57MB，这时kernel以及整个程序可以正常运行。如果分的少于20块，就不行了，应该是kernel挂掉吧。
除了系统界面显示，没有其他用到现存的了，应该不会占用显存而剩60MB的情况。
错误提示是cuda error unknown。
LZ您好：
您的问题显然可以分为两个环节：1：申请空间和存储数据；2：kernel使用数据。
根据一般的经验，1：应该问题不大，所以建议您着重检查一下2：，也就是您的kernel的情况，这个一般是您kernel访存越界了，导致kernel挂掉。
大致分析如上，供您参考。
祝您好运！
cuda error unknown是典型的kernel越界，而不是cudaMalloc分配失败。
所以您的问题（1）不在于无法分配大约60MB的显存。
而在于（2）您的kernel在处理超过一定规模的数据的时候存在越界。
建议使用nsight调试，可以直接定位到越界位置，然后查看此行是否符合预期，进行修改。
感谢两位版主周末还来答疑！
以下是代码，我看不出来哪里有问题，还请版主大人指教。
__global__ void GPUkernel( int height, int width, void *Pin_data, double *d_Dgauss, int Isize, void
*Pout_data)
{
int Iinterval = Isize/2; //Isize = 51；
int col = threadIdx.x + blockIdx.x*blockDim.x; //请问我这样访问可以吗？
int row = threadIdx.y + blockIdx.y*blockDim.y;
unsigned short sum = 0;
for(int i=-Iinterval; i<=Iinterval;i++)
{
for(int j=-Iinterval; j<=Iinterval;j++)
{
sum += d_Dgauss[(i+Iinterval)*Isize + (j+Iinterval)]*(*((unsigned short *)Pin_data +
(row+i)*width + col+j));
}
}
*((unsigned short *)Pout_data + row*width + col ) =*((unsigned short *)Pin_data + Iinterval*gridDim.x +
Iinterval) - sum;
}
extern "C"
void GPU( int div_height, int div_width, void *h_in_data, double *h_gauss, int size, double d_nodata, void
*h_out_data)
{
cudaSetDevice(0);
int MemSize = div_height * div_width; //原图像的h和w都在20000+，这里传进来的值分别大约为4000+和5000+
int bdx = 32;
int bdy = 32;
int gdx = (div_width + bdx-1)/bdx;
int gdy = (div_height + bdy -1)/bdy;
dim3 grd_dim( gdx, gdy);
dim3 blk_dim( bdx, bdy);
checkCudaErrors(cudaDeviceSynchronize());
void *d_in_data;
checkCudaErrors(cudaMalloc((void**)&d_in_data, MemSize*sizeof(unsigned short)));
checkCudaErrors(cudaMemset(d_in_data, 0, MemSize*sizeof(unsigned short)));
checkCudaErrors(cudaMemcpy(d_in_data, h_in_data, MemSize*sizeof(unsigned short), cudaMemcpyHostToDevice));
void *d_out_data;
checkCudaErrors(cudaMalloc((void**)&d_out_data, MemSize*sizeof(unsigned short)));
checkCudaErrors(cudaMemset(d_out_data, 0, MemSize*sizeof(unsigned short)));
double *d_gauss;
checkCudaErrors(cudaMalloc((void**)&d_gauss, size*size*sizeof(double)));
checkCudaErrors(cudaMemset(d_gauss, 0, size*size*sizeof(double)));
checkCudaErrors(cudaMemcpy(d_gauss, h_gauss, size*size*sizeof(double), cudaMemcpyHostToDevice));
GPUkernel<<< grd_dim, blk_dim >>>( div_height, div_width, d_in_data, d_gauss, I_size, d_out_data );
getLastCudaError("Kernel execution failed!");
checkCudaErrors(cudaDeviceSynchronize());
checkCudaErrors(cudaMemcpy(h_out_data, d_out_data, MemSize*sizeof(unsigned short), cudaMemcpyDeviceToHost));
checkCudaErrors(cudaFree(d_in_data));
checkCudaErrors(cudaFree(d_out_data));
checkCudaErrors(cudaFree(d_gauss));
cudaDeviceReset();
}
楼主您看了我的建议了吗？
（1）我们肉眼调试上百行，找到越界的位置，相当折腾，而且没个几个小时找不出来。
（2）您如果请请移动您的鼠标到nsight菜单，选择start cuda debugging (以及选择cuda memory check), 只需要1秒钟找到问题。
为何不使用此建议呢？
感谢来访，周末愉快。
非常感谢横扫版主中肯的建议，只是我现在没有合适的机器，明天就试试。
也祝您周末愉快。
感谢您的来访，祝您调试顺利。
版主你好！
nsight定位越界出现在这行:
sum += d_Dgauss[(i+Iinterval)*Isize + (j+Iinterval)]*(*((unsigned short *)Pin_data + (row+i)*width + col+j));
而这个算法是在CPU上运行时可行的，为什么到CUDA里就不行了？
还有就是，如何处理越界？这个和<<<>>>里的设置有关吗?
LZ您好：
1：“算法可行”与“您的GPU实现是正确的”这是两个不同的问题。您只能用您CPU上的实现可以正常运作来佐证您前者，却不能同样地用此来佐证后者。
2：您需要在您的代码正确地表达您算法意图的前提下，保证您访问的地址空间都是在您申请的地址空间范围之内，如此可以避免越界。
以及，根据您的代码，您的row和col都是和线程编号相关的并在您的代码中用于寻址，以及nsight报告了此处越界，而线程编号和<<<>>>中给定的发布线程的形状和规模是相关的，因此您<<<>>>中给定的参数和您的故障很可能是有关系的。
但是请您注意的是，这多半是您实现上的BUG引起的，比如您没有正确地将您的算法意图变成您的代码实现。而不是说<<<>>>中线程规模在某些情况下一定会出现问题，您需要保证您的代码的处理和您的<<<>>>中的线程发布规格，以及您kernel函数个各个参数（包括申请的空间）之间是相互配合的，是相容的。
大致如此，供您参考。
祝您调试愉快~
首先恭喜楼主成功定位到越界的行，这证明我们之前的讨论是正确的：
您的确是kernel自身有问题，而不是贵卡不能使用大于64MB内存。
那么第二个问题，“我该怎么修改啊”–楼主原文。
既然您这行有1个下标访问和1个指针访问，同时越界了，您首先需要定位是两者的哪一个：
（1）将此行拆分成独立的下标访问和指针访问两行。
（2）查看哪行越界。
然后您需要思考：
如果是下标访问挂了：我这个下标的写法为何不合理？合理的写法是什么？
如果是指针挂了：我的这个指针用的对吗？应该指向哪里？
这些都需要根据您的算法要求，进行思考，而不是让版主直接指出是哪里错误。（这显然不可能的）
例如，如果您发现您的下标是[88888], 而您正确的值最大是[77777](都是举例，非真实数据），那么您应该考虑为何您的下标计算会得到88888，是哪里写错了。
这样才能解决。
请按上述步骤考虑。谢谢。
版主您好！
我修改了指针访问，但程序一运行还是提示cudaErrorUnkwown，现在的问题是，启用nsight调试，然后就像卡死一样没有任何反应了，但还可以设置断点、结束调试等操作。但不像原来那样一下定位到越界的地方（也许这次是其他bug，但调试器没有输出任何信息）。等待半小时也没反应。请问这又是什么情况？
（1）您是说，如果直接运行，提示cuda error unknown么？这表明您还有越界。
（2）您继续说，是在这个基础上运行nsight会运行半个小时未结束么？这个真心不知道了。按理说nsight不会导致您半个小时连续卡住的。这个不知道原因。
(1) 是的，继续越界，继续改。
（2）是的。那不就是没法调试了？或者说是nsight的bug？
嗯嗯。第一个继续改是好的，请仔细检查每处访存的场合。
关于第二个，我从来没遇见过直接运行可以执行（并返回错误），但是在nsight下死循环的情况。这个不能提供建议。
又来麻烦版主了，非常感谢之前斑竹热情细致地回答。
程序启动不了，很奇怪nsight出现了这个提示：
CUDA grid launch failed: CUcontext: 3057776 CUmodule: 263549264 Function: Z9GPUkerneliiPvPdidS
CUDA（线程）网格启动失败？
我的配置是这样的：
int bdx = 32;
int bdy = 32;
int gdx = div_width; //div_width = 6129
int gdy = div_height; //div_height= 4915
dim3 grd_dim( gdx, gdy);
dim3 blk_dim( bdx, bdy);
GPUkernel<<< grd_dim, blk_dim >>>(……）；
说奇怪是因为之前这样的配置我用过，没问题，现在再次用却出问题。不知为何。因为我改了核函数，会不会是因为核函数内部不合理导致的？
LZ您好：
单纯32*32=1024线程，这个参数是无问题的，grid尺寸您这里也正常的。
但是在1024线程的情况下，如果您kernel的资源用超了，导致一个block所需的资源超过一个SM所能提供的最大值，此时是无法执行的。在缩小blcok线程数量的时候或许可以。
您可以先检查一下。
大致如此，祝您好运~
斑竹您好，
我现在似乎走入一个怪圈，始终跳不出去，情况就如13#描述的那样，不管我现在如何修改数组访问还是指针访问都这样。
请斑竹以及各位支招，谢谢大家！
LZ您好：
这个只能您自己慢慢调试了，别人无法猜出问题所在。
祝您调试顺利~