我有GeForce GTX460 SE,所以它是:6 SM x 48 CUDA核心= 288 CUDA核心.众所周知,在一个Warp中包含32个线程,并且在一个块中同时(一次)只能执行一个Warp.也就是说,在单个多处理器(SM)中,即使有48个可用核心,也可以同时只执行一个Block,一个Warp和32个线程？
此外,可以使用threadIdx.x和blockIdx.x来分发具体的Thread和Block的示例.要分配它们,请使用内核<<< Blocks,Threads >>>().但是如何分配特定数量的Warp-s并分发它们,如果不可能那么为什么还要去了解Warps呢？
如果我使用其网格具有维度的网格启动我的内核:
dim3 block_dims(16,16);
网格块现在如何分成经线？这样一个块的前两行是形成一个warp,还是前两列,还是这个任意排序？
假设GPU计算能力为2.0.
注意:此问题特定于nVIDIA Compute Capability 2.1设备.以下信息可从CUDA编程指南v4.1获得:
在计算能力2.1设备中,每个SM具有48个SP(核心),用于整数和浮点运算.每个warp由32个连续的线程组成.每个SM都有2个warp调度程序.在每个指令发布时,一个warp调度程序选择一个准备好的线程warp并发出2个内核上的warp 指令.
我的疑惑:
在概括将二维数组的值向右移动一个空间（环绕行边界）的内核时，我遇到了扭曲同步问题。完整的代码附在下面。
该代码适用于任意数组宽度、数组高度、线程块数和每个块的线程数。选择33的线程大小（即比完整扭曲更长的线程），第33个线程不会__syncthreads()调用。这会导致输出数据出现问题。该问题仅在存在多个经纱且阵列的宽度大于线程数（例如，宽度 = 35 和 34 线程）时出现。
下面是一个缩小的例子，说明发生了什么（实际上，数组需要有更多的元素才能让内核产生错误）。
初始数组：
0 1 2 3 4
5 6 7 8 9
预期结果：
4 0 1 2 3
9 5 6 7 8
内核产生：
4 0 1 2 3
8 5 6 7 8
第一行正确完成（如果有多个块，则为每个块），所有后续行都具有重复的倒数第二个值。我已经测试了这两种不同的卡（8600GT 和 GTX280）并得到相同的结果。我想知道这是否只是我的内核的错误，还是无法通过调整我的代码来修复的问题？
完整的源文件包含在下面。
谢谢你。
#include <cstdio>
#include <cstdlib>
// A method to ensure all reads use the same logical layout.
inline __device__ __host__ int loc(int x, int y, int width)
{
return y*width + x;
} …以下代码将32数组中的每个32元素与每个元素组的第一个元素相加:
int i = threadIdx.x;
int warpid = i&31;
if(warpid < 16){
s_buf[i] += s_buf[i+16];__syncthreads();
s_buf[i] += s_buf[i+8];__syncthreads();
s_buf[i] += s_buf[i+4];__syncthreads();
s_buf[i] += s_buf[i+2];__syncthreads();
s_buf[i] += s_buf[i+1];__syncthreads();
}
我以为我可以消除__syncthreads()代码中的所有内容,因为所有操作都是在同一个warp中完成的.但如果我消除它们,我会得到垃圾结果.它不会对性能产生太大影响,但我想知道为什么我需要__syncthreads()这里.
我想知道CUDA应用程序的warp调度顺序是否是确定性的.
具体来说,我想知道在相同设备上具有相同输入数据的同一内核的多次运行时,warp执行的顺序是否保持不变.如果没有,是否有任何可能强制执行warp执行的命令(例如在调试依赖于顺序的算法时)？
在 cuda 示例中，例如这里，__match_all_sync__match_any_sync 用来。
这是一个例子，其中一个经线被分成多个（一个或多个）组，每个组都跟踪自己的原子计数器。
// increment the value at ptr by 1 and return the old value
__device__ int atomicAggInc(int* ptr) {
int pred;
//const auto mask = __match_all_sync(__activemask(), ptr, &pred); //error, should be any_sync, not all_sync
const auto mask = __match_any_sync(__activemask(), ptr, &pred);
const auto leader = __ffs(mask) - 1; // select a leader
int res;
const auto lane_id = ThreadId() % warpSize;
if (lane_id == leader) { // leader does the update
res …来自 CUDA 编程指南：
[Warp shuffle 函数] 在 warp 内的线程之间交换变量。
我知道这是共享内存的替代方案，因此它被用于扭曲中的线程来“交换”或共享值。但它背后的直觉是什么（它是如何工作的）？与使用共享内存相比，它有什么好处？
在CUDA 9中,nVIDIA似乎有了这种"合作团体"的新概念; 由于某些原因我不完全清楚,__ballot()现在(= CUDA 9)被弃用赞成__ballot_sync().这是别名还是语义改变了？
...对于其他内置组件的类似问题,现在已__sync()添加到他们的名字中.
我有关于以下问题的查询:
假设我们有一个9*7图片(x方向7像素,y方向9像素),假设4*4线程块和每个warp 8个线程,有多少warp将有控制偏差？
这里如何组织块和经线？对于x或水平方向,我可以假设每行2个块.类似地,对于垂直方向,每列3个块.但是,经线将如何组织？有人可以指出warp的线程id,以及控制分歧发生的情况(Thread ids等).
谢谢