魏晨多大在前面的两文章中，我们研究了如何在主机和设备之间高效地移动数据。在我们的 CUDA C / C ++系列的第六篇文章中，我们将讨论如何有效地从内核中访问设备存储器，特别是

　　在 CUDA 设备上有几种内存，每种内存的作用域、生存期和缓存行为都不同。到目前为止，在本系列中，我们已经使用了驻留在设备DRAM中的全局内存，用于主机和设备之间的传输，以及内核的数据输入和输出。这里的名称global是指作用域，因为它可以从主机和设备访问和修改。全局内存可以像下面代码片段的第一行那样使用__device__de Clara 说明符在全局（变量）范围内声明，或者使用cudaMalloc()动态分配并分配给一个常规的 C 指针变量，如第 7 行所示。全局内存分配可以在应用程序的生命周期内保持。根据设备的计算能力，全局内存可能被缓存在芯片上，也可能不在芯片上缓存。

　　}在讨论全局内存访问性能之前，我们需要改进对 CUDA 执行模型的理解。我们已经讨论了如何将线程被分组为线程块分配给设备上的多处理器。在执行过程中，有一个更精细的线程分组到warps。GPU上的多处理器以 SIMD （单指令多数据）方式为每个扭曲执行指令。所有当前支持 CUDA – 的 GPUs 的翘曲尺寸（实际上是 SIMD 宽度）是 32 个线程。

　　将线程分组为扭曲不仅与计算有关，而且与全局内存访问有关。设备coalesces全局内存加载并存储由一个 warp 线程发出的尽可能少的事务，以最小化 DRAM 带宽（在计算能力小于 2 . 0 的旧硬件上，事务合并在 16 个线程的一半扭曲内，而不是整个扭曲中）。为了弄清楚 CUDA 设备架构中发生聚结的条件，我们在三个Tesla 卡上进行了一些简单的实验： a Tesla C870 （计算能力 1 . 0 ）、 Tesla C1060 （计算能力 1 . 3 ）和 Tesla C2050 （计算能力 2 . 0 ）。

　　我们运行两个实验，使用如下代码（GitHub 上也有）中所示的增量内核的变体，一个具有数组偏移量，这可能导致对输入数组的未对齐访问，另一个是对输入数组的跨步访问。

　　}此代码可以通过传递“ fp64 ”命令行选项以单精度（默认值）或双精度运行偏移量内核和跨步内核。每个内核接受两个参数，一个输入数组和一个表示访问数组元素的偏移量或步长的整数。内核在一系列偏移和跨距的循环中被称为。

　　设备内存中分配的数组由 CUDA 驱动程序与 256 字节内存段对齐。该设备可以通过 32 字节、 64 字节或 128 字节的事务来访问全局内存。对于 C870 或计算能力为 1 . 0 的任何其他设备，半线程的任何未对齐访问（或半扭曲线程不按顺序访问内存的对齐访问）将导致 16 个独立的 32 字节事务。由于每个 32 字节事务只请求 4 个字节，因此可以预期有效带宽将减少 8 倍，这与上图（棕色线）中看到的偏移量（不是 16 个元素的倍数）大致相同，对应于线程的一半扭曲。

　　对于计算能力为 1 . 2 或 1 . 3 的 Tesla C1060 或其他设备，未对准访问的问题较少。基本上，通过半个线程对连续数据的未对齐访问在几个“覆盖”请求的数据的事务中提供服务。由于未请求的数据正在传输，以及不同的半翘曲所请求的数据有些重叠，因此相对于对齐的情况仍然存在性能损失，但是这种损失远远小于 C870 。

　　计算能力为 2 . 0 的设备，如 Tesla C250 ，在每个多处理器中都有一个 L1 缓存，其行大小为 128 字节。该设备将线程的访问合并到尽可能少的缓存线中，从而导致对齐对跨线程顺序内存访问吞吐量的影响可以忽略不计。

　　对于快速的全局内存访问，我们有不同的看法。对于大步进，无论架构版本如何，有效带宽都很差。这并不奇怪：当并发线程同时访问物理内存中相距很远的内存地址时，硬件就没有机会合并这些访问。从上图中可以看出，在 Tesla C870 上，除 1 以外的任何步幅都会导致有效带宽大幅降低。这是因为 compute capability 1 . 0 和 1 . 1 硬件需要跨线程进行线性、对齐的访问以进行合并，因此我们在 offset 内核中看到了熟悉的 1 / 8 带宽。 Compute capability 1 . 2 及更高版本的硬件可以将访问合并为对齐的段（ CC 1 . 2 / 1 . 3 上为 32 、 64 或 128 字节段，在 CC 2 . 0 及更高版本上为 128 字节缓存线），因此该硬件可以产生平滑的带宽曲线。

　　当访问多维数组时，线程通常需要索引数组的更高维，因此快速访问是不可避免的。我们可以使用一种名为共享内存的 CUDA 内存来处理这些情况。共享内存是一个线程块中所有线程共享的片上内存。共享内存的一个用途是将多维数组的 2D 块以合并的方式从全局内存提取到共享内存中，然后让连续的线程跨过共享内存块。与全局内存不同，对共享内存的快速访问没有惩罚。我们将在下一篇文章中详细介绍共享内存。

　　在这篇文章中，我们讨论了如何从 CUDA 内核代码中有效地访问全局内存的一些方面。设备上的全局内存访问与主机上的数据访问具有相同的性能特征，即数据局部性非常重要。在早期的 CUDA 硬件中，内存访问对齐和跨线程的局部性一样重要，但在最近的硬件上，对齐并不是什么大问题。另一方面，快速的内存访问会损害性能，使用片上共享内存可以减轻这种影响。在下一篇文章中，我们将详细探讨共享内存，之后的文章中，我们将展示如何使用共享内存来避免在矩阵转置过程中出现跨步全局内存访问。

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

　　UID_BASE (0x1FF1E800) 寄存器，但我遇到了 HardFault。我查看了参考手册，第 137 页说

　　物联网 （IoT） 描述了嵌入传感器、软件和其他技术的物联网，目的是通过互联网与其他

　　的服务。此过程应该是无缝的，并且应尽可能少地完成，从而减少用户为实际使用设置

　　变量、局部变量是变量自身的身份。身份的不同是靠出生地决定的。而能否被全能局

　　为128KiB，可以提供多种轻量级网络协议，轻量级的图形框架，以及丰富的IOT总线读写部件等。可支撑的产品如智能家居领域的连接类模组、传感器

　　的服务。此过程应该是无缝的，并且应该花费尽可能少的时间来完成，从而减少用户为实际使用设置

　　鉴于KAIST的HPC根源，将DirectCXL原型放在一起的研究人员专注于使用远程直接

　　(通常是芯片内部的各个I2C，SPI, USB等控制器的寄存器或者外部

　　流程 /

　　碳化硅二极管多为肖特基二极管。第一个商用 SiC 肖特基二极管是在 10 多年前推出的。从那时起，这些

　　测量SiC MOSFET /

　　虚拟页的时候，产生缺页异常。如果是文件映射，那么分配物理页，把文件指定区间的数据读到物理页

　　实现的每个系统组件的性能，从而在不增加功耗的情况下实现更快的设计。所有电子

　　对计算机系统来说是一项非常重要的资源，直接影响着系统运行的性能。最初的时候，系统是直接运行在物理

　　管理详解 /

　　对象初始化链表介绍链表是一种数据结构，跟其他的结构体类似，初始化后本身占用

　　引言在使用stm32或者其他单片机的时候，会经常使用到串口通讯，那么如何

　　接收数据呢？假如这段数据是不定长的有如何高效接收呢？同学A：数据来了就会进入串口中断，在中断

　　在物联网(IoT)日渐普及下，边缘运算(Edge Computing)越来越重要，这种分布式计算的架构将数据处理由网络中心移往边缘节点，能更快更

　　泄漏？ /

　　空间用作特殊用途（给安全模块使用，给其它处理器使用，或是给特定的驱动程序使用等），在 Device Tree 中有提供两种方法对预留

　　空间？—必须发送系统调用；4、如何发生系统调用？—调用操作系统提供的函数接口AP...

　　映射 /

　　获取用来提高效率，当然，这种改变不能影响指令顺序的行为处理有多个总线接口在

　　创建的，而后手工做初始化。以下代码是一个关于静态线程和动态线程的例子：/* 线 的对象和

　　使用可编程时钟器件？设计SERDES参考时钟源面临的挑战有哪些？如何去实现XAUI参考时钟源？

　　作用于定义它的文件里，不能作用到其他文件里，即被static关键字修饰过的变量具有文件作用域。这样即使两个不同的源文件都定义了相同的静态

　　创建的，而后手工做初始化。 以下代码是一个关于静态线程和动态线程的例子： /* 线 的对象和运行时用到

　　革命 /

　　使用AD628精密增益模块 /

　　华天电力专业生产局部放电测试仪（又称局部放电成套装置），接下来为大家分享如何更

　　使用局部放电监控局部放电的大多数在线监视或定期监视都是使用未经校准的测量方法

　　Block Started by Symbol 的简称，BSS 段属于静态

　　分析 /

　　管理部分，本章源码超级多，很烧脑，但笔者关键处都加了注释。废话不多说，开始吧。初始化整个

　　的死死的.按不住奴才那这主子就不合格,就不是一个稳定的运作系统. 试着想想 实际

　　管理部分，本章源码超级多，很烧脑，但笔者关键处都加了注释。废话不多说，开始吧。初始化整个

　　成本高、无法监督巡检人员工作、巡检数据信息化程度低等缺陷。 一.系统背景: 如何

　　的安全隐患？ 苏州新导RFID电力资产巡检管理系统的推出，彻底改变了电

　　将巡检数据信息化 /

　　)：在虚拟地址空间中，段内可以使用的最大偏移量2) 分段实现· 逻辑地址的段寄存器

　　创新的角度来看，通过物联网技术传输的庞大数据量是令人难以置信的，而正是这些数据使世界构建成为人们现在所采用的互联生态系统，而所有这些制约因素对物联网

　　时虽然可以进行临界区保护（原子操作），但是貌似还是会出问题！比如A任务在

　　，然后运行计算。因为后者会占用大量矩阵的大量空间。 RTL设计如何直接

　　嗨，大家好，我发现了一些线索，但我看到的任何一个都没有指引我。有可能直接

　　SRAM吗？如果不是，有可能用DMA来做吗？如果有人能提供一些指示，所以我可以开始调查它吗？我感兴趣

　　，则在时钟上升沿传输数据;主从双方可以分别通过使IRDY#或TRDY#无效，在数据期中插入等待周期

　　dev_add说起 /

　　资源的方式有两种：动态映射(ioremap)和静态映射(map_desc)。

　　驱动中使用它. 最后, __GFP_NORETRY 告知分配器立即放弃如果得不到请求的

　　区段__GFP_DMA和__GFP_HIGHMEM的使用与平台相关，Linux把

　　static这个说明符在不同的地方所起的作用是不同的。应予以注意。Tips:A.若

　　关闭或驱动被卸载时释放IO端口范围。 in、out、ins和outs汇编语言指令都可以

　　溢出的解决办法：1、等比例缩小图片。2、对图片采用软引用，及时进行 recycle( ) 操作。3、使用加载图片框架处理图片，如专业处理图片的 ImageLoader 图片加载框架，还有XUtils 的 BitMapUtils 来处理。

　　的管理 /

　　详解 /

　　址和偏移地址的区别 /

　　C style=cursor:pointer;color：#D05C38;text-decoration:underline;》C

　　PLL器件的一个组成部分。要确定偏移的值，我们必须考虑到影响信号的布线延迟和任何外部

　　关闭控制回路 /

　　数组的确比较方便，但是在加上RTOS后就是另一种情况了。相比消息队列，使用

　　变量相比事件标志组主要有如下三个问题：1、使用事件标志组可以让RTOS

　　代码– 刚开始学习的时候最重要的就是看代码，但是学习Linux系统，不能一头扎进

　　上的寄存器来进行的，通常包括控制寄存器、状态寄存器和数据寄存器三大类，外设的寄存器通常被连续地编址。根据CPU体系结构

　　（外设寄存器）资源了）。 为了配置寄存器，我们需要知道寄存器在操作系统

　　的正确使用将防止90%的ESD有关的问题，但是为了达到最后10%，需要另一种保护：离子化。中和那些可产生静电电荷的装配

　　未来5年，全球智能电表应用将明显持续增长。本文中将探讨如何才能使智能电网更

　　运行? /

　　。我们可以把它看成是最大容量为4K的普通文件。我们要能够像打开普通文件

　　使用AD628精密增益模块 /

　　保养维护好硬盘避免硬盘的损坏，我们需要提醒您注意以下事项，尽量避免硬盘出现问题。1、为您的机器提供不间断电源（UPS）当硬盘开始工作时，一般都处于高速旋转之中，如果硬盘

　　把握采购机遇背景 ： 在全球性的金融危机蔓延的当下，您是否在为打开国际、国内配套之门而呕心沥血？配套采购的未来发展趋势如何？机会何在？是否在为开拓国际售后市场而绞尽脑汁

　　NVIDIA Jetson Xavier NX世上最小嵌入式AI超级计算机

　　骁龙XR2助力奇遇MIX率先实现双目全彩透视MR！华为专利收入持续攀升！热点科技新闻点评

　　西安理工大学-物联网培训第一讲-ESP32的VSCode开发环境搭建7-6#物联网