找回密码
 立即注册

QQ登录

只需一步,快速开始

查看: 92|回复: 2

关于cuda核函数循环问题

[复制链接]
发表于 2018-12-25 17:44:12 | 显示全部楼层 |阅读模式
GTC
在cuda的核函数中使用循环,加上__syncthreads()栅栏函数,使得块之间能够同步,但是仍然没有办法解决:让所有线程结束一个循环之后再一起开始下一个循环。代码实例如下:
__global__ void bitonic_sort(int* arr){
        const unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x;
        for (unsigned int i = 2; i <= 4; i <<= 1){
                for (unsigned int j = i>>1; j>0;j>>=1)
                {
                        unsigned int tid_comp = tid ^ j;
                        if (tid_comp > tid){
                                if ((tid & i) == 0){ //ascending
                                        if (arr[tid]>arr[tid_comp]){
                                                swap(arr[tid], arr[tid_comp]);
                                        }
                                }
                                else{ //desending
                                        if (arr[tid]<arr[tid_comp]){
                                                swap(arr[tid], arr[tid_comp]);
                                        }
                                }
                        }
                        __syncthreads();
                }
        }
}

当i==2,j==0时一切正常,但当i==4,j==2时就会出现问题,有的线程没有等其他线程全部完成后,就已经开始了下一个循环:i==4,j==1。但是我已经设置了栅栏函数了。我的理解是应当等所有的i=4,j==2循环结束后在开始j==1的循环。。。作为一个小白,我实在不太清楚为什么。请各位帮助我弄清楚这个逻辑。万分感谢!!!
回复

使用道具 举报

发表于 2018-12-26 13:40:32 | 显示全部楼层
Jetson TX2
任何情况下,用__syncthreads()都只能block内部同步,试图用它来进行block之间能够同步,这个不用尝试了,做不到的。它本身就不是设计成能够全局同步的。

最简单的解决方案:
重新开一次kernel即可“全局同步”。 死亡和重新开始是最整齐的同步了。
回复 支持 反对

使用道具 举报

发表于 2018-12-26 14:39:34 | 显示全部楼层
Tesla P100
或者你可以选择将范围限定在一个block内部的threads之间,则也是可以的(如果这么小的规模适用的话)。
回复 支持 反对

使用道具 举报

您需要登录后才可以回帖 登录 | 立即注册

本版积分规则

关闭

站长推荐上一条 /1 下一条

快速回复 返回顶部 返回列表