CUDA并行线程

Question

运行CUDA程序时，我遇到了以下问题：

1. 我调用与有2个线程

   CUDAkernel<<<1,2>>>

2. 里面的单块一个简单的内核内核我做到以下几点：

   int i = threadIdx.x; if (i==0){ waitabit(); }

   if (i==1){ waitabit(); }

所以，无论是内核线程调用相同的功能waitabit()这非常浪费一些 时钟周期：

__device__ void waitabit(){ 
    clock_t start = clock(); 
     clock_t now; 
     for (;;) { 
     now = clock(); 
     clock_t cycles = now > start ? now - start : now + (0xffffffff - start); 
      if (cycles >= 10000000 ) 
      {break;} 
     }   
} 

现在的问题：功能waitabit()通过延迟线程0.008秒。我自然认为这些线程并行运行，所以他们两个都会将平行于平均延迟0.008秒（大致），整个内核的延迟将大致为0.008秒。

但是，情况并非如此。内核依次执行它们并且延迟为0.016，即2 * 0.008

并行性是不正确地完成的吗？

在此先感谢！

Answer 1

这是一台SIMT机器。在任何给定的时间，只有一条指令是由一个warp处理的。在控制流发散的情况下，if路径和else路径的处理按顺序处理，而不是并行处理。当warp的所有线程到达你的第一个if语句时，线程0处理if路径，而所有其他线程什么都不做。然后，warp在该构造结束时重新同步并开始并行处理。然后他们击中了第二条if语句，并且只有线程1继续而其他人在等待。然后它们在第二个if结构的末尾再次重新同步并开始以锁步方式处理。

所以你的例子的净效果是两个if语句被顺序处理。这是预料之中的。