1. GPU和CUDA
GPU,即图形处理器,是一种专门用于处理图像和视频的计算设备。GPU由许多小型处理器核心组成,可以同时处理许多简单的计算任务。与CPU不同,GPU具有数千个处理器核心,这使得它们非常适合进行并行计算。
CUDA是一种并行计算平台和编程模型,用于利用GPU进行高性能计算。CUDA是NVIDIA推出的技术,是GPU计算的事实标准。
2. CUDA编程基础
#include <stdio.h>
__global__ void helloCUDA()
{
printf("Hello CUDA from GPU!\n");
}
int main()
{
helloCUDA<<<1,1>>>();
cudaDeviceSynchronize();
return 0;
}
2.1 核函数
CUDA的核心概念是核函数。核函数是在GPU上执行的函数,可以并行地处理多个数据元素。在CUDA中,核函数被称为“kernel”。
在Kernel中,可以使用特殊的语法来访问GPU内存、线程和块,并使用“<<<…>>>”运算符在GPU上调用它。这些语法包括__global__、shared、__device__等。在上面的示例中,定义了一个名为“helloCUDA”的Kernel,并使用“<<<1,1>>>”运算符在GPU上调用它。
2.2 线程网格和线程块
每个核函数由一个线程网格和一个线程块组成。线程网格由多个线程块组成,线程块由多个线程组成。在编写CUDA程序时,需要明确线程网格和线程块之间的关系。
线程块和线程网格之间的关系非常重要。线程块中的线程可以共享本地内存和同步操作,而线程网格中的线程可以共享全局内存。在GPU中,读取全局内存需要较长时间,因此使用本地内存可以显著提高性能。在CUDA中,线程块和线程被称为“Block”和“Thread”。
一个Block是由多个Thread组成的,一个Kernel可以启动多个Block。在上面的示例中,使用“<<<1,1>>>”运算符启动了一个Block,并使用“blockIdx.x”和“threadIdx.x”来访问当前Block和Thread的ID。
CUDA内存模型
CUDA使用一种特殊的内存模型来管理GPU内存。在CUDA程序中,可以使用以下四种类型的内存:
- Global memory:全局内存是GPU上的主要内存池。它可以被所有的Block和Thread访问,并可以在主机和设备之间传输数据。
- Shared memory:共享内存是Block内的线程之间共享的内存。它可以用于加速数据访问。
- Constant memory:常量内存是只读的内存区域,可以被所有Block和Thread访问。它通常用于存储常量数据。
- Texture memory:纹理内存是一种高速缓存,用于存储图像数据。它具有一些高级特性,例如自动缩放和插值。
- 在CUDA程序中,可以使用以下函数来分配、释放和访问GPU内存:
cudaMalloc():用于在GPU上分配全局内存。
cudaFree():用于释放GPU上的全局内存。
cudaMemcpy():用于在主机和设备之间复制数据。
global void():用于在GPU上定义Kernel。
shared void():用于在共享内存中声明变量。
参数 | 解释 |
---|---|
cudaMalloc() | 用于在GPU上分配全局内存 |
cudaFree() | 用于释放GPU上的全局内存 |
cudaMemcpy() | 用于在主机和设备之间复制数据 |
cudaMalloc() | 用于在GPU上定义Kernel |
cudaMalloc() | 用于在共享内存中声明变量 |
示例程序:向量加法
现在,将介绍一个向量加法的CUDA程序,在这个程序中,将在GPU上对两个向量进行加法运算,并将结果存储在第三个向量中。以下是程序的主要部分:
#include <stdio.h>
__global__ void vecAdd(float *a, float *b, float *c, int n)
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) {
c[i] = a[i] + b[i];
}
}
int main()
{
int n = 1000;
float *a, *b, *c;
float *d_a, *d_b, *d_c;
int size = n * sizeof(float);
a = (float*)malloc(size);
b = (float*)malloc(size);
c = (float*)malloc(size);
cudaMalloc(&d_a, size);
cudaMalloc(&d_b, size);
cudaMalloc(&d_c, size);
for (int i = 0; i < n; i++) {
a[i] = i;
b[i] = i * 2;
}
cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice);
int blockSize = 256;
int numBlocks = (n + blockSize - 1) / blockSize;
vecAdd<<<numBlocks, blockSize>>>(d_a, d_b, d_c, n);
cudaMemcpy(c, d_c, size, cudaMemcpyDeviceToHost);
for (int i = 0; i < n; i++) {
printf("%f\n", c[i]);
}
free(a);
free(b);
free(c);
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
return 0;
}
在这个程序中,首先在主机上分配三个向量(a、b、c)的内存,并使用cudaMalloc()在GPU上分配三个向量(d_a、d_b、d_c)的内存。然后使用cudaMemcpy()函数将主机上的向量a和b复制到GPU上的向量d_a和d_b中。
接下来,计算需要启动的Block数量和每个Block中需要启动的Thread数量。在本例中,将Block大小设置为256,将向量大小设置为1000,因此需要启动4个Block。然后使用<<<…>>>运算符启动Kernel,并在GPU上执行向量加法运算。
最后使用cudaMemcpy()函数将GPU上的向量d_c复制回主机上,然后释放所有主机和GPU内存。
这个程序中的Kernel函数是向量加法,其定义如下:
__global__ void vecAdd(float *a, float *b, float *c, int n)
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) {
c[i] = a[i] + b[i];
}
}
在这个Kernel函数中,每个线程将计算一个向量元素的和。i是线程的唯一标识符,计算方式是通过块号和线程号计算得到。这里使用了CUDA内置的blockIdx和threadIdx变量。首先计算线程的全局索引i,如果i小于向量大小n,则将a和b向量中的元素相加,并将结果存储在c向量中。
请注意,该程序在启动Kernel时使用了<<<…>>>运算符。该运算符指定了Kernel应该在多少个Block和每个Block中有多少个Thread上运行。
此外,该程序使用了cudaMemcpy()函数,该函数用于在主机和GPU之间复制数据。该函数具有四个参数:源指针、目标指针、复制的字节数以及复制方向(从主机到GPU或从GPU到主机)。
总结
本文简要介绍了CUDA的基本概念和编程模型,包括核函数、线程网格和线程块、内存管理和编译器。此外,还介绍了一个简单的CUDA程序,可以将两个向量相加。