0
  • 聊天消息
  • 系统消息
  • 评论与回复
登录后你可以
  • 下载海量资料
  • 学习在线课程
  • 观看技术视频
  • 写文章/发帖/加入社区
会员中心
创作中心

完善资料让更多小伙伴认识你,还能领取20积分哦,立即完善>

3天内不再提示

CUDA并行计算平台的C/C++接口的简单介绍

星星科技指导员 来源:NVIDIA 作者:Mark Harris 2022-04-11 10:13 次阅读

本文是 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 及其存储器。在主机上运行的代码可以管理主机和设备上的内存,还可以启动在设备上执行的函数 kernels 。这些内核由许多 GPU 线程并行执行。

鉴于 CUDA 编程模型的异构性, CUDA C 程序的典型操作序列是:

声明并分配主机和设备内存。

初始化主机数据。

将数据从主机传输到设备。

执行一个或多个内核。

将结果从设备传输到主机。

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

第一个 CUDA C 程序

在最近的一篇文章中,我演示了 萨克斯比的六种方法 ,其中包括一个 CUDA C 版本。 SAXPY 代表“单精度 A * X + Y ”,是并行计算的一个很好的“ hello world ”示例。在这篇文章中,我将剖析 CUDA C SAXPY 的一个更完整的版本,详细解释它的作用和原因。完整的 SAXPY 代码是:

#include 

__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
", 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运行时API使用cudaMalloc函数分配的设备数组。CUDA中的主机和设备有独立的内存空间,这两个空间都可以从主机代码进行管理(CUDAC内核也可以在支持它的设备上分配设备内存)。

然后,主机代码初始化主机数组。在这里,我们设置了一个 1 数组,以及一个 2 数组。

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

为了初始化设备数组,我们只需使用cudaMemcpy将数据从xy复制到相应的设备数组d_xd_y,它的工作方式与标准的 Cmemcpy函数一样,只是它采用了第四个参数,指定了复制的方向。在本例中,我们使用cudaMemcpyHostToDevice指定第一个(目标)参数是设备指针,第二个(源)参数是主机指针。

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

在运行内核之后,为了将结果返回到主机,我们使用cudaMemcpycudaMemcpyDeviceToHost,从d_y指向的设备数组复制到y指向的主机数组。

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

启动内核

cord [EZX13 内核由以下语句启动:

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

三个 V 形符号之间的信息是 执行配置 ,它指示有多少设备线程并行执行内核。在 CUDA 中,软件中有一个线程层次结构,它模仿线程处理器在 GPU 上的分组方式。在 CUDA 编程模型中,我们谈到启动一个 grid 为 螺纹块 的内核。执行配置中的第一个参数指定网格中线程块的数量,第二个参数指定线程块中的线程数。

线程块和网格可以通过为这些参数传递 dim3 (一个由 CUDA 用 x 、 y 和 z 成员定义的简单结构)值来生成一维、二维或三维的线程块和网格,但是对于这个简单的示例,我们只需要一维,所以我们只传递整数。在本例中,我们使用包含 256 个线程的线程块启动内核,并使用整数算术来确定处理数组( (N+255)/256 )的所有 N 元素所需的线程块数。

对于数组中的元素数不能被线程块大小平均整除的情况,内核代码必须检查内存访问是否越界。

清理

完成后,我们应该释放所有分配的内存。对于使用 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__ de __global__ 说明符定义诸如 Clara 这样的内核。设备代码中定义的变量不需要指定为设备变量,因为假定它们驻留在设备上。在这种情况下, n 、 a 和 i 变量将由每个线程存储在寄存器中,指针 x 和 y 必须是指向设备内存地址空间的指针。这确实是真的,因为当我们从宿主代码启动内核时,我们将 d_x 和 d_y 传递给了内核。但是,前两个参数 n 和 a 没有在主机代码中显式传输到设备。因为函数参数在 C / C ++中是默认通过值传递的,所以 CUDA 运行时可以自动处理这些值到设备的传输。 CUDA 运行时 API 的这一特性使得在 GPU 上启动内核变得非常自然和简单——这几乎与调用 C 函数一样。

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

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

生成用于访问数组元素的全局索引。我们在这个例子中没有使用它,但是还有一个 gridDim ,它包含在启动的第一个执行配置参数中指定的网格维度。

在使用该索引访问数组元素之前,将根据元素的数量 n 检查其值,以确保没有越界内存访问。如果一个数组中的元素数不能被线程块大小平均整除,并且结果内核启动的线程数大于数组大小,则需要进行此检查。内核的第二行执行 SAXPY 的元素级工作,除了边界检查之外,它与 SAXPY 主机实现的内部循环相同。

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

编译和运行代码

CUDA C 编译器 nvcc 是 NVIDIA CUDA 工具箱 的一部分。为了编译我们的 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__ de Clara 说明符;启动内核时使用的执行配置;内置的设备变量 blockDim 、 blockIdx 和 threadIdx 用来识别和区分并行执行内核的 GPU 线程。

异类 CUDA 编程模型的一个优点是,将现有代码从 C 移植到 CUDA C 可以逐步完成,一次只能移植一个内核。

在本系列的下一篇文章中,我们将研究一些性能度量和度量。

关于作者

Mark Harris 是 NVIDIA 杰出的工程师,致力于 RAPIDS 。 Mark 拥有超过 20 年的 GPUs 软件开发经验,从图形和游戏到基于物理的模拟,到并行算法和高性能计算。当他还是北卡罗来纳大学的博士生时,他意识到了一种新生的趋势,并为此创造了一个名字: GPGPU (图形处理单元上的通用计算)。

审核编辑:郭婷

声明:本文内容及配图由入驻作者撰写或者入驻合作网站授权转载。文章观点仅代表作者本人,不代表电子发烧友网立场。文章及其配图仅供工程师学习之用,如有内容侵权或者其他违规问题,请联系本站处理。 举报投诉
  • 存储器
    +关注

    关注

    38

    文章

    7484

    浏览量

    163769
  • cpu
    cpu
    +关注

    关注

    68

    文章

    10855

    浏览量

    211602
  • gpu
    gpu
    +关注

    关注

    28

    文章

    4729

    浏览量

    128897
收藏 人收藏

    评论

    相关推荐

    同样是函数,在CC++中有什么区别

    同样是函数,在 CC++ 中有什么区别? 第一个返回值。 C语言的函数可以不写返回值类型,编译器会默认为返回 int。 但是 C++ 的函数,除了构造和析构这两个特殊的函数,必须
    的头像 发表于 11-29 10:25 274次阅读

    C7000 C/C++优化指南用户手册

    电子发烧友网站提供《C7000 C/C++优化指南用户手册.pdf》资料免费下载
    发表于 11-09 15:00 0次下载
    <b class='flag-5'>C</b>7000 <b class='flag-5'>C</b>/<b class='flag-5'>C++</b>优化指南用户手册

    TMS320C6000优化C/C++编译器v8.3.x

    电子发烧友网站提供《TMS320C6000优化C/C++编译器v8.3.x.pdf》资料免费下载
    发表于 11-01 09:35 0次下载
    TMS320<b class='flag-5'>C</b>6000优化<b class='flag-5'>C</b>/<b class='flag-5'>C++</b>编译器v8.3.x

    TMS320C28x优化C/C++编译器v22.6.0.LTS

    电子发烧友网站提供《TMS320C28x优化C/C++编译器v22.6.0.LTS.pdf》资料免费下载
    发表于 10-31 10:10 0次下载
    TMS320<b class='flag-5'>C</b>28x优化<b class='flag-5'>C</b>/<b class='flag-5'>C++</b>编译器v22.6.0.LTS

    C语言和C++中结构体的区别

    同样是结构体,看看在C语言和C++中有什么区别?
    的头像 发表于 10-30 15:11 199次阅读

    C7000优化C/C++编译器

    电子发烧友网站提供《C7000优化C/C++编译器.pdf》资料免费下载
    发表于 10-30 09:45 0次下载
    <b class='flag-5'>C</b>7000优化<b class='flag-5'>C</b>/<b class='flag-5'>C++</b>编译器

    GPU加速计算平台是什么

    GPU加速计算平台,简而言之,是利用图形处理器(GPU)的强大并行计算能力来加速科学计算、数据分析、机器学习等复杂计算任务的软硬件结合系统。
    的头像 发表于 10-25 09:23 245次阅读

    C++语言基础知识

    电子发烧友网站提供《C++语言基础知识.pdf》资料免费下载
    发表于 07-19 10:58 7次下载

    C++中实现类似instanceof的方法

    函数,可实际上C++中没有。但是别着急,其实C++中有两种简单的方法可以实现类似Java中的instanceof的功能。 在 C++ 中,确定对象的类型是编程中实际需求,使开发人员
    的头像 发表于 07-18 10:16 574次阅读
    <b class='flag-5'>C++</b>中实现类似instanceof的方法

    C/C++代码动态测试工具VectorCAST插桩功能演示#代码动态测试 #C++

    C++代码
    北汇信息POLELINK
    发布于 :2024年04月18日 11:57:45

    鸿蒙OS开发实例:【Native C++

    使用DevEco Studio创建一个Native C++应用。应用采用Native C++模板,实现使用NAPI调用C标准库的功能。使用C标准库hypot
    的头像 发表于 04-14 11:43 2594次阅读
    鸿蒙OS开发实例:【Native <b class='flag-5'>C++</b>】

    使用 MISRA C++:2023® 避免基于范围的 for 循环中的错误

    在前两篇博客中,我们 向您介绍了新的 MISRA C++ 标准 和 C++ 的历史 。在这篇博客中,我们将仔细研究以 C++ 中 for 循环为中心的特定规则。
    的头像 发表于 03-28 13:53 787次阅读
    使用 MISRA <b class='flag-5'>C++</b>:2023® 避免基于范围的 for 循环中的错误

    c语言,c++,java,python区别

    C语言、C++、Java和Python是四种常见的编程语言,各有优点和特点。 C语言: C语言是一种面向过程的编程语言。它具有底层的特性,能够对计算
    的头像 发表于 02-05 14:11 2370次阅读

    vb语言和c++语言的区别

    Microsoft开发的一种面向对象的事件驱动编程语言。它的设计目标是简化编程过程,让初学者也能快速上手。与之相比,C++语言是一种通用的、面向对象的编程语言,其设计目标是提供高性能的系统级编程。 语法: VB语言的语法较为简单,使用了很多可读性强的关键词,如“
    的头像 发表于 02-01 10:20 2263次阅读

    C++简史:C++是如何开始的

    MISRA C++:2023,MISRA® C++ 标准的下一个版本,来了!为了帮助您做好准备,我们介绍了 Perforce 首席技术支持工程师 Frank van den Beuken 博士撰写
    的头像 发表于 01-11 09:00 582次阅读
    <b class='flag-5'>C++</b>简史:<b class='flag-5'>C++</b>是如何开始的