• [交流分享] 安培架构的MIG及其应用案例分享
    转载自:https://blog.csdn.net/weixin_39626181/article/details/111123999介绍 — 什么是 MIG引言:今年5月14日,NVIDIA 发布了最新的 GPU 架构: 安培,以及基于安培架构的最新的 GPU : A100。安培提供了许多新的特性,MIG 是其中一项非常重要的新特性。MIG的全名是 Multi-Instance GPU,它最多可以把 A100 GPU 切成七个 GPU 实例。每一个 GPU 实例都有各自的流处理器 (SMs) 和内存系统。因此,使用MIG后,每一个 GPU 实例都能够保证各个使用者的工作的时延和吞吐量是可预期的、提供服务品质,并且提高 GPU 的使用率。如图表 1 所示,我们最多可以将 A100 切分成 7 个实例,每一个实例可以提供给不同的使用者运行不同的程序。在使用者的程序对 GPU 的计算量要求比较少的情况下,MIG 最多能够提供高达 7 倍的吞吐量。此外,由于每一个实例都有各自的计算资源和内存资源,因此不必担心因为其中一位使用者的计算量过大或是程式运行错误而干扰到其他的使用者。图片1:MIG 的一个重要的应用,是让提供云服务的公司能够将 GPU 切分成数个实例,并将这些实例提供给不同的承租者使用。由于不同实例之间的计算资源与内存资源都是透过硬件分离,因此能够确保资料的安全性、错误的隔离,以及提供稳定可靠的服务。MIG v.s. MPS v.s. multi-stream在 MIG 推出之前,我们也能够透过 CUDA MPS (Multi-Process Service) 来提高 GPU 使用率。但 MPS 的缺点在于,多个使用者会使用共同的内存,因此使用者的程序会互相影响,除了无法保证推理的速度和吞吐量之外,也有可能因为其中一位使用者的程序出错而导致其他使用者受到干扰。而 MIG 克服了 MPS 面临的问题。MIG 藉由硬件上的分离,保证了使用者的程序不会互相干扰,进而能够让程序的时延和吞吐量能符合预期。在表格 1当中,我们比较了多流 (multi-stream)、MPS 以及 MIG 的优缺点。其中,多流的使用限制较少,同时也很灵活,但对代码更动的需求大,并且无法避免使用者之间的互相干扰,使用者必须小心的使用以避免产生错误。MPS 则不需要更动代码即可使用,可以同时执行的程序也较 MIG 多 (48 与 7),内存的使用与分配也是自动处理的,不需要人工的介入。缺点在于,无法避免多个用户对于 GPU 资源的竞争;最后,MIG 虽然可以同时执行的程序数量最少,但和 MPS 一样不需要使用者另外更动代码,同时在安全性与可靠性上面也是三者中最佳的。这三样技术并不互相冲突,使用者可以根据使用的情境与场景选择与搭配使用。表格1:多流、MPG、MIG 的比较GPU 实例与计算实例要深入地了解 MIG 之前,需要先介绍两个 MIG 引入的重要名词与观念,GPU 实例 (GPU Instance) 和计算实例 (Compute Instance)。在这之前,我们提到的都是 GPU 实例。而在实际的使用上,GPU 实例还可以再切分成多个计算实例。同一个 GPU 实例当中的计算实例是共享内存但拥有独立的流处理器。因此若使用者有多项互相关联的工作,并且使用共享内存较为便利的话,可以选择在一个 GPU 实例当中使用多个计算实例来处理。图表2 GPU 实例与计算实例图表 2 展示了 GPU 实例和计算实例的区别。不同的 GPU 实例之前,他们的内存系统是分开的,以此能够保证内存的服务品质 (QoS, Quality of Service),并且避免受到其他使用者的错误干扰。而在一个 GPU 实例当中,可以有多个计算实例,例如图表 2 中最左侧的 GPU 实例包含了四个计算实例,这四个计算实例各自负责了不同的工作,但他们的内存是共用的。透过这样子的设计,MIG 提供使用者灵活的调度 GPU 的资源。表格2:GPU 实例的设定图表3 GPU 实例的组合表格 2 展示了不同大小的 GPU 实例他们具备的流处理器比例、内存比例、以及可以分配的数量。从图表 3,我们可以更清楚地看到MIG是如何对整个 GPU 去进行切割。在 A100 中,我们共有 8 份内存以及 7 份流处理器可以使用。其中,最小的实例 MIG 1g.5gb,具备一份内存和一份流处理器。而 MIG 2g.10gb 具备两份内存以及两份流处理器。需要特别注意的是,MIG 3g.20gb 和 MIG 4g.20gb虽然具备的流处理器数量不同,但同样拥有四份的内存。最大的实例 MIG 7g.40gb 则具备所有的流处理器以及内存。在使用时,我们可以灵活的分配实例的大小。举例来说,我们可以使用一个 4g.20gb 实例、一个 2g.10gb 实例,以及一个 1g.5gb 实例。图表 4 展示了我们使用的 GPU 实例。这里需要注意的是,虽然我们使用了所有的流处理器,但只有用到 7 份的内存,有一份内存会浪费掉。图表4 MIG GPU 实例分配的例子 当然,目前 MIG 的使用上还是有限制的。首先,MIG 只支援 Linux 系统,并且要求CUDA 11 以及450.36.06 以后的 NVIDIA 驱动软体。另外,在开启 MIG 时,不支援图像的接口(像是 OpenGL)、不支援 GPU to GPU P2P 的传输 (包括 PCIe 和 NVLINK)、GPU instance 之间的 CUDA IPC 也不支援,但支援计算实例之间的 CUDA IPC。如何管理 MIG 实例至此,我们已经展示了 MIG 带来的好处。但管理者要如何去管理这些 MIG 实例呢?要如何避免使用者之间只能使用被分配到的 MIG 实例,而不会去影响到其他的 MIG 实例呢?在 CUDA 11 里头,管理者可以透过赋予使用者访问特定檔案的权限来限制使用者对 MIG 实例的使用权限。 这些管理的檔案放在 “/proc/driver/nvidia/capabilities” 这个资料夹下,如图表 5 所示。对于 MIG 的管理,可以透过 “mig/config” 和 “mig/monitor” 来控制权限。其中,”mig/config” 一般只有 root 权限有访问的权限。拥有访问这个檔案权限的使用者能够管理 MIG 的实例,例如实例的创建和删除。而 “mig/monitor” 的权限则能让使用者看到整个 GPU 的资讯,例如目前 GPU 内存的使用量、是否开启 MIG 模式等等。 除了上述两个管理和监控 MIG 支援的权限之外,MIG 也提供针对特定 GPU 实例或是计算实例的使用权限。例如若要让使用者能够只能使用图表 5 当中的 gpu0 中 gi0 GPU实例下的ci0计算实例,则赋予该使用者访问 gpu0/gi0/access” 和 “gpu0/gi0/ci0/access” 这两个檔案的访问权限,并且移除该使用者访问其他实例的权限。如此一来,此位使用者便只能使用该实例而无法接触到其他实例 。图表5 MIG 管理树状图如何使用 MIG在这个章节当中,我们会一步一步的展示如何使用 MIG。开启 MIGMIG 的相关指令都整合进了 NVIDIA System Management Interface(nvidia-smi)当中。要开启 MIG,首先需要关闭所有和 GPU 相关的程序,因为启动 MIG 的时候会对 GPU 进行重启。其次,开启 MIG 一般需要 root 权限。确认符合上述的条件之后,就可以透过 sudo nvidia-smi mig 1 来开启 MIG。需要特别注意的是,开启 MIG 之后,在还没有建立 GPU 实例和计算实例之前,是不能使用 GPU 的。如果使用者在这个情况下直接去执行程序,会返回找不到相关装置的错误。如图表 6 所示。图表6 未建立 GPU 实例与计算实例时,执行 cuda 程序的错误在开启 MIG 之后,nvidia-smi 展示的界面会有所改变,增加了一栏针对 MIG 的说明,如图表 7 所示。图表7 开启 MIG 之后的 nvidia-smi创建 GPU 实例与计算实例接下来,我们可以透过 sudo nvidia-smi mig -lgip 来确认我们能够创建的 GPU 实例。图表 8 是我们将看到的样子。第一列是 GPU 的号码、第二列是各个实例设定的名称、第三列是创建该实例时使用的 ID、第四列是该实例剩余可创建数目,以及该实例最多可创建数目、第五列是该实例的内存大小、第七列是流处理器的数目。图表8 GPU 实例的资料-1 创建 GPU 实例的指令为 “sudo nvidia-smi mig -cgi ”。我们先透过下列的指令创建一个 1g.5gb 的实例,并且再次确认实例的资料 sudo nvidia-smi mig -cgi 19 sudo nvidia-smi mig -lgip图表9 GPU 实例的资料-2 从图表 9 我们可以发现,1g.5gb、3g.20gb、7g.39gb这三个实例剩余可创建数目减少了1。这是因为我们在创建一个 1g.5gb 实例之后,我们剩下 6 份流处理器和 7 份内存,因此无法再创建 7g.39gb 的实例,3g.20gb 的实例由于需要四份的内存,因此也只能再创建一个。 下一步,我们一次再创建两个 1g.5gb 实例以及一个 4g.20gb 实例。 sudo nvidia-smi mig -cgi 19,19,5 sudo nvidia-smi mig -lgip图表10 GPU 实例的资料-3 从图表 10 我们可以看到,所有 GPU 实例的剩余可创建数目都为 0。接下来,我们可以透过 “sudo nvidia-smi mig -lcip”看到我们现在拥有的 GPU 实例,以及可创建的计算实例。图表 11 展示了我们可以创建的计算实例有哪些。第一列代表的是所属的 GPU、第二列是所属的 GPU 实例的 ID、第三列是该计算实例的名称、第四列是创建该计算实例的 ID、第五列是该计算实例剩余可创建数目及最大可创建数目、第六列是该计算实例拥有的流处理器数目。另外,我们可以发现,表格中没有关于内存的资讯,这是因为一个 GPU 实例里头的所有内存是共用的,因此在创建计算实例时不需要再进行分配。 例如 GPU 实例 ID 7、8、9 他们都是 1g.5gb 的 GPU 实例,因此都只能创建一份 1g.5gb 的计算实例。而 GPU 实例 ID 2 是一个 4g.20gb 的 GPU 实例,可以创建 4 个拥有 14 个流处理器的 1c.4g.20gb 的计算实例、或者 2 个拥有 28 个流处理器的 2c.4g.20gb 的计算实例、或者 1 个拥有 56 个流处理器的 4g.20gb 的计算实例。图表11 计算实例资料-1 下一步,我们透过 ”sudo nvidia-smi mig -gi -cci ”在 GPU 实例 ID 7、8、9 上面各创建一个 1g.5gb 的计算实例,并在 GPU 实例 ID 2 上面创建两个 1c.4g.20gb 的计算实例与一个 2c.4g.20gb 的计算实例。 sudo nvidia-smi mig -gi 7,8,9 -cci 0 sudo nvidia-smi mig -gi 2 -cci 0,0,1 sudo nvidia-smi mig -lcip在创建完之后,我们可以会发现计算实例的资料会有所改变,如图表 12所示,可以发现所有 GPU 实例的可创建计算实例数目都为 0。接下来,我们再回头来看一下 “nvidia-smi” 目前的情况。从图表 13我们可以看见”MIG devices” 这一栏增加了新的资讯,里头包括了各个 GPU 实例的内存大小以及 ID,还有在 GPU 实例当中的计算实例的 ID 以及流处理器的数量。图表12 计算实例资料-1图表13 创建完计算实例后的 nvidia-smi在指定的计算实例上运行程序创建完计算实例之后,我们要如何在特定的实例上面执行我们的程序呢?目前,我们要在指定的 GPU 上面执行程序时,是透过 “CUDA_VISIBLE_DEVICES” 这一环境变数。而 CUDA 11 在这个环境变数上进行了扩充,除了能够指定第几个 GPU 之外,也能透过计算实例的 UUID 来直接指定使用哪个计算实例。首先,我们透过 “nvidia-smi -L” 来得到所有计算实例的 UUID。从图表 14,我们可以看到各个计算实例对应的UUID。这边 UUID 的命名格式为MIG-//。以图表 14 为例,所有实例的GPU-UUID都是“38b3962a-109c-d69d-c3d3-3c2e8cff25cb”,而 MIG 1c.4g.20gb Device 0 的 GPU 实例 ID 为 2,计算实例 ID 为 0。图表14 nvidia-smi -L 的结果 在得到计算实例的 UUID 之后,我们便可以透过 “CUDA_VISIBLE_DEVICES” 这个环境变量来指定。例如下面的指令可以在 2c.4g.20gbDevice2 上面执行程序: CUDA_VISIBLE_DEVICES=MIG-GPU-38b3962a-109c-d69d-c3d3-3c2e8cff25cb/2/2./bin/encoder_sample 32 12 32 12 64 1 0 图表15为运行时 “nvidia-smi” 的结果。图表15 在 2c.4g.20gb Device2 上运行程序删除计算实例与 GPU 实例最后,在计算完之后,我们要删除目前现有的计算实例和 GPU 实例。首先,我们需要先删除 GPU 实例当中的计算实例,才能删除 GPU 实例。而删除计算实例的方式为 ”sudo nvidia-smimig -gi -ci -dci”。例如下面的指令能够删除 GPU 实例 ID 7 上的计算实例 ID 0: sudo nvidia-smi mig -gi 7 -ci 0 -dci 下一步,我们可以透过 “sudo nvidia-smi mig-gi -dgi” 来删除 GPU 实例。例如下面的指令能够删除 GPU 实例 ID 7: sudo nvidia-smi mig -gi 7 -dgi图表16 删除 GPU 实例 ID 7之后的 nvidia-smi 透过 "nvidia-smi",我们能看到 GPU 实例 ID 7已经被删除,如图表 16 所示。 最后,我们透过以下的指令删除剩余的计算实例和 GPU 实例: sudo nvidia-smi mig -gi 2 -ci 0,1,2 -dci sudo nvidia-smi mig -gi 8,9 -ci 0 -dci sudo nvidia-smi mig -gi 2,8,9 -dgi 透过 "nvidia-smi",我们能看到所有的 GPU 实例和计算实例都已经被删除,如图表 17所示。图表17 删除剩余计算实例与 GPU 实例后的 nvidia-smi 最后,我们透过以下的指令关闭 MIG 模式: sudo nvidia-smi -mig 0 透过 nvidia-smi 我们可以看到已经将 MIG 模式关闭,如图表 18 所示。图表18 关闭 MIG 之后的 nvidia-smi应用案例分享 — 以 Faster Transformer 为例在这个章节,我们透过一个实际的例子来展示如何使用 MIG,以及 MIG 带来的效果。这里,我们使用的例子是 FasterTransformer (FT),这是 NVIDIA 发布的开源代码,针对 transformer 相关的推理进行了优化。FT 最新的代码在 https://github.com/NVIDIA/DeepLearningExamples/tree/master/FasterTransformer/v2.1 ,使用者可以透过 git clone 下载,并按照说明檔来将 FT 建立起来。建立的流程在此不再赘述。需要特别注意的是,由于 A100 是 compute capability 是 80,在建立 FT时需要将-DSM 设为 80。在建立完 FT 之后,我们接下来创建 7 个 1g.5gb 的 GPU 实例,并且在各个 GPU 实例上面创建一个计算实例。nvidia-smi mig -cgi 19,19,19,19,19,19,19nvidia-smi mig -gi 7,8,9,11,12,13,14 -cci 0接下来透过nvidia-smi -L得到所有计算实例的名字之后,透过 CUDA_VISIBLE_DEVICES 这个环境变量来指定。透过 nvidia-smi,我们可以看到7个实例同时在运行,如图表 19 所示。图表19 7个计算实例同时运行MIG 的性能接下来,我们展示 A100 在使用 MIG 前后的计算速度以及吞吐量的比较,并且和前一代的 GPU,NVIDIA T4 和 NVIDIA V100 进行比较。其中,NVIDIA T4 和 NVIDIA V100 的 CUDA 版本是 CUDA 10.2,NVIDIA A100 的 CUDA 版本是 CUDA 11。另外,这边只比较使用 FP16 的速度。当 A100 使用 MIG 时,若没有特别说明,我们的 GPU 实例和计算实例大小相同,并且每个 GPU 实例只使用一个计算实例。表格3:T4、V100、A100 的比较* 模型: BERT Base on FasterTransformer 2.1, sequencelength 128 表格 3 比较了 V100、T4、A100 不使用 MIG,A100 使用 7 个 MIG 计算实例、以及 A100 不使用 MIG 但使用 MPS 同时跑七个程式的时延以及吞吐量。这里需要特别注意的是,A100在batchsize 1时,会比 V100 还慢,这是由于 CUDA11 的新特性,FT 在使用 cuBLAS 时无法手动选择算法导致。从表格当中我们可以看到,在 batch size = 1 这种计算量比较小的情境下,MIG 可以带来 4 倍以上的吞吐量,而代价是时延会提升百分之五十。MPS 的效果则稍微比 MIG 要好一些。和上一代的 V100 与 T4 对比,A100 在开启 MIG 之后,吞吐量也提升成 3~4 倍。而在 batch size=128 这种计算量较大的情境下,开启 MIG 并不会带来什么效益,由于使用的流处理器较少(在使用 MIG 时,最多只能使用 98 个流处理器,而 A100 有 108 个流处理器),因此吞吐量反而会些微的下降。在这种情境下,A100 的吞吐量为 T4 的 6 倍以上,V100 的两倍左右。表格4:V100 与 A100 7个计算实例在不同 batch size 下的对比* 模型: BERT Base on FasterTransformer 2.1, sequencelength 128 表格 4 比较了 V100 与 A100 使用 7 个 MIG GPU 实例在不同 batch size 下的吞吐量与时延。我们可以看到在 batch size 为 8 的时候,A100 的吞吐量已经很接近峰值,距离峰值不到百分之十;另一方面,V100 在 batch size 8 的食后,吞吐量距离峰值还有百分之三十左右,在 batch size 32 时,吞吐量距离峰值也还有百分之十左右。这代表和上一代的 GPU 相比,MIG 在提升 GPU 的使用率上有很大的进步。表格5:A100、A100 使用 1 个计算实例与 A100 使用 7 个计算实例的比较* 模型: BERT Base on FasterTransformer 2.1, sequence length 128表格6:A100、A100 使用 1 个计算实例与 A100 使用 7 个计算实例的比较 -2* 模型: BERT Large on FasterTransformer 2.1, sequencelength 384 表格 5 和表格 6 比较了A100、A100 使用一个 7g.40gb 的计算实例、以及 A100 使用 7 个 1g.5gb 的计算实例,在 BERT Base 和 BERT Large 上的差别。从这两张表格当中,我们可以得到几个结论。首先,A100 在不使用 MIG 下,和使用一个 7g.40gb 的实例下,其吞吐量和时延都差不多。其次,在 BERT Base 这个较小的模型上,MIG 在 batch size 1 时能带来四倍以上的吞吐量、在 batch size 4 时能带来两倍以上的吞吐量。只有到 batch size 64 以上才没有什么效益。而在 BERT Large 上,在 batch size 1 时,MIG 能带来百分之五十以上的提升,但在更大的 batch size 上,MIG 就没有什么明显的笑意。结 论MIG 是新 GPU 架构-安培,以及 CUDA 11 推出的新特性。它最多能将一个 A100 GPU 切分成 7 个 GPU 实例,大幅提升 GPU 的使用率,同时藉由硬件上的分离提供可靠的服务品质与错误的分离。 更多 A100 与 MIG 的相关资讯,可以访问 https://www.nvidia.com/en-us/data-center/a100/ https://docs.nvidia.com/datacenter/tesla/mig-user-guide/index.html#cuda-visible-devices 欢迎大家前往参考。本文来自微信公众号NVIDIA开发者社区
  • [算子使用] GPU异构算子全流程开发指导
    Tips① 此文档详细介绍了MindSpore GPU异构算子开发流程,与官方文档相比本文档更加侧重于开发文件的解读以及常用开发方法的讲解。同时本文档用词相对简单,主要帮助大家了解GPU算子开发需要写什么,各种文件的作用是什么以及应该怎么写这些文件,而官方文档中则更偏向于基础概念和框架的介绍,建议大家将两个文档结合起来阅读,这样更能够加深理解。② 本文篇幅较长,如果有自己熟悉的内容可以直接跳过,但是建议大家能够仔细阅读第三章,充分了解GPU算子需要开发哪些文件、每个开发文件的作用以及常用的方法,这样可以更快的入门,开始开发算子。③ MindSpore GPU异构算子调用流程如下④ 此文档也有网页版:https://bbs.huaweicloud.com/blogs/364623,大家愿意的话可以支持一下,刷一刷阅读量。⑤ 本文主要为前期准备和开发流程,附加了两个关于接口文档测试的内容,后续进入测试阶段我也会写一个GPU算子测试指导。⑥ 如有遗漏或错误,欢迎指出与修改。常用网址① GPU 算子全流程开发指导录屏 ② BartlettWindow 算子 PR ③ MindSpore 算子 Issue 查询网址 ④ MindSpore 官方安装网址 ⑤ MindSpore 官方算子查询网址 ⑥ 谭升 -GPU 编程一、 环境配置1. 连接服务器(1) ssh跳转连接使用服务器链接软件在同一个连接会话中,依次连续输入以下命令:① 跳转服务器1:ssh jump@xxx.xxx.xxx.xxx passwd: xxxxxxxx② 跳转服务器2:ssh test@xxx.xxx.xxx.xxx passwd: xxxxxxxx③ GPU算子服务器:ssh user14@xxx.xxx.xxx.xxx -p xxxx pwd:xxxxxxxx输入ls,到达此界面即为连接成功!(2) VPN连接方法① 点击右下角网络(WiFi那个标识),点击网络和Internet设置。② 点击VPN,添加VPN(我已经设置过一个,请忽略)。③ 设置GPU算子开发网络连接,输入如下用户名和密码,点击保存。VPN信息:服务器地址:xxx.xxx.xxx.xxx用户名:xxxxxxxx密码:xxxxxxxx④ 打开控制面板-网络连接,并选择刚创建的vpn右键属性⑤ 选择网络-IPV4-属性-高级,取消勾选“在远程网络上使用默认网关”⑥ 选择安全,选择允许使用这些协议,点击确定⑦ 连接VPN后直接连接服务器ssh user14@xxx.xxx.xxx.xxx -p xxx pwd:xxxxxxxx输入ls,到达此界面即为连接成功!2. 下载mindspore包(1) 注册gitee账号在官网注册自己的gitee账号用于后续加入团队开发仓库和提交代码。https://gitee.com (2) 加入Owner仓库创建好gitee账户后联系团队owner,让他在自己的账户上邀请团队成员成为仓库开发者,这样成员的账号才有权限往owner仓库直接提交代码。将邀请链接或二维码发送给团队成员邀请其加入仓库:成员接收邀请后,owner可以在私信栏找到成员的申请,进入后点击同意即可。(3) 创建分支成员加入团队后登录自己的账号可以看到owner的主仓,点击进入点击分支点击新建分支,输入自己想取的分支名(一个算子使用一个单独的分支名):创建完成后回到owner的mindspore仓库,可以查询到自己的分支:(4) 下载mindspore连接服务器后创建自己的文件夹:mkdier wzb_SSSqG激活公共环境ci3.7source /home/Public/env.sh创建个人conda环境conda create -n wzb_SSSqG --clone ci3.7复制测试环境脚本cp /home/Public/env.sh wzb_SSSqG/cd wzb_SSSqG/sed -i 's/ci3.7/wzb_SSSqG/g' env.sh每次登陆后都需要执行以下脚本,开启测试环境source env.sh下载团队owner的mindspore仓库git clone https://gitee.com/EJaven/mindspore.git 进入mindspore文件夹,切换到自己创建的分支cd mindsporegit checkout SSSqG二、 开发流程本章将大致介绍本次GPU异构算子开发的整体流程,并对需要开发的文件进行简要介绍。大家最好能够先详细阅读一遍华为官方GPU算子开发指南,了解一些基本概念后结合此文档进行学习和理解。华为 GPU 算子开发指南 1. MindSpore GPU异构算子调用流程2. 开发文件清单依据以上GPU异构算子调用流程,我们可以总结本次GPU异构算子开发所需开发文件以及文件功能如下:算子Python侧前端定义文件正向单算子:mindspore/python/mindspore/ops/operations/yyy_ops.py(yyy为算子所属类别:[array、math、other、random、sparse]等)反向单算子:mindspore/python/mindspore/ops/operations/_grad_ops.py功能描述:编写算子Python侧前端接口,包括算子声明、校验可写性、注册算子属性算子C++侧前端推理文件正向单算子:mindspore/core/ops/xxx.hmindspore/core/ops/xxx.cc反向单算子:mindspore/core/ops/grad/xxx_grad.hmindspore/core/ops/grad/xxx_grad.cc(xxx为算子名)算子PrimitivePtr定义:mindspore/core/ops/core_ops.h算子infer推理实际值注册:mindspore/core/abstract/ops/primitive_infer_map.cc功能描述:编写算子C++侧前端接口以及推理函数,包括算子类声明、输入shape/type校验、输出shape/type的infer推理算子C++侧后端适配文件正向单算子:mindspore/ccsrc/plugin/device/gpu/kernel/yyy/xxx_gpu_kernel.hmindspore/ccsrc/plugin/device/gpu/kernel/yyy/xxx_gpu_kernel.cc反向单算子:mindspore/ccsrc/plugin/device/gpu/kernel/yyy/xxx_grad_gpu_kernel.hmindspore/ccsrc/plugin/device/gpu/kernel/yyy/xxx_grad_gpu_kernel.cc(yyy为算子所属类别:[array、math、other、random、sparse], xxx为算子名)功能描述:编写算子C++侧后端接口以及适配函数,包括算子注册、初始化参数校验、内存计算、调用cuda核函数、数据类型注册算子cuda核函数开发正向单算子:mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/xxx_impl.cuhmindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/xxx_impl.cumindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_class/xxx_helper.h反向单算子:mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/xxx_grad_impl.cuhmindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/xxx_grad_impl.cumindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_class/xxx_grad_helper.h(xxx_helper.h和xxx_grad_helper.h文件为可选文件,如果算子逻辑复杂就写,不复杂可以不用写)功能描述:编写算子cuda核函数,包括核函数模板声明、函数模板定义、函数模板实例化、实现多线程核函数算子Python反向实现文件正向单算子:mindspore/python/mindspore/ops/_grad_experimental/grad_yyy_ops.py(yyy为算子所属类别:[array、math、other、random、sparse]等)反向单算子:反向算子一般无反向,所以不用写功能描述:编写算子反向实现函数,包括算子反向注册、实现算子反向逻辑算子ST测试文件正向单算子:tests/st/ops/gpu/test_xxx_op.py反向单算子:tests/st/ops/gpu/test_xxx_grad_op.py(xxx为算子名,反向单算子视作正向单算子来测试)功能描述:测试算子在PYNATIVE_MODE和GRAPH_MODE两种模式下的功能和精度是否符合标准3. 算子分类介绍(1) 正向单算子平常大家使用的基本上都是正向算子,比如pow算子、mean算子、max算子等等,这些算子的功能就是其数学推导,也即正向计算过程。(2) 正向算子的反向实现正向单算子的反向实现逻辑,往往只需要在mindspore/python/mindspore/ops/_grad_experimental/grad_yyy_ops.py文件中使用反向单算子或其他一些正向单算子的组合来实现求算子反向的逻辑即可。反向单算子样例:正向算子组合样例:(3) 反向单算子反向单算子实际上就是把正向单算子求反向的逻辑单独封装成一个算子,这样在实现“正向单算子的反向实现时”就可以直接调用这个算子的反向单算子。在测试的时候反向单算子可以按照正向单算子的方式来调用算子进行输出的校验,但是在确保没问题之后,正常情况下必须通过mindspore的自动求导来隐式调用反向单算子,不可显示调用。(4) 动态shape算子动态shape算子指的是算子输出y的shape会依据输入x的数值的变化而变化,比如BartlettWindow算子需要实现BartlettWindow函数:该算子其中一个输入为window_length,公式中表示为N,即输出y的元素个数。这时,我们输出y的shape = [N],即输入window_length发生变化时,输出y的shape会随之变化,这种算子就称为动态shape算子。一定要注意是输出y的shape会依据输入x的“数值”的变化而变化,如果是输入x的shape变化导致输出y的shape变化,这种不算动态shape算子。(5) 支持动态shape测试大家依据第(4)条判断完自己的算子是否为动态shape算子之后,无论结果,GPU异构算子验收时要求所有算子支持“动态shape输入”测试,动态shape输入和动态shape算子不是一个概念,实际上动态shape输入就是对原先静态的输入添加一步gather,只需要参照模板编写测试用例和网络即可,动态shape输入和普通静态输入差别如下:三、 算子开发阅读完第二章,大家应该已经了解本次GPU异构算子的调用流程以及需要开发的文件,即已经知道这个项目是要做什么。因此本章将依据第二章中罗列的需要开发的文件进行逐一讲解,主要介绍每个需要开发的文件中需要写哪些方法,这些方法的作用是什么。另外本章仅讲解正向单算子的开发过程,反向单算子的开发过程实际上一样的,仅仅是把原先正向单算子的输出y以及其对应的梯度dy也作为输入来依据逻辑进行计算罢了。为了便于理解,我将以之前开发过的一个BartlettWindow算子为例,分析该算子中各个文件中编写的代码的含义,但为了保证阅读的连贯性和排版的美观。我不会在本章节直接放出该算子的代码截图,而是将该算子的详细解析放在附录中并以超链接的方式链接到文中。除此之外大家还可以直接前往BartlettWindow算子的PR:https://gitee.com/mindspore/mindspore/pulls/35601 查看全部完整代码。1. 算子Python侧前端定义(1) mindspore/python/mindspore/ops/operations/yyy_ops.py文件名中的yyy代表算子所属类型[array、math、other、random、sparse],此文件功能为编写算子python侧前端接口,之后我们调用算子的时候实际上就是从这个地方开始执行的。其主要需要编写的内容和方法如下:① 算子接口注释编写关于算子的描述,包括:功能、数学表达式、参数、输入、输出、异常处理、支持的平台和算子样例等信息。这一部分很容易出现遗漏,同时还需要注意异常处理的报错信息要清晰明确,要让测试人员读得懂,不能模棱两可,否则很大概率会被直接打回。算子接口注释编写完之后需要按照如何做 doctest 和如何做接口网页自验证 进行注释文档的校验以及接口网页的自验证,以确保所写注释是正确并且符合规范的。BartlettWindow 算子接口注释 BartlettWindow 算子接口网页 ② Init()函数编写算子的初始化函数,主要进行算子参数的校验,包括属性类型的校验、输入类型的校验、数据类型的校验等。常见方法及其作用如下:• self.add_prim_attr("xxx", 1000)给算子添加一个属性xxx,其数值为1000• validator.check_value_type("xxx", xxx, [a,b,c], self.name):检验xxx参数的类型是否为a,b,c,如果都不是则报错。• validator.check_type_name("xxx", xxx, [float16,float32] self.name)检验xxx参数的数据类型是否为float16,float32,如果都不是则报错其他还有很多检验的方法,这些方法的命名都非常的清晰,大家看一下方法名就能大概猜出来它的作用,所以如果大家有想要校验的参数可以在库里现有的代码中搜一下用法。BartlettWindow 算子 接口 i nit() 函数 ③ 推理函数现有MindSpore算子库中的部分算子同时有Python侧与C++侧的推理函数,包括InferType()和InferShape()。为了更好地性能以及其他一些功能上的需求,默认优先加载C++层的推理函数。我们只需要定义C++侧的推理函数即可。因此如见到Python侧的推理函数,包括infer_dtype(), infer_shape(), infer()等函数时请直接忽略。2. 算子C++侧前端推理(1) mindspore/core/ops/xxx.h此文件用于声明C++侧算子类、前端接口函数以及用于输入shape/type校验以及输出shape/type推理的infer函数,另外如果后续需要用到算子的某个属性值则需要增加该属性的set和get函数。此文件开发难度不大,参照一下模板和库中其他算子写就行了:BartlettWindow 算子 bartlett_window.h 文件 (2) mindspore/core/ops/xxx.cc此文件用于实现xxx.h文件中声明的接口函数、输入type校验和输出type推导的infertype函数、输入shape校验和输出shape推导的infershape函数和属性相关的初始化、set和get函数。① MyOpsInferType函数校验输入的type(是否为tuple、tensor等等)、dtype(是否为float32、int32等等)以及存在约束关系的变量数据类型是否相同,并最终推理出输出的dtype,编码模板如下:图例外的其余校验请自行查找相应函数的用法。② MyOpsInferShape函数校验输入的shape是否满足要求(维度要求,大小要求)、存在约束关系的变量之间shape是否满足要求(相等、更大、更小),并最终推理出输出的shape,编码模板如下:Infershape部分代码需要开发者针对自己的算子进行完整仔细的分析,考虑到每个参数自身shape的限制,更要考虑到不同参数之间的相互约束,最后需要通过逻辑推理计算出输出的shape值。③ 属性init()、set()、get()函数如果在之后的后端适配和核函数实现部分需要用到算子的某些属性数值的话就需要在此处编写这三个函数,用来给该属性初始化,赋值和取值,以MyOps的属性myattr为例,代码模板如下:大家在编写完编译的时候,可能会发现自己的kMyAttr这个变量没有声明,此时我们只需要在mindspore/core/ops/op_name.h文件中新增一个kMyAttr变量的声明即可,代码模板如下:④ MyOpsInfer函数此函数为整个InferShape和InferType的组合函数,作用为判断空值、调用InferShape和InferType函数,写法也相对固定,代码模板如下:⑤ 宏定义 此处代码形式完全固定,调用MIND_API_BASE_IMPL用于注册算子继承关系,REGISTER_PRIMITI VE_EVAL_IMPL 用于注册算子推理函数,依据以下代码模板编写即可:BartlettWindow 算子 bartlett_window.cc 文件 (3) mindspore/core/ops/core_ops.h此文件用于添加算子PrimitivePtr定义,对算子进行注册,使得算子可以调用C++侧接口,此处代码形式也是固定的,参考以下模板即可:BartlettWindow 算子 core_ops.h 文件 (4) mindspore/core/abstract/ops/primitive_infer_map.cc首先我们需要了解如下信息:这一块简单来说就是我们的算子在动态shape测试时会调用两次infershape函数,而第一次调用的时候处于图编译状态,此时输出的shape无法确定,并且输入参数BuildValue得到的是AnyValue,无法获取真实值,而第二次调用的时候是出于运行时,这时可以获取到输入参数的真实数值。而想要获取到输入参数的真实数值,我们就需要在mindspore/core/abstract/ops/primitive_infer_map.cc文件中对算子infershape阶段需要获取真实数值的输入参数进行注册,注册代码也是固定形式的,参考以下模板即可:ShapeSet{0,2}的含义就是MyOps算子在推理时,需要用到第1、3两个输入的真实值。BartlettWindow 算子 primitive_infer_map.cc 文件 3. 算子C++侧后端适配(1) mindspore/ccsrc/plugin/device/gpu/kernel/yyy/xxx_gpu_kernel.h此处用于适配框架接口,由于对mindspore框架适配的需要, 在定义算子类时,都需要继承自NativeGpuKernelMod基类。根据最新框架要求,不可以继承DeprecatedNativeGpuKernelMod基类。否则算子打回重新适配 。而根据 NativeGpuKernelMod基类定义,所有算子都必须实现如下几个虚函数接口 :Init()、Resize()、GetOpSupport()、Launch(),除此之外为了适配动态shape测试,我们还需要编写内存清理ResetResource()函数以及定义一些类私有变量。此部分代码形式固定,大部分只需要参照模板编写即可:① Init()函数声明② Resize()函数声明③ GetOpSupport()函数声明④ Launch()函数声明⑤ RestResource()函数实现此函数用于在动态shape过程中清除之前输入、输出以及工作空间所申请的内存,避免之前的计算结果影响新一轮计算,代码模板如下:⑥ 类私有变量声明BartlettWindow 算子 bartlett_window_gpu_kernel.h 文件 (2) mindspore/ccsrc/plugin/device/gpu/kernel/yyy/xxx_gpu_kernel.cc此文件主要需要实现xxx_gpu_kernel.cc中声明的各种方法,用于校验输入、初始化参数、重计算并更新内存空间、核函数调用以及数据类型注册。① Init()函数实现此函数需要完成对于输入的一些简单校验、依据输入的数据类型选择正确的类模板、最后初始化一些参数并计算与内存相关的一些变量值,代码模板如下:② Resize()函数实现此函数用于判断是否为动态shape测试、清空内存、判断输出是否为空并申请新内存,代码模板如下:③ 核函数调用此函数获取输入、输出和工作空间的内存地址并传入核函数接口即可,代码模板如下:④ 数据类型注册此函数进行数据类型注册,参照代码模板编写即可:⑤ GetOpSupport()函数声明直接复制库里现有代码即可:BartlettWindow 算子 bartlett_window_gpu_kernel.cc 文件 4. 算子cuda核函数开发(1) mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/xxx_impl.cuh此文件用于函数模板声明,确定好自己需要用到的输入参数后参照库中已有代码进行修改即可:B artlettWindow 算子 b artlett_window_impl.cuh 文件 (2) mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/xxx_impl.cu此文件包含三大部分,第一部分是函数模板定义、第二部分为函数模板显示实例化、第三部分为多线程核函数实现。① 函数模板定义依据xxx_impl.cuh文件中声明的函数写相应的定义即可,形式固定,参照模板写即可:② 函数模板显示实例化依据算子支持的数据类型将函数模板进行实例化,其目的是为了加快实际计算时的运行速率,代码形式固定,参照模板写即可:③ 多线程核函数实现这一块才是真正用来实现算子计算逻辑的地方,也是整个GPU异构算子开发中最重要的部分,想要写好这一块的代码,大家首先需要了解自己算子的功能、计算逻辑,同时还需要有一定的Cuda编程基础,这边推荐大家看一下华为官网上推荐的cuda编程教程:cuda 编程(谭升) 大家主要了解一下host、device、grid、block和thread的概念就可以大致理解cuda是如何进行多线程计算的,之后就可以参照友商或者自己构建多线程计算逻辑,实现算子在GPU上的并行计算。以加法算子Add为例,假设有一个输入x和一个输入y,我们需要求z = Add(x, y),其线性计算和并行计算逻辑如下:BartlettWindow 算子 bartlett_window_impl.cu 文件 (3) mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_class/xxx_helper.hBartlettWindow算子逻辑不复杂,因此没有写helper.h文件,故此处无法给出解析,大家可以参考官方开发文档中关于cudaKernel 封装 的介绍,也可以参考库中其他算子该文件的写法。5. 算子Python反向实现文件(1) mindspore/python/mindspore/ops/_grad_experimental/grad_yyy_ops.pyBartlettWindow算子无反向,故无需修改此文件。大家可以参考华为官方文档中关于注册 算子反向 的描述以及正向算子的反向实现 。算子反向实现的模板如下图所示,而实际计算逻辑则需要大家自己进行求导推理或者参照友商反向实现逻辑进行实现:6. 算子ST测试文件(1) tests/st/ops/gpu/test_xxx_op.py此文件用于门禁ST测试的用例,需要包括GRAPH_MODE和PYNATIVE_MODE两种模式的测试用例,同时需要注意这里的输入必须使用固定值,不可以随机生成,具体写法参照库中已有代码即可。BartlettWindow 算子 test_bartlett_window_op.py 文件 四、 如何做doctest用于检查接口注释是否规范1、进入GPU服务器中自己conda环境中的mindspore包:cd /disk1/user14/.conda/envs/xxx/lib/python3.7/site-packages/mindspore2、在当前路径传入conftest.py文件(见附件)3、修改conftest.py文件中导入的operation类型:我的是array,改成自己对应的4、进入算子python侧前端接口定义目录:cd /disk1/user14/.conda/envs/xxx/lib/python3.7/site -packages/mindspore/ops/operations5、在当前目录下传入需要doctest的算子接口文件array_ops_xxx.py:前面的array改成自己的算子类型,后面的xxx为算子名称6、增添算子接口将自己算子的前端定义复制粘贴到上面创建的array_ops_xx.py文件中,并且要增加支持的平台类型,就是在目前已经支持的平台前加上GPU,不仅要在此mindspore包中修改,自己提交的代码中也需要加上:从哪里复制?从原本算子对应的前端接口处复制,就是在此路径下的array_ops.py(array换成自己算子的类型)文件中搜自己的算子前端定义。7、在此路径下执行doctest命令(array换成自己的算子类型):pytest --disable-warnings -vra --doctest-modules -o doctest_optionflags=NORMALIZE_WHITESPACE --tb=long array_ops_xxx.py8、如果通过结果如下图,如果报错根据报错信息修改注释五、如何做接口网页自验证用于检查网页前端接口是否正确,主要依照官方文档接口注释网页自验部分:算子前端定义 wiki 需要注意四点:① clone docs代码时选择老的分支:git clone -b r1.6 https://gitee.com/mindspore/docs.git ② 编辑operation.rstoperation.rst文件可以用电脑的记事本打开,并且注意添加到自己算子所属类型下面③ 报错信息生成网页过程中会报一些没有图片的错误,不需要管。生成完毕后需要将/disk1/user14/wzb/docs/docs/mindspore/api/build_zh_cn路径下的整个html文件夹下载下来④ 网页截图html文件夹全部下载之后,打开index.html文件,在其中找自己算子的前端接口网页截图即可。附录1:BartlettWindow算子开发详解本文采取代码+注释方式解析BartlettWindow算子的开发过程。1. 算子Python侧前端定义(1) mindspore/python/mindspore/ops/operations/other_ops.py① 算子接口注释写完算子接口注释后可以通过()接口网页自验证生成如下网页文件:② Init()函数③ 推理函数BartlettWindow算子无Python侧Infer推理函数2. 算子C++侧前端推理(1) mindspore/core/ops/bartlett_window.h(2) mindspore/core/ops/bartlett_window.cc① BartlettWindowInferType函数图例外的其余校验请自行查找相应函数的用法。② BartlettWindowInferShape函数③ 属性init()、set()、get()函数mindspore/core/ops/op_name.h ④ BartlettWindowInfer函数⑤ 宏定义 (3) mindspore/core/ops/core_ops.h(4) mindspore/core/abstract/ops/primitive_infer_map.cc3. 算子C++侧后端适配(1) mindspore/ccsrc/plugin/device/gpu/kernel/other/bartlett_window _gpu_kernel.h① Init()函数声明② Resize()函数声明③ GetOpSupport()函数声明④ Launch()函数声明⑤ RestResource()函数实现⑥ 类私有变量声明(2) mindspore/ccsrc/plugin/device/gpu/kernel/other/bartlett_window _gpu_kernel.cc① Init()函数实现② Resize()函数实现③ 核函数调用④ 数据类型注册⑤ GetOpSupport()函数声明4. 算子cuda核函数开发(1) mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/bartlett_window _impl.cuh(2) mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/bartlett_window _impl.cu① 函数模板定义② 函数模板显示实例化③ 多线程核函数实现(3) mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_class/bartlett_window _helper.hBartlettWindow算子无需编写helper.h文件5. 算子Python反向实现文件(1) mindspore/python/mindspore/ops/_grad_experimental/grad_other_ops.pyBartlettWindow算子无反向6. 算子ST测试文件(1) tests/st/ops/gpu/test_ bartlett_window _op.py
  • [算子编译] 【GPU算子开发】测试用例单个执行正确,一起执行报错 cudaStreamSynchronize failed
    【功能模块】ScaleAndTranslate算子开发:https://gitee.com/mindspore/mindspore/issues/I5EWTN?from=project-issue【操作步骤&问题现象】1、总共3个测试用例,依次第一个入参image的类型为float32、int8、int16.2、这三个用例单独执行是正确的,一起执行时,仅第一个float32会通过,经过其他测试发现,就是images的类型变化,导致后续的用例不通过。3、经过打印gpu_kernel以及cu文件的调用过程,均没有报错。报错可能发生在计算出结果后,框架从device端取数据产生的错误(个人猜测)。不知道是在哪里的原因导致的。【截图信息】1、自己打印的gpu_kernel以及cu文件日志(一次完整流程)。2、截图3、截图4、截图【日志信息】(可选,上传日志内容或者附件)PR链接:https://gitee.com/mindspore/mindspore/pulls/39777
  • [其他干货] CUDA编程(九)原子操作
    CUDA的原子操作针对的是Global Memory或者是Shared Memory。为什么要引入原子操作这个概念。我们从前几天的训练营课程得知:Shared Memory是可被同一个block的所有thread访问(读写)的。Global Memory相当于显存,可以被所有thread访问(读写)的。那么,这两种Memory,就很可能会遇到多个thread同时读写同一块内存区域的问题。假如两个线程都在做“读取-修改-写入"操作,如果在这个操作中,出现互相交错的情况,就会出现混乱。举个例子,比如有块内存里面的值是10,A、B两个用途为”加一“的线程同时读该块内存,然后各自都加1,A将值变为11,再写回去;B也将值改为11,也写了回去。这个时候,结果就变成了11。但是显然我们要求的结果应为12。我们只好要求将“读取-修改-写入"捆绑成一个逻辑上的单体操作,不可拆分,逻辑上顺序进行,保证一次性成功。这样才能确保任何一次的操作对变量的操作结果的正确性。常用的原子操作函数如下:这些函数大多会返回原子操作前的变量值。原子操作的函数存在多态,适用于不同数据类型和精度的版本,以atomicAdd为例:我们来实战吧!(a)实战1:对1000万的整型数组求和关于对向量所有元素求和这个事情,讲师何老师提供了一个框架。他通过ppt介绍了这个框架的原理。看起来比较复杂。他以只有32个数据的求和为例图示了这个过程:具体的代码如下:sum.cu#include<stdio.h> #include<stdint.h> #include<time.h> //for time() #include<stdlib.h> //for srand()/rand() #include<sys/time.h> //for gettimeofday()/struct timeval #include"error.cuh" #define N 10000000 #define BLOCK_SIZE 256 #define BLOCKS ((N + BLOCK_SIZE - 1) / BLOCK_SIZE) __managed__ int source[N]; //input data __managed__ int final_result[1] = {0}; //scalar output __global__ void _sum_gpu(int *input, int count, int *output) { __shared__ int sum_per_block[BLOCK_SIZE]; int temp = 0; for (int idx = threadIdx.x + blockDim.x * blockIdx.x; idx < count; idx += gridDim.x * blockDim.x ) { temp += input[idx]; } sum_per_block[threadIdx.x] = temp; //the per-thread partial sum is temp! __syncthreads(); //**********shared memory summation stage*********** for (int length = BLOCK_SIZE / 2; length >= 1; length /= 2) { int double_kill = -1; if (threadIdx.x < length) { double_kill = sum_per_block[threadIdx.x] + sum_per_block[threadIdx.x + length]; } __syncthreads(); //why we need two __syncthreads() here, and, if (threadIdx.x < length) { sum_per_block[threadIdx.x] = double_kill; } __syncthreads(); //....here ? } //the per-block partial sum is sum_per_block[0] if (blockDim.x * blockIdx.x < count) //in case that our users are naughty { //the final reduction performed by atomicAdd() if (threadIdx.x == 0) atomicAdd(output, sum_per_block[0]); } } int _sum_cpu(int *ptr, int count) { int sum = 0; for (int i = 0; i < count; i++) { sum += ptr[i]; } return sum; } void _init(int *ptr, int count) { uint32_t seed = (uint32_t)time(NULL); //make huan happy srand(seed); //reseeding the random generator //filling the buffer with random data for (int i = 0; i < count; i++) ptr[i] = rand(); } double get_time() { struct timeval tv; gettimeofday(&tv, NULL); return ((double)tv.tv_usec * 0.000001 + tv.tv_sec); } int main() { //********************************** fprintf(stderr, "filling the buffer with %d elements...\n", N); _init(source, N); //********************************** //Now we are going to kick start your kernel. cudaDeviceSynchronize(); //steady! ready! go! fprintf(stderr, "Running on GPU...\n"); double t0 = get_time(); _sum_gpu<<<BLOCKS, BLOCK_SIZE>>>(source, N, final_result); CHECK(cudaGetLastError()); //checking for launch failures CHECK(cudaDeviceSynchronize()); //checking for run-time failurs double t1 = get_time(); int A = final_result[0]; fprintf(stderr, "GPU sum: %u\n", A); //********************************** //Now we are going to exercise your CPU... fprintf(stderr, "Running on CPU...\n"); double t2 = get_time(); int B = _sum_cpu(source, N); double t3 = get_time(); fprintf(stderr, "CPU sum: %u\n", B); //******The last judgement********** if (A == B) { fprintf(stderr, "Test Passed!\n"); } else { fprintf(stderr, "Test failed!\n"); exit(-1); } //****and some timing details******* fprintf(stderr, "GPU time %.3f ms\n", (t1 - t0) * 1000.0); fprintf(stderr, "CPU time %.3f ms\n", (t3 - t2) * 1000.0); return 0; } 由于其原理略有复杂,张小白是这么想的:以上的代码其实是提供了一个GPU遍历所有字段的框架,这是一个分而治之的思路:block中的多个线程负责多个数据点,这些点被规约(reduce/缩减)到一个标量。这样每个block中都有一个标量的结果。但blocks有很多,这些变量组成的数组/向量,还需要二次缩减到最终的1个标量值。以上过程存在两步reduce,第一步用并行折半缩减(规约),第二步直接用原子操作函数atomicAdd规约。两步完成后,得到了单一点。我们运行下试试:可见,CPU和GPU求和的结果是一致的,说明这个遍历所有字段的框架是没问题的。看下性能:(b)实战2:对1000万的整型数组求出最大值和最小值基于上面实战1分析的原理,我们接着分析本题的解题思路:同样使用两步reduce,第一步用并行折半缩减(规约),第二步直接用原子操作atomicMax和atomicMin规约。两步完成后,得到了单一点(最大值/最小值)。于是我们就像搭积木那样,将一个sum改为一个max和一个min,代码变动如下:min_or_max.cu#include<stdio.h> #include<stdint.h> #include<time.h> //for time() #include<stdlib.h> //for srand()/rand() #include<sys/time.h> //for gettimeofday()/struct timeval #include"error.cuh" #define N 10000000 #define BLOCK_SIZE 256 #define BLOCKS ((N + BLOCK_SIZE - 1) / BLOCK_SIZE) __managed__ int source[N]; //input data //__managed__ int final_result[2] = {INT_MIN,INT_MAX}; //scalar output __managed__ int final_result_max = INT_MIN; //scalar output __managed__ int final_result_min = INT_MAX; //scalar output __global__ void _sum_min_or_max(int *input, int count, int *max_output, int *min_output) { __shared__ int max_per_block[BLOCK_SIZE]; __shared__ int min_per_block[BLOCK_SIZE]; int max_temp = 0; int min_temp = 0; for (int idx = threadIdx.x + blockDim.x * blockIdx.x; idx < count; idx += gridDim.x * blockDim.x ) { //temp += input[idx]; max_temp = (input[idx] > max_temp) ? input[idx] :max_temp; min_temp = (input[idx] < min_temp) ? input[idx] :min_temp; } max_per_block[threadIdx.x] = max_temp; //the per-thread partial max is temp! min_per_block[threadIdx.x] = min_temp; //the per-thread partial max is temp! __syncthreads(); //**********shared memory summation stage*********** for (int length = BLOCK_SIZE / 2; length >= 1; length /= 2) { int max_double_kill = -1; int min_double_kill = -1; if (threadIdx.x < length) { max_double_kill = (max_per_block[threadIdx.x] > max_per_block[threadIdx.x + length]) ? max_per_block[threadIdx.x] : max_per_block[threadIdx.x + length]; min_double_kill = (min_per_block[threadIdx.x] < min_per_block[threadIdx.x + length]) ? min_per_block[threadIdx.x] : min_per_block[threadIdx.x + length]; } __syncthreads(); //why we need two __syncthreads() here, and, if (threadIdx.x < length) { max_per_block[threadIdx.x] = max_double_kill; min_per_block[threadIdx.x] = min_double_kill; } __syncthreads(); //....here ? } //the per-block partial sum is sum_per_block[0] if (blockDim.x * blockIdx.x < count) //in case that our users are naughty { //the final reduction performed by atomicAdd() //if (threadIdx.x == 0) atomicAdd(output, max_per_block[0]); if (threadIdx.x == 0) atomicMax(max_output, max_per_block[0]); if (threadIdx.x == 0) atomicMin(min_output, min_per_block[0]); } } int _max_min_cpu(int *ptr, int count, int *max1, int *min1) { int max = INT_MIN; int min = INT_MAX; for (int i = 0; i < count; i++) { //sum += ptr[i]; max = (ptr[i] > max)? ptr[i]:max; min = (ptr[i] < min)? ptr[i]:min; } //printf(" CPU max = %d\n", max); //printf(" CPU min = %d\n", min); *max1 = max; *min1 = min; return 0; } void _init(int *ptr, int count) { uint32_t seed = (uint32_t)time(NULL); //make huan happy //srand(seed); //reseeding the random generator //filling the buffer with random data for (int i = 0; i < count; i++) { //ptr[i] = rand() % 100000000; ptr[i] = rand() ; if (i % 2 == 0) ptr[i] = 0 - ptr[i] ; } } double get_time() { struct timeval tv; gettimeofday(&tv, NULL); return ((double)tv.tv_usec * 0.000001 + tv.tv_sec); } int main() { //********************************** fprintf(stderr, "filling the buffer with %d elements...\n", N); _init(source, N); //********************************** //Now we are going to kick start your kernel. cudaDeviceSynchronize(); //steady! ready! go! fprintf(stderr, "Running on GPU...\n"); double t0 = get_time(); _sum_min_or_max<<<BLOCKS, BLOCK_SIZE>>>(source, N, &final_result_max, &final_result_min); CHECK(cudaGetLastError()); //checking for launch failures CHECK(cudaDeviceSynchronize()); //checking for run-time failures double t1 = get_time(); //int A = final_result[0]; fprintf(stderr, " GPU max: %d\n", final_result_max); fprintf(stderr, " GPU min: %d\n", final_result_min); //********************************** //Now we are going to exercise your CPU... fprintf(stderr, "Running on CPU...\n"); double t2 = get_time(); int cpu_max=0; int cpu_min=0; int B = _max_min_cpu(source, N, &cpu_max, &cpu_min); printf(" CPU max = %d\n", cpu_max); printf(" CPU min = %d\n", cpu_min); double t3 = get_time(); //fprintf(stderr, "CPU sum: %u\n", B); //******The last judgement********** if ( final_result_max == cpu_max && final_result_min == cpu_min ) { fprintf(stderr, "Test Passed!\n"); } else { fprintf(stderr, "Test failed!\n"); exit(-1); } //****and some timing details******* fprintf(stderr, "GPU time %.3f ms\n", (t1 - t0) * 1000.0); fprintf(stderr, "CPU time %.3f ms\n", (t3 - t2) * 1000.0); return 0; } 这里需要指出几点:(1)初始化最大值变量final_result_max的时候,给它赋最小值INT_MIN;初始化最小值变量final_result_min的时候,给它赋最大值INT_MAX,这样在它比较的时候,就一定会被比下去,换成最新的值。如果有人不小心写反了,那么就完蛋了。不信大家可以试试。(2)在产生1000万个随机数的时候,张小白采纳了何老师的建议,每两个数就有一个正数,有一个负数。这样不会导致原来取最小值永远是0的情况。编译运行:看起来CPU和GPU算出的结果都是一致的。怎么样?简单吧?上面的代码,张小白偷懒,使用了两个managed变量记录结果,张小白看了看后面的作业,还有一道“找到1000万数据中前10个最大值”的题目,感觉还是用 数组会更合适点。也许可以无缝的升级解决后面这道题,于是张小白又做了以下改动:#include<stdio.h> #include<stdint.h> #include<time.h> //for time() #include<stdlib.h> //for srand()/rand() #include<sys/time.h> //for gettimeofday()/struct timeval #include"error.cuh" #define N 10000000 #define BLOCK_SIZE 256 #define BLOCKS ((N + BLOCK_SIZE - 1) / BLOCK_SIZE) __managed__ int source[N]; //input data __managed__ int final_result[2] = {INT_MIN,INT_MAX}; //scalar output //__managed__ int final_result_max = INT_MIN; //scalar output //__managed__ int final_result_min = INT_MAX; //scalar output //__global__ void _sum_min_or_max(int *input, int count, int *max_output, int *min_output) __global__ void _sum_min_or_max(int *input, int count,int *output) { __shared__ int max_per_block[BLOCK_SIZE]; __shared__ int min_per_block[BLOCK_SIZE]; int max_temp = 0; int min_temp = 0; for (int idx = threadIdx.x + blockDim.x * blockIdx.x; idx < count; idx += gridDim.x * blockDim.x ) { //temp += input[idx]; max_temp = (input[idx] > max_temp) ? input[idx] :max_temp; min_temp = (input[idx] < min_temp) ? input[idx] :min_temp; } max_per_block[threadIdx.x] = max_temp; //the per-thread partial max is temp! min_per_block[threadIdx.x] = min_temp; //the per-thread partial max is temp! __syncthreads(); //**********shared memory summation stage*********** for (int length = BLOCK_SIZE / 2; length >= 1; length /= 2) { int max_double_kill = -1; int min_double_kill = -1; if (threadIdx.x < length) { max_double_kill = (max_per_block[threadIdx.x] > max_per_block[threadIdx.x + length]) ? max_per_block[threadIdx.x] : max_per_block[threadIdx.x + length]; min_double_kill = (min_per_block[threadIdx.x] < min_per_block[threadIdx.x + length]) ? min_per_block[threadIdx.x] : min_per_block[threadIdx.x + length]; } __syncthreads(); //why we need two __syncthreads() here, and, if (threadIdx.x < length) { max_per_block[threadIdx.x] = max_double_kill; min_per_block[threadIdx.x] = min_double_kill; } __syncthreads(); //....here ? } //the per-block partial sum is sum_per_block[0] if (blockDim.x * blockIdx.x < count) //in case that our users are naughty { //the final reduction performed by atomicAdd() //if (threadIdx.x == 0) atomicAdd(output, max_per_block[0]); //if (threadIdx.x == 0) atomicMax(max_output, max_per_block[0]); //if (threadIdx.x == 0) atomicMin(min_output, min_per_block[0]); if (threadIdx.x == 0) atomicMax(&output[0], max_per_block[0]); if (threadIdx.x == 0) atomicMin(&output[1], min_per_block[0]); } } int _max_min_cpu(int *ptr, int count, int *max1, int *min1) { int max = INT_MIN; int min = INT_MAX; for (int i = 0; i < count; i++) { //sum += ptr[i]; max = (ptr[i] > max)? ptr[i]:max; min = (ptr[i] < min)? ptr[i]:min; } //printf(" CPU max = %d\n", max); //printf(" CPU min = %d\n", min); *max1 = max; *min1 = min; return 0; } void _init(int *ptr, int count) { uint32_t seed = (uint32_t)time(NULL); //make huan happy srand(seed); //reseeding the random generator //filling the buffer with random data for (int i = 0; i < count; i++) { //ptr[i] = rand() % 100000000; ptr[i] = rand() ; if (i % 2 == 0) ptr[i] = 0 - ptr[i] ; } } double get_time() { struct timeval tv; gettimeofday(&tv, NULL); return ((double)tv.tv_usec * 0.000001 + tv.tv_sec); } int main() { //********************************** fprintf(stderr, "filling the buffer with %d elements...\n", N); _init(source, N); //********************************** //Now we are going to kick start your kernel. cudaDeviceSynchronize(); //steady! ready! go! fprintf(stderr, "Running on GPU...\n"); double t0 = get_time(); //_sum_min_or_max<<<BLOCKS, BLOCK_SIZE>>>(source, N, &final_result_max, &final_result_min); _sum_min_or_max<<<BLOCKS, BLOCK_SIZE>>>(source, N,final_result); CHECK(cudaGetLastError()); //checking for launch failures CHECK(cudaDeviceSynchronize()); //checking for run-time failures double t1 = get_time(); //int A = final_result[0]; //fprintf(stderr, " GPU max: %d\n", final_result_max); //fprintf(stderr, " GPU min: %d\n", final_result_min); fprintf(stderr, " GPU max: %d\n", final_result[0]); fprintf(stderr, " GPU min: %d\n", final_result[1]); //********************************** //Now we are going to exercise your CPU... fprintf(stderr, "Running on CPU...\n"); double t2 = get_time(); int cpu_max=0; int cpu_min=0; int B = _max_min_cpu(source, N, &cpu_max, &cpu_min); printf(" CPU max = %d\n", cpu_max); printf(" CPU min = %d\n", cpu_min); double t3 = get_time(); //fprintf(stderr, "CPU sum: %u\n", B); //******The last judgement********** //if ( final_result_max == cpu_max && final_result_min == cpu_min ) if ( final_result[0] == cpu_max && final_result[1] == cpu_min ) { fprintf(stderr, "Test Passed!\n"); } else { fprintf(stderr, "Test failed!\n"); exit(-1); } //****and some timing details******* fprintf(stderr, "GPU time %.3f ms\n", (t1 - t0) * 1000.0); fprintf(stderr, "CPU time %.3f ms\n", (t3 - t2) * 1000.0); return 0; } 分别在定义:__managed__ int final_result[2] = {INT_MIN,INT_MAX}; //scalar output核函数定义:__global__ void _sum_min_or_max(int *input, int count,int *output)核函数操作:if (threadIdx.x == 0) atomicMax(&output[0], max_per_block[0]); if (threadIdx.x == 0) atomicMin(&output[1], min_per_block[0]); 以及核函数调用:_sum_min_or_max<<<BLOCKS, BLOCK_SIZE>>>(source, N,final_result);这几个地方做了改动。开始编译,运行:(Quardo P1000上运行)(Nano上运行)运行没问题,但是貌似GPU运行时间(81ms)比CPU运行时间(22ms)要慢一些。比较在Nano上GPU运行时间(154ms)比CPU运行时间(126ms),好像结果中确实GPU的速度并不占优势。这是什么原因呢?计算包括访存密集型还是计算密集型等类型。无论是加法,还是max/min,都是访存密集的计算。除非独立显卡,且提前预取或者传输数据到显存,否则GPU无论是managed数据自动迁移,或者GPU和CPU一样的享受同样的带宽(Jetson上),都不会占据优势。那么,将过程泛化到怎样的f(a,b)操作,才能让GPU具有显著的优势呢?哪怕是在Jetson这种CPU和GPU有同样的访存带宽,或者哪怕是强制走了慢速的PCI-E传输的带宽,GPU依然能比CPU的运算快得多呢?这个问题,就留给大家思索了!听说阅读 樊哲勇老师的小红书《CUDA 编程:基础与实践》可以找到解决之路哦~~
  • [其他干货] CUDA编程(八)统一内存
    从前面的矩阵乘的代码中可以看出,要写好一个CUDA的代码,需要分配HOST内存(malloc或cudaMallocHost),需要分配DEVICE内存(cudaMalloc),需要将HOST内存数据复制到DEVICE(cudaMemcpy),需要完成GPU核函数的调用,需要把核函数的调用结果在复制回HOST(cudaMemcpy),还需要对前面的各种内存做释放工作(free,cudaFreeHost,cudaFree)。这些工作,虽然是套路,显然还是太繁琐了。于是,聪明的Nvidia在CUDA 6.0以上的版本提出了一个叫做Unified Memory(统一内存)的概念。它把GPU内存、CPU内存在编码层面屏蔽起来:它是可以从系统的任何CPU、GPU访问的单个内存地址空间。它允许应用程序分配可以从CPUs或GPUs上允许的代码读取或者写入数据。具体的方式如下所示:它把原来CPU上的malloc改为cudaMallocManaged,并且分配好的内存地址可以直接被GPU的核函数(图中的 qsort)使用(还记得原来的代码需要先cudaMallocHost/malloc,在cudaMemcpy吗?这里统统不要了。统一内存除了上面使用的 cudaMallocManaged函数来定义变量以外,还可以使用 __managed__ 标识符来表示这是一块统一内存。(这个前面可能还需要再加上 __device__ 标识符供 核函数使用。统一内存使用的时候要借助于 cudaDeviceSynchronize() 来确保CPU和GPU同步。统一内存不显式的区分HOST还是DEVICE的memory,它简化了代码,增强了代码的通用性。统一内存只能在HOST申请。这里面有几个误区需要澄清下:(1)张小白原来以为,只有 Nvidia Jetson Orin那种显存和内存合二为一的设备才有统一内存的概念。但其实并不是——满足 SM架构大于3.0(Kepler架构以上)都可以使用统一内存的方式来编程。逻辑上任何GPU卡或者设备都可以使用统一内存,但是从效果上来看,只有真正的融合为一体的设备(如Jetson AGX Orin),才有最好的统一内存的效果。(2)对于矩阵乘的代码而言,统一内存相当于对Global Memory的一个等效版本,而共享内存则是对SM内部的一种速度优化方式。两者是无关的。也就是说,你在使用统一内存的代码中可以同时使用共享内存。(2)使用了 __managed__ 标识符或 cudaMallocManaged 之后,确实代码中不需要 cudaMalloc,cudaMemcpy这些代码了。但是系统底层其实还会根据情况,决定自己是否需要执行相关的GPU内存分配和 HOST和DEVICE内存的互相拷贝的动作。举个例子,对于张小白的Nvidia Quardo P1000的显卡而言,HOST内存在自己的笔记本内存上(大概有64G),DEVICE内存在GPU显卡(大概有4G)。在这样的环境运行代码,系统仍然会做 申请HOST内存,申请DEVICE内存,HOST内存与DEVICE内存复制等动作。但是对于张小白新购置的了不起的Nvidia AGX Orin而言,HOST内存就是DEVICE内存(大概有32G)。两者不仅仅叫做统一内存,其实还叫做同一内存(张小白自创的)。也就是说,ARM CPU和Nvidia GPU共享一个物理内存。具体的说明可参见:https://zhuanlan.zhihu.com/p/486130961同一内存最大的好处就是:下面典型的三个动作,1、3都可以省略了:所以典型的代码就从左边的模式变成了右边的模式:(1)定义变量:仅需要定义unified memory的变量。节省了空间。(2)HOST->DEVICE:步骤省略(3)执行核函数:跟原来一样(4)DEVICE->HOST:步骤省略(5)显式同步:只是统一内存比原来的方式多一个CPU等待GPU完成的动作。注:上述图片(含代码)来自于上面链接中的文章。那么,统一内存到底是怎么实现的呢?这里借助了下图的做法:CUDA在现有内存池的结构上增加了一个 统一内存系统。开发人员可以直接访问任何内存或者显存资源。当CUDA发现需要访问GPU内存时,如果一开始定义在HOST侧,并且对其进行了初始化,CUDA会自动执行数据拷贝,所以,仍然会受制于PCI-E的带宽和延迟。我们可以看到在这个情况下,代码和运行时变量前后的变迁:好了,概念好像整理得差不多了。下面开始实战:我们把昨天的矩阵乘的代码(包含共享内存优化部分)拿过来,然后看看该怎么优化。原来的代码是这样的:matrix_mul.cuh#pragma once #include <stdio.h> #define CHECK(call) \ do \ { \ const cudaError_t error_code = call; \ if (error_code != cudaSuccess) \ { \ printf("CUDA Error:\n"); \ printf(" File: %s\n", __FILE__); \ printf(" Line: %d\n", __LINE__); \ printf(" Error code: %d\n", error_code); \ printf(" Error text: %s\n", \ cudaGetErrorString(error_code)); \ exit(1); \ } \ } while (0)matrix_mul_old.cu#include <stdio.h> #include <math.h> #include "error.cuh" #include "matrix_mul.cuh" #define BLOCK_SIZE 32 __global__ void gpu_matrix_mult(int *a,int *b, int *c, int m, int n, int k) { int row = blockIdx.y * blockDim.y + threadIdx.y; int col = blockIdx.x * blockDim.x + threadIdx.x; int sum = 0; if( col < k && row < m) { for(int i = 0; i < n; i++) { sum += a[row * n + i] * b[i * k + col]; } c[row * k + col] = sum; } } __global__ void gpu_matrix_mult_shared(int *d_a, int *d_b, int *d_result, int m, int n, int k) { __shared__ int tile_a[BLOCK_SIZE][BLOCK_SIZE]; __shared__ int tile_b[BLOCK_SIZE][BLOCK_SIZE]; int row = blockIdx.y * BLOCK_SIZE + threadIdx.y; int col = blockIdx.x * BLOCK_SIZE + threadIdx.x; int tmp = 0; int idx; for (int sub = 0; sub < gridDim.x; ++sub) { idx = row * n + sub * BLOCK_SIZE + threadIdx.x; tile_a[threadIdx.y][threadIdx.x] = row<n && (sub * BLOCK_SIZE + threadIdx.x)<n? d_a[idx]:0; idx = (sub * BLOCK_SIZE + threadIdx.y) * n + col; tile_b[threadIdx.y][threadIdx.x] = col<n && (sub * BLOCK_SIZE + threadIdx.y)<n? d_b[idx]:0; __syncthreads(); for (int k = 0; k < BLOCK_SIZE; ++k) { tmp += tile_a[threadIdx.y][k] * tile_b[k][threadIdx.x]; } __syncthreads(); } if(row < n && col < n) { d_result[row * n + col] = tmp; } } void cpu_matrix_mult(int *h_a, int *h_b, int *h_result, int m, int n, int k) { for (int i = 0; i < m; ++i) { for (int j = 0; j < k; ++j) { int tmp = 0.0; for (int h = 0; h < n; ++h) { tmp += h_a[i * n + h] * h_b[h * k + j]; } h_result[i * k + j] = tmp; } } } int main(int argc, char const *argv[]) { int m=100; int n=100; int k=100; //声明Event cudaEvent_t start, stop, stop2, stop3 , stop4 ; //创建Event CHECK(cudaEventCreate(&start)); CHECK(cudaEventCreate(&stop)); CHECK(cudaEventCreate(&stop2)); int *h_a, *h_b, *h_c, *h_cc; CHECK(cudaMallocHost((void **) &h_a, sizeof(int)*m*n)); CHECK(cudaMallocHost((void **) &h_b, sizeof(int)*n*k)); CHECK(cudaMallocHost((void **) &h_c, sizeof(int)*m*k)); CHECK(cudaMallocHost((void **) &h_cc, sizeof(int)*m*k)); for (int i = 0; i < m; ++i) { for (int j = 0; j < n; ++j) { h_a[i * n + j] = rand() % 1024; } } for (int i = 0; i < n; ++i) { for (int j = 0; j < k; ++j) { h_b[i * k + j] = rand() % 1024; } } int *d_a, *d_b, *d_c; CHECK(cudaMalloc((void **) &d_a, sizeof(int)*m*n)); CHECK(cudaMalloc((void **) &d_b, sizeof(int)*n*k)); CHECK(cudaMalloc((void **) &d_c, sizeof(int)*m*k)); // copy matrix A and B from host to device memory CHECK(cudaMemcpy(d_a, h_a, sizeof(int)*m*n, cudaMemcpyHostToDevice)); CHECK(cudaMemcpy(d_b, h_b, sizeof(int)*n*k, cudaMemcpyHostToDevice)); unsigned int grid_rows = (m + BLOCK_SIZE - 1) / BLOCK_SIZE; unsigned int grid_cols = (k + BLOCK_SIZE - 1) / BLOCK_SIZE; dim3 dimGrid(grid_cols, grid_rows); dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE); //开始start Event cudaEventRecord(start); //非阻塞模式 cudaEventQuery(start); //gpu_matrix_mult<<<dimGrid, dimBlock>>>(d_a, d_b, d_c, m, n, k); gpu_matrix_mult_shared<<<dimGrid, dimBlock>>>(d_a, d_b, d_c, m, n, k); //开始stop Event cudaEventRecord(stop); //由于要等待核函数执行完毕,所以选择阻塞模式 cudaEventSynchronize(stop); //计算时间 stop-start float elapsed_time; CHECK(cudaEventElapsedTime(&elapsed_time, start, stop)); printf("start-》stop:Time = %g ms.\n", elapsed_time); CHECK(cudaMemcpy(h_c, d_c, (sizeof(int)*m*k), cudaMemcpyDeviceToHost)); //cudaThreadSynchronize(); //开始stop2 Event CHECK(cudaEventRecord(stop2)); //非阻塞模式 //CHECK(cudaEventSynchronize(stop2)); cudaEventQuery(stop2); //计算时间 stop-stop2 float elapsed_time2; cudaEventElapsedTime(&elapsed_time2, stop, stop2); printf("stop-》stop2:Time = %g ms.\n", elapsed_time2); //销毁Event CHECK(cudaEventDestroy(start)); CHECK(cudaEventDestroy(stop)); CHECK(cudaEventDestroy(stop2)); //CPU函数计算 cpu_matrix_mult(h_a, h_b, h_cc, m, n, k); int ok = 1; for (int i = 0; i < m; ++i) { for (int j = 0; j < k; ++j) { if(fabs(h_cc[i*k + j] - h_c[i*k + j])>(1.0e-10)) { ok = 0; } } } if(ok) { printf("Pass!!!\n"); } else { printf("Error!!!\n"); } // free memory cudaFree(d_a); cudaFree(d_b); cudaFree(d_c); cudaFreeHost(h_a); cudaFreeHost(h_b); cudaFreeHost(h_c); return 0; }先执行一下:没啥问题。我们来分析一下:上面的代码用到了 h_a, h_b, h_c, h_cc 4个HOST内存,还用到了 d_a, d_b, d_c 三个DEVICE内存。其中,abc是对应的。而cc是放CPU运算结果专用的。其实我们可以把h_cc直接改为malloc的内存就行了。但是为了好看,也可以将这4个HOST内存都改为统一内存。我们将统一内存起名为 u_a, u_b, u_c, u_cc吧!魔改开始:将代码中 h_a->u_a,h_b->u_b,h_c->u_c,h_cc->u_cc,其他变量做相应的适当修改。matrix_mul.cu#include <stdio.h> #include <math.h> #include "error.cuh" #include "matrix_mul.cuh" #define BLOCK_SIZE 32 __managed__ int u_a[100*100]; __managed__ int u_b[100*100]; __managed__ int u_c[100*100]; __managed__ int u_cc[100*100]; __global__ void gpu_matrix_mult(int *a,int *b, int *c, int m, int n, int k) { int row = blockIdx.y * blockDim.y + threadIdx.y; int col = blockIdx.x * blockDim.x + threadIdx.x; int sum = 0; if( col < k && row < m) { for(int i = 0; i < n; i++) { sum += u_a[row * n + i] * u_b[i * k + col]; } u_c[row * k + col] = sum; } } __global__ void gpu_matrix_mult_shared(int *u_a, int *u_b, int *u_result, int m, int n, int k) { __shared__ int tile_a[BLOCK_SIZE][BLOCK_SIZE]; __shared__ int tile_b[BLOCK_SIZE][BLOCK_SIZE]; int row = blockIdx.y * BLOCK_SIZE + threadIdx.y; int col = blockIdx.x * BLOCK_SIZE + threadIdx.x; int tmp = 0; int idx; for (int sub = 0; sub < gridDim.x; ++sub) { idx = row * n + sub * BLOCK_SIZE + threadIdx.x; tile_a[threadIdx.y][threadIdx.x] = row<n && (sub * BLOCK_SIZE + threadIdx.x)<n? u_a[idx]:0; idx = (sub * BLOCK_SIZE + threadIdx.y) * n + col; tile_b[threadIdx.y][threadIdx.x] = col<n && (sub * BLOCK_SIZE + threadIdx.y)<n? u_b[idx]:0; __syncthreads(); for (int k = 0; k < BLOCK_SIZE; ++k) { tmp += tile_a[threadIdx.y][k] * tile_b[k][threadIdx.x]; } __syncthreads(); } if(row < n && col < n) { u_result[row * n + col] = tmp; } } void cpu_matrix_mult(int *u_a, int *u_b, int *u_result, int m, int n, int k) { for (int i = 0; i < m; ++i) { for (int j = 0; j < k; ++j) { int tmp = 0.0; for (int h = 0; h < n; ++h) { tmp += u_a[i * n + h] * u_b[h * k + j]; } u_result[i * k + j] = tmp; } } } int main(int argc, char const *argv[]) { int m=100; int n=100; int k=100; //声明Event cudaEvent_t start, stop, stop2, stop3 , stop4 ; //创建Event CHECK(cudaEventCreate(&start)); CHECK(cudaEventCreate(&stop)); CHECK(cudaEventCreate(&stop2)); //int *h_a, *h_b, *h_c, *h_cc; //CHECK(cudaMallocHost((void **) &h_a, sizeof(int)*m*n)); //CHECK(cudaMallocHost((void **) &h_b, sizeof(int)*n*k)); //CHECK(cudaMallocHost((void **) &h_c, sizeof(int)*m*k)); //CHECK(cudaMallocHost((void **) &h_cc, sizeof(int)*m*k)); for (int i = 0; i < m; ++i) { for (int j = 0; j < n; ++j) { u_a[i * n + j] = rand() % 1024; } } for (int i = 0; i < n; ++i) { for (int j = 0; j < k; ++j) { u_b[i * k + j] = rand() % 1024; } } //int *d_a, *d_b, *d_c; //CHECK(cudaMalloc((void **) &d_a, sizeof(int)*m*n)); //CHECK(cudaMalloc((void **) &d_b, sizeof(int)*n*k)); //CHECK(cudaMalloc((void **) &d_c, sizeof(int)*m*k)); // copy matrix A and B from host to device memory //CHECK(cudaMemcpy(d_a, h_a, sizeof(int)*m*n, cudaMemcpyHostToDevice)); //CHECK(cudaMemcpy(d_b, h_b, sizeof(int)*n*k, cudaMemcpyHostToDevice)); unsigned int grid_rows = (m + BLOCK_SIZE - 1) / BLOCK_SIZE; unsigned int grid_cols = (k + BLOCK_SIZE - 1) / BLOCK_SIZE; dim3 dimGrid(grid_cols, grid_rows); dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE); //开始start Event cudaEventRecord(start); //非阻塞模式 cudaEventQuery(start); //gpu_matrix_mult<<<dimGrid, dimBlock>>>(d_a, d_b, d_c, m, n, k); gpu_matrix_mult_shared<<<dimGrid, dimBlock>>>(u_a, u_b, u_c, m, n, k); //开始stop Event cudaEventRecord(stop); //由于要等待核函数执行完毕,所以选择阻塞模式 cudaEventSynchronize(stop); //计算时间 stop-start float elapsed_time; CHECK(cudaEventElapsedTime(&elapsed_time, start, stop)); printf("start-》stop:Time = %g ms.\n", elapsed_time); //CHECK(cudaMemcpy(h_c, d_c, (sizeof(int)*m*k), cudaMemcpyDeviceToHost)); //cudaThreadSynchronize(); //开始stop2 Event CHECK(cudaEventRecord(stop2)); //非阻塞模式 //CHECK(cudaEventSynchronize(stop2)); cudaEventQuery(stop2); //计算时间 stop-stop2 float elapsed_time2; cudaEventElapsedTime(&elapsed_time2, stop, stop2); printf("stop-》stop2:Time = %g ms.\n", elapsed_time2); //销毁Event CHECK(cudaEventDestroy(start)); CHECK(cudaEventDestroy(stop)); CHECK(cudaEventDestroy(stop2)); //CPU函数计算 cpu_matrix_mult(u_a, u_b, u_cc, m, n, k); int ok = 1; for (int i = 0; i < m; ++i) { for (int j = 0; j < k; ++j) { if(fabs(u_cc[i*k + j] - u_c[i*k + j])>(1.0e-10)) { ok = 0; } } } if(ok) { printf("Pass!!!\n"); } else { printf("Error!!!\n"); } // free memory //cudaFree(d_a); //cudaFree(d_b); //cudaFree(d_c); //cudaFreeHost(h_a); //cudaFreeHost(h_b); //cudaFreeHost(h_c); return 0; }执行一下:额,好像速度没啥变化。查看下性能:
  • [调试调优] 使用mindspore复现segmenter时,在GPU上速度精度均正常,但是在Ascend910上推理速度出奇的慢
    【功能模块】mindspore复现语义分割算法segmenter【操作步骤&问题现象】1、使用mindspore复现segmenter时,在GPU上速度精度均正常,但是在Ascend910上推理速度出奇的慢, 一张图片要几十秒,非常不正常请专业人员帮忙解决一下,可以提供代码
  • [安装] 【mindelec模块】【安装后代码测试问题】
    【功能模块】wsl下conda环境中安装的mindelec【操作步骤&问题现象】1、在wsl中完成安装mindspore-gpu,并完成安装测试2、利用wget https://ms-release.obs.cn-north-4.myhuaweicloud.com/1.5.0/MindScience/x86_64/mindscience_mindelec_ascend-0.1.0-cp37-cp37m-linux_x86_64.whl  下载然后pip安装mindelec模块3.python -c “import mindelec”测试通过4.进行样例代码测试时报错如下。【截图信息】进行电磁频域solve.py代码测试时,更改Ascend为GPU代码报错如下:现实确认GPU显存和内存均能够满足需求的情况下仍然报错。希望官方可以将mindelec利用gpu进行测试,有利于致力于开发mindspore电磁计算工作平台的工作者更好开发。或者可以放出一个经过测试的mindelec的gpu版本的whl文件便于开发者安装。十分感谢蹲个官方答复。【日志信息】(可选,上传日志内容或者附件)
  • [问题求助] 【MindElec的GPU编译版本】【MindElec源码编译过程中出现问题】Could NOT find Python3
    【功能模块】电磁计算平台MindElec【操作步骤&问题现象】需求:在RTX3060 Gpu上运行mindelec电磁计算平台,由于官方只支持ascend,只能用源码编译的方式1、在wsl中完成minds pore1.7-gpu版本的安装及测试后,尝试在Ubuntu-20.04的wsl中利用源码编译的方式安装mindelec敲出命令 bash build.sh -e gpu 出现报错-- Could NOT find Python3 (missing: Python3_INCLUDE_DIRS Development) (found version "3.8.10")CMake Error at CMakeLists.txt:35 (find_python_package):  Unknown CMake command "find_python_package".-- Configuring incomplete, errors occurred!See also "/home/lihongji/mindscience/MindElec/build/mindelec/CMakeFiles/CMakeOutput.log".2、详细操作信息:(base) lihongji@LAPTOP-P7NAUT00:~/mindscience/MindElec$ bash build.sh -e gpumkdir: created directory '/home/lihongji/mindscience/MindElec/output'---------------- MindElec: build start ----------------start build mindelec project.-DDEBUG_MODE=off -DBUILD_PATH=/home/lihongji/mindscience/MindElec/build/ -DENABLE_GPU=ONMD LIBS CACHE PATH:  /home/lihongji/mindscience/MindElec/build/mindelec/.mdlibset make thread num: 8pkg name:pybind11,pybind11pybind11 config hash: 5253018179879c380bb2d33c476f2438-- Found pybind11: /home/lihongji/mindscience/MindElec/build/mindelec/.mdlib/pybind11_5253018179879c380bb2d33c476f2438/include (found version "2.6.1" )Found pkg: pybind11pkg name:glog,glogglog config hash: 6f581a3db68ffe9e125bd4c496da9aaa_FIND:/home/lihongji/mindscience/MindElec/build/mindelec/.mdlib/glog_6f581a3db68ffe9e125bd4c496da9aaafound /home/lihongji/mindscience/MindElec/build/mindelec/.mdlib/glog_6f581a3db68ffe9e125bd4c496da9aaa/lib/libglog.soFound libs: glog::glogpkg name:nlohmann_json,nlohmann_jsonnlohmann_json config hash: e0a3765b49b7f4747a26bc3db79d0028========== External libs built successfully ==========-- Could NOT find Python3 (missing: Python3_INCLUDE_DIRS Development) (found version "3.8.10")CMake Error at CMakeLists.txt:35 (find_python_package):  Unknown CMake command "find_python_package".-- Configuring incomplete, errors occurred!See also "/home/lihongji/mindscience/MindElec/build/mindelec/CMakeFiles/CMakeOutput.log".之前已经完成过一部分的编译,重新编译只显示该部分信息【截图信息】【日志信息】(可选,上传日志内容或者附件)
  • [问题求助] 【华为云GPU产品】【tensorflow的版本】P型弹性云服务器具体支持tensorflow的版本
    【功能模块】 请问P型弹性云服务器具体支持tensorflow哪个版本呢?因为tensorflow1.5的版本和tensorflow2版本相差还是比较大的,不同项目需要不同的tensorflow的版本
  • [问题求助] 【众智产品】【npu迁移问题】1p npu性能与gpu相同,8p差了四倍
    【问题现象】1、1p性能达到gpu效果,但是8p npu差距很大2、大多数Npu的利用率为0【截图信息】【日志信息】(可选,上传日志内容或者附件)
  • [算子编译] 请问Graph Kernel Fusion(图算融合)在mindspore1.7.0下会生成融合后的mindIR的.dot文件吗
    【功能模块】图算融合,GPU (NVIDIA-RTX3080) 验证【操作步骤&问题现象】1、参考(基于mindspore0.5.0)链接1: https://gitee.com/mindspore/course/tree/master/06_distributed/graph_kernel2、参考(基于mindspore1.0.0)链接2: https://bbs.huaweicloud.com/forum/forum.php?mod=viewthread&tid=78817&page=1&replytype=13、按照教程和另一位大佬的帖子开启图算融合后,对基本组合算子和自组合算子进行试验,发现mindspore1.7.0生成的ir和.dot文件和之前版本的mindir文件的命名形式不一样,发现没有生成例子当中的文件(hwopt_d_fuse_basic_opt_end_graph_0.dot/hwopt_d_composite_opt_end_graph_0.dot)。未找到如下图红框中类似的图算融合后生成的ir/dot文件.minspore1.0.0生成的ir文件如下图(源自链接2楼主)请问是否有图算融合的整体流程图(细化到生成每个IR的小阶段)参考,多谢!【截图信息】本人基于mindspore1.7.0生成的mindir文件信息中,未找到图算融合后的.dot。烦请指导~多谢【日志信息】(可选,上传日志内容或者附件)
  • [问题求助] 【Ascend 910】【模型训练】代码可以跑通,但是npu计算速度比1080的gpu还要慢几十倍
    【功能模块】模型更新参数慢【操作步骤&问题现象】1、模型的参数量不大,仅2.6M2、一个batch,在npu上耗时2分07秒,然而在1080的卡上,一个batch总共也只需要1秒左右3、更新一次参数就需要耗时30s,这部分代码应该还比较标准,就这三行【截图信息】【日志信息】(可选,上传日志内容或者附件)2022-07-18 22:23:50.891451 ,idx: 12022-07-18 22:23:51.057450 hpn_learner2022-07-18 22:23:51.057613 Squeezing building blocks2022-07-18 22:24:49.295080 Mixing building blocks2022-07-18 22:25:07.476452 Decode the encoded 4D-tensor2022-07-18 22:25:07.484426 ,update model parameters2022-07-18 22:25:37.292860 ,Evaluate classify_prediction2022-07-18 22:25:57.887156 ,Evaluate update2022-07-18 22:25:57.888581 ,Evaluate write_process2022-07-18 22:25:57.889203 ,idx: 2
  • [安装] RTX3060 Laptop GPU,CUDA11.6.134,win11要怎么安装昇思GPU版?
    如题
  • [其他干货] CUDA编程(四)Global Memory
    在GPU上,on-board memory包含以下类型:local memory 每个thread一个。线程私有。global memory 每个grid一个。每个thread都可以读。constant memory 每个grid一个。只读。每个thread都可以读。texture memory 每个grid一个。只读。每个thread都可以读。on-chip memory包含以下类型:registers 每个thread一个。线程私有。shared memory 每个block一个,一个block下所有线程都可以访问。HOST内存函数malloc 申请memset 初始化free 释放DEVICE内存函数cudaMalloc 申请cudaMemset 初始化cudaFree 释放请注意,这里函数只返回状态。所以分配的内存地址作为函数参数。HOST《-》DEVICE互相拷贝cudaMemcpy( 目的内存地址,源内存地址,内存大小,cudaMemcpyHostToDevice/cudaMemcpyDeviceToHost/cudaMemcpyDeviceToDevice/cudaMemcpyHostToHost)以矩阵乘为例:CPU的做法是嵌套循环,如上图所示。GPU的做法应该是使用 index( blockIdx和 threadIdx的组合公式)替换原来的下标i,j。这也是一般CUDA程序的套路——把for loop展开成每个线程处理其中的一步。那么,如何使用CUDA将坐标拆开呢?将二维坐标(矩阵)改为 在全局中的索引:需要找到每个线程需要处理元素的位置。ty=线程在y方向的坐标tx=线程在x方向的坐标ty=blockIdx.y*blockDim.y + threadIdx.ytx=blockIdx.x*blockDim.x + threadIdx.xnx=x方向有多少数据。index = ty * nx + tx目的是将高维降为低维。矩阵乘的每个核函数的算法如下:典型的核函数算法代码如下:需要注意:矩阵乘 矩阵M是 mXn,矩阵N是 nXk,这里面需要 矩阵M和矩阵N都有n。否则无法相乘。上代码:matrix_mul.cu#include <stdio.h> #include <math.h> #define BLOCK_SIZE 16 //使用GPU进行矩阵计算 __global__ void gpu_matrix_mult(int *a,int *b, int *c, int m, int n, int k) { int row = blockIdx.y * blockDim.y + threadIdx.y; int col = blockIdx.x * blockDim.x + threadIdx.x; int sum = 0; if( col < k && row < m) { for(int i = 0; i < n; i++) { sum += a[row * n + i] * b[i * k + col]; } c[row * k + col] = sum; } } //使用CPU进行矩阵计算 void cpu_matrix_mult(int *h_a, int *h_b, int *h_result, int m, int n, int k) { for (int i = 0; i < m; ++i) { for (int j = 0; j < k; ++j) { int tmp = 0.0; for (int h = 0; h < n; ++h) { tmp += h_a[i * n + h] * h_b[h * k + j]; } h_result[i * k + j] = tmp; } } } int main(int argc, char const *argv[]) { /* 矩阵A mXn,矩阵B nXk --》矩阵乘计算的结果是 mXk */ int m=3; int n=4; int k=5; int *h_a, *h_b, *h_c, *h_cc; //分配原矩阵的内存 h是host memory cudaMallocHost((void **) &h_a, sizeof(int)*m*n); cudaMallocHost((void **) &h_b, sizeof(int)*n*k); //分配 CPU结果内存 cudaMallocHost((void **) &h_c, sizeof(int)*m*k); //分配 GPU结果内存 cudaMallocHost((void **) &h_cc, sizeof(int)*m*k); //初始化矩阵A(mxn) srand(time(0)); printf("---------------h_a------------------\n"); for (int i = 0; i < m; ++i) { for (int j = 0; j < n; ++j) { h_a[i * n + j] = rand() % 1024; printf("%d", h_a[i * n + j] ); printf(" "); } printf("\n"); } //初始化矩阵B(nxk) printf("---------------h_b------------------\n"); for (int i = 0; i < n; ++i) { for (int j = 0; j < k; ++j) { h_b[i * k + j] = rand() % 1024; printf("%d", h_b[i * k + j] ); printf(" "); } printf("\n"); } int *d_a, *d_b, *d_c; //分配 原矩阵的GPU内存 d是device memory cudaMalloc((void **) &d_a, sizeof(int)*m*n); cudaMalloc((void **) &d_b, sizeof(int)*n*k); //分配 目的矩阵的GPU内存 cudaMalloc((void **) &d_c, sizeof(int)*m*k); // copy matrix A and B from host to device memory cudaMemcpy(d_a, h_a, sizeof(int)*m*n, cudaMemcpyHostToDevice); cudaMemcpy(d_b, h_b, sizeof(int)*n*k, cudaMemcpyHostToDevice); unsigned int grid_rows = (m + BLOCK_SIZE - 1) / BLOCK_SIZE; unsigned int grid_cols = (k + BLOCK_SIZE - 1) / BLOCK_SIZE; dim3 dimGrid(grid_cols, grid_rows); dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE); //GPU计算,结果放入h_c gpu_matrix_mult<<<dimGrid, dimBlock>>>(d_a, d_b, d_c, m, n, k); cudaMemcpy(h_c, d_c, sizeof(int)*m*k, cudaMemcpyDeviceToHost); //cudaThreadSynchronize(); //CPU计算,结果直接放入h_cc cpu_matrix_mult(h_a, h_b, h_cc, m, n, k); int ok = 1; for (int i = 0; i < m; ++i) { for (int j = 0; j < k; ++j) { // 比较大小的时候使用 a-b<0.0000000001 if(fabs(h_cc[i*k + j] - h_c[i*k + j])>(1.0e-10)) { ok = 0; } } } printf("---------------h_c cpu result------------------\n"); for(int i=0;i<m;i++) { for(int j=0;j<k;j++) { //矩阵小的时候还可以打印,大的时候就别打了 printf("%d",h_c[i*k + j] ); printf(" "); } printf("\n"); } printf("---------------h_cc gpu result----------------\n"); for(int i=0;i<m;i++) { for(int j=0;j<k;j++) { //矩阵小的时候还可以打印,大的时候就别打了 printf("%d",h_cc[i*k + j] ); printf(" "); } printf("\n"); } if(ok) { printf("Pass!!!\n"); } else { printf("Error!!!\n"); } // free memory cudaFree(d_a); cudaFree(d_b); cudaFree(d_c); cudaFreeHost(h_a); cudaFreeHost(h_b); cudaFreeHost(h_c); return 0; }代码中张小白加上了注释,已经介绍得比较清楚了。我们执行下看看:代码以 3X4和4X5的矩阵相乘,得到了3X5的矩阵结果。这个结果跟CPU计算的结果做了对比。显示Pass表示结果是一致的(其实张小白把两个结果都打印的出来,当然也是一致的)这里面有个小TIPS,就是在调用rand()生成随机数的时候,可以使用srand(time(0)) 做随机数种子,这样下次调用的时候跟这次生成的内容就会不一样。如果去掉这句话,每次执行的结果都是一样的。当然,如果在同一秒同时执行,srand(time(0)) 也会导致同时生成的随机数是一样的。这点需要注意。
  • [其他干货] CUDA编程(二)CUDA代码执行原理
    我们把 CPU,内存这块区域叫做“主机(HOST)”,把GPU,显存这块区域叫做“设备(DEVICE)”。(是不是跟昇腾有点类似?)CUDA的代码执行包含以下几步:简述一下,就是 host_to_device-》在device上并行计算-》device_to_host。cuda程序其实是一个对C的扩展程序。其后缀名为.cu,如果头文件则为.cuh。这个.cu 程序除了C程序的语法外,还有一些cuda的特有部分,比如它在函数前面加了前缀,分为 __global__, __host__,__device__ 三种。对于__global__:所谓的“执行配置”,下面会提到,比如说是 <<< >>>中间的内容。这个标识符将一个C函数声明成一个 核函数。它只能在设备(device)上执行。对于__host__:对于__device_:这几个前缀定义了这些代码运行的设备,这会让程序决定在哪个设备上运行。对于一个简单的Hello World代码而言:#include <stdio.h> void hello_from_cpu() { printf("Hello World from the CPU!\n"); } int main(void) { hello_from_cpu(); return 0;如果我们想让它在GPU上运行,仅需要做两步:(1)将被调用的函数 hello_from_cpu 改为 hello_from_gpu ,前面加上 __global__ 将其定义为核函数。(2)在main主函数调用的时候,加上执行配置<<< >>>部分,如加上<<<1,1>>>则为并行1次,如加上<<<2,4>>>则运行 2X4次。我们看看实际代码修改后的效果:#include <stdio.h> __global__ void hello_from_gpu() { printf("Hello World from the GPU!\n"); } int main(void) { hello_from_gpu<<<1,1>>>(); return 0; }cu代码必须使用nvcc编译,编译的时候要根据GPU架构的不同填不同的参数。其中,arch参数如下:code参数如下:举个简单的例子,张小白这台笔记本的显卡是Quardo P1000,是Pascal架构,那么参数是compute_61和sm_61.我们执行以下语句:/usr/local/cuda/bin/nvcc -arch=compute_61 -code=sm_61 hello_cuda.cu -o hello_cuda./hello_cuda如果将执行配置改为2,4:可以发现这个核函数被执行了8次。
总条数:340 到第
上滑加载中