CUDA如何调试死循环

今天调试全局同步的函数,如下
157 //GPU lock-free synchronization function
158 device
159 void __gpu_sync(unsigned goalVal, volatile unsigned *Arrayin, volatile unsigned *Arrayout) {
160 // thread ID in a block
161 unsigned tid_in_blk = threadIdx.x * blockDim.y + threadIdx.y;//caculate the thread id in the block.
162 unsigned nBlockNum = gridDim.x * gridDim.y;
163 unsigned bid = blockIdx.x * gridDim.y + blockIdx.y;//caculate the block id in the grid
164 // only thread 0 is used for synchronization
165 if (tid_in_blk == 0) {
166 Arrayin[bid] = goalVal;
167 __threadfence();
168 }
169 if (bid == 1) {
170 if (tid_in_blk < nBlockNum) {//every thread whose id < block number in block 0 is hooked by one block, means to sysnc all threads of all blocks
171 while (Arrayin[tid_in_blk] != goalVal){//wait for accomplishmet of Array[bid] = goalVal from other block
172 //Do nothing here
173 }
174 }
175 __syncthreads();
176 if (tid_in_blk < nBlockNum) {
177 Arrayout[tid_in_blk] = goalVal;
178 __threadfence();
179 }
180 }
181 if (tid_in_blk == 0) {
182 while (Arrayout[bid] != goalVal) {
183 //Do nothing here
184 }
185 }
186 // printf(“waiting for set arrayout[%d] = %d, goalVal = %d, bid: %d, tid_in_blk : %d”, bid, arrayout[bid], goalVal, bid, tid_in_blk);
187 __syncthreads();
188 }

然后在183行死循环,在CUDA-GDB中ctrl+c停止后,无法print变量(没有上下文),printf也没有显示,怎么办,谢谢

无法printf很奇怪,两种可能,一种是kernel内部有严重错误,一种是kerel调用完没有与cpu同步。你可以在kernel<<<>>>();后面加上cudaDeviceSynchronize(); printf(“%s”,cudaGetErrorString(cudaGetLastError()));检查下是不是no error。

如果CUDA-GDB无法打印出变量,可能的原因比较多,在这种状况下建议使用printf。

无法print的原因是kernel死循环了,没有调用完,我理解的cuda的printf只是把数据放到缓冲区里,不是真正print。即使后面加上cudaDeviceSynchronize(); printf(“%s”,cudaGetErrorString(cudaGetLastError()));
也无法执行。

http://cudazone.nvidia.cn/forum/forum.php?mod=viewthread&tid=6740
这个帖子和我看的是一篇文章,原来是BLOCK多于SM数的缘故。

但CUDA-GDB不能print还是件很蛋疼的事儿。

LZ您好:

建议使用安全的实现方式,而不是先使用不安全的实现方式,再去纠结因此带来的各种问题。

祝您编码顺利~

谢谢版主,能否给讲一下,我列出的同步函数为什么不安全,谢谢。因为要GPU没有全局的同步机制,为解决这个问题,有使用如上的同步函数的,也有划分为多个kernel,kernel之间同步,然后我需要比较这两种方法的优劣,还请赐教。

LZ您好:

仅建议使用结束kernel的方法实现全局同步。

使用原子累加标记的方法仅适用于总block数量小于该GPU最大驻留blocks数目的情况,在您引用的论坛帖子的2#已经有较为详细的解说。而不同的GPU在运行该程序时,因为硬件的不同,所允许的最大驻留blocks数目是不同的。这样,您的在较强的GPU上可以运行的代码跑在一个较弱的GPU上,可能会直接卡死。因而,这样是不安全的做法。

同时需要说明的是,一般习惯上grid内的blocks数量都是大于GPU最大驻留blocks数量的,并以此来提升GPU的使用效率。而这种同步的方法的blocks数目的要求与惯常习惯背道而驰。

大致如此,供您参考。

祝您编码顺利~