当前位置: 首页 > news >正文

CUDA系列-Kernel Launch-8

这里写目录标题

  • kernel launch

本章主要追踪一下kernel launch的流程,会不断完善。


kernel launch

先抛出一个问题,如果在一个循环中不断的发送kernel(kernel 内部while死循环),会是什么结果。

// kernel 函数
__global__ void kernel(float *a, int n) {int id = threadIdx.x + blockIdx.x * blockDim.x;while(1) {//a[id] = sqrt(a[id] + 1);//这句注释掉对结果没有影响}
}// 持续不断的把kernelfun送入某一个具体stream
int main() {
//1. 声明变量(略)//2. 设置cudaLimitDevRuntimePendingLaunchCount为128/1000等cudaDeviceSetLimit(cudaLimitDevRuntimePendingLaunchCount, 128)//3. 创建stream
StreamCreate(&stream);//4. launch kernel ctrl+C 退出while (1) {
//grid_dim, block_dim一次性占满所有资源或者<<<1,1>>>kernel<<<grid_dim, block_dim, 0, stream>>>(buffer, size);}
...
//5. 销毁资源sync();StreamDestroy(&stream1);
}上面345可以改为多线程,一个线程一个stream.
其中还有一个简单的办法,首先在stream中发射一个阻塞的hostfun,然后发送空kernel也能计算到其大小,参考部分有相关代码

结果:

持续的发送一个小kernel到1个stream中,在1022次kernal launch 后,host出现block。3个stream中现象和1个stream一样,也是在1022次后被阻塞住。

详细参数如下(无论cudaLimitDevRuntimePendingLaunchCount设置为多少,下面结果没有变化)

indexgriddimblockdimstremresult
1111从1022次开始阻塞
2113从1022次开始阻塞
317281281从1022次开始阻塞
417281283从1022次开始阻塞
617281288从1022次开始阻塞
71/17281/12812约~743次开始阻塞
81/17281/12816约~550次开始阻塞
91/17281/12848约~224次开始阻塞
101/17281/128128约~12-33次开始阻塞
  1. cudaLimitDevRuntimePendingLaunchCount 的设置对结果没有影响,上面表格中,cudaLimitDevRuntimePendingLaunchCount无论设置为128,256,1000等,最后结果都是一样的。因为它是CUDA Dynamic Parallelism 嵌套launch的一个控制参数(后面会有证明)。

2)grid_dim, block_dim大小对结果也没有影响,因为此刻限制issue item的数量的是cuda runtimes 中的stream queue(实际上该queue 被map到另外一个叫做channel的对象上),对于正在执行或者pending状态的item,都会在queue中占用资源,直到其完成。上面都说明限制在某个stream上launch kernel的瓶颈点是一个host端侧的资源,所以排除cudaLimitDevRuntimePendingLaunchCount,因为它是一个device侧的资源控制量。

3)stream用户可以创建很多个,但是stream queue最后都是被map到channel上,channel的数量是有限的,并且channel又分为很多类型,不同类型其capacity也不一样,其中看实验结果,其中CU_CHANNEL_COMPUTE类型的只有8个,并且每个channel容纳1022个kernel func。

  1. 如果stream数量少于channel的数量,那么每个stream对应一个channel,如果stream的数量大于channel,distributes work evenly across all channels。
cuiLaunchstreamBeginPushWithFlagsstreamBeginPushWithDescstreamBeginPushOnChannelWithFlagschannelBeginPushInternalchannelBeginPushInternal_UnderLockchannelMustAdvance_UnderlockchannelMustAdvance_WaitForGPFIFOchannelCanAdvanceGPFIFO(在这里判断pushbuf和fifo entry)gpfifoHasPushbufferSpacegpfifoAdvanceGpuGetpushbufferHasSpace

整个调用链是这样的:
当向stream中下发kerneL的时候,stream会找到一个channel,该channel中有一个gpfifo的queue,其内部有一个ring_buffer(4M),另外还维护着一个semaphore queue(Max:1024),我们下发的每个kernel都会写道对应的ring_buffer中,并且每个kernel对专门对应一个semaphore entry放在semaphore queue中,当GPU开始执行kernel的时候,ring_buffer中的kernel data是不能删除的,只有当kernel执行完后,GPU 发送一个semaphore signal给CPU,CPU收到后会找对应的semaphore entry,让其释放资源,因为fifo有顺序要求,所以如果前面的kernel没有执行完,后面的kernel执行完,那么依然会block.也就是说只要第一个kernel在执行,即使后面全部是empty kernel 那么依然会block.

(继续完善)

相关文章:

CUDA系列-Kernel Launch-8

这里写目录标题 kernel launch 本章主要追踪一下kernel launch的流程&#xff0c;会不断完善。 kernel launch 先抛出一个问题&#xff0c;如果在一个循环中不断的发送kernel&#xff08;kernel 内部while死循环&#xff09;&#xff0c;会是什么结果。 // kernel 函数 __glo…...

# 消息中间件 RocketMQ 高级功能和源码分析(四)

消息中间件 RocketMQ 高级功能和源码分析&#xff08;四&#xff09; 一、 消息中间件 RocketMQ 源码分析&#xff1a;回顾 NameServer 架构设计。 1、RocketMQ 架构设计 消息中间件的设计思路一般是基于主题订阅发布的机制&#xff0c;消息生产者&#xff08;Producer&…...

如何通过数据库与AI实现以图搜图?OceanBase向量功能详解

OceanBase支持向量数据库的基础能力 当前&#xff0c;数据库存储系统与人工智能技术的结合&#xff0c;可以体现在两个主要的应用方向上。 一、近似搜索。它利用大语言模型&#xff08;LLM&#xff0c;简称大模型&#xff09;的嵌入&#xff08;embedding&#xff09;技术&am…...

Kafka内外网分流配置listeners和advertised.listeners

问题背景&#xff1a; Kafka部署在内网&#xff0c;内网Java服务会使用Kafka收发消息&#xff0c;另外&#xff0c;Java服务会与其他第三方系统使用kafka实现数据同步&#xff0c;也就是外网也会发送消息到kafka&#xff0c;外网IP做了端口映射到了内网&#xff0c;advertised…...

Linux系统编程——网络编程

目录 一、对于Socket、TCP/UDP、端口号的认知&#xff1a; 1.1 什么是Socket&#xff1a; 1.2 TCP/UDP对比&#xff1a; 1.3 端口号的作用&#xff1a; 二、字节序 2.1 字节序相关概念&#xff1a; 2.2 为什么会有字节序&#xff1a; 2.3 主机字节序转换成网络字节序函数…...

信息安全技术基础知识-经典题目

【第1题】 1.在信息安全领域&#xff0c;基本的安全性原则包括机密性(Confidentiality)、完整性(Integrity)和 可用性(Availability)。机密性指保护信息在使用、传输和存储时 (1) 。信息加密是保证系统机密性的常用手段。使用哈希校验是保证数据完整性的常用方法。可用性指保证…...

nextjs(持续学习中)

return ( <p className{${lusitana.className} text-xl text-gray-800 md:text-3xl md:leading-normal}> Welcome to Acme. This is the example for the{’ } Next.js Learn Course , brought to you by Vercel. ); } 在顶级 /public 文件夹下提供静态资产 **默认 /…...

数据预处理与特征工程、过拟合与欠拟合

数据预处理与特征工程 常用的数据预处理步骤 向量化&#xff1a;将数据转换成pytorch张量值归一化&#xff1a;将特定特征的数据表示成均值为0&#xff0c;标准差为1的数据的过程&#xff1b;取较小的值&#xff1a;通常在0和1之间&#xff1b;相同值域处理缺失值特征工程&am…...

甲辰年五月十四风雨思

甲辰年五月十四风雨思 夜雨消暑气&#xff0c;远光归家心。 ​只待万窗明&#xff0c;朝夕千家勤。 ​苦乐言行得&#xff0c;酸甜日常品。 宫商角徵羽&#xff0c;​仁义礼智信。...

java分别使用 iText 7 库和iText 5 库 将excel转成PDF导出,以及如何对excel转PDF合并单元格

第一种 package com.junfun.pms.report.util;import com.itextpdf.kernel.font.PdfFontFactory; import com.itextpdf.layout.Document; import com.itextpdf.layout.element.Paragraph; import com.itextpdf.layout.property.TextAlignment; import com.itextpdf.layout.prop…...

Java特性之设计模式【访问者模式】

一、访问者模式 概述 在访问者模式&#xff08;Visitor Pattern&#xff09;中&#xff0c;我们使用了一个访问者类&#xff0c;它改变了元素类的执行算法。通过这种方式&#xff0c;元素的执行算法可以随着访问者改变而改变。这种类型的设计模式属于行为型模式。根据模式&…...

【教师资格证考试综合素质——法律专项】未成年人保护法笔记以及练习题

《中华人民共和国未成年人保护法》 目录 第一章 总 则 第二章 家庭保护 第三章 学校保护 第四章 社会保护 第五章 网络保护 第六章 政府保护 第七章 司法保护 第八章 法律责任 第九章 附 则 介一&#xff0e;首次颁布&#xff1a;第一部《中华人民共和国未成年人保护法…...

6.19作业

TCP服务器 #include <stdio.h> #include <sys/types.h> #include <sys/socket.h> #include <unistd.h> #include <arpa/inet.h> #include <netinet/in.h> #include <string.h>#define PORT 8888 #define IP "192.168.124.39&q…...

java 线程之间通信-volatile 和 synchronized

你好&#xff0c;我是 shengjk1&#xff0c;多年大厂经验&#xff0c;努力构建 通俗易懂的、好玩的编程语言教程。 欢迎关注&#xff01;你会有如下收益&#xff1a; 了解大厂经验拥有和大厂相匹配的技术等 希望看什么&#xff0c;评论或者私信告诉我&#xff01; 文章目录 一…...

资源宝库网站!人人必备的神器!

面对网络中海量的内容&#xff0c;一个高效、便捷的网络导航工具&#xff0c;可以帮助我们快速查找使用网络资源。无论是职场精英还是学生党&#xff0c;使用导航网站都可以帮助我们提升效率。下面小编就来和大家分享一款资源宝库网站-办公人导航-实用的办公生活导航网站&#…...

Redis实战—优惠卷秒杀(锁/事务/代理对象的应用)

本博客为个人学习笔记&#xff0c;学习网站与详细见&#xff1a;黑马程序员Redis入门到实战 P50 - P54 目录 优惠卷秒杀下单功能实现 超卖问题 悲观锁与乐观锁 实现CAS法乐观锁 一人一单功能实现 代码优化 代码细节分析 优惠卷秒杀下单功能实现 ​ ​ Controller层…...

HTML星空特效

目录 写在前面 完整代码 代码分析 运行效果 系列文章 写在后面 写在前面 100行代码实现HTML星空特效。 完整代码 全部代码如下。 <!DOCTYPE html PUBLIC "-//W3C//DTD HTML 4.01 Transitional//EN" "http://www.w3.org/TR/html4/loose.dtd"&g…...

银行数仓项目实战(四)--了解银行业务(存款)

文章目录 项目准备存款活期定期整存整取零存整取存本取息教育储蓄定活两便通知存款 对公存款对公账户协议存款 利率 项目准备 &#xff08;贴源层不必写到项目文档&#xff0c;因为没啥操作没啥技术&#xff0c;只是数据。&#xff09; 可以看到&#xff0c;银行的贴源层并不紧…...

MySQL版本发布模型

MySQL 8.0 之后使用了新的版本控制和发布模型&#xff0c;分为两个主线&#xff1a;长期支持版&#xff08;LTS&#xff09;以及创新版。这两种版本都包含了缺陷修复和安全修复&#xff0c;都可以用于生产环境。 下图是 MySQL 的版本发布计划&#xff1a; 长期支持版 MySQL…...

java: 不兼容的类型: org.apache.xmlbeans.XmlObject无法转换为x2006.main.CTRow

我使用的xmlbeans版本是5.0&#xff0c;使用xmlbeans包做转换时&#xff0c;报错&#xff0c;正如标题显示得那样 解决办法 额外再引入下面的jar包 <dependency><groupId>org.apache.xmlbeans</groupId><artifactId>xmlbeans</artifactId><…...

KubeSphere 容器平台高可用:环境搭建与可视化操作指南

Linux_k8s篇 欢迎来到Linux的世界&#xff0c;看笔记好好学多敲多打&#xff0c;每个人都是大神&#xff01; 题目&#xff1a;KubeSphere 容器平台高可用&#xff1a;环境搭建与可视化操作指南 版本号: 1.0,0 作者: 老王要学习 日期: 2025.06.05 适用环境: Ubuntu22 文档说…...

Docker 离线安装指南

参考文章 1、确认操作系统类型及内核版本 Docker依赖于Linux内核的一些特性&#xff0c;不同版本的Docker对内核版本有不同要求。例如&#xff0c;Docker 17.06及之后的版本通常需要Linux内核3.10及以上版本&#xff0c;Docker17.09及更高版本对应Linux内核4.9.x及更高版本。…...

7.4.分块查找

一.分块查找的算法思想&#xff1a; 1.实例&#xff1a; 以上述图片的顺序表为例&#xff0c; 该顺序表的数据元素从整体来看是乱序的&#xff0c;但如果把这些数据元素分成一块一块的小区间&#xff0c; 第一个区间[0,1]索引上的数据元素都是小于等于10的&#xff0c; 第二…...

django filter 统计数量 按属性去重

在Django中&#xff0c;如果你想要根据某个属性对查询集进行去重并统计数量&#xff0c;你可以使用values()方法配合annotate()方法来实现。这里有两种常见的方法来完成这个需求&#xff1a; 方法1&#xff1a;使用annotate()和Count 假设你有一个模型Item&#xff0c;并且你想…...

C++.OpenGL (10/64)基础光照(Basic Lighting)

基础光照(Basic Lighting) 冯氏光照模型(Phong Lighting Model) #mermaid-svg-GLdskXwWINxNGHso {font-family:"trebuchet ms",verdana,arial,sans-serif;font-size:16px;fill:#333;}#mermaid-svg-GLdskXwWINxNGHso .error-icon{fill:#552222;}#mermaid-svg-GLd…...

短视频矩阵系统文案创作功能开发实践,定制化开发

在短视频行业迅猛发展的当下&#xff0c;企业和个人创作者为了扩大影响力、提升传播效果&#xff0c;纷纷采用短视频矩阵运营策略&#xff0c;同时管理多个平台、多个账号的内容发布。然而&#xff0c;频繁的文案创作需求让运营者疲于应对&#xff0c;如何高效产出高质量文案成…...

使用Spring AI和MCP协议构建图片搜索服务

目录 使用Spring AI和MCP协议构建图片搜索服务 引言 技术栈概览 项目架构设计 架构图 服务端开发 1. 创建Spring Boot项目 2. 实现图片搜索工具 3. 配置传输模式 Stdio模式&#xff08;本地调用&#xff09; SSE模式&#xff08;远程调用&#xff09; 4. 注册工具提…...

【Linux】自动化构建-Make/Makefile

前言 上文我们讲到了Linux中的编译器gcc/g 【Linux】编译器gcc/g及其库的详细介绍-CSDN博客 本来我们将一个对于编译来说很重要的工具&#xff1a;make/makfile 1.背景 在一个工程中源文件不计其数&#xff0c;其按类型、功能、模块分别放在若干个目录中&#xff0c;mak…...

uniapp 小程序 学习(一)

利用Hbuilder 创建项目 运行到内置浏览器看效果 下载微信小程序 安装到Hbuilder 下载地址 &#xff1a;开发者工具默认安装 设置服务端口号 在Hbuilder中设置微信小程序 配置 找到运行设置&#xff0c;将微信开发者工具放入到Hbuilder中&#xff0c; 打开后出现 如下 bug 解…...

JDK 17 序列化是怎么回事

如何序列化&#xff1f;其实很简单&#xff0c;就是根据每个类型&#xff0c;用工厂类调用。逐个完成。 没什么漂亮的代码&#xff0c;只有有效、稳定的代码。 代码中调用toJson toJson 代码 mapper.writeValueAsString ObjectMapper DefaultSerializerProvider 一堆实…...