cuda¶

NVIDIA®CUDA®工具包为创建高性能GPU加速应用程序提供了一个开发环境。有了它，您可以在GPU加速的嵌入式系统、台式工作站、企业数据中心、基于云的平台和超级计算机上开发、优化和部署应用程序。该工具包包括GPU加速库、调试和优化工具、C/C++编译器和运行库。

安装列表¶

序号	集群	版本	模块名	位置
1	hpckapok1	12.3	cuda/12.3	/share/software/cuda/local/cuda-12.3
2	hpckapok1	11.8	cuda/11.8	/share/software/cuda/local/cuda-11.8
3	hpckapok1	11.7	cuda/11.7	/share/software/cuda/local/cuda-11.7
4	hpckapok1	11.3.0	cuda/11.3.0	/share/software/cuda/local/cuda-11.3.0
5	hpckapok1	10.2	cuda/10.2	/share/software/cuda/local/cuda-10.2
6	hpckapok2	12.3	cuda/12.3.0	/public/software/cuda/local/cuda-12.3
7	hpckapok2	11.8.0	cuda/11.8.0	/public/software/cuda/local/cuda-11.8.0
8	hpckapok2	11.6	cuda/11.6.2	/public/software/cuda/local/cuda-11.6
9	hpckapok2	10.2.0	cuda/10.2.0	/public/software/cuda/local/cuda-10.2.0

cuDNN¶

NVIDIA CUDA®深度神经网络库（cuDNN）是用于深度神经网络的GPU加速基元库。cuDNN为标准例程提供了高度调优的实现，如前向和后向卷积、注意力(attention)、matmul、池化(pooling)和规范化(normalization)。

安装环境¶

序号	cuda版本	cuDNN版本
1	cuda-10.2	8.7.0
2	cuda-11.x,cuda-12.x	8.9.7

注：cuDNN库文件位于cuda文件夹的lib64目录下，包含文件位于cuda文件夹的include文件目录下。

一个简单的例子¶

我们将从一个简单的C++程序开始，该程序添加两个数组的元素，每个数组有一百万个元素，在cpu上运行。

#include <iostream>
#include <math.h>

// function to add the elements of two arrays
void add(int n, float *x, float *y)
{
  for (int i = 0; i < n; i++)
      y[i] = x[i] + y[i];
}

int main(void)
{
  int N = 1<<20; // 1M elements

  float *x = new float[N];
  float *y = new float[N];

  // initialize x and y arrays on the host
  for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
  }

  // Run kernel on 1M elements on the CPU
  add(N, x, y);

  // Check for errors (all values should be 3.0f)
  float maxError = 0.0f;
  for (int i = 0; i < N; i++)
    maxError = fmax(maxError, fabs(y[i]-3.0f));
  std::cout << "Max error: " << maxError << std::endl;

  // Free memory
  delete [] x;
  delete [] y;

  return 0;
}

编译并运行这个C++程序。将上面的代码放在一个文件中，并将其保存为add.cpp，然后使用C++编译器进行编译，在Linux上使用g++：

$ g++ add.cpp -o add_cpu

然后运行：

$ ./add_cpu
Max error: 0

不出所料，它打印出总和中没有错误，然后退出。

现在我想让这个计算（并行）在GPU的许多核心上运行。

首先，我只需要将我们的add函数转换为GPU可以运行的函数，称为CUDA中的内核。要做到这一点，我所要做的就是将说明符__global__添加到函数中，这告诉CUDA C++编译器这是一个在GPU上运行的函数，可以从CPU代码中调用。

// CUDA Kernel function to add the elements of two arrays on the GPU
__global__
void add(int n, float *x, float *y)
{
  for (int i = 0; i < n; i++)
      y[i] = x[i] + y[i];
}

这些__global__函数被称为内核，在GPU上运行的代码通常被称为设备代码，而在CPU上运行的是主机代码。

然后，要在GPU上进行计算，需要分配GPU可访问的内存。CUDA中的统一内存通过提供可供系统中所有GPU和CPU访问的单个内存空间，使这一点变得容易。要在统一内存中分配数据，请调用cudaMallocManaged（），它返回一个指针，您可以从主机（CPU）代码或设备（GPU）代码访问该指针。要释放数据，只需将指针传递给cudaFree（）。

只需要将上面代码中对new的调用替换为对cudaMallocManaged（）的调用，并将对delete[]的调用替换成对cudaFree的调用。

// Allocate Unified Memory -- accessible from CPU or GPU
float *x, *y;
cudaMallocManaged(&x, N*sizeof(float));
cudaMallocManaged(&y, N*sizeof(float));

...

// Free memory
cudaFree(x);
cudaFree(y);

最后，需要启动add（）内核，它在GPU上调用它。CUDA内核启动是使用三角括号语法<<>>>指定的。我只需要将它添加到要添加到参数列表之前的调用中。

add<<<1, 1>>>(N, x, y);

现在，您只需要知道这一行启动一个GPU线程来运行add（）。

还有一件事：需要CPU等待内核完成后再访问结果（因为CUDA内核启动不会阻止调用CPU线程）。要做到这一点，只需在对CPU进行最后的错误检查之前调用cudaDeviceSynchronize（）。

以下是完整的代码：

#include <iostream>
#include <math.h>
// Kernel function to add the elements of two arrays
__global__
void add(int n, float *x, float *y)
{
  for (int i = 0; i < n; i++)
    y[i] = x[i] + y[i];
}

int main(void)
{
  int N = 1<<20;
  float *x, *y;

  // Allocate Unified Memory – accessible from CPU or GPU
  cudaMallocManaged(&x, N*sizeof(float));
  cudaMallocManaged(&y, N*sizeof(float));

  // initialize x and y arrays on the host
  for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
  }

  // Run kernel on 1M elements on the GPU
  add<<<1, 1>>>(N, x, y);

  // Wait for GPU to finish before accessing on host
  cudaDeviceSynchronize();

  // Check for errors (all values should be 3.0f)
  float maxError = 0.0f;
  for (int i = 0; i < N; i++)
    maxError = fmax(maxError, fabs(y[i]-3.0f));
  std::cout << "Max error: " << maxError << std::endl;

  // Free memory
  cudaFree(x);
  cudaFree(y);

  return 0;
}

CUDA文件的文件扩展名为.cu。因此，将此代码保存在名为add.cu的文件中，并使用CUDA C++编译器nvcc进行编译。

$ module load cuda/12.3
$ nvcc add.cu -o add_cuda

需要放在gpu节点上执行，创建脚本add.slurm：

#!/bin/bash
#SBATCH -J test
#SBATCH -p gpuA800
#SBATCH -o %j.out
#SBATCH -e %j.err
#SBATCH -N 1
#SBATCH --ntasks-per-node=1
#SBATCH --gres=gpu:1

module load cuda/12.3
./add_cuda

提交脚本：

$ sbatch add.slurm

查看输出结果%j.out，可知道结果为：Max error: 0.000000

这只是第一步，因为正如所写的，这个内核只适用于单个线程，因为运行它的每个线程都将对整个数组执行加法。此外，还存在竞争条件，因为多个并行线程将同时读取和写入相同的位置。

既然您已经运行了一个内核，其中有一个线程可以进行一些计算，那么如何使其并行呢？关键在于CUDA的<<1，1>>>语法。这被称为执行配置，它告诉CUDA运行时在GPU上启动要使用多少并行线程。这里有两个参数，但让我们从更改第二个参数开始：线程块中的线程数。CUDA GPU使用大小为32的倍数的线程块来运行内核，因此256个线程是一个合理的选择。

add<<<1, 256>>>(N, x, y);

如果我只运行这个更改的代码，它将在每个线程中进行一次计算，而不是将计算分散到并行线程中。为了正确地执行此操作，我需要修改内核。CUDA C++提供了关键字，让内核获取正在运行的线程的索引。具体来说，threadIdx.x包含块中当前线程的索引，blockDim.x包含块中的线程数。我只需修改循环，使其使用并行线程遍历数组。

__global__
void add(int n, float *x, float *y)
{
  int index = threadIdx.x;
  int stride = blockDim.x;
  for (int i = index; i < n; i += stride)
      y[i] = x[i] + y[i];
}

add函数没有太大变化。事实上，将index设置为0，将步幅设置为1，使其在语义上与第一个版本相同。

将文件保存为add_block.cu，然后再次用nvcc编译并运行。可以看到运行速度有很大提升。

参考资料¶

Contributor：mzliu

最后更新: 2025 年 06 月 30 日