删除线程后可以使用__syncthreads()吗？

Question

__syncthreads()在我有意使用线程丢弃的块中使用是否安全return？

文档说明__syncthreads() 必须由块中的每个线程调用,否则它将导致死锁,但实际上我从未经历过这样的行为.

示例代码:

__global__ void kernel(float* data, size_t size) {
    // Drop excess threads if user put too many in kernel call.
    // After the return, there are `size` active threads.
    if (threadIdx.x >= size) {
        return;
    }

    // ... do some work ...

    __syncthreads(); // Is this safe?

    // For the rest of the kernel, we need to drop one excess thread
    // After the return, there are `size - 1` active threads
    if (threadIdx.x + 1 == size) {
        return;
    }

     // ... do more work ...

    __syncthreads(); // Is this safe?
}

Answer 1

简短问题的答案是"不".围绕__syncthreads()指令的Warp级分支分歧将导致死锁并导致内核挂起.您的代码示例不保证安全或正确.实现代码的正确方法是这样的:

__global__ void kernel(...)

    if (tidx < N) {
        // Code stanza #1
    }

    __syncthreads();


    if (tidx < N) {
        // Code stanza #2
    }

    // etc
}

以便__syncthreads()无条件执行指令.

编辑:只是添加一些确认此断言的附加信息,__syncthreads()调用将编译到bar.sync所有体系结构的PTX 指令中.PTX2.0指南(p133)记录bar.sync并包含以下警告:

障碍是在每个warp的基础上执行的,就好像warp中的所有线程都是活动的一样.因此,如果warp中的任何线程执行bar指令,则好像warp中的所有线程都执行了bar指令.warp中的所有线程都会停止,直到屏障完成,并且屏障的到达计数会增加warp大小(而不是warp中活动线程的数量).在有条件执行的代码中,只有在已知所有线程以相同方式评估条件(warp不发散)时才应使用bar指令.由于障碍是在每个warp的基础上执行的,因此可选的线程数必须是warp大小的倍数.

因此,尽管有任何相反的断言,在__syncthreads()调用周围进行条件分支是不安全的,除非您可以100%确定任何给定warp中的每个线程都遵循相同的代码路径并且不会发生warp散布.

Answer 2

Compute Capability 7.x(Volta)更新:

随着warp中线程之间引入独立线程调度,CUDA最终在实践中更加严格,现在匹配记录的行为.从编程指南:

虽然__syncthreads()一直被记录为同步线程块中的所有线程,但Pascal和先前的体系结构只能在warp级别强制执行同步.在某些情况下,只要每个warp中至少有一些线程到达屏障,这就允许屏障成功而不会被每个线程执行.从Volta开始,每个线程强制执行CUDA内置的__syncthreads()和PTX指令bar.sync(及其派生词),因此在块中所有未退出的线程达到之前不会成功.利用先前行为的代码可能会死锁,必须进行修改以确保所有未退出的线程都到达屏障.

下面是前一个答案,其中讨论了伏尔塔前的行为.

更新:这个答案可能不会在talonmies之上添加任何东西(取决于你对这个主题的理解,我想),但冒着过于冗长的风险我会提供帮助我更好地理解这些信息的信息.此外,如果你对"幕后"的工作方式不感兴趣,或者除了官方文档之外可能有什么可能,那么这里没有什么可看的.尽管如此,我仍然不建议做出超出正式记录的假设,特别是在希望支持多种或未来架构的环境中.我主要想指出的是,虽然这被CUDA编程指南明确地称为不良实践,但实际行为__syncthreads()可能与描述它的方式有些不同,对我来说有趣.我想要的最后一件事就是传播错误信息,所以我愿意讨论并修改我的答案!

一些重要的事实

这个答案没有TL; DR因为有太多误解的可能性,但这里有一些相关的事实:

• __syncthreads() 对于块中的warp而不是块中的所有线程来说,行为就像是一个障碍,尽管按照建议使用时它也是相同的.
• 如果warp中的任何线程执行PTX bar指令(例如from _syncthreads),则好像warp中的所有线程都有.
• 当a bar.sync被调用时(由instrinsic生成__syncthreads()),该块和屏障的到达计数增加了warp大小.这就是以前的要点.
• 线程分歧(多个路径)通过序列化分支的执行来处理.序列化的顺序是一个可能导致麻烦的因素.
• warp中的线程不同步__syncthreads().该指令不会导致warp停止并等待发散路径上的线程.分支执行是序列化的,因此只有当分支重新加入或代码终止时,warp中的线程才会重新同步.在此之前,分支机构按顺序独立运行.同样,块的每个warp中只有一个线程需要命中__syncthreads()执行才能继续.

官方文档和其他来源支持这些声明.

口译和文件

由于__syncthreads()在块中而不是块中的所有线程中作为warp的屏障,如编程指南中所述,如果每个warp中至少有一个线程碰到屏障,那么似乎可以很简单地提前退出.(但这并不是说你不能用内在函数引起死锁!)这也假设它__syncthreads()总是生成一个简单的bar.sync a;PTX指令,并且它的语义也不会改变,所以不要在生产中这样做.

我遇到的一项有趣的研究实际上调查了当你违反CUDA编程指南的建议时会发生什么,他们发现尽管通过滥用__syncthreads()条件块确实可能导致死锁,但并非所有在条件块中使用内在函数代码会这样做.从论文的D.1节:

编程指南建议仅当条件在整个线程块中进行相同的求值时,才能在条件代码中使用syncthreads().本节的其余部分将研究syncthreads()在违反此建议时的行为.我们证明了syncthreads()作为warp的屏障,而不是线程.我们表明,当warp的线程由于分支发散而被序列化时,一条路径上的任何syncthreads()都不会等待来自另一条路径的线程,而只等待在同一线程块内运行的其他warp.

此声明与talonmies引用的PTX文档的位置一致.特别:

障碍是在每个warp的基础上执行的,就好像warp中的所有线程都是活动的一样.因此,如果warp中的任何线程执行bar指令,则好像warp中的所有线程都执行了bar指令.warp中的所有线程都会停止,直到屏障完成,并且屏障的到达计数会增加warp大小(而不是warp中活动线程的数量).

正是从这个很清楚,为什么可选的线程数b的bar.sync a{, b};指令必须是经大小的倍数-只要在经单个线程执行bar指令的到来计数由经大小,而不是线程的经纱数递增这实际上触及了障碍.尽早终止的线程(遵循不同的路径)无论如何都被有效地计为到达.现在,引用段落中的下一句话确实表示不要__syncthreads()在条件代码中使用,除非"已知所有线程都以相同的方式评估条件(warp不会发散)".这似乎是一个过于严格的建议(对于当前的架构),旨在确保到达计数实际上反映了触及障碍的实际线程数.如果击中屏障的至少一个线程增加了整个扭曲的到达次数,那么您可能真的有一点灵活性.

在PTX文档中,bar.sync a;通过__syncthreads()等待当前协作线程阵列(块)中的所有线程到达屏障而生成的指令没有歧义a.然而,重点在于,每当屏障被击中时(默认情况下b未指定),通过以每个经线大小的倍数递增到达计数来确定"所有线程"的当前状态.这部分不是未定义的行为,至少不是并行线程执行ISA版本4.2.

请记住,即使没有条件,warp中也可能存在非活动线程 - "线程数不是warp大小的倍数的块的最后一个线程".(SIMT架构说明).然而__syncthreads()在这样的街区中并没有被禁止.

例子

提前退出版本1:

__global__ void kernel(...)

    if (tidx >= N)
        return;      // OK for <32 threads to hit this, but if ALL
                     // threads in a warp hit this, THEN you are deadlocked
                     // (assuming there are other warps that sync)

    __syncthreads(); // If at least one thread on this path reaches this, the 
                     // arrival count for this barrier is incremented by 
                     // the number of threads in a warp, NOT the number of 
                     // threads that reach this in the current warp.
}

如果每个warp中至少有一个线程命中同步,则不会死锁,但可能的问题是发散代码路径执行的序列化顺序.您可以更改上面的内核以有效地交换分支.

提前退出版本2:

__global__ void kernel(...)

    if (tidx < N) {
        // do stuff

        __syncthreads();
    }
    // else return;
}

如果你在warp中至少有一个线程击中了障碍,那么仍然没有死锁,但在这种情况下,分支执行的顺序是否重要？我不这么认为,但要求特定的执行订单可能是个坏主意.

本文在一个更为复杂的例子中证明了这一点,与一个微不足道的早期退出相比,这也提醒我们在扭曲分歧时要谨慎.这里warp的前半部分(tid[0,15]上的线程id )写入一些共享内存并执行__syncthreads(),而另一半(tid[16,31]上的线程id )也执行__syncthreads()但现在从写入的共享内存位置读取经线的上半部分.首先忽略共享内存测试,您可能会在任一障碍处遇到死锁.

// incorrect code to demonstrate behavior of __syncthreads
if (tid < 16 ) {
  shared_array[tid] = tid;
  __syncthreads();
}
else {
  __syncthreads();
  output[tid] =
    shared_array[tid%16];
}

没有死锁,表示__syncthreads()不会在warp中同步分叉线程. 不同的代码路径在warp中被序列化,并且它只需要代码路径中的一个线程来使调用__syncthreads()在每个warp级别工作.

但是,共享内存位显示某些不可预测的行为可以进入此位置.warp的后半部分没有从上半部分获得更新值,因为首先执行了warp和else块的分支发散序列化执行.因此该函数不能正常工作,但它也表明__syncthreads()不会在warp中同步不同的线程.

摘要

__syncthreads()不等待warp中的所有线程,并且单个线程在warp中的到达有效地将整个warp计为已到达屏障.(现在的架构).

__syncthreads()在条件代码中使用会很危险,因为序列化的线程执行有多么不同.

仅当您了解条件代码的工作方式以及如何处理分支分歧(在 warp 中发生)时才使用条件代码中的内在函数.

请注意,我没有说继续使用__syncthreads()的方式与记录方式不一致.