为什么在这个例子中 PyCUDA 比 C CUDA 快

问题描述

我正在探索从 OpenCL 迁移到 CUDA，并做了一些测试来衡量 CUDA 在各种实现中的速度。令我惊讶的是，在下面的示例中，PyCUDA 实现比 C CUDA 示例快约 20%。

我读了很多关于 C CUDA 代码“发布版本”的帖子。我确实尝试在 makefile 中包含 -Xptxas -O3 并且这确实没有任何区别。我还尝试调整执行内核的块大小。不幸的是，它也无助于提高速度。

我的问题是：

• 导致 C CUDA 和 PYCUDA 之间速度差异的原因可能是什么？
• 如果 PYCUDA 中的“高级”（缺乏更好的词）编译是原因之一，我该如何优化 C CUDA 代码的编译？
• 在这种情况下，还有其他方法可以提高 C CUDA 的速度吗？

虽然我很欣赏一般性评论，但我正在寻找可以在我的机器上验证的可行建议。谢谢！

import pycuda.autoinit
import pycuda.driver as drv
import numpy as np

from pycuda.compiler import SourceModule
import time


mod = SourceModule(
    """
__global__ void saxpy(int n,const float a,float *x,float *y)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n){
        y[i] = a * x[i] + y[i];
    }
}
"""
)

saxpy = mod.get_function("saxpy")

N = 1 << 25
time_elapse = 0.0

for i in range(100):
    # print(i)
    # print(N)

    x = np.ones(N).astype(np.float32)
    y = 2 * np.ones(N).astype(np.float32)
    start = time.time()
    saxpy(
        np.int32(N),np.float32(2.0),drv.In(x),drv.InOut(y),block=(512,1,1),grid=(int(N / 512) + 1,)
    time_elapse += (time.time() - start)


print(time_elapse )
print(y[-100:-1])
print(y.sum())
print(N * 4.0)

#include <stdio.h>
#include <time.h>
#define DIM 512



__global__ void saxpy(int n,float a,float *y)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n)
        y[i] = a * x[i] + y[i];
}

int main(int num_iterations)
{
    double start;
    double cputime;
    int N = 1 << 25;
    float *x,*y,*d_x,*d_y;
    int i,j;
    for (j = 0; j < num_iterations; j++)
    {
        x = (float *)malloc(N * sizeof(float));
        y = (float *)malloc(N * sizeof(float));

        cudaMalloc(&d_x,N * sizeof(float));
        cudaMalloc(&d_y,N * sizeof(float));

        for (i = 0; i < N; i++)
        {
            x[i] = 1.0f;
            y[i] = 2.0f;
        }

        cudaMemcpy(d_x,x,N * sizeof(float),cudaMemcpyHostToDevice);
        cudaMemcpy(d_y,y,cudaMemcpyHostToDevice);

        // Perform SAXPY on 1M elements
        start = clock();
        saxpy<<<(N + DIM) / DIM,DIM>>>(N,2.0f,d_x,d_y);
        cputime += ((double)(clock() - start) / CLOCKS_PER_SEC);
        cudaMemcpy(y,d_y,cudaMemcpyDeviceToHost);

        // float maxError = 0.0f;
        // for (int i = 0; i < N; i++){
        //     maxError = max(maxError,abs(y[i] - 4.0f));
        //     //printf("y[%d]: %f\n",i,y[i]);
        // }
        // printf("Max error: %f\n",maxError);

        cudaFree(d_x);
        cudaFree(d_y);
        free(x);
        free(y);
    }

 
    printf("cpu time is %f\n",cputime);
    return 0;
}

我将上述文件保存为 cuda_example.cu 并在 makefile 中使用以下命令编译它：

nvcc -arch=sm_61 -Xptxas -O3,-v -o main cuda_example.cu

解决方法

如果我按原样执行您的 CUDA-C 代码，并将 num_iterations 设置为 300，如下所示：

#include <stdio.h>
#include <time.h>
#define DIM 512

__global__ void saxpy(int n,float a,float *x,float *y)
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n)
    y[i] = a * x[i] + y[i];
}

int main()
{
double start = clock();
int N = 1 << 25;
float *x,*y,*d_x,*d_y;
int i,j;

int num_iterations = 300;
x = (float *)malloc(N * sizeof(float));
y = (float *)malloc(N * sizeof(float));

cudaMalloc(&d_x,N * sizeof(float));
cudaMalloc(&d_y,N * sizeof(float));

for (i = 0; i < N; i++)
{
   x[i] = 1.0f;
   y[i] = 2.0f;
}
cudaMemcpy(d_x,x,N * sizeof(float),cudaMemcpyHostToDevice);
cudaMemcpy(d_y,y,cudaMemcpyHostToDevice);

for (j = 0; j < num_iterations; j++){
    saxpy<<<(N + DIM) / DIM,DIM>>>(N,2.0f,d_x,d_y);
    cudaDeviceSynchronize();
}
cudaMemcpy(y,d_y,cudaMemcpyDeviceToHost);
cudaFree(d_x);
cudaFree(d_y);
free(x);
free(y);

double cputime = ((double)(clock() - start) / CLOCKS_PER_SEC);
printf("cpu time is %f\n",cputime);
return 0;
}

那么在 Geforce GTX 1650 上执行您的程序大约需要 60 秒。您的代码效率极低，因为您每次迭代都在 GPU 和设备之间来回复制数据。 因此，让我们将循环限制为内核执行：

sudo ln -s /snap/dotnet-sdk/current/dotnet /usr/local/bin/dotnet

如果我这样做，那么执行时间变为 1.36 秒。执行类似于 PyCUDA 代码的操作，我得到了大约 19 秒的执行时间。