Nvidia CUDA

NVIDIA CUDA(Compute Unified Device Architecture)是由NVIDIA开发的并行计算平台和编程模型。CUDA使开发者能够使用高性能NVIDIA GPU(图形处理器)执行通用计算任务。通过CUDA,开发者可以利用GPU强大的计算能力进行大规模并行计算,显著提高应用程序性能,特别是在科学计算、机器学习、图像处理和物理模拟等领域。CUDA支持多种编程语言,如C、C++和Python,并提供丰富的库和工具集,帮助开发者轻松实现GPU加速计算。

cuda logo

CUDA编程基础

我们如何在程序中调用GPU资源来加速代码执行?通过使用CUDA,我们可以轻松管理GPU内的数千个计算核心。CUDA编程模型是一种异构模型,需要CPU和GPU协同工作。在CUDA编程中,使用"host"和"device"这两个术语来区分代码执行的设备。

  • Host:CPU和主机内存
  • Device:GPU和显存

通常,主程序加载在CPU上,然后将初始化数据复制到GPU。GPU完成计算后,将结果复制回CPU。

我们可以在CUDA程序中使用kernel来实现上述过程。在CUDA中,kernel指的是一个函数,当调用时,会导致GPU同时启动许多线程来执行这个kernel,从而实现并行性。每个线程使用其线程ID来对应输入数据的索引,确保每个线程执行相同的kernel但处理不同的数据。

GPU拥有大量计算核心,CUDA使用两级组织结构来管理这些核心。以下是CUDA编程中需要介绍的三个概念:Thread、Block和Grid。

  • Thread:CUDA并行程序由许多线程执行。
  • Block:若干线程组合在一起形成一个块。
  • Grid:多个块进一步组织成一个网格。

注意:同一块中的线程可以通过共享内存进行同步和通信。

cpu and gpu

在CUDA中,每个线程都有一个唯一的标识符称为ThreadIdx,它根据Grid和Block的划分方式而变化。

此外,下面是CUDA内存模型的简要介绍,如下图所示。每个线程都有自己私有的本地内存,而每个线程块包含共享内存,块内的所有线程都可以访问,其生命周期与线程块相同。此外,所有线程都可以访问全局内存。还有一些只读内存空间:常量内存和纹理内存。内存结构涉及程序优化,这里不再详细讨论。

cpu and gpu

CUDA演示案例

现在,让我们通过一个简单的例子来体验GPU加速。在这个例子中,我们将在CPU和GPU上计算两个向量的和,并比较CPU和GPU之间的计算速度。

准备运行环境

JetPack操作系统已经包含了CUDA运行时环境。我们可以使用jtop工具检查CUDA版本。打开终端并运行jtop。

bash
jtop

jtop

注意:如果您的系统未安装jtop,请参阅本指南

下面我们将使用Python编写CUDA代码。因此,需要在终端中运行以下命令来安装Numba

bash
pip install numba

运行示例代码

步骤1: 在根目录中创建本课程的Python脚本:

bash
mkdir -p ~/reComputer_J/3.4_CUDA
touch ~/reComputer_J/3.4_CUDA/add_compare.py

步骤2: 在脚本所在目录中打开VSCode:

bash
cd ~/reComputer_J/3.4_CUDA
code .

步骤3: 在左侧导航栏中,我们可以看到hello_world.py。在python文件中输入以下代码:

点击展开代码
python
from numba import cuda, float32
import numpy as np
import time


@cuda.jit
def add_gpu(a, b, out):
  tx = cuda.threadIdx.x
  ty = cuda.blockIdx.x
  block_size = cuda.blockDim.x
  grid_size = cuda.gridDim.x
  start = tx + ty * block_size
  stride = block_size * grid_size

  for i in range(start, a.shape[0], stride):
      out[i] = a[i] + b[i]

def add_cpu(a, b, out):
  for i in range(a.shape[0]):    
      out[i] = a[i] + b[i]

def main():
  # Prepare test data
  n = 10000000
  a = np.arange(n).astype(np.float32)
  b = np.arange(n).astype(np.float32)
  out = np.empty_like(a)

  # 1. Test CPU computation time
  start = time.time()
  add_cpu(a, b, out)
  print("CPU Time Taken:", time.time()-start)

  # 2. Test GPU computation time
  # Copy data from host memory to device memory
  d_a = cuda.to_device(a)
  d_b = cuda.to_device(b)
  d_out = cuda.device_array_like(out)
  # Define the number of threads per block and the number of blocks
  threads_per_block = 256  
  blocks_per_grid = (n + threads_per_block - 1) // threads_per_block
  # Compute the result
  start_gpu = time.time()
  add_gpu[blocks_per_grid, threads_per_block](d_a, d_b, d_out)
  end_gpu = time.time()
  # Copy the result from device memory back to host memory
  d_out.copy_to_host(out)

  print("GPU Time Taken:", end_gpu-start_gpu)


if __name__ == "__main__":
  main()

步骤4: 点击运行按钮,在终端中查看Python脚本的结果。

launch vscode

从终端打印的结果可以看出,GPU的计算速度更快。

更多CUDA学习教程

教程类型描述
Getting Started with CUDA Programming(zh)文档中国人写的CUDA编程入门极简教程
CUDA C++ Programming Guide文档CUDA模型和接口的编程指南