CUDA - CUDA C&C++简介

pla66

已于 2023-06-12 21:41:05 修改

阅读量426

点赞数 1

分类专栏： CUDA 文章标签： c++ c语言

于 2023-06-05 23:14:10 首次发布

原文链接：https://developer.nvidia.com/blog/easy-introduction-cuda-c-and-c/

版权

CUDA 专栏收录该内容

8 篇文章 2 订阅

订阅专栏

原文链接：An Easy Introduction to CUDA C and C++

更新（2017年1月）：查看新的更简单的CUDA介绍！

文章目录

CUDA编程模型基础
第一个CUDA C程序
编译和运行代码
总结与结论

这篇文章是CUDA C&C++系列文章中的第一篇，CUDA C&C++是CUDA并行计算平台的C/C++接口。本系列文章假设您熟悉C语言编程。我们也将针对Fortran程序员推出一系列关于CUDA Fortran的对应文章。这两个系列将涵盖CUDA平台上并行计算的基本概念。从这里开始，除非我另有说明，否则我将使用术语“CUDA C”作为“CUDA C&C++”的简写。CUDA C本质上是C/C++加入一些扩展，允许使用多个线程并行地在GPU上执行函数。

CUDA编程模型基础

在我们开始CUDA C编程之前，一个CUDA编程模型和所使用的一些术语的概述将使CUDA新手从中受益。

CUDA编程模型是同时使用CPU和GPU的异构模型。在CUDA中，主机（host） 指的是CPU及其内存，而设备（device） 指的是GPU及其内存。在主机上运行的代码可以同时管理主机和设备上的内存，还可以启动内核（kernel）——在设备上执行的函数。这些内核由许多GPU线程并行执行。

鉴于CUDA编程模型的异构性，一个CUDA C程序的典型操作顺序为：

声明并分配主机和设备内存
初始化主机数据
将数据从主机传输到设备
执行一个或多个内核
将结果从设备传输到主机

记住这个操作顺序，让我们看一个CUDA C示例。

第一个CUDA C程序

在最近的一篇文章中，我展示了SAXPY的六种方法，其中包括一个CUDA C版本。SAXPY（Single-Precision A·X Plus Y）代表“单精度A·X加Y”，是并行计算的一个很好的“hello world”示例。在这篇文章中，我将剖析CUDA C SAXPY的一个更完整的版本，详细解释要做什么以及为什么要做。完整的SAXPY代码是：

#include <stdio.h>

__global__
void saxpy(int n, float a, float *x, float *y)
{
  int i = blockIdx.x*blockDim.x + threadIdx.x;
  if (i < n) y[i] = a*x[i] + y[i];
}

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

  cudaMalloc(&d_x, N*sizeof(float)); 
  cudaMalloc(&d_y, N*sizeof(float));

  for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
  }

  cudaMemcpy(d_x, x, N*sizeof(float), cudaMemcpyHostToDevice);
  cudaMemcpy(d_y, y, N*sizeof(float), cudaMemcpyHostToDevice);

  // Perform SAXPY on 1M elements
  saxpy<<<(N+255)/256, 256>>>(N, 2.0f, d_x, d_y);

  cudaMemcpy(y, d_y, N*sizeof(float), cudaMemcpyDeviceToHost);

  float maxError = 0.0f;
  for (int i = 0; i < N; i++)
    maxError = max(maxError, abs(y[i]-4.0f));
  printf("Max error: %f\n", maxError);

  cudaFree(d_x);
  cudaFree(d_y);
  free(x);
  free(y);
}

函数saxpy是在GPU上并行运行的内核，main函数是主机代码。让我们从主机代码开始讨论这个程序。

主机代码

main函数声明了两对数组：

  float *x, *y, *d_x, *d_y;
  x = (float*)malloc(N*sizeof(float));
  y = (float*)malloc(N*sizeof(float));

  cudaMalloc(&d_x, N*sizeof(float)); 
  cudaMalloc(&d_y, N*sizeof(float));

指针x和y指向使用malloc分配的主机数组，而d_x和d_y数组指针指向使用CUDA runtime API中的cudaMalloc函数分配的设备数组。CUDA中的主机和设备具有独立的内存空间，两者都可以通过主机代码进行管理（CUDA C内核也可以在支持它的设备上分配设备内存）。

接下来主机代码初始化主机数组。这里，我们将x设置为一个值全部为1的数组，将y设置为值全部为2的数组。

  for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
  }

为了初始化设备数组，我们只需使用cudaMemcpy将数据从x和y复制到相应的设备数组d_x和d_y，cudaMemcpy的工作原理与标准的C memcpy函数类似，只是它需要第四个参数来指明复制的方向。在此，我们使用cudaMemcpyHostToDevice来指明第一个（目的）参数是设备指针，第二个（源）参数是主机指针。

  cudaMemcpy(d_x, x, N*sizeof(float), cudaMemcpyHostToDevice);
  cudaMemcpy(d_y, y, N*sizeof(float), cudaMemcpyHostToDevice);

运行内核后，为了将结果返回到主机，我们使用cudaMemcpy和cudaMemcpyDeviceToHost选项将d_y指向的设备数组复制到y指向的主机数组。

cudaMemcpy(y, d_y, N*sizeof(float), cudaMemcpyDeviceToHost);

启动内核

saxpy内核通过以下代码启动：

saxpy<<<(N+255)/256, 256>>>(N, 2.0, d_x, d_y);

三个V形符号之间的信息是执行配置（execution configuration），它规定了有多少设备线程并行地执行内核。在CUDA中，软件层面存在一个线程层次结构，它模拟了GPU上线程处理器（processor）的分组方式。在CUDA编程模型中，我们谈到用一个由thread blocks构成的grid启动内核。执行配置中的第一个参数指定grid中thread block的数量，第二个参数指定一个thread block中的thread数量。

通过传递dim3（一个由CUDA定义的具有x、y和z成员的简单结构体）值，thread block和grid可以定义为一维、二维或三维的，但对于这个简单的示例，我们只需要一个维度，因此我们传递整数。此处我们使用包含256个线程的thread block启动内核，并使用整数运算来确定处理数组的所有N个元素所需的thread blocks的数量（(N+255)/256）。

对于数组中元素的数量不能被thread block大小整除的情况，内核代码必须检查内存访问是否越界。

整理

计算完成后，我们应该释放所有分配的内存。对于使用cudaMalloc()分配的设备内存，只需调用cudaFree()即可。对于主机内存，请像往常一样使用free()。

  cudaFree(d_x);
  cudaFree(d_y);
  free(x);
  free(y);

设备代码

现在我们转到内核代码。

__global__
void saxpy(int n, float a, float *x, float *y)
{
  int i = blockIdx.x*blockDim.x + threadIdx.x;
  if (i < n) y[i] = a*x[i] + y[i];
}

在CUDA中，我们使用__global__声明说明符（declaration specifier）定义诸如saxpy之类的内核。设备代码中定义的变量不需要指定为设备变量，因为程序已经假定它们位于设备上。在这里，n、a和i变量将由每个线程存储在寄存器（register）中，指针x和y必须是指向设备内存地址空间的指针。显然，我们从主机代码启动内核时，将d_x和d_y传递给了内核。然而，前两个参数n和a并没有在主机代码中显式地传输到设备。因为函数参数在C/C++中默认是按值传递的，所以CUDA runtime可以自动处理这些值到设备的传输。CUDA runtime API的这一特性使得在GPU上启动内核变得非常自然和容易——这几乎与调用C函数相同。

在我们的saxpy内核中只有两行。如前所述，内核由多个线程并行执行。如果我们希望每个线程处理数组中的一个元素，那么我们需要一种区分和识别每个线程的方法。CUDA为此定义了变量blockDim、blockIdx和threadIdx。这些预定义变量的类型为dim3，类似于主机代码中的执行配置参数。预定义变量blockDim为每个thread block的维度（内核启动时的第二个执行配置参数）。预定义变量threadIdx和blockIdx分别为线程在thread block内的索引和thread block在grid内的索引。表达式

    int i = blockDim.x * blockIdx.x + threadIdx.x

生成用于访问数组元素的全局索引。此外还有gridDim（grid的维度（内核启动时的第一个执行配置参数）），我们在这个例子中没有使用它。

在使用此索引访问数组元素之前，将根据元素数量n检查该值，以确保没有越界的内存访问。如果数组中的元素数量不能被thread block大小整除，因此内核启动的线程数量大于数组大小，则需要进行此检查。内核的第二行执行SAXPY的逐元素操作，除了边界检查之外，它与在主机实现的SAXPY的内层循环相同。

if (i < n) y[i] = a*x[i] + y[i];

编译和运行代码

CUDA C编译器nvcc是NVIDIA CUDA Toolkit的一部分。为了编译我们的SAXPY示例，我们将代码保存在一个扩展名为.cu的文件中，比如SAXPY.cu。然后我们可以使用nvcc进行编译。

nvcc -o saxpy saxpy.cu

然后我们可以运行代码：

% ./saxpy
Max error: 0.000000

总结与结论

通过SAXPY这个简单CUDA C实现的演练，您现在了解了CUDA C编程基础。将C代码“移植”到CUDA C只需要对C进行几个扩展：设备内核函数使用的__global__声明说明符；启动内核时使用的执行配置；以及用于识别和区分并行执行内核的GPU线程的内置设备变量blockDim、blockIdx和threadIdx。

异构CUDA编程模型的一个优点是，可以增量地（一次一个内核）将现有代码从C移植到CUDA C。

在本系列的下一篇文章中，我们将了解一些性能测量方法和指标。