0

我试图了解下一段代码中翘曲发散的惩罚是什么。我了解原则上如何处理扭曲分歧（小分支的预测指令，大分支的扭曲投票和分支 - 如果所有扭曲都同意，否则预测指令并且没有分支，与小分支相同）。但是，我不了解具体细节 - 如何处理带有中断/继续的 while 循环。

在下面的示例中，当通道 X 的 scrapEverythingCondition() 计算结果为 true 时，将发生以下哪种情况：

  1. 评估在内部循环中进行，直到 n == N_N，通道 X 始终评估 nops，i 递增，所有通道现在一起工作。
  2. 除了车道 X 之外的每个人都执行 someMoreWork() 而车道 X 评估 nops，车道 X 执行 someCostlyInitialization() 然后 n=0 而其他所有人都评估 nops，所有车道继续一起评估内部循环（显然具有不同的 n 值）。
  3. 还有一些我没有想到的。

代码：

__global__ void chainKernel() {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    while (i < N_I) {
        someCostlyInitialization();
        for(int n = 0; n < N_N; ++n) {
            someStatisticsComputations(n);
            if (scrapEverythingCondition(n)) {
                // Everything we did for current i is no good. Scrap and begin again
                i -= BLOCKS*THREADS;
                break;
            }
            someMoreWork();
        }
        i += BLOCKS*THREADS;
    }
}

我尝试编译到 PTX 并查看生成的代码，但它对我来说太复杂了:(

编辑：感谢 Maku 的回答。我还能够使用散布在代码周围的老式 printf() 来验证答案。我能够看到哪些线程以什么顺序到达哪里，并且确实选项 1 是正确的（通道 X 被暂停，直到内部 for 循环耗尽）。

4

2 回答 2

0

我在这个问题上找到了一个有趣的文档：pdf

据我了解，控制流语句（包括break）定义线程的同步点。在您的情况下，它将在 i += BLOCKS*THREADS; So lane X 离开for循环并等待其他线程到达上述行。

于 2013-07-02T09:02:26.217 回答
0

据我了解，所有scrapEverythingCondition(n) 为真的线程都在执行if 块内的内容并退出for 循环。在执行此块之前，所有其他线程都被阻塞。当这些线程退出 for 循环时，其他 trheads 将执行 someMoreWork();

试用 NVidia Visual Profiler。它确实有助于分析此类问题。

这里还有一些关于此的信息（第 13 - 18 页）： http:
//mc.stanford.edu/cgi-bin/images/3/34/Darve_cme343_cuda_3.pdf

于 2013-07-02T09:29:17.460 回答