CUDA 的性能取决于声明变量

Question

在这种情况下是否有任何提高 CUDA 性能的技巧，例如声明全局/局部变量、参数传递、内存复制。

在下面的示例中，我试图找出 sum_gpu_FAST 和 sum_gpu_SLOW 两种性能差异太大的原因。

在这里你可以看到整个示例代码。

#include <iostream>
#include <chrono>
#define N 10000000
__global__
void sum_gpu_FAST(int (&data)[N][2], int& sum, int n) {  // runtime : 2.42342s
    int s = 0;
    for (int i = 0; i < n; i++) 
        s += data[i][0] * 10 + data[i][1];
    sum = s;
}
__global__
void sum_gpu_SLOW(int (&data)[N][2], int& sum, int n) {  // runtime : 436.64ms
    sum = 0;
    for (int i = 0; i < n; i++) {
        sum += data[i][0] * 10 + data[i][1];
    }
}
void sum_cpu(int (*data)[2], int& sum, int n) {
    for (int i = 0; i < n; i++) {
        sum +=  data[i][0] * 10 + data[i][1];
    }
}
int main()
{
    int (*v)[2] = new int[N][2];
    for (int i = 0; i < N; i++)
        v[i][0] = 1, v[i][1] = 3;
    printf ("-CPU------------------------------------------------\n");
    {
        int sum = 0;
        auto start = std::chrono::system_clock::now();
        sum_cpu(v, sum, N);
        auto end   = std::chrono::system_clock::now();
        // print output
        std::cout << sum << " / " << (end-start).count() / 1000000 << "ms" << std::endl;
    }
    printf ("-GPU-Ready------------------------------------------\n");
    int *dev_sum       = nullptr;
    int (*dev_v)[N][2] = nullptr;
    cudaMalloc((void **)&dev_v,   sizeof(int[N][2]));
    cudaMalloc((void **)&dev_sum, sizeof(int));
    cudaMemcpy(dev_v, v, sizeof(int[N][2]), cudaMemcpyHostToDevice);
    printf("-GPU-FAST-------------------------------------------\n");
    {
        int sum = 0;
        auto start = std::chrono::system_clock::now();
        sum_gpu_FAST<<<1, 1>>> (*dev_v, *dev_sum, N);
        cudaDeviceSynchronize(); // wait until end of kernel
        auto end   = std::chrono::system_clock::now();
        // print output
        cudaMemcpy( &sum, dev_sum, sizeof(int), cudaMemcpyDeviceToHost );
        std::cout << sum << " / " << (end-start).count() / 1000000 << "ms" << std::endl;
    }
    printf("-GPU-SLOW-------------------------------------------\n");
    {
        int sum = 0;
        auto start = std::chrono::system_clock::now();
        sum_gpu_SLOW<<<1, 1>>> (*dev_v, *dev_sum, N);
        cudaDeviceSynchronize(); // wait until end of kernel
        auto end   = std::chrono::system_clock::now();
        // print output
        cudaMemcpy( &sum, dev_sum, sizeof(int), cudaMemcpyDeviceToHost );
        std::cout << sum << " / " << (end-start).count() / 1000000 << "ms" << std::endl;
    }
    printf("----------------------------------------------------\n");
    return 0;
}

Answer 1

在下面的示例中，我试图找出 sum_gpu_FAST 和 sum_gpu_SLOW 两种性能差异太大的原因。

在快速情况下，您正在创建一个包含（可能）在寄存器中的局部变量：

int s = 0;

在循环迭代期间，从全局内存中进行读取，但唯一的写入操作是对寄存器：

for (int i = 0; i < n; i++) 
    s += data[i][0] * 10 + data[i][1];

在慢速情况下，运行总和包含在驻留在全局内存中的变量中：

sum = 0;

因此，在每次循环迭代中，更新的值都会写入全局内存：

for (int i = 0; i < n; i++) {
    sum += data[i][0] * 10 + data[i][1];

因此，循环在每次迭代时都会有额外的开销来写入全局内存，这比将总和保存在寄存器中要慢。

我不会完全剖析 SASS 代码来比较这两种情况，因为编译器会在快速情况下围绕循环展开和可能的其他因素做出其他决定，但我的猜测是不需要在循环迭代期间将结果存储到全局内存也大大有助于循环展开。但是我们可以根据每个案例的 SASS 代码的尾部做一个简单的推论：

                Function : _Z12sum_gpu_FASTRA10000000_A2_iRii
        .headerflags    @"EF_CUDA_SM70 EF_CUDA_PTX_SM(EF_CUDA_SM70)"
        /*0000*/                   MOV R1, c[0x0][0x28] ;                        /* 0x00000a0000017a02 */
                                                                                 /* 0x000fd00000000f00 */
...
        /*0b00*/                   STG.E.SYS [R2], R20 ;                         /* 0x0000001402007386 */
                                                                                 /* 0x000fe2000010e900 */
        /*0b10*/                   EXIT ;                                        /* 0x000000000000794d */
                                                                                 /* 0x000fea0003800000 */

在上面的快速案例中，我们看到在内核末尾、返回语句 (EXIT) 之前以及内核中的任何循环之外有一个全局存储 (STG) 指令.虽然我没有全部展示出来，但在快速内核中确实没有其他的STG 指令，除了最后的那个。我们在慢内核的尾部看到了一个不同的故事：

        code for sm_70
                Function : _Z12sum_gpu_SLOWRA10000000_A2_iRii
        .headerflags    @"EF_CUDA_SM70 EF_CUDA_PTX_SM(EF_CUDA_SM70)"
        /*0000*/                   IMAD.MOV.U32 R1, RZ, RZ, c[0x0][0x28] ;       /* 0x00000a00ff017624 */
                                                                                 /* 0x000fd000078e00ff */
...
        /*0460*/                   STG.E.SYS [R2], R7 ;                          /* 0x0000000702007386 */
                                                                                 /* 0x0005e2000010e900 */
        /*0470*/              @!P0 BRA 0x2f0 ;                                   /* 0xfffffe7000008947 */
                                                                                 /* 0x000fea000383ffff */
        /*0480*/                   EXIT ;                                        /* 0x000000000000794d */
                                                                                 /* 0x000fea0003800000 */

慢速内核以循环内的STG 指令结束循环。慢速内核在整个内核中也有许多 STG 指令实例，可能是因为编译器展开。