我的理解是，warp 是通过任务调度程序在运行时定义的一组线程，CUDA 的一个性能关键部分是 warp 内线程的分歧，有没有办法很好地猜测硬件将如何构造 warp在线程块内？
例如，我启动了一个线程块中包含 1024 个线程的内核，扭曲是如何排列的，我可以从线程索引中看出（或至少做出一个很好的猜测）吗？
因为通过这样做，可以最大限度地减少给定经纱内线程的发散。
假设我有一个带有local_size = 8*8*8的OpenGL计算着色器.调用如何映射到nVidia GPU warp？同样的调用是否gl_LocalInvocationID.x会在同一个warp中？还是？还是z？我并不是指所有的调用,我只是指一般的聚合.
我问这个是因为在一个时刻进行了优化,并非所有的调用都有工作要做,所以我希望它们处于相同的变形中.
在CUDA Developer Blog上阅读了这篇文章后，我努力了解何时可以安全\正确使用__activemask()代替__ballot_sync()。
在活动掩码查询部分中，作者写道：
这是不正确的，因为它将导致部分和而不是总计。
之后，在“ 机会扭曲级编程”部分中，他们使用此函数__activemask()是因为：
如果要在库函数中使用扭曲级编程，但是不能更改函数接口，则可能会很困难。
nvcc设备代码可以访问内置值,warpSize该值设置为执行内核的设备的warp大小(即在可预见的将来为32).通常你不能把它区分为常数 - 但是如果你试图声明一个长度为warpSize的数组,你就会抱怨它是非常量的...(使用CUDA 7.5)
所以,至少为了这个目的,你有动力去做(编辑):
enum : unsigned int { warp_size = 32 };
在你的标题中的某个地方.但是现在 - 我应该选择哪个,何时？:warpSize,或warp_size？
编辑: warpSize显然是PTX中的编译时常量.问题仍然存在.