对于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

另一方面，我的测试表明，启动足够多的块，以便每个线程仅处理一个元素，总是胜过“最佳”块数。以下是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

对于GPU上与数据无关的问题，每个元素启动1个线程是否始终是最佳选择？

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

1. 每个线程一个元素。每个线程执行1个元素加载，工作和1个存储。GPU喜欢每个可用线程具有很多公开的具有并行功能的指令，以隐藏延迟。您的示例由一个负载和一个线程一个存储组成，而忽略了其他指令（如索引算术等）。在示例GPU中，您有4个SM，每个SM最多可支持2048个线程（对于当今几乎所有GPU都是如此） ，因此最大飞行补码为8192个线程。因此最多可以将8192个负载发送到内存管道，然后我们要使机器停顿，直到数据从内存中返回，以便可以发出相应的存储指令。此外，在这种情况下，我们还有与退出线程块和启动新线程块相关的开销，

2. 每个线程有多个元素，在编译时未知。在这种情况下，我们有一个循环。编译器在编译时不知道循环范围，因此它可能会展开循环，也可能不会展开循环。如果它没有展开循环，则每个循环迭代都有一个加载，然后是一个存储。这并没有给编译器一个很好的机会来重新排序（独立）指令，因此实际效果可能与情况1相同，除了我们在处理循环方面有一些额外的开销。

3. 每个线程有多个元素，在编译时已知。您尚未真正提供此示例，但这通常是最佳方案。在parallelforall Blog 矩阵转置示例中，本质上是复制内核的作者选择让每个线程执行复制“工作”的8个元素。然后，编译器会看到一个循环：

     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条指令都可以发出-它们都是独立的。发出这些指令后，线程将在飞行中的第一条指令处停滞。这有好处吗？在某些情况下，确实如此。收益会因您所运行的GPU而异。ST指令条指令上-直到实际从全局内存返回相应的数据为止。在非展开情况下，机器可以发出第一LD条指令，但是会立即命中从属ST指令，因此它可能会停在那里。最终的结果是，在前两种情况下，我只能进行8192次LD操作，而到内存子系统，但是在第三种情况下，我可以进行65536次操作LD

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

为什么会有速度差异？

如果不仔细分析代码，可能很难回答。但是，通常情况下，启动足够多的线程来完全满足GPU的瞬时承载能力并不是一个好的策略，这可能是由于“尾部效应”或其他类型的效率低下所致。块受某些其他因素（例如寄存器或共享内存使用率）的限制也可能是这种情况。通常有必要仔细分析并研究生成的机器代码，以完全回答此类问题。但是，循环开销可能会明显影响您的比较，这基本上就是我的情况2与上面的情况1。

（请注意，在我的“伪”机器代码示例中的内存索引并不是编写良好的网格跨越式复制循环所期望的-它们只是出于示例目的，以展示展开以及通过编译器指令重新排序可以带来的好处）。