我正在寻找可以为 CPU(使用 g++)和 GPU(使用 nvcc)编写的最简洁的代码量,GPU 的性能始终优于 CPU。任何类型的算法都是可以接受的。
澄清一下:我实际上是在寻找两个短代码块,一个用于 CPU(在 g++ 中使用 C++),另一个用于 GPU(在 nvcc 中使用 C++),GPU 的性能优于该代码块。最好以秒或毫秒为单位。可能的最短代码对。
首先,我将重申我的评论:GPU 是高带宽、高延迟的。试图让 GPU 在纳秒的工作(甚至是毫秒或秒的工作)上击败 CPU 完全失去了做 GPU 工作的意义。下面是一些简单的代码,但要真正体会到 GPU 的性能优势,你需要一个很大的问题规模来分摊启动成本……否则,它就毫无意义。我可以在两英尺的比赛中击败法拉利,仅仅是因为转动钥匙、启动发动机和踩踏板需要一些时间。这并不意味着我在任何有意义的方面都比法拉利更快。
在 C++ 中使用类似的东西:
#define N (1024*1024)
#define M (1000000)
int main()
{
float data[N]; int count = 0;
for(int i = 0; i < N; i++)
{
data[i] = 1.0f * i / N;
for(int j = 0; j < M; j++)
{
data[i] = data[i] * data[i] - 0.25f;
}
}
int sel;
printf("Enter an index: ");
scanf("%d", &sel);
printf("data[%d] = %f\n", sel, data[sel]);
}
在 CUDA/C 中使用类似的东西:
#define N (1024*1024)
#define M (1000000)
__global__ void cudakernel(float *buf)
{
int i = threadIdx.x + blockIdx.x * blockDim.x;
buf[i] = 1.0f * i / N;
for(int j = 0; j < M; j++)
buf[i] = buf[i] * buf[i] - 0.25f;
}
int main()
{
float data[N]; int count = 0;
float *d_data;
cudaMalloc(&d_data, N * sizeof(float));
cudakernel<<<N/256, 256>>>(d_data);
cudaMemcpy(data, d_data, N * sizeof(float), cudaMemcpyDeviceToHost);
cudaFree(d_data);
int sel;
printf("Enter an index: ");
scanf("%d", &sel);
printf("data[%d] = %f\n", sel, data[sel]);
}
如果这不起作用,请尝试使 N 和 M 更大,或将 256 更改为 128 或 512。
一个非常非常简单的方法是计算前 100,000 个整数的平方,或者一个大型矩阵运算。Ita 易于实现,并且通过避免分支、不需要堆栈等来利用 GPU 的优势。我不久前使用 OpenCL 与 C++ 进行了此操作,并获得了一些非常惊人的结果。(2GB GTX460 的性能大约是双核 CPU 的 40 倍。)
您是在寻找示例代码,还是只是想法?
编辑
40x 与双核 CPU 相比,而不是四核。
一些指示:
正如我在对@Paul R 的评论回复中所说,考虑使用 OpenCL,因为它可以让您轻松地在 GPU 和 CPU 上运行相同的代码,而无需重新实现它。
(回想起来,这些可能非常明显。)
作为参考,我用时间测量做了一个类似的例子。使用 GTX 660,GPU 加速为 24 倍,其中除了实际计算之外,它的操作还包括数据传输。
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <time.h>
#define N (1024*1024)
#define M (10000)
#define THREADS_PER_BLOCK 1024
void serial_add(double *a, double *b, double *c, int n, int m)
{
for(int index=0;index<n;index++)
{
for(int j=0;j<m;j++)
{
c[index] = a[index]*a[index] + b[index]*b[index];
}
}
}
__global__ void vector_add(double *a, double *b, double *c)
{
int index = blockIdx.x * blockDim.x + threadIdx.x;
for(int j=0;j<M;j++)
{
c[index] = a[index]*a[index] + b[index]*b[index];
}
}
int main()
{
clock_t start,end;
double *a, *b, *c;
int size = N * sizeof( double );
a = (double *)malloc( size );
b = (double *)malloc( size );
c = (double *)malloc( size );
for( int i = 0; i < N; i++ )
{
a[i] = b[i] = i;
c[i] = 0;
}
start = clock();
serial_add(a, b, c, N, M);
printf( "c[0] = %d\n",0,c[0] );
printf( "c[%d] = %d\n",N-1, c[N-1] );
end = clock();
float time1 = ((float)(end-start))/CLOCKS_PER_SEC;
printf("Serial: %f seconds\n",time1);
start = clock();
double *d_a, *d_b, *d_c;
cudaMalloc( (void **) &d_a, size );
cudaMalloc( (void **) &d_b, size );
cudaMalloc( (void **) &d_c, size );
cudaMemcpy( d_a, a, size, cudaMemcpyHostToDevice );
cudaMemcpy( d_b, b, size, cudaMemcpyHostToDevice );
vector_add<<< (N + (THREADS_PER_BLOCK-1)) / THREADS_PER_BLOCK, THREADS_PER_BLOCK >>>( d_a, d_b, d_c );
cudaMemcpy( c, d_c, size, cudaMemcpyDeviceToHost );
printf( "c[0] = %d\n",0,c[0] );
printf( "c[%d] = %d\n",N-1, c[N-1] );
free(a);
free(b);
free(c);
cudaFree( d_a );
cudaFree( d_b );
cudaFree( d_c );
end = clock();
float time2 = ((float)(end-start))/CLOCKS_PER_SEC;
printf("CUDA: %f seconds, Speedup: %f\n",time2, time1/time2);
return 0;
}
我同意 David 关于 OpenCL 是测试这一点的好方法的评论,因为在 CPU 和 GPU 上运行的代码之间切换是多么容易。如果您能够在 Mac 上工作,Apple 提供了一些使用 OpenCL 进行 N 体模拟的示例代码,内核在 CPU、GPU 或两者上运行。您可以在它们之间实时切换,并在屏幕上显示 FPS 计数。
对于一个更简单的情况,他们有一个“hello world”OpenCL 命令行应用程序,它以类似于 David 所描述的方式计算平方。这很可能可以毫不费力地移植到非 Mac 平台上。要在 GPU 和 CPU 使用率之间切换,我相信您只需要更改
int gpu = 1;
在 hello.c 源文件中的行中,CPU 为 0,GPU 为 1。
Apple 在他们的主要 Mac 源代码列表中有更多的 OpenCL 示例代码。
David Gohara 博士在有关该主题的介绍性视频会议的最后(大约 34 分钟)执行分子动力学计算时,举了一个 OpenCL 的 GPU 加速示例。在他的计算中,他看到从在 8 个 CPU 内核上运行的并行实现到单个 GPU 的加速大约提高了 27 倍。同样,这不是最简单的示例,但它展示了一个真实世界的应用程序以及在 GPU 上运行某些计算的优势。
我还使用 OpenGL ES 着色器对移动空间进行了一些修改,以执行基本计算。我发现,在 GPU 上作为着色器运行时,在图像上运行的简单颜色阈值着色器比在 CPU 上为该特定设备执行的相同计算快大约 14-28 倍。