CUDA学习笔记（九）Dynamic Parallelism

news2024/12/29 15:32:21

本篇博文转载于https://www.cnblogs.com/1024incn/tag/CUDA/，仅用于学习。

Dynamic Parallelism

到目前为止，所有kernel都是在host端调用，CUDA Dynamic Parallelism允许GPU kernel在device端创建调用。Dynamic Parallelism使递归更容易实现和理解，由于启动的配置可以由device上的thread在运行时决定，这也减少了host和device之间传递数据和执行控制。我们接下来会分析理解使用Dynamic Parallelism。

Nested Execution

在host调用kernel和在device调用kernel的语法完全一样。kernel的执行则被分为两种类型：parent和child。一个parent thread，parent block或者parent grid可以启动一个新的grid，即child grid。child grid必须在parent 之前完成，也就是说，parent必须等待所有child完成。

当parent启动一个child grid时，在parent显式调用synchronize之前，child不保证会开始执行。parent和child共享同一个global和constant memory，但是有不同的shared 和local memory。不难理解的是，只有两个时刻可以保证child和parent见到的global memory完全一致：child刚开始和child完成。所有parent对global memory的操作对child都是可见的，而child对global memory的操作只有在parent进行synchronize操作后对parent才是可见的。

Nested Hello World on the GPU

为了更清晰的讲解Dynamic Parallelism，我们改编最开始写的hello world程序。下图显示了使用Dynamic Parallelism的执行过程，host调用parent grid（每个block八个thread）。thread 0调用一个child grid（每个block四个thread），thread 0 的第一个thread又调用一个child grid（每个block两个thread），依次类推。

下面是具体的代码，每个thread会先打印出Hello World；然后，每个thread再检查自己是否该停止。

__global__ void nestedHelloWorld(int const iSize,int iDepth) {
    int tid = threadIdx.x；
    printf("Recursion=%d: Hello World from thread %d block %d\n",iDepth,tid,blockIdx.x);
    // condition to stop recursive execution
    if (iSize == 1) return;
    // reduce block size to half
    int nthreads = iSize>>1;
    // thread 0 launches child grid recursively
    if(tid == 0 && nthreads > 0) {
        nestedHelloWorld<<<1, nthreads>>>(nthreads,++iDepth);
        printf("-------> nested execution depth: %d\n",iDepth);
    }
}

编译：

$ nvcc -arch=sm_35 -rdc=true nestedHelloWorld.cu -o nestedHelloWorld -lcudadevrt

-lcudadevrt是用来连接runtime库的，跟gcc连接库一样。-rdc=true使device代码可重入，这是DynamicParallelism所必须的，至于原因则将是一个比较大的话题，以后探讨。

代码的输出为：

./nestedHelloWorld Execution Configuration: grid 1 block 8
Recursion=0: Hello World from thread 0 block 0
Recursion=0: Hello World from thread 1 block 0
Recursion=0: Hello World from thread 2 block 0
Recursion=0: Hello World from thread 3 block 0
Recursion=0: Hello World from thread 4 block 0
Recursion=0: Hello World from thread 5 block 0
Recursion=0: Hello World from thread 6 block 0
Recursion=0: Hello World from thread 7 block 0
-------> nested execution depth: 1
Recursion=1: Hello World from thread 0 block 0
Recursion=1: Hello World from thread 1 block 0
Recursion=1: Hello World from thread 2 block 0
Recursion=1: Hello World from thread 3 block 0
-------> nested execution depth: 2
Recursion=2: Hello World from thread 0 block 0
Recursion=2: Hello World from thread 1 block 0
-------> nested execution depth: 3
Recursion=3: Hello World from thread 0 block 0

这里的01234….输出顺序挺诡异的，太规整了，我们暂且认为CUDA对printf做过修改吧。还有就是，按照CPU递归程序的经验，这里的输出顺序就更怪了，当然，肯定不是编译器错误或者CUDA的bug，大家可以在调用kernel后边加上cudaDeviceSynchronize，就可以看到“正常”的顺序了，原因也就清楚了。

使用nvvp可以查看执行情况，空白说明parent在等待child执行结束：

$nvvp ./nesttedHelloWorld

接着，我们尝试使用两个block而不是一个：

$ ./nestedHelloWorld 2

输出是：

./nestedHelloWorld 2Execution Configuration: grid 2 block 8
Recursion=0: Hello World from thread 0 block 1
Recursion=0: Hello World from thread 1 block 1
Recursion=0: Hello World from thread 2 block 1
Recursion=0: Hello World from thread 3 block 1
Recursion=0: Hello World from thread 4 block 1
Recursion=0: Hello World from thread 5 block 1
Recursion=0: Hello World from thread 6 block 1
Recursion=0: Hello World from thread 7 block 1
Recursion=0: Hello World from thread 0 block 0
Recursion=0: Hello World from thread 1 block 0
Recursion=0: Hello World from thread 2 block 0
Recursion=0: Hello World from thread 3 block 0
Recursion=0: Hello World from thread 4 block 0
Recursion=0: Hello World from thread 5 block 0
Recursion=0: Hello World from thread 6 block 0
Recursion=0: Hello World from thread 7 block 0
-------> nested execution depth: 1
-------> nested execution depth: 1
Recursion=1: Hello World from thread 0 block 0
Recursion=1: Hello World from thread 1 block 0
Recursion=1: Hello World from thread 2 block 0
Recursion=1: Hello World from thread 3 block 0
Recursion=1: Hello World from thread 0 block 0
Recursion=1: Hello World from thread 1 block 0
Recursion=1: Hello World from thread 2 block 0
Recursion=1: Hello World from thread 3 block 0
-------> nested execution depth: 2
-------> nested execution depth: 2
Recursion=2: Hello World from thread 0 block 0
Recursion=2: Hello World from thread 1 block 0
Recursion=2: Hello World from thread 0 block 0
Recursion=2: Hello World from thread 1 block 0
-------> nested execution depth: 3
-------> nested execution depth: 3
Recursion=3: Hello World from thread 0 block 0
Recursion=3: Hello World from thread 0 block 0

从上面结果来看，首先应该注意到，所有child的block的id都是0。下图是调用过程，parent有两个block了，但是所有child都只有一个blcok：

nestedHelloWorld<<<1, nthreads>>>(nthreads, ++iDepth);

注意：Dynamic Parallelism只有在CC3.5以上才被支持。通过Dynamic Parallelism调用的kernel不能执行于不同的device（物理上实际存在的）上。调用的最大深度是24，但实际情况是，kernel要受限于memory资源，其中包括为了同步parent和child而需要的额外的memory资源。

Nested Reduction

学过算法导论之类的算法书应该知道，因为递归比较消耗资源的，所以如果可以的话最好是展开，而这里要讲的恰恰相反，我们要实现递归，这部分主要就是再次证明DynamicParallelism的好处，有了它就可以实现像C那样写递归代码了。

下面的代码就是一份实现，和之前一样，每个child的有一个block，block中第一个thread调用kernel，不同的是，parent的grid有很多的block。第一步还是讲global memory的地址g_idata转化为每个block本地地址。然后，if判断是否该退出，退出的话，就将结果拷贝回global memory。如果不该退出，就进行本地reduction，一般的线程执行in-place（就地）reduction，然后，同步block来保证所有部分和的计算。thread0再次产生一个只有一个block和当前一半数量thread的child grid。

__global__ void gpuRecursiveReduce (int *g_idata, int *g_odata,
unsigned int isize) {
// set thread ID
unsigned int tid = threadIdx.x;
// convert global data pointer to the local pointer of this block
int *idata = g_idata + blockIdx.x*blockDim.x;
int *odata = &g_odata[blockIdx.x];
// stop condition
if (isize == 2 && tid == 0) {
g_odata[blockIdx.x] = idata[0]+idata[1];
return;
}
// nested invocation
int istride = isize>>1;
if(istride > 1 && tid < istride) {
// in place reduction
idata[tid] += idata[tid + istride];
}
// sync at block level
__syncthreads();
// nested invocation to generate child grids
if(tid==0) {
gpuRecursiveReduce <<<1, istride>>>(idata,odata,istride);
// sync all child grids launched in this block
cudaDeviceSynchronize();
}
// sync at block level again
__syncthreads();
}

编译运行，下面结果是运行在Kepler K40上面：

$ nvcc -arch=sm_35 -rdc=true nestedReduce.cu -o nestedReduce -lcudadevrt
./nestedReduce starting reduction at device 0: Tesla K40c
array 1048576 grid 2048 block 512
cpu reduce elapsed 0.000689 sec cpu_sum: 1048576
gpu Neighbored elapsed 0.000532 sec gpu_sum: 1048576<<<grid 2048 block 512>>>
gpu nested elapsed 0.172036 sec gpu_sum: 1048576<<<grid 2048 block 512>>>

相较于neighbored，nested的结果是非常差的。

从上面结果看，2048个block被初始化了。每个block执行了8个recursion，16384个child block被创建，__syncthreads也被调用了16384次。这都是导致效率很低的原因。

当一个child grid被调用后，他看到的memory是和parent完全一样的，因为child只需要parent的一部分数据，block在每个child grid的启动前的同步操作是不必要的，修改后：

__global__ void gpuRecursiveReduceNosync (int *g_idata, int *g_odata,unsigned int isize) {
// set thread ID
unsigned int tid = threadIdx.x;
// convert global data pointer to the local pointer of this block
int *idata = g_idata + blockIdx.x * blockDim.x;
int *odata = &g_odata[blockIdx.x];
// stop condition
if (isize == 2 && tid == 0) {
g_odata[blockIdx.x] = idata[0] + idata[1];
return;
}
// nested invoke
int istride = isize>>1;
if(istride > 1 && tid < istride) {
idata[tid] += idata[tid + istride];
if(tid==0) {
gpuRecursiveReduceNosync<<<1, istride>>>(idata,odata,istride);
}
}
}

运行输出，时间减少到原来的三分之一：

./nestedReduceNoSync starting reduction at device 0: Tesla K40c
array 1048576 grid 2048 block 512
cpu reduce elapsed 0.000689 sec cpu_sum: 1048576
gpu Neighbored elapsed 0.000532 sec gpu_sum: 1048576<<<grid 2048 block 512>>>
gpu nested elapsed 0.172036 sec gpu_sum: 1048576<<<grid 2048 block 512>>>
gpu nestedNosyn elapsed 0.059125 sec gpu_sum: 1048576<<<grid 2048 block 512>>>

不过，性能还是比neighbour-paired要慢。接下来在做点改动，主要想法如下图所示，kernel的调用增加了一个参数iDim，这是因为每次递归调用，child block的大小就减半，parent 的blockDim必须传递给child grid，从而使每个thread都能计算正确的global memory偏移地址。注意，所有空闲的thread都被移除了。相较于之前的实现，每次都会有一半的thread空闲下来而被移除，也就释放了一半的计算资源。

__global__ void gpuRecursiveReduce2(int *g_idata, int *g_odata, int iStride,int const iDim) {
// convert global data pointer to the local pointer of this block
int *idata = g_idata + blockIdx.x*iDim;
// stop condition
if (iStride == 1 && threadIdx.x == 0) {
g_odata[blockIdx.x] = idata[0]+idata[1];
return;
}
// in place reduction
idata[threadIdx.x] += idata[threadIdx.x + iStride];
// nested invocation to generate child grids
if(threadIdx.x == 0 && blockIdx.x == 0) {
gpuRecursiveReduce2 <<<gridDim.x,iStride/2>>>(
g_idata,g_odata,iStride/2,iDim);
}
}

编译运行：

./nestedReduce2 starting reduction at device 0: Tesla K40c
array 1048576 grid 2048 block 512
cpu reduce elapsed 0.000689 sec cpu_sum: 1048576
gpu Neighbored elapsed 0.000532 sec gpu_sum: 1048576<<<grid 2048 block 512>>>
gpu nested elapsed 0.172036 sec gpu_sum: 1048576<<<grid 2048 block 512>>>
gpu nestedNosyn elapsed 0.059125 sec gpu_sum: 1048576<<<grid 2048 block 512>>>
gpu nested2 elapsed 0.000797 sec gpu_sum: 1048576<<<grid 2048 block 512>>>

从这个结果看，数据又好看了不少，可以猜测，大约是由于调用了较少的child grid，我们可以用nvprof来验证下：

$ nvprof ./nestedReduce2

部分输出结果如下，第二列上显示了dievice kernel 的调用次数，第一个和第二个创建了16384个child grid。gpuRecursiveReduce2八层nested Parallelism只创建了8个child。

Calls (host) Calls (device) Avg Min Max Name
1 16384 441.48us 2.3360us 171.34ms gpuRecursiveReduce
1 16384 51.140us 2.2080us 57.906ms gpuRecursiveReduceNosync
1 8 56.195us 22.048us 100.74us gpuRecursiveReduce2
1 0 352.67us 352.67us 352.67us reduceNeighbored

对于一个给定的算法，我们可以有很多种实现方式，避免大量的nested 调用可以提升很多性能。同步对算法的正确性至关重要，但也是一个消耗比较大的操作，block内部的同步操作倒是可以去掉。因为在device上运行nested程序需要额外的资源，nested调用是有限的。

本文来自互联网用户投稿，该文观点仅代表作者本人，不代表本站立场。本站仅提供信息存储空间服务，不拥有所有权，不承担相关法律责任。如若转载，请注明出处：http://www.coloradmin.cn/o/1123400.html

如若内容造成侵权/违法违规/事实不符，请联系多彩编程网进行投诉反馈，一经查实，立即删除！

相关文章

信钰证券：长江电力180亿市值，招商证券、摩根大通等浮盈超一成

信钰证券：长江电力180亿市值，招商证券、摩根大通等浮盈超一成

本周A股限售股解禁规划环比有所上升。 Wind数据核算闪现&#xff0c;除去新上市公司&#xff0c;本周共有64家公司限售股解禁&#xff0c;解禁数量51.52亿股&#xff0c;以最新收盘价核算&#xff08;下同&#xff09;&#xff0c;解禁市值776.21亿元。 本周解禁市值跨越10亿…
阅读更多...
C++二分算法的应用：寻找峰值原理、源码及测试用例

C++二分算法的应用：寻找峰值原理、源码及测试用例

说明 此文是课程https://edu.csdn.net/course/detail/38771 的讲义。 源码下载&#xff1a;https://download.csdn.net/download/he_zhidan/88458478 题目 长度为n的数组nums&#xff0c;请返回任意一峰值的索引。符合以下条件之一i便是峰值的索引。 n等于1 i等于0 n>…
阅读更多...
构建并训练单层神经网络模型

构建并训练单层神经网络模型

1. 单层神经网络 单层神经网络由一个层次中的多个神经元组成，总体来看，单层神经网络的结构是：一个输入层，一个隐含层，一个输出层。 图中描述的是前向网络，但其反向传播网络的结构也是一样的。蓝色方框表示输入层，绿色圆圈表示隐含层，输出层没有绘制出来。 1.1 加载…
阅读更多...
Linux安装Redis（这里使用Redis6，其它版本类似）

Linux安装Redis（这里使用Redis6，其它版本类似）

目录 一、选择需要安装的Redis版本二、下载并解压Redis三、编译安装Redis四、启动Redis4.1、修改配置文件4.2、启动 五、测试连接5.1、本地连接使用自带客户端redis-cli连接操作redis5.2、外部连接使用RedisDesktopManager操作redis 六、关闭Redis七、删除Redis 一、选择需要安…
阅读更多...
【吞噬星空】战神宫全体投票，为罗峰脱罪，徐欣补办婚礼，洪成功恢复脑电波

【吞噬星空】战神宫全体投票，为罗峰脱罪，徐欣补办婚礼，洪成功恢复脑电波

【侵权联系删除】【文/郑尔巴金】 吞噬星空动画第90集即将更新&#xff0c;官方相当给力&#xff0c;提前曝光了图文情报与先行预告。虽然罗峰与巴巴塔尚未正式开始闯荡宇宙&#xff0c;但却是斩杀阿特金三大巨头的平稳生活。不但有战神宫为罗峰脱罪&#xff0c;而且还给徐欣补…
阅读更多...
python如何创建自己的对冲交易算法

python如何创建自己的对冲交易算法

在这篇文章中&#xff0c;我解释了如何创建一个人工智能来每天为我进行自动交易。 随着机器学习的现代进步和在线数据的轻松访问&#xff0c;参与量化交易变得前所未有的容易。为了让事情变得更好&#xff0c;AWS 等云工具可以轻松地将交易想法转化为真正的、功能齐全的交易机器…
阅读更多...
重组件的优化和页面渲染十万条数据

重组件的优化和页面渲染十万条数据

重组件的优化和页面渲染十万条数据的优化 重组件的优化vue2写法vue3写法 页面渲染十万条数据的优化使用虚拟列表的方式 重组件的优化 以下代码原理是使用requestAnimationFrame(callback) 方法 vue2写法 Test01.vue <template><div class"container">&…
阅读更多...
第十届山东省大学生网络安全技能大赛【神秘的base】【小试牛刀】

第十届山东省大学生网络安全技能大赛【神秘的base】【小试牛刀】

神秘的base 题目描述 EvAzEwo6E9RO4qSAHq42E9KvEv5zHDt34GtdHGJaHD7NHG42bwd神奇密码&#xff1a; xbQTZqjN8ERuwlzVfUIrPkeHd******LK697o2pSsGDncgm3CBh/Xy1MF4JAWta解题思路 这个题&#xff0c;上午一直零解&#xff0c;后来放出了hint&#xff0c;提示了base64换表。 这…
阅读更多...
iOS 配置通用链接(Universal Link)服务端和开发者后台都配置好了，还是跳转不到App

iOS 配置通用链接(Universal Link)服务端和开发者后台都配置好了，还是跳转不到App

目录 一、什么是 Universal Link&#xff1f; 1.背景介绍 2.特点 3.运行机制原理&流程图 二、配置教程 1.第一步&#xff1a;开启 Associated Domains 服务 1.1 开通 Associated Domains 2.第二步&#xff1a;服务器配置 apple-app-site-association&#xff08;AAS…
阅读更多...
binlog 和 redolog 有什么区别

binlog 和 redolog 有什么区别

binlog 和 redolog 都是 Mysql 里面用来记录数据库数据变更操作的日志. binlog 其中 binlog 主要用来做数据备份、数据恢复和数据同步&#xff0c;在Mysql 的主从数据同步的场景中&#xff0c;master 节点的数据变更&#xff0c;会写入到 binlog 中&#xff0c;然后再把 binl…
阅读更多...
会声会影2024永久破解中文版下载

会声会影2024永久破解中文版下载

会声会影2024中文版是一款功能强大的视频编辑软件、大型视频制作软件、专业视频剪辑软件。会声会影专业视频编辑处理软件&#xff0c;可以用于剪辑合并视频&#xff0c;制作视频&#xff0c;屏幕录制&#xff0c;光盘制作&#xff0c;视频后期编辑、添加特效、字幕和配音等操作…
阅读更多...
Allegro在走线时如何隐藏其它网络的飞线

Allegro在走线时如何隐藏其它网络的飞线

Allegro在走线时，自动隐藏其它网络的飞线，在走线的过程中不会被其它飞线干扰，从而方便快速的走线。 在走线过程中，其它飞线自动隐藏。 操作方法如下： 先激活走线命令，然后在Options选项卡上勾选Auto-blank other rats 则在走线时，其它网络的预拉线全部自动隐藏。
阅读更多...
Mybatis plus中的逻辑删除源码跟踪

Mybatis plus中的逻辑删除源码跟踪

网上搜了一堆文章&#xff0c;都没有贴出源码&#xff1a; 疑问&#xff1a; 逻辑删除&#xff0c;到底是在哪改变的sql&#xff1f;&#xff1f;&#xff1f; 跟踪方法——DeleteById 说明&#xff1a;逻辑删除&#xff0c;实质用的update语句&#xff1b; 组装逻辑删除字段…
阅读更多...
分享画PAD图的软件-PADFlowChart

分享画PAD图的软件-PADFlowChart

软件的可执行文件下载&#xff1a;PADFlowChart-exe.zip 如果有帮助望三联
阅读更多...
第63讲：MySQL存储过程变量传参的核心概念与案例

第63讲：MySQL存储过程变量传参的核心概念与案例

文章目录 1.存储过程传参2.存储过程传参的语法结构3.存储过程传参经典案例3.1.IN类型的传参3.2.IN和OUT同时使用的传参3.3.INOUT类型的传参 1.存储过程传参 在创建存储过程时&#xff0c;存储过程名称后面的括号里是可以传入一些参数的&#xff0c;例如传入一个变量&#xff0…
阅读更多...
毕业三年，月薪30K，我想跟你聊聊

毕业三年，月薪30K，我想跟你聊聊

大家好&#xff0c;我是冰河~~ 很多读者私信问我&#xff0c;自己工作三年多了&#xff0c;随着工作年限的不断增长&#xff0c;感觉自己的技术水平与自己的工作年限严重不符。想跳槽出去换个新环境吧&#xff0c;又感觉自己的能力达不到心仪公司的标准&#xff0c;即使投了简…
阅读更多...
蔚蓝科技推BabyAlpha阿尔法机器狗，表象之下是否会重蹈虚假覆辙

蔚蓝科技推BabyAlpha阿尔法机器狗，表象之下是否会重蹈虚假覆辙

最近笔者在朋友圈又刷到这样一条广告&#xff1a;关于一个叫蔚蓝阿尔法机器狗的推广。颇有刷着刷着刷出一条机器狗来的既视感。 经过了解&#xff0c;这是蔚蓝科技公司为旗下新机器狗BabyAlpha推出&#xff0c;投放的朋友圈销售广告。为什么说又呢&#xff0c;因为印象中这家公…
阅读更多...
单链表（第二部分）单链表的实现！！！

单链表（第二部分）单链表的实现！！！

1.单链表的头文件 #pragma once #define _CRT_SECURE_NO_WARNINGS #include<stdio.h> #include<stdlib.h> #include<assert.h> #include<string.h> #include"contact.h"typedef struct PersonInfo SLTDataType; typedef struct SListNode {…
阅读更多...
技术文档工具『Writerside』抢鲜体验

技术文档工具『Writerside』抢鲜体验

前言 2023 年 10 月 16 日&#xff0c;JetBrains 宣布以早期访问状态推出 Writerside&#xff0c;基于 IntelliJ 平台的 JetBrains IDE&#xff0c;开发人员可使用它编写、构建、测试和发布技术文档&#xff0c;可以作为 JetBrains IDE 中的插件使用&#xff0c;也可以作为独立…
阅读更多...
2023年中国汽车塑料模具市场规模、竞争格局及行业趋势分析[图]

2023年中国汽车塑料模具市场规模、竞争格局及行业趋势分析[图]

汽车注塑模具主要用来制造汽车内外饰件以及座椅等其他塑料零部件&#xff0c;其中又以汽车内外饰件模具最多。汽车内外饰件主要由各类塑料、表皮、织物或复合材料制成&#xff0c;用到的模具主要是塑料模具。从现代汽车使用的材料来看&#xff0c;无论是外装饰件、内装饰件&…
阅读更多...
推荐文章
最新文章

Copyright @ 2022~2023