stream overlap 有没有对launch的线程数限制,或者kernel限制呢?

我用的是 单个 stream overlap 的操作。下边是部分代码
cudaStreamCreate(&stream);
int N = 1000000; //分块大小

uint *gpuWords ,gpuH = NULL; //gpuWords 为待处理的数据
cudaMalloc((void **)&gpuWords, N
64);

cudaMalloc((void **)&gpuH, 4sizeof(uint)); //结果的存放地址,所有的线程的结果 。
uint tmp[4] = {0};
cudaMemcpy(gpuH, tmp, 4
sizeof(uint), cudaMemcpyHostToDevice);

int dynShmemPerThread = 64;
int staticShmemPerBlock = 32;

int tpb =128; // tpb is number of threads per block

int gridDim[3];
calculate_grid_parameters(gridDim, tpb, N, dynShmemPerThread, staticShmemPerBlock);//计算任务数为N 的block 的维度

for (int i=0 ; i < SUM; i+=N) //SUM为总任务数
{
if(SUM-i < N) //剩下的不够 分配, 重新计算block 维度
{
N=SUM-i;
calculate_grid_parameters(gridDim, tpb, N, dynShmemPerThread, staticShmemPerBlock);
}
cudaMemcpyAsync(gpuWords, &paddedWords[0+i64], N64, cudaMemcpyHostToDevice, stream);
execute_kernel(gridDim[0], gridDim[1], tpb, tpb*dynShmemPerThread, N, gpuWords, gpuH, stream, i);//, stream, i);
std::cerr<< " I =" << i << "N = " << N << “\n”;
}
cudaStreamSynchronize(stream);//cudaThreadSynchronize();

uint ret[4];
cudaMemcpy(ret, gpuHashes, sizeof(uint)*4, cudaMemcpyDeviceToHost); //结果copy 会主机

出现段错误的 stack[attach]3280[/attach]

cudaMemcpyAsync中 的 paddedWords 我在主机内存中 用malloc 和 mallocHostAlloc 都尝试过,对100+M 数据都运行成功 ,都与 500M 的数据运行,都会出现下边的错误。
这样分块 overlap 是因为我的显存不够一次存放。
这样我运行 数据量小的数据的时候 100+M ,运行成功。
但是当我运行 500+M 的数据的时候,就不成功(当我一次将500M 的数据装入显存来处理可以处理成功)

本该 ,500M 的数据 运行到 i = 67 000 000 ,N =108864
可是,我的运行结果是 i=33 000 000 , N =1 000 000,就出现了段错误。(我想前边的都一样,出错也该到结尾出错呀)

于是,我想是不是 对overlap 登录的kernel 或者 thread 有要求么?
不知道 ,那为能给点意见。。为什么在这或出现段错误呢? 我该这么找到这个原因?

为什么会在这边出错了?

对于CUDA 3.2(包含)以后的用户来说,要使用compute和transfer的overlap, 对线程和kernel均无要求。

您导致的segmentation fault, (以及windows上的access violation), 更大的是您host code中的错误。而不是CUDA无法支持。

(但依然无法排除CUDA runtime内部bug的可能)

举个例子说,您有2个指针p和q,
p是指向一段任意的host memory地址的,(例如int *p = (int *)0x11223344;)
q是指向一段有效的device memory地址的,
您直接上个cudaMemcpy, 就有相当大的可能导致Access Violation(windows上)和Segmentation fault(Linux上), 但这个不是cudaMemcpy的错误,而是您提供的host指针提领的时候本来就无法有效。

请检查您的malloc成功(host上).

建议立刻尝试此建议。

(以及依然无法排除cuda runtime突然内部BUG了,挂了,这也是有可能的)
(以及同样无法突然排除楼主机器突然内存有问题,挂了,这依然也是有可能的)
(以及同样无法排除突然楼主RP下降,导致程序出错,这也是有可能的)

LZ您好:

您的代码不完整,以及注释还有很多错字,我只能就现有的信息加以简要说明。

1:单个stream是无法实现计算和copy overlap的,用cudaMemcpyAsync()也不行。无论在默认stream还是您申请的单个stream里面都是顺序执行的。不知您如何要单个stream实现overlap。

2:要实现overlap还需要使用页锁定内存,以及最好是telsa卡,geforce系列根据驱动版本不同,行为有所不同。

3:无论是默认stream还是其他您创建的stream,这里面启动kernel的限制都是一样的。

综上有两点结论:
1)您用一个stream是无法overlap的
2)您当前的程序错误应该是您代码的问题,而不是其他限制。

大致如上,请您参考。

祝您编码顺利~

谢谢斑竹!
是内存的溢出了。。

谢谢斑竹您,我找到错误了是内存的溢出 。。。我用一个流,可以使用 cudaMemcpyAsync() 。这样可以解决我的显存不足的问题。。。。过会再改成overlap的形式。。

嗯嗯。恭喜楼主,就估计是这个原因,否则cudaMemcpy*真心不那么容易挂的。

建议楼主以后多多养成对malloc以及cudaMalloc之类的分配是否成功的检测(以及及时释放),这样可以避免很多问题。:slight_smile:

感谢您的来访,但ICE的一些建议也不错,不妨看看?