对于 GPU 上的数据独立问题，每个元素启动 1 个线程是否总是最优的？

Question

我正在编写一个简单的 memcpy 内核来测量我的 GTX 760M 的内存带宽并将其与 cudaMemcpy() 进行比较。看起来像这样：

template<unsigned int THREADS_PER_BLOCK>
__global__ static
void copy(void* src, void* dest, unsigned int size) {
    using vector_type = int2;
    vector_type* src2 = reinterpret_cast<vector_type*>(src);
    vector_type* dest2 = reinterpret_cast<vector_type*>(dest);

    //This copy kernel is only correct when size%sizeof(vector_type)==0
    auto numElements = size / sizeof(vector_type);

    for(auto id = THREADS_PER_BLOCK * blockIdx.x + threadIdx.x; id < numElements ; id += gridDim.x * THREADS_PER_BLOCK){
        dest2[id] = src2[id];
    }
}

我还计算了达到 100% 占用率所需的块数，如下所示：

THREADS_PER_BLOCK = 256 
Multi-Processors: 4 
Max Threads per Multi Processor: 2048 
NUM_BLOCKS = 4 * 2048 / 256 = 32

另一方面，我的测试表明，启动足够多的块以便每个线程只处理一个元素总是优于 "optimal" 块计数。以下是 400mb 数据的时间：

bandwidth test by copying 400mb of data.
cudaMemcpy finished in 15.63ms. Bandwidth: 51.1838 GB/s
thrust::copy finished in 15.7218ms. Bandwidth: 50.8849 GB/s
my memcpy (195313 blocks) finished in 15.6208ms. Bandwidth: 51.2137 GB/s
my memcpy (32 blocks) finished in 16.8083ms. Bandwidth: 47.5956 GB/s

所以我的问题是：

为什么会有速度差异？

当每个元素都可以完全独立于所有其他元素进行处理时，为每个元素启动一个线程是否有任何缺点？

Answer 1

Is starting 1 thread per element always optimal for data independent problems on the GPU?

不总是。让我们考虑 3 种不同的实现。在每种情况下，我们都假设我们正在处理一个简单的可并行化问题，每个线程涉及一个元素加载、一些 "work" 和一个元素存储。在您的复制示例中，基本上没有工作 - 只是加载和存储。

每个线程一个元素。每个线程执行 1 个元素加载、工作和 1 个存储。 GPU 喜欢在每个可用线程中有大量公开的可并行发布的指令，以隐藏延迟。您的示例由每个线程一个加载和一个存储组成，忽略索引算法等其他指令。在您的示例 GPU 中，您有 4 个 SM，每个 SM 最多能够补充 2048 个线程（当今几乎所有 GPU 都是如此），所以最大的飞行补充是 8192 个线程。所以最多可以向内存管道发出 8192 次加载，然后我们将停止机器直到数据从内存返回，以便可以发出相应的存储指令。此外，对于这种情况，我们有与退出线程块和启动新线程块相关的开销，因为每个块只处理 256 个元素。
每个线程多个元素，编译时未知。在这种情况下，我们有一个循环。编译器在编译时不知道循环范围，因此它可能会也可能不会展开循环。如果它不展开循环，那么我们在每次循环迭代中都有一个加载，然后是一个存储。这不会为编译器提供重新排序（独立）指令的好机会，因此净效果可能与情况 1 相同，只是我们有一些与处理循环相关的额外开销。
每个线程多个元素，在编译时已知。你还没有真正提供这个例子，但它通常是最好的场景。在 parallelforall 博客 matrix transpose example 中，本质上是复制内核的作者选择让每个线程执行 8 个复制元素 "work"。然后编译器看到一个循环：
```
  LOOP:  LD R0, in[idx];
         ST out[idx], R0;
         ...
         BRA  LOOP;
```
它可以展开（比方说）8 次：
```
     LD R0, in[idx];
     ST out[idx], R0;
     LD R0, in[idx+1];
     ST out[idx+1], R0;
     LD R0, in[idx+2];
     ST out[idx+2], R0;
     LD R0, in[idx+3];
     ST out[idx+3], R0;
     LD R0, in[idx+4];
     ST out[idx+4], R0;
     LD R0, in[idx+5];
     ST out[idx+5], R0;
     LD R0, in[idx+6];
     ST out[idx+6], R0;
     LD R0, in[idx+7];
     ST out[idx+7], R0;
```
然后它可以重新排序指令，因为操作是独立的：
```
     LD R0, in[idx];
     LD R1, in[idx+1];
     LD R2, in[idx+2];
     LD R3, in[idx+3];
     LD R4, in[idx+4];
     LD R5, in[idx+5];
     LD R6, in[idx+6];
     LD R7, in[idx+7];
     ST out[idx], R0;
     ST out[idx+1], R1;
     ST out[idx+2], R2;
     ST out[idx+3], R3;
     ST out[idx+4], R4;
     ST out[idx+5], R5;
     ST out[idx+6], R6;
     ST out[idx+7], R7;
```
以增加套准压力为代价。与非展开循环情况相比，这里的好处是前 8 LD 条指令都可以发出——它们都是独立的。发出这些指令后，线程将在第一个 ST 指令处停止 - 直到相应的数据实际从全局内存中返回。在非展开的情况下，机器可以发出第一条 LD 指令，但会立即命中相关的 ST 指令，因此它可能会停在那里。这样做的目的是，在前两种情况下，我只能对内存子系统执行 8192 LD 次操作，但在第三种情况下，我能够在其中执行 65536 LD 条指令航班。这有好处吗？在某些情况下，确实如此。好处会因您运行使用的 GPU 而异。

我们在这里所做的是有效地（与编译器一起工作）增加每个线程可以发出的指令数，然后线程就会陷入停顿.这也称为增加 exposed 并行度，基本上是通过这种方法中的 ILP。它是否有任何好处将取决于您的实际代码、您的实际 GPU 以及当时 GPU 中的其他内容。但是使用诸如此类的技术来增加暴露的并行性始终是一个很好的策略，因为发出指令的能力是 GPU 隐藏其必须处理的各种形式延迟的方式，因此我们有效地提高了 GPU 隐藏延迟的能力, 用这种方法。

Why is there a speed difference?

如果不仔细分析代码，这可能很难回答。然而通常情况下，启动足够个线程来完全满足GPU的瞬时承载能力并不是一个好的策略，可能是由于"tail effect"或其他类型的低效率.块也可能受到其他一些因素的限制，例如寄存器或共享内存的使用。通常需要仔细分析并可能研究生成的机器代码才能完全回答此类问题。但循环开销可能会显着影响您的比较，这基本上是我上面的案例 2 与我的案例 1。

（请注意，我的 "pseudo" 机器代码示例中的内存索引并不是您对编写良好的跨网格复制循环所期望的 - 它们只是为了演示展开及其好处的示例目的通过编译器指令重新排序）。

Answer 2

一句话回答：当每个元素有一个线程时，您需要为每个元素支付线程设置成本——至少，将参数从常量内存复制到寄存器——这是一种浪费。

对于 GPU 上的数据独立问题，每个元素启动 1 个线程是否总是最优的？

Is starting 1 thread per element always optimal for data independent problems on the GPU?

cuda

gpu

gpgpu