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

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

3天内不再提示

如何在CUDA C/C++中实现主机和设备同步执行

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

在 本系列文章的第一篇 中,我们通过检查 CUDA C/C++ SAXPY 来研究 CUDA C / C ++的基本元素。在第二篇文章中,我们将讨论如何分析这个和其他 CUDA C / C ++代码的性能。我们将依赖于这些性能测量技术在未来的职位,性能优化将变得越来越重要。

CUDA 性能度量通常是从主机代码中完成的,可以使用 CPU 计时器或 CUDA 特定计时器来实现。在讨论这些性能度量技术之前,我们需要讨论如何在主机和设备之间同步执行。

主机设备同步

让我们看看数据传输和来自上一篇文章的 SAXPY 主机代码的内核启动:

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

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



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



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

使用cudaMemcpy()在主机和设备之间的数据传输是synchronous(或blocking)传输。同步数据传输在之前发出的所有 CUDA 调用完成之前不会开始,后续的 CUDA 调用在同步传输完成之前无法开始。因此,第三行的saxpy内核启动在第二行从yd_y的传输完成后才会发出。另一方面,内核启动是异步的。一旦内核在第三行启动,控制权立即返回到 CPU ,而不是等待内核完成。而 MIG ht 似乎为设备在最后一行主机数据传输设置了一个竞争条件,数据传输的阻塞性质确保了内核在传输开始之前完成。

用 CPU 计时器计时内核执行

现在让我们来看看如何使用 CPU 计时器为内核执行计时。

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

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



t1 = myCPUTimer();

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

cudaDeviceSynchronize();

t2 = myCPUTimer();



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

除了对通用主机时间戳函数myCPUTimer()的两次调用外,我们还使用显式同步屏障cudaDeviceSynchronize()来阻止 CPU 的执行,直到设备上以前发出的所有命令都已完成。如果没有这个屏障,这段代码将测量内核发射时间,而不是内核执行时间

使用 CUDA 事件计时

使用主机设备同步点(如cudaDeviceSynchronize()的一个问题是它们会暂停 GPU 管道。因此, CUDA 通过CUDA 事件 API为 CPU 定时器提供了一个相对轻量级的替代方案。 CUDA 事件 API 包括在两个记录的事件之间调用create破坏事件、record事件和以毫秒为单位计算已用时间

CUDA 事件利用 CUDA streams. CUDA 流只是按顺序在设备上执行的操作序列。在某些情况下[vx3 . 4 可以交叉使用 vx3 . 4]的流。到目前为止, GPU 上的所有操作都发生在默认流或流 0 (也称为“空流”)中。

在下面的清单中,我们将 CUDA 事件应用于 SAXPY 代码。

cudaEvent_t start, stop;

cudaEventCreate(&start);

cudaEventCreate(&stop);



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

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



cudaEventRecord(start);

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

cudaEventRecord(stop);



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



cudaEventSynchronize(stop);

float milliseconds = 0;

cudaEventElapsedTime(&milliseconds, start, stop);

CUDA 事件属于cudaEvent_t类型,使用cudaEventCreate()cudaEventDestroy()创建和销毁事件。在上面的代码中cudaEventRecord()将启动和停止事件放入默认流 stream 0 。当事件到达流中的事件时,设备将记录事件的时间戳。函数cudaEventSynchronize()会阻止 CPU 的执行,直到记录指定的事件为止。cudaEventElapsedTime()函数在第一个参数中返回录制startstop之间经过的毫秒数。该值的分辨率约为半微秒。

内存带宽

现在我们有了一种精确计时内核执行的方法,我们将使用它来计算带宽。在评估带宽效率时,我们同时使用理论峰值带宽和观察到的或有效的内存带宽。

理论带宽

理论带宽可以使用产品文献中提供的硬件规格计算。例如, NVIDIA Tesla M2050 GPU 使用内存时钟速率为 1546 MHz 的 DDR (双数据速率) RAM 和 384 位宽的内存接口。使用这些数据项, NVIDIA Tesla M2050 的峰值理论内存带宽为 148 GB / s ,如下所示。

BWTheoretical= 1546 * 106* (384 / 8) * 2 / 109= 148 GB / s

在这个计算中,我们将内存时钟速率转换为赫兹,乘以接口宽度(除以 8 ,将位转换为字节),再乘以 2 ,这是由于数据速率加倍。最后,我们除以 109将结果转换为 GB / s 。

有效带宽

我们通过计时特定的程序活动和了解程序如何访问数据来计算有效带宽。我们用下面的等式。

BWEffective=(RB+WB( VZX50]* 109)

这里,BWEffective有效带宽,单位为 GB / s ,RB是每个内核读取的字节数,WB是每个内核写入的字节数,t是以秒为单位的运行时间。下面是完整的代码。

#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 = 20 * (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;

  }



  cudaEvent_t start, stop;

  cudaEventCreate(&start);

  cudaEventCreate(&stop);



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

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



  cudaEventRecord(start);



  // Perform SAXPY on 1M elements

  saxpy<<<(N+511)/512, 512>>>(N, 2.0f, d_x, d_y);



  cudaEventRecord(stop);



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



  cudaEventSynchronize(stop);

  float milliseconds = 0;

  cudaEventElapsedTime(&milliseconds, start, stop);



  float maxError = 0.0f;

  for (int i = 0; i < N; i++) {

    maxError = max(maxError, abs(y[i]-4.0f));

  }



  printf("Max error: %fn", maxError);

  printf("Effective Bandwidth (GB/s): %fn", N*4*3/milliseconds/1e6);

}

在带宽计算中,N*4是每个数组读或写传输的字节数, 3 的因子表示x的读取和y的读写。经过的时间存储在变量milliseconds中,以明确单位。请注意,除了添加带宽计算所需的功能外,我们还更改了数组大小和线程块大小。在 Tesla M2050 上编译并运行此代码:

$ ./saxpy

Max error: 0.000000

Effective Bandwidth (GB/s): 110.374872

测量计算吞吐量

我们刚刚演示了如何测量带宽,带宽是数据吞吐量的度量。另一个对性能非常重要的指标是计算吞吐量。计算吞吐量的常用度量是 GFLOP / s ,它代表“每秒千兆浮点运算”,其中 Giga 是 10 的前缀9. 我们通常测量 SAXPY 的吞吐量,因为每一个 SAXPY 运算都是有效的

GFLOP/s Effective== 2 N /( t :《* 109)

N 是 SAXPY 操作中的元素数, t 是以秒为单位的运行时间。与理论峰值带宽一样,理论峰值 GFLOP / s 可以从产品文献中获得(但是计算它可能有点棘手,因为它与体系结构非常相关)。例如, Tesla M2050 GPU 的单精度浮点吞吐量理论峰值为 1030 GFLOP / s ,双倍精度的理论峰值吞吐量为 515 GFLOP / s 。

SAXPY 为计算的每个元素读取 12 个字节,但是只执行一个乘法加法指令( 2 个浮点运算),因此很明显它是带宽受限的,因此在这种情况下(实际上在许多情况下),带宽是衡量和优化的最重要的指标。在更复杂的计算中,在 FLOPs 级别测量性能可能非常困难。因此,更常见的是使用分析工具来了解计算吞吐量是否是一个瓶颈。应用程序通常提供特定于问题(而不是特定于体系结构)的吞吐量指标,因此对用户更有用。例如,天文 n 体问题的“每秒十亿次相互作用”,或分子动力学模拟的“每天纳秒”。

总结

这篇文章描述了如何使用 CUDA 事件 API 为内核执行计时。 CUDA 事件使用 GPU 计时器,因此避免了与主机设备同步相关的问题。我们提出了有效带宽和计算吞吐量性能指标,并在 SAXPY 内核中实现了有效带宽。很大一部分内核是内存带宽限制的,因此计算有效带宽是性能优化的第一步。在以后的文章中,我们将讨论如何确定带宽、指令或延迟是性能的限制因素。

CUDA 事件还可以用于确定主机和设备之间的数据传输速率,方法是在 cudaMemcpy() 调用的任一侧记录事件。

如果你在这个设备上运行一个关于内存不足的错误[ZC9],你可能会得到一个更小的错误。实际上,到目前为止,我们的示例代码还没有费心检查运行时错误。在[VZX337]中,我们将学习如何在 CUDA C / C ++中执行错误处理以及如何查询当前设备以确定它们可用的资源,以便我们可以编写更健壮的代码。

关于作者

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

审核编辑:郭婷

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

    关注

    28

    文章

    4673

    浏览量

    128558
  • API
    API
    +关注

    关注

    2

    文章

    1471

    浏览量

    61745
  • 计时器
    +关注

    关注

    1

    文章

    416

    浏览量

    32602
收藏 人收藏

    评论

    相关推荐

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

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

    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>编译器

    ostream在c++的用法

    ostream 是 C++ 标准库中一个非常重要的类,它位于 头文件(实际上,更常见的是通过包含 头文件来间接包含 ,因为 包含了 和 )。 ostream 类及其派生类(如 std::cout
    的头像 发表于 09-20 15:11 413次阅读

    C++实现类似instanceof的方法

    函数,可实际上C++没有。但是别着急,其实C++中有两种简单的方法可以实现类似Java的instanceof的功能。 在
    的头像 发表于 07-18 10:16 496次阅读
    <b class='flag-5'>C++</b><b class='flag-5'>中</b><b class='flag-5'>实现</b>类似instanceof的方法

    FX2 CY7C68013A如何在C++环境中使用LoadEEPROM函数?

    我使用的是 FX2 CY7C68013A 芯片。 我知道 CyUSB.NET 库中有我需要的 LoadEEPROM 函数。 请问如何在 C++ 环境而不是 C#/CLR 环境中使用该函
    发表于 05-31 06:59

    何在FX3 SuperSpeed explorer等电路板上使用openOCD调试C++项目?

    配置与文档的完全相同。 因此,我想请教如何在 FX3 SuperSpeed explorer 等电路板上使用 openOCD 调试我的 C++ 项目? 回到纯 C 项目并不是一个真正
    发表于 05-23 08:16

    C/C++两种宏实现方式

    #ifndef的方式受C/C++语言标准支持。它不仅可以保证同一个文件不会被包含多次,也能保证内容完全相同的两个文件(或者代码片段)不会被不小心同时包含。
    的头像 发表于 04-19 11:50 526次阅读

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

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

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

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

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

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

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

    的 MISRA C++:2023 博客系列的第二部分。 在这篇博客,我们将深入探讨 C++ 的历史、编程语言多年来的发展历程以及它的下一步发展方向。
    的头像 发表于 01-11 09:00 512次阅读
    <b class='flag-5'>C++</b>简史:<b class='flag-5'>C++</b>是如何开始的

    C语言和C++那些不同的地方

    ++11标准。根据不同的标准,它们的功能也会有所不同,但是越新的版本支持的编译器越少,所以本文在讨论的时候使用的C语言标准是C89,C++标准是C++99.我们来介绍
    的头像 发表于 12-07 14:29 889次阅读
    <b class='flag-5'>C</b>语言和<b class='flag-5'>C++</b><b class='flag-5'>中</b>那些不同的地方

    c++怎么开始编程

    应用程序、嵌入式系统和网络应用程序等各种领域。 在开始编程之前,你需要安装C++的编程环境。首先,你需要下载并安装一个编译器,比如微软的Visual Studio、GNU的GCC或者Clang。这些编译器可以将你的C++代码编译成可执行
    的头像 发表于 11-27 15:56 873次阅读

    c++多行注释快捷键

    C++,多行注释(也称为块注释)是一种用于注释大段代码或多个语句的方法。当你希望暂时禁用一些代码或者解释特定部分代码的作用时,多行注释是非常有用的。 在C++,多行注释以 /*
    的头像 发表于 11-22 10:24 7914次阅读

    何在非Autosar应用执行I2C

    我是NXP产品的新产品,我试图了解它的生态系统。我需要在非Autosar应用执行I2C,所以我需要了解它是如何工作的,所以我有一些问题: 1) 在 RTD_I2C_UM.pdf
    发表于 11-13 07:04