问题标签 [gpu-warp]

In generalizing a kernel thats shifts the values of a 2D array one space to the right (wrapping around the row boundaries), I have come across a warp synchronization problem. The full code is attached and included below.

The code is meant to work for arbitrary array width, array height, number of thread blocks, and number of threads per block. When choosing a thread size of 33 (i.e. one more thread than a full warp), the 33rd thread doesn't synchronize with __syncthreads() is called. This causes problems with the output data. The problem is only present when there is more than one warp, and the width of the array is more than the number of threads (e.g. with width=35 and 34 threads).

The following is a downsized example of what happens (in reality the array would need to have more elements for the kernel to produce the error).

Initial array:

Expected Result:

Kernel Produces:

The first line is done correctly (for each block if there are more than one), with all subsequent lines having the second last value repeated. I have tested this one two different cards (8600GT and GTX280) and get the same results. I would like to know if this is just a bug with my kernel, or a problem that can't be fixed by adjusting my code?

The full source file is included below.

Thank you.

一个经线是 32 个线程。32 个线程是否在多处理器中并行执行？如果 32 个线程没有并行执行，那么 warp 中就没有竞争条件。在经历了一些例子后，我有了这个疑问。

如果我用一个块有尺寸的网格开始我的内核：

网格块现在是如何分裂成扭曲的？这种块的前两行形成一个扭曲，还是前两列，或者这是任意排序的？

假设 GPU 计算能力为 2.0。

注意：此问题特定于 nVIDIA Compute Capability 2.1 设备。以下信息来自 CUDA Programming Guide v4.1：

在计算能力 2.1 设备中，每个SM有 48 个SP（内核）用于整数和浮点运算。每个经线由 32 个连续的线组成。每个 SM 有 2 个warp 调度器。在每个指令发布时间，一个warp 调度程序选择一个准备好的线程warp，并为核心上的warp 发出2 条指令。

我的疑惑：

• 一个线程将在一个核心上执行。设备如何在单个时钟周期或单个多周期操作中向线程发出 2 条指令？
• 这是否意味着这两条指令应该相互独立？
• 这两条指令可以在内核上并行执行，可能是因为它们在内核中使用了不同的执行单元？这是否也意味着只有在 2 条指令执行完毕后，或者在其中一条指令之后，warp 才会准备就绪？

以下代码将32数组中的每个元素与每个32元素组的第一个元素相加：

我以为我可以消除__syncthreads()代码中的所有内容，因为所有操作都是在同一个经线中完成的。但是如果我消除它们，我会得到垃圾结果。它不会对性能产生太大影响，但我想知道为什么我需要__syncthreads()这里。

我有 GeForce GTX460 SE，所以它是：6 SM x 48 CUDA Cores = 288 CUDA Cores。众所周知，一个 Warp 中包含 32 个线程，并且在一个块中同时（一次）只能执行一个 Warp。也就是说，在单个多处理器（SM）中只能同时执行一个 Block、一个 Warp 和只有 32 个线程，即使有 48 个内核可用？

此外，可以使用threadIdx.x和blockIdx.x来分配具体的Thread和Block的示例。要分配它们，请使用内核 <<< Blocks, Threads >>> ()。但是如何分配特定数量的 Warp 并分配它们，如果不可能，那么为什么还要费心去了解 Warp 呢？

当一个 warp 完成一个内核，但同一块的另一个 warp 仍在运行时，完成的 warp 是否会被阻塞，直到同一块的其他 warp 完成，或者完成的 warp 是否可供另一个块立即重用，而另一个当前块的扭曲仍在运行？

我的理解是warp是在运行时通过任务调度程序定义的一组线程，CUDA的一个性能关键部分是warp中线程的分歧，有没有办法很好地猜测硬件将如何构建warp在线程块内？

例如，我在一个线程块中启动了一个具有 1024 个线程的内核，如何安排扭曲，我可以从线程索引中判断（或至少做出一个好的猜测）吗？

因为通过这样做，可以最大限度地减少给定warp中线程的分歧。

在在线竞赛检查文档中，严重性级别具有对危险级别警告的描述：
这方面的一个示例是由于翘曲级别编程导致的危险，假设线程正在分组进行。

该语句令人困惑，因为线程是按组处理的。（SM 跨warp 执行代码。）如果它们没有分组处理，那么它们是如何处理的？
“翘曲级别编程”是什么意思？（非扭曲级编程是什么？）

我对 GPU 中的扭曲有疑问。

我使用了以下配置：

• GeForce 210
• Cuda 能力主要/次要：1.2
• 2 个多处理器，8 个 CUDA 核心/MP：16 个 CUDA 核心
• 经纱尺寸：32

以下是运行时间（我使用了 nsight）：

Warps（=32 个线程）同时运行，并且有 2 个 MP。所以我认为这个 GPU 的最大能力是 64 线程，但是 16*32 线程几乎同时运行。考虑到warp scheduler，我无法理解这个结果。

我的问题是：

1. 为什么 16*32 线程与 32 线程几乎同时运行？
2. 为什么 64*32 运行时间不是 32*32 的两倍
3. 我听说全局内存访问速度和寄存器一样快。这样对吗？（包括 3.5 GPU 或昂贵的 GPU）