近年来,随着人工智能的发展,“卖铲人”NVIDIA在人工智能行业的地位越来越重要。CUDA编程也是人工智能重要知识点,本系列通过学习三章文章对CUDA编程有个快速理解和入门:第一章介绍基础知识点,理解CUDA编程基础概念;第二章通过2个简单程序的编程、编译和执行对CUDA编程有个初步掌握;第三章通过CUDA并行编程的实操掌握入门。
CUDA 与 C 的快速比较
下表并排比较了 C 和 CUDA 中的 hello world 程序。C
void c_hello(){
printf("Hello World!\n");
}
int main() {
c_hello();
return 0;
}
CUDA
__global__ void cuda_hello(){
printf("Hello World from GPU!\n");
}
int main() {
cuda_hello<<<1,1>>>();
return 0;
}
C 和 CUDA 实现的主要区别在于global说明符和<<<…>>>语法。global说明符表示在设备(GPU)上运行的函数。此类函数可以通过主机代码调用,例如main()示例中的函数,也称为“内核”。
调用内核时,其执行配置通过<<<…>>>语法提供,例如cuda_hello<<<1,1>>>()。在 CUDA 术语中,这称为“内核启动(1,1)”。我们将在第三章讨论该参数。
编译 CUDA 程序
编译 CUDA 程序与编译 C 程序类似。NVIDIA nvcc 在 CUDA 工具包中提供了一个名为 CUDA 编译器来编译 CUDA 代码,通常存储在扩展名为 的文件中 .cu 。例如$> nvcc hello.cu -o hello
使用上述命令编译 CUDA 程序时,你可能会看到以下警告
nvcc warning : The 'compute_20', 'sm_20', and 'sm_21' architectures are deprecated, and may be removed in a future release (Use -Wno-deprecated-gpu-targets to suppress warning).
现在可以忽略此警告。
执行
CUDA hello world 示例不执行任何操作,即使程序已编译,屏幕上也不会显示任何内容。为了使操作生效,我们将研究向量加法。以下是用 C 语言实现的向量加法示例( )。该示例计算存储在数组和./vector_add.c中的两个向量的加法,并将结果放入数组 中。a b out
#define N 10000000
void vector_add(float *out, float *a, float *b, int n) {
for(int i = 0; i < n; i++){
out[i] = a[i] + b[i];
}
}
int main(){
float *a, *b, *out;
// Allocate memory
a = (float*)malloc(sizeof(float) * N);
b = (float*)malloc(sizeof(float) * N);
out = (float*)malloc(sizeof(float) * N);
// Initialize array
for(int i = 0; i < N; i++){
a[i] = 1.0f; b[i] = 2.0f;
}
// Main function
vector_add(out, a, b, N);
}
练习:将向量加法转换为 CUDA
在第一个练习中,我们将使用 hello world 作为示例转换 vector_add.c 为 CUDA 程序。 vector_add.cu复制vector_add.c到****vector_add.cu
$> cp vector_add.c vector_add.cu
转换vector_add()为 GPU 内核
__global__ void vector_add(float *out, float *a, float *b, int n) {
for(int i = 0; i < n; i++){
out[i] = a[i] + b[i];
}
}
vector_add()将调用更改main()****为内核调用
vector_add<<<1,1>>>(out, a, b, N);
编译并运行程序
$> nvcc vector_add.c -o vector_add
$> ./vector_add
你会注意到程序无法正常工作。原因是 CPU 和 GPU 是独立的实体。两者都有自己的内存空间。CPU 不能直接访问 GPU 内存,反之亦然。在 CUDA 术语中,CPU 内存称为主机内存,GPU 内存称为设备内存。指向 CPU 和 GPU 内存的指针分别称为主机指针和设备指针。
设备内存管理
CUDA 提供了几个用于分配设备内存的函数。最常见的是 cudaMalloc() 和 cudaFree() 。这两个函数的语法如下cudaMalloc(void **devPtr, size_t count);
cudaFree(void *devPtr);
cudaMalloc() count在设备内存中分配大小的内存,并将设备指针更新devPtr为分配的内存。cudaFree()释放设备指针指向的设备内存区域<font style="color:rgb(231, 76, 60);background-color:rgb(252, 252, 252);">devPtr</font>
。它们分别相当于 C 中的malloc()和free()
内存传输
主机和设备内存之间的数据传输可以通过函数完成 cudaMemcpy ,类似于 memcpy C 语言。语法 cudaMemcpy 如下cudaMemcpy(void *dst, void *src, size_t count, cudaMemcpyKind kind)
函数将大小为 的内存count从src复制到dst。kind表示方向。,对于典型用法, kind的值是cudaMemcpyHostToDevice或cudaMemcpyDeviceToHost。 还有其他可能的值,但在本教程中我们不会涉及它们。
练习(续):完成向量加法
1. 为数组 a 、 b 和分配和释放设备内存 out 。 2. 在主机和设备内存之间传输 a 、 b 和 out 。 - 测验:内核执行之前和之后必须传输哪个数组?示例:数组“a”的解决方案
```plain void main(){ float *a, *b, *out; float *d_a;a = (float*)malloc(sizeof(float) * N);
// Allocate device memory for a
cudaMalloc((void**)&d_a, sizeof(float) * N);
// Transfer data from host to device memory
cudaMemcpy(d_a, a, sizeof(float) * N, cudaMemcpyHostToDevice);
…
vector_add<<<1,1>>>(out, d_a, b, N);
…
// Cleanup after kernel execution
cudaFree(d_a);
free(a);
}
1. <font style="color:rgb(64, 64, 64);background-color:rgb(252, 252, 252);">编译并测量性能。(请参阅(</font>[<font style="color:rgb(64, 64, 64);background-color:rgb(252, 252, 252);">./solutions/vector_add.cu</font>](https://cuda-tutorial.readthedocs.io/en/latest/tutorials/tutorial01/solutions/vector_add.cu)<font style="color:rgb(64, 64, 64);background-color:rgb(252, 252, 252);">)中的解决方案)</font>
```plain
$> nvcc vector_add.cu -o vector_add
$> time ./vector_add
性能分析
使用 time 不会提供有关程序性能的太多信息。NVIDIA 提供了一个名为 的命令行分析器工具 nvprof ,它提供了有关 CUDA 程序性能的更多深入信息。要分析向量加法,请使用以下命令
$> nvprof ./vector_add
以下是 Tesla M2050 上的分析结果示例
==6326== Profiling application: ./vector_add
==6326== Profiling result:
Time(%) Time Calls Avg Min Max Name
97.55% 1.42529s 1 1.42529s 1.42529s 1.42529s vector_add(float*, float*, float*, int)
1.39% 20.318ms 2 10.159ms 10.126ms 10.192ms [CUDA memcpy HtoD]
1.06% 15.549ms 1 15.549ms 15.549ms 15.549ms [CUDA memcpy DtoH]
总结
在本章中,我们演示了如何在 CUDA 中编写一个简单的向量加法。我们介绍了 GPU 内核及其从主机代码的执行。此外,我们还介绍了 CPU 和 GPU 之间分离内存空间的概念。我们还演示了如何管理设备内存。但是,我们仍然没有并行运行程序。内核执行配置<<<1,1>>>表明内核仅以 1 个线程启动。在下一章中,我们将修改向量加法以并行运行。