al-folio v1.x 博客迁移与深度定制实践总结

<!-- 动态聚合标签 -->
<div class="tag-category-list">
  <ul class="p-0 m-0">
    
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/ai">AI</a>
        <span class="count">(6)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/algorithm">Algorithm</a>
        <span class="count">(4)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/ardupilot">Ardupilot</a>
        <span class="count">(3)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/bash">Bash</a>
        <span class="count">(9)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/cad">CAD</a>
        <span class="count">(2)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/cam">CAM</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/cmake">CMake</a>
        <span class="count">(6)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/cpu">CPU</a>
        <span class="count">(5)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/cuda">CUDA</a>
        <span class="count">(21)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/clang">Clang</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/computer">Computer</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/cpp">Cpp</a>
        <span class="count">(58)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/cute">CuTe</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/dart">Dart</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/designpattern">DesignPattern</a>
        <span class="count">(2)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/ffmpeg">FFmpeg</a>
        <span class="count">(2)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/fast-dds">Fast-DDS</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/flutter">Flutter</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/gcc">GCC</a>
        <span class="count">(4)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/gdb">GDB</a>
        <span class="count">(2)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/gpu">GPU</a>
        <span class="count">(2)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/git">Git</a>
        <span class="count">(2)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/hpc">HPC</a>
        <span class="count">(3)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/helios">Helios</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/intel">Intel</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/jetson">Jetson</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/llvm">LLVM</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/libevent">Libevent</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/linux">Linux</a>
        <span class="count">(18)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/mavlink">MAVLink</a>
        <span class="count">(2)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/mpi">MPI</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/markdown">Markdown</a>
        <span class="count">(2)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/mathjax">MathJax</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/memory">Memory</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/nvidia">NVIDIA</a>
        <span class="count">(2)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/networking">Networking</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/nginx">Nginx</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/occt">OCCT</a>
        <span class="count">(11)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/office">Office</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/opencl">OpenCL</a>
        <span class="count">(6)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/opencv">OpenCV</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/opengl">OpenGL</a>
        <span class="count">(2)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/openmaic">OpenMAIC</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/performance">Performance</a>
        <span class="count">(20)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/plugin">Plugin</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/powershell">PowerShell</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/python">Python</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/qemu">QEMU</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/qgc">QGC</a>
        <span class="count">(11)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/qml">QML</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/qt">Qt</a>
        <span class="count">(5)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/rtps">RTPS</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/redis">Redis</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/shell">Shell</a>
        <span class="count">(2)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/tbb">TBB</a>
        <span class="count">(7)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/template">Template</a>
        <span class="count">(2)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/tools">Tools</a>
        <span class="count">(2)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/vscode">VSCode</a>
        <span class="count">(6)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/vtk">VTK</a>
        <span class="count">(6)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/vtune">VTune</a>
        <span class="count">(7)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/vulkan">Vulkan</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/website">Website</a>
        <span class="count">(3)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/al-folio">al-folio</a>
        <span class="count">(2)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/blog">blog</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/docker">docker</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/frontend">frontend</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/jekyll">jekyll</a>
        <span class="count">(3)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/oh-my-posh">oh-my-posh</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/oh-my-zsh">oh-my-zsh</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/posh-git">posh-git</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/ubuntu">ubuntu</a>
        <span class="count">(2)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/zed">zed</a>
        <span class="count">(1)</span>
      </li>
    
      <li>
        <i class="fa-solid fa-hashtag fa-sm"></i>
        <a href="/blog/tag/zsh">zsh</a>
        <span class="count">(1)</span>
      </li>
    
  </ul>
</div>

---
layout: none
permalink: /assets/js/data/search.json
---
[

  { "title": "al-folio v1.x 博客迁移与深度定制实践总结", "url": "/blog/2026/al-folio-v1-customization-guide/", "categories": "tech, blog",
  "tags": "jekyll, al-folio, frontend, blog", "date": "2026-08-30", "content":
  "## 1. 概述与背景本站基于流行的学术个人主页与技术博客模板 [al-folio](https://github.com/alshedivat/al-folio) 搭建。近期模板进行了 v1.x 大版本架构重构（核心组件与 SASS 抽离为 Ruby Gem `al_folio_core`，并引入 Tailwind CSS 预编译体系）。为了在享受新版模板持续更新与现代技术栈优势的同时，完美保留 170+ 篇博文资产与个性化阅读体验，我们对新仓库进行了一系列系统性的迁移、问题修复与深度定制。```┌──────────────────────────────────────────────────────────┐│ al-folio v1.x 站点架构 │├─────────────────┬──────────────────────┬─────────────────┤│ Ruby Gem 核心 │ 本地覆盖层 │ 客户端增强 ││ (al_folio_core) │ (Shadowing Overrides)│ (Client Hooks) │├─────────────────┼──────────────────────┼─────────────────┤│ • 基础 Layouts │ • _layouts/ │ • Medium-Zoom ││ • 基础 SCSS 变量 │ • _sass/ │ • Tocbot 3-Level││ • 插件机制 │ • _includes/ │ • FullTextSearch││ • 构建工具链 │ • assets/js/ │ • Theme Toggle │└─────────────────┴──────────────────────┴─────────────────┘```---## 2. 本地开发与 Docker 容器环境### 2.1. 推荐本地运行工作流在本地采用 Docker 容器化环境运行，无需在本机安装复杂的 Ruby/Bundler 依赖：```bash# 启动本地开发服务（后台/前台）docker compose up# 若修改了依赖或配置需重新构建docker compose up --build# 停止容器并释放端口docker compose down```### 2.2. 本地环境常见问题与解决1. **端口 8080 冲突**： - 本地已有其他开发服务占用了 `8080` 端口。 - **解决**：在 `docker-compose.yml` 与 `docker-compose-slim.yml` 中将宿主机端口映射修改为 `8081:8080`，通过 `http://localhost:8081` 访问。2. **构建时外部 RSS 源连接超时（medium.com）**： - 模板默认配置了外部 Medium 源抓取，国内本地网络环境下因超时导致构建中断。 - **解决**：在 `_config.yml` 中注释掉 `external_sources` 相关配置。---## 3. CI/CD 与 GitHub Actions 部署修复在将已有博文推送到 GitHub 进行 Pages 部署时，遇到了一系列自动化检查拦截，主要解决措施如下：1. **Prettier 格式化检查失败**： - GitHub Actions 的 `prettier.yml` 强制执行 `npx prettier . --check`。 - **解决**：本地统一执行 `npx prettier . --write` 完成全量格式化；并在 `.prettierignore` 中配置忽略规则。2. **图片丢失导致 `broken-links-site.yml` 报错**： - 之前误删除了模板自带的部分示例图片资源（如 `assets/img/1.jpg`~`12.jpg`）。 - **解决**：完整补齐项目示例图片资源，确保所有内置示例页面的外链与静态文件检查全绿通过。3. **标签大小写冲突（Tag Case Collisions）**： - 原博文中存在大小写混用（如 `Ubuntu` 与 `ubuntu`），在 Linux/macOS 大小写不敏感文件系统交叉构建时会导致生成路径冲突。 - **解决**：统一规范博文 Front Matter 中的标签为小写格式。---## 4. 栏目精简与个人信息定制### 4.1. 导航栏结构重构为使站点保持极简聚焦的技术博客风格，评估并保留了 3 个核心栏目，将其余学术专属页面（论文、授课等）在导航栏中隐藏：| 栏目 | 状态 | 导航顺序 (`nav_order`) | 说明 || :------------------------------------- | :--- | :--------------------: | :------------------------- || **About** | 开启 | — | 个人主页与简介 || **Blog** | 开启 | 1 | 技术博客主页 || **Projects** | 开启 | 2 | 个人项目与开源作品 || _CV / Publications / Repos / Teaching_ | 隐藏 | `nav: false` | 隐藏冗余页面，保留文件备用 |### 4.2. About 页面与社交链接- **个人信息**：在 `_pages/about.md` 中更新了个性化 Slogan、技术方向简介与个人头像（`starry-sky.jpg`），同时关闭了 `announcements`（新闻）与 `selected_papers`。- **社交链接**：在 `_data/socials.yml` 中配置了 GitHub、博客园（cnblogs）、知乎、邮箱与 RSS 订阅源。---## 5. 动态标签（Tags）与分类（Categories）系统### 5.1. 原生问题al-folio 默认的 `_pages/blog.md` 中使用的是配置文件中预定义的静态标签数组（`site.display_tags` / `site.display_categories`），导致博文实际使用的标签无法自动出现在博客主页，且会展示很多空标签。### 5.2. 改造方案在 `_pages/blog.md` 中将静态循环替换为 Liquid 动态聚合语法：```liquid {% assign sorted_tags = site.tags | sort %} {% for tag in sorted_tags %} {{ tag[0] }} ({{ tag[1].size }}) {% endfor %} ```**效果**：新增、修改博文标签时，博客首页标签栏实时同步更新，自动统计文章篇数，无任何空标签。---## 6. 博文排版布局与三级目录（TOC）深度扩展### 6.1. 调整内容显示宽度与压缩留白- **全局版心**：在 `_config.yml` 中将 `max_width` 从默认的 `930px` 调整为 **`1400px`**。- **排版栅格**：在 `_layouts/default.liquid` 中采用稳定可靠的 `col-sm-9`（正文 75%，宽约 `1020px`）与 `col-sm-3`（目录 25%，宽约 `350px`）。- **效果**：大幅减少了宽屏显示器下两侧过宽的无意义空白，正文阅读视野开阔。### 6.2. 目录（TOC）列表框高度与平滑滚动在 `_sass/_utilities.scss` 中优化 `#toc-sidebar` 样式：- 设置 `max-height: calc(100vh - 6rem);` 充分利用整个屏幕高度；- 开启 `overflow-y: auto;` 并配置轻量滚动条 `scrollbar-width: thin;`，长文章目录平滑滚动。### 6.3. 支持三级及以上深度目录（H2 / H3 / H4）1. **JS 标题扫描与 Tocbot 升级**（`assets/js/common.js`）： ```javascript window.tocbot.init({ tocSelector: \"#toc-sidebar\", contentSelector: '[role=\"main\"]', headingSelector: \"h2, h3, h4\", // 扩展解析 H2、H3、H4 ignoreSelector: \"[data-toc-skip]\", hasInnerContainers: true, collapseDepth: resolveTocCollapseDepth(), orderedList: false, activeLinkClass: \"is-active-link\", scrollSmooth: true, scrollSmoothOffset: -80, headingsOffset: 80, }); ```2. **SCSS 层级缩进增强**（`_sass/_utilities.scss`）： 为三级嵌套 `.toc-list .toc-list .toc-list` 添加了 `padding-left: 0.8rem;` 缩进，形成清晰的树状视觉层次。---## 7. 图片宽度自适应与 Medium-Zoom 灯箱增强### 7.1. 大图宽度自适应（防截断）在 `_sass/_typography.scss` 中为内容区所有图片添加响应式约束：```scss.post-content img,#markdown-content img,article.post-content img,.post img { max-width: 100% !important; height: auto !important; display: block; margin: 0.75rem auto; cursor: zoom-in;}```### 7.2. 点击弹出全屏大图（带白底衬托与单例防重）1. **白底衬底（解决透明 SVG / 黑色文字看不清）**： 在 `_sass/_utilities.scss` 中为 `.medium-zoom-image--opened` 赋予白底卡片衬底，使透明背景图与暗色蒙层形成高对比度： ```scss .medium-zoom-overlay { z-index: 1050; backdrop-filter: blur(3px); } .medium-zoom-image--opened { z-index: 1051; background-color: #ffffff !important; /* 保证透明图自带白色画布底 */ border-radius: 6px; padding: 8px; box-shadow: 0 12px 40px rgba(0, 0, 0, 0.6) !important; } ```2. **单例模式绑定（解决重复弹窗与鬼影残留）**： 在 `assets/js/common.js` 中使用全局单例 `window.__alFolioZoomInstance` 并配合 `dataset.zoomAttached` 标记，严格保证每张图片仅绑定一次实例。---## 8. 全站全文检索系统（Full-Text Search）### 8.1. 方案选型与入口替换模板自带的 `ninja-keys` 是轻量命令面板，仅索引标题，无法搜索正文。我们采用了**按需轻量索引 + 客户端全文检索**的架构替换原有入口：- **顶部导航入口**：将顶部放大镜图标直接指向 `/search/` 独立全文搜索页；- **全局快捷键**：在全站任意位置按下 **`Ctrl + K`** 或 **`Cmd + K`**，即可一键直达搜索页并自动聚焦输入框。### 8.2. 构建期索引生成（`assets/js/data/search.json`）通过 Liquid 模板在编译时输出全站 170+ 篇博文的轻量 JSON 索引（含标题、URL、分类、标签、日期及经过纯文本剥离的正文）：```liquid---layout: nonepermalink: /assets/js/data/search.json---[{% for post in site.posts %} { \"title\": {{ post.title | jsonify }}, \"url\": {{ post.url | relative_url | jsonify }}, \"categories\": {{ post.categories | join: ', ' | jsonify }}, \"tags\": {{ post.tags | join: ', ' | jsonify }}, \"date\": {{ post.date | date: '%Y-%m-%d' | jsonify }}, \"content\": {{ post.content | strip_html | strip_newlines | normalize_whitespace | jsonify }} }{% unless forloop.last %},{% endunless %}{% endfor %}]```### 8.3. 搜索页交互与动态反馈（`_pages/search.md`）- **索引加载状态**：页面初始加载时显示 `⏳ 正在加载站内全文索引库...`，就绪后转为绿色提示 `✅ 索引库已加载（共 176 篇博文）`；- **检索中反馈**：输入关键词时实时显示 `⏳ 正在检索 “xxx”...`；- **多词与全文匹配**：支持中英文多词 AND 检索，结果卡片展示**正文上下文摘要片段（Snippet）**与**高亮关键词（``）**；- **一键清空**：输入框右侧提供 `✕` 快速清空按钮；- **URL 查询联动**：支持 `/search/?q=关键词` 链接直达自动搜索。---## 9. Prettier 规则优化与 CI 构建提速在全站博文与媒体文件较多时，Prettier 若遍历大量 PDF、图片和压缩包会导致本地检查与 GitHub Actions 构建变慢。我们在 `.prettierignore` 中添加了针对性忽略规则：```gitignore# 二进制媒体文件**/*.png**/*.jpg**/*.jpeg**/*.gif**/*.webp**/*.svg**/*.pdf**/*.zip**/*.rar**/*.7z**/*.doc**/*.docx# 编译缓存与带 Liquid 的 JSON_site/**.jekyll-cache/**assets/js/data/**```**优化效果**：Prettier 检查从扫描全量二进制文件压缩至 **8 秒内极速完成**，CI/CD 流程大幅加速。---## 10. 核心改动文件索引| 文件路径 | 作用与定制内容 || :--------------------------- | :-------------------------------------------------------------------------- || `_config.yml` | 全局版心 `max_width: 1400px`、博客名称、禁用 external_sources、精简 plugins || `_layouts/default.liquid` | 覆盖核心布局，稳定配置 `col-sm-9` 与 `col-sm-3` 栅格，挂载 TOC 侧边栏 || `_includes/header.liquid` | 将顶部放大镜按钮改造为直达 `/search/` 的全文搜索入口 || `_pages/about.md` | 个人主页简介、头像与社交配置 || `_pages/blog.md` | 改造为 Liquid 动态提取全站 Tags 与 Categories || `_pages/search.md` | 独立全文检索主页，集成状态加载指示器与高亮摘要 || `assets/js/data/search.json` | 全站博文全文检索轻量化 JSON 索引库 || `assets/js/common.js` | Tocbot 扩展支持 H2/H3/H4、Medium-Zoom 单例管理、`Ctrl+K` 全局快捷键 || `_sass/_layout.scss` | 容器尺寸约束 || `_sass/_utilities.scss` | TOC 容器高度扩展与滚动、图片弹窗白底卡片衬底与暗色遮罩 || `_sass/_typography.scss` | 正文图片 `max-width: 100%` 自适应约束与表格斑马纹边框优化 || `.prettierignore` | 过滤静态媒体与压缩包，提升代码检查与 Actions 部署速度 || `docker-compose.yml` | 本地容器端口调整为 `8081:8080`，避免本地端口冲突 |" },

  { "title": "在Ubuntu上部署OpenMAIC", "url": "/blog/2026/OpenMAIC%E9%83%A8%E7%BD%B2/", "categories": "AI",
  "tags": "AI, OpenMAIC", "date": "2026-05-22", "content":
  "1. 安装WSL2使用管理员启动PowerShell，执行启用WSL2：# 启用WSL功能：/all -- 所有用户启用，/norestart -- 不重启系统dism.exe /online /enable-feature /featurename:Microsoft-Windows-Subsystem-Linux /all /norestart# WSL2 依赖 Windows 虚拟机功能，需要额外启用dism.exe /online /enable-feature /featurename:VirtualMachinePlatform /all /norestart# 将WSL2设为默认WSLwsl --set-default-version 2执行完成之后，重启win10/win11。接下来在WSL2中安装Ubuntu系统。# 查看可用的WSL发行版本wsl --list --online# 安装过程中，在用户目录下如果有.wslconfig引起的错误，可用直接上这个配置文件# 安装过程中，会提示输入用户名和密码，输入完成之后就安装完成了wsl --install -d Ubuntu-26.04 --location D:\\WSL\\Ubuntu-26.04其他命令列表：# 退出 WSLwsl --shutdown# 验证安装wsl -l -v# 启动wsl -d Ubuntu-26.04# 删除命令wsl --unregister Ubuntu-26.04Remove-Item -Recurse -Force \"D:\\WSL\\Ubuntu-26.04\"如果Windows系统上安装了Nvidia CUDA套件，WSL2会自动安装Nvidia驱动程序，并且可以直接使用GPU进行计算。使用命令nvidia-smi可以验证GPU是否可用。安装并启动之后，可以在Windows中访问WSL2中的Ubuntu系统，比如在文件管理器中输入路径\\\\wsl$\\Ubuntu-26.04\\home\\\u0026lt;用户名\u0026gt;\\。也可以从WSL/Ubuntu系统中访问Windows文件系统，比如输入路径/mnt/d/访问D盘。1.1. WSL2/Ubuntu 26.04 系统配置Ubuntu 26.04系统更换软件源。备份之后清空文件/etc/apt/sources.list.d/ubuntu.sources，并粘贴如下内容：Types: debURIs: https://mirrors.aliyun.com/ubuntuSuites: resolute resolute-updates resolute-backportsComponents: main restricted universe multiverseSigned-By: /usr/share/keyrings/ubuntu-archive-keyring.gpgTypes: debURIs: https://mirrors.aliyun.com/ubuntuSuites: resolute-securityComponents: main restricted universe multiverseSigned-By: /usr/share/keyrings/ubuntu-archive-keyring.gpg安装必要的软件：sudo apt install -y git cmake build-essential wget unzip vim htopsudo apt install python3 python3-pip -y资料: WSL2 完全指南：在 Windows 上安装和使用 Linux 的终极教程：备注：不要参考文中关于docker的安装部分 Ubuntu26.04 可以稳定使用的国内源1.2. WSL清理及备份首先查看Ubuntu占用空间分配情况：usdo apt install -y ncdusudo ncdu --exclude=\"mnt*\" --exclude=\"proc*\" /###################### docker 清理docker volume prune # 清理 无用的volumesdocker system prune # 删除无用的镜像和容器及构建缓存# 删除不必要的dockerdocker imagesdocker rmi xxx执行WSL压缩（在windows/powrshell环境下）：wsl --shutdowndiskpartDISKPART\u0026gt; select vdisk file=\"D:\\WSL\\Ubuntu-26.04\\ext4.vhdx\"DISKPART\u0026gt; compact vdiskDISKPART\u0026gt; detach vdisk# 退回到Powershell环境备份及恢复命令：# 备份指定系统到指定位置wsl --export Ubuntu-26.04 D:\\Ubuntu-26.04.tar# 还原指定系统wsl --import Ubuntu-26.04 D:\\Ubuntu-26.04.tar2. 安装Nodejs以及pnpm需要安装Node.js和pnpm，并更换源，命令如下：# 可以使用nvm在本地安装Node.js以及npm，并选择Node.js/npm 的版本# nvm是一个Node.js版本管理工具：https://github.com/nvm-sh/nvmcurl -o- https://raw.githubusercontent.com/nvm-sh/nvm/v0.40.4/install.sh | bashnvm install 24 # 安装Node.js 24版本nvm use 24 # 切换到Node.js 24版本node -v # 验证Node.js版本# 也可以管理员安装Node.jssudo apt install nodejs npmnode -v # 验证Node.js版本npm -v # 验证npm版本# 安装pnpmnpm install -g pnpm# 更换源为阿里云，阿里云速度快，同步及时npm config set registry https://registry.npmmirror.compnpm config set registry https://registry.npmmirror.com3. 安装 docker 以及 NVIDIA Container Toolkit需要在带Nvidia GPU的服务器上部署MinerU。不带GPU的服务器，使用CPU运行MinerU，性能会非常差。MinerU可以使用本地部署，或者使用docker部署。使用docker部署+Nvidia GPU，需要安装NVIDIA Container Toolkit。首先在WSL2/Ubuntu系统上安装docker，命令如下：sudo apt-get updatesudo apt-get install -y ca-certificates curl util-linux-extrasudo install -m 0755 -d /etc/apt/keyrings# 从阿里云下载 GPG keysudo curl -fsSL https://mirrors.aliyun.com/docker-ce/linux/ubuntu/gpg -o /etc/apt/keyrings/docker.ascsudo chmod a+r /etc/apt/keyrings/docker.asc# 使用阿里云仓库，原始仓库国内连接不上echo \\ \"deb [arch=$(dpkg --print-architecture) signed-by=/etc/apt/keyrings/docker.asc] https://mirrors.aliyun.com/docker-ce/linux/ubuntu \\ $(. /etc/os-release \u0026amp;\u0026amp; echo \"$VERSION_CODENAME\") stable\" | \\ sudo tee /etc/apt/sources.list.d/docker.list \u0026gt; /dev/nullsudo apt-get updatesudo apt-get install -y docker-ce docker-ce-cli containerd.io docker-buildx-plugin docker-compose-plugin# 将当前用户添加到docker用户组中，以便不使用sudo运行docker命令sudo usermod -aG docker $USERnewgrp docker再执行安装NVIDIA Container Toolkit的命令：# Install the prerequisites for NVIDIA Container Toolkitsudo apt-get update \u0026amp;\u0026amp; sudo apt-get install -y --no-install-recommends ca-certificates curl gnupg2# Configure the production repository and the GPG keycurl -fsSL https://nvidia.github.io/libnvidia-container/gpgkey | sudo gpg --dearmor -o /usr/share/keyrings/nvidia-container-toolkit-keyring.gpg \\ \u0026amp;\u0026amp; curl -s -L https://nvidia.github.io/libnvidia-container/stable/deb/nvidia-container-toolkit.list | \\ sed 's#deb https://#deb [signed-by=/usr/share/keyrings/nvidia-container-toolkit-keyring.gpg] https://#g' | \\ sudo tee /etc/apt/sources.list.d/nvidia-container-toolkit.list# Optional: use experimental packagessudo sed -i -e '/experimental/ s/^#//g' /etc/apt/sources.list.d/nvidia-container-toolkit.listsudo apt update# Install the NVIDIA Container Toolkit packagesexport NVIDIA_CONTAINER_TOOLKIT_VERSION=1.19.1-1sudo apt-get install -y \\ nvidia-container-toolkit=${NVIDIA_CONTAINER_TOOLKIT_VERSION} \\ nvidia-container-toolkit-base=${NVIDIA_CONTAINER_TOOLKIT_VERSION} \\ libnvidia-container-tools=${NVIDIA_CONTAINER_TOOLKIT_VERSION} \\ libnvidia-container1=${NVIDIA_CONTAINER_TOOLKIT_VERSION}安装完成之后，执行如下命令配置docker使用NVIDIA Container Toolkit：sudo nvidia-ctk runtime configure --runtime=dockersudo systemctl restart docker验证 NVIDIA Container Toolkit 是否安装成功：docker run --rm --gpus all nvidia/cuda:12.2.0-base-ubuntu22.04 nvidia-smi3.1. 配置 docker Hub 仓库镜像sudo tee /etc/docker/daemon.json \u0026lt;\u0026lt;EOF{ \"runtimes\": { \"nvidia\": { \"args\": [], \"path\": \"nvidia-container-runtime\" } }, \"registry-mirrors\": [ \"https://docker.m.daocloud.io\", \"https://mirror.baidubce.com\" ]}EOFsudo systemctl restart docker参考资料： Windows + WSL2 + Ubuntu 22.04 + Docker: 完整安装与配置教程 (自定义路径 ext4 镜像数据盘方案)4. 安装/部署/运行 MinerU下载仓库：git clone https://github.com/opendatalab/MinerU.git构建镜像：cd MinerU/docker/chinamkdir -p ~/mineru/pdfsmkdir -p ~/mineru/outputcp ../../demo/pdfs/demo1.pdf ~/mineru/pdfs/docker run --rm --gpus=all -v ~/mineru/pdfs:/data/pdfs -v ~/mineru/output:/data/output mineru:latest mineru -p /data/pdfs/demo1.pdf -o /data/output其中： -v ~/mineru/pdfs:/data/pdfs 将本地的pdfs目录挂载到容器的/data/pdfs目录，-v ~/mineru/output:/data/output 将本地的output目录挂载到容器的/data/output目录。 mineru -p /data/pdfs/demo1.pdf -o /data/output 是在容器中运行MinerU命令，指定输入pdf文件和输出目录。 -p、-o两个参数，表示运行一次MinerU，处理一个pdf文件，输出结果到指定目录。随后就退出了，导致docker容器也随之退出。 --gpus=all去掉这个参数，则使用CPU进行计算，速度会非常慢。 使用docker images命令可以查看本地的docker镜像列表，确认是否成功构建了mineru:latest镜像。MinerU的另外一种方式：docker run -it --name mineru --gpus=all mineru:latest /bin/bash这种方式，会在容器中启动一个交互式的bash终端，用户可以在这个终端中执行MinerU命令，处理多个pdf文件，输出结果到指定目录。用户可以在这个交互式终端中执行多次MinerU命令，处理多个pdf文件，输出结果到指定目录。用户处理完成之后，可以输入exit命令退出交互式终端，随之docker容器也会退出。MinerU还提供web方式访问，使用Gradio Web进行部署，见参考博客。 MinerU 安装部署完全指南：CPU、GPU、Docker 保姆级教程：参考 五、Docker 部署 使用如上命令运行docker+MinerU，构建镜像的Dockerfile文件中，是使用pip安装的方式安装MinerU的依赖包，与本地代码没有关系，即不会部署本地代码。4.1. Docker compose 服务在生产环境下，使用docker compose部署服务会更方便，运行命令如下：# CDN加速下载compose.yaml文件，避免下载失败# compose.yml文件中，是使能了GPU的配置wget https://gcore.jsdelivr.net/gh/opendatalab/MinerU@master/docker/compose.yaml# 运行命令，启动服务。docker在后台运行，使用`-d`参数docker compose -f compose.yaml --profile api up -d# 验证服务是否就绪curl http://localhost:8000/health# 停止并移除容器（推荐）docker compose -f compose.yaml --profile api down# 只停止不删除容器docker compose -f compose.yaml --profile api stop其他关联命令用法：# 查看容器状态docker compose -f compose.yaml --profile api ps# 查看日志docker logs -f mineru-apidocker logs -f -t mineru-api # 带时间戳的日志# 查看日志，另外一种方式docker compose -f compose.yaml --profile api logs -f # 实时日志docker compose -f compose.yaml --profile api logs --tail=100 # 查看最近100行日志MinerU服务的所有配置，都在compose.yaml文件中，如端口、镜像名称等等。 MinerU启动之后，在OpenMAIC的配置文件.env.local中，将MINERU_API_URL=http://localhost:8000，即可使用本地部署的MinerU服务。有关MinerU后端参数，见MinerU仓库的issue：使用docker compose启动mineru-api，如何选择hybrid-auto-engine和vlm-auto-engine。MinerU不能很好的识别pdf中的目录结构，转换得到的markdown文件，目录似乎会都识别成二级目录##。需要借助LLM进行辅助识别，可以使用本地部署Ollama，或者第三方LLM（付费，需要API Key）。见MinerU代码仓库的Discussions：5003、4413。4.2. MinerU 环境变量部分实用可配置的 MinerU 环境变量：# 限制单卡GPU显存export MINERU_VIRTUAL_VRAM_SIZE=8# 使用modelscope源加速国内下载，默认HuggingFace# 官方文档：https://opendatalab.github.io/MinerU/zh/usage/model_source/export MINERU_MODEL_SOURCE=modelscope# 指定使用pipeline后端#export MINERU_BACKEND=pipeline4.3. MinerU　测试如下测试命令，将 MinerU 返回的json结果保存到mineru_response.json文件中，并将响应头保存到mineru_headers.txt文件中：curl.exe -F \"files=@E:\\义务教育教科书·物理八年级全一册.pdf\" -F \"parse_method=auto\" -F \"backend=hybrid-auto-engine\" -F \"return_content_list=true\" -F \"return_images=true\" http://172.22.191.184:8000/file_parse -o mineru_response.json --max-time 1800 -D mineru_headers.txt -sS4.4. 部分资料 使用指南 issue mineru-vl-client 如何配置可用显存等参数呢 issue vlm以及pipeline识别某pdf出现内存泄露占用大量内存导致进程被操作系统kill5. 本地部署 QWen TTSTODO5.1. 资料 MOSS-TTS：MOSS-TTS主打高保真，相对QWen TTS，Voice Clone时间比较长（QWen TTS似乎限制3秒），但QWen TTS时延小（\u0026lt;96ms）。另外，MOSS-TTS nano版本占用资源很小，可以在CPU上运行。6. 部署及使用 OpenMAICdocker部署、运行命令，以及清理命令如下：# 在OpenMAIC代码目录中docker build --no-cache -t openmaic .# docker build --no-cache --build-arg APK_MIRROR=mirrors.aliyun.com --build-arg NPM_REGISTRY=https://registry.npmmirror.com -t openmaic .docker run --rm -p 3000:3000 --env-file .env.local openmaicdocker stop $(docker ps -q --filter \"ancestor=openmaic\") 2\u0026gt;/dev/nulldocker rm $(docker ps -aq --filter \"ancestor=openmaic\") 2\u0026gt;/dev/null# 清理所有项目的 BuildKit 缓存docker builder prune -f.env.local配置文件基于.env.example修改如下：# 本地部署的MinerU服务地址，远程则还需要远程地址及API KeyPDF_MINERU_BASE_URL=http://172.22.191.184:8000# PR提交的新增的参数，设置OpenMAIC等待MinerU的时间，单位秒，默认5分钟MINERU_TIMEOUT_MS=1800000# 使用DeepSeek模型，配置API KeyDEEPSEEK_API_KEY=\u0026lt;....\u0026gt;DEEPSEEK_BASE_URL=# Example: deepseek-v4-pro,deepseek-v4-flashDEEPSEEK_MODELS=deepseek-v4-pro# web search引擎的API KeyTAVILY_API_KEY=\u0026lt;....\u0026gt;# 配置PorxyHTTP_PROXY=http://192.168.11.139:7890HTTPS_PROXY=http://192.168.11.139:7890# 允许局域网访问，如localhost，192.168.x.x等ALLOW_LOCAL_NETWORKS=true6.1. 平台链接 DeepSeek开放平台 阿里云平台 Tavily" },

  { "title": "AI工具大全", "url": "/blog/2026/AI%E5%B7%A5%E5%85%B7%E5%A4%A7%E5%85%A8/", "categories": "AI",
  "tags": "AI", "date": "2026-05-16", "content":
  "IDE Cursor Zed Kiro Antigravity VS Code Trae GooseADE CodexVS Code extension GitHub Copilot Cline Roo Code Amp Code (free with ads)Cloud agent Codex Jules Cursor background agent GitHub Copilot Claude Code on the webLocal parallel agent Vibe Kanban Crystal by Stravu Conductor Sculptor by imbue emdashCLIHere are the hyperlinks to the official websites (or documentation pages) for each of the three CLI tools: Codex CLI Claude Code Gemini CLI Kimi CLI Qwen3-Coder CLI Amp Code CLI OpenCode Pi Mistral Vibe CLI GitHub Copilot CLI Goose CLI Warp CLI Jules ToolsAgents framework/library OpenAI Agents SDK Google ADK (Agent Development Kit) Microsoft Agent Framework Claude Agent SDKMCP Playwright Chrome Dev Tools Context7Web Search API Brave Exa PerplexityImage models Nano Banana and Nano Banana Pro by Google Seedream by BytedanceOpen weight models Z-Image by Alibaba (HF: Turbo, Base, Edit) Qwen-Image by Alibaba (HF: Qwen-Image and Qwen-Image-Edit) LongCat-Image-Edit by MeituanOCR models GLM-OCRAudio models SAM-Audio by Meta Fun-Audio-Chat by Alibaba PersonaPlex by NvidiaTTS Gemini 3.1 Flash TTSOpen-weight Chatterbox-Turbo Qwen3-TTS: HFTranscription (STT) Parakeet v3: HF Qwen3-ASR: HF Voxtral Transcribe 2: HFTranslate models TranslateGemma by Google (HF: TranslateGemma) HY-MT by Tencent (HF: HY-MT15)Evaluation OpenAI evals Stax by GoogleVibe coding platforms Google AI Studio Apps (frontend only) Google Firebase Studio (full-stack) Cloudflare VibeSDK Lovable Bolt AI Replit v0 by Vercel Base44 GitHub Spark" },

  { "title": "Fast DDS入门（On-Going）", "url": "/blog/2026/Fast-DDS%E5%85%A5%E9%97%A8/", "categories": "FastDDS",
  "tags": "Fast-DDS, RTPS", "date": "2026-04-26", "content":
  "1. Fast DDS功能模块TODO2. 发现协议发现协议类型枚举定义(include/fastdds/rtps/attributes/RTPSParticipantAttributes.hpp)：enum class DiscoveryProtocol{ NONE, /*!\u0026lt; NO discovery whatsoever would be used. Publisher and Subscriber defined with the same topic name would NOT be linked. All matching must be done manually through the addReaderLocator, addReaderProxy, addWriterProxy methods. */ SIMPLE, /*!\u0026lt; Discovery works according to 'The Real-time Publish-Subscribe Protocol(RTPS) DDS Interoperability Wire Protocol Specification' */ EXTERNAL, /*!\u0026lt; A user defined PDP subclass object must be provided in the attributes that deals with the discovery. Framework is not responsible of this object lifetime. */ CLIENT, /*!\u0026lt; The participant will behave as a client concerning discovery operation. Server locators should be specified as attributes. */ SERVER, /*!\u0026lt; The participant will behave as a server concerning discovery operation. Discovery operation is volatile (discovery handshake must take place if shutdown). */ BACKUP, /*!\u0026lt; The participant will behave as a server concerning discovery operation. Discovery operation persist on a file (discovery handshake wouldn't repeat if shutdown). */ SUPER_CLIENT /*!\u0026lt; The participant will behave as a client concerning all internal behaviour. Remote servers will treat it as a server and will share every discovery information. */};TODO3. Fast DDS中的环境变量Fast DDS提供了一些运行时环境变量，用于配置和调整DDS实体的行为。这些环境变量可以在系统环境中设置，也可以通过一个环境文件（json）中加载，这个环境文件路径通过环境变量FASTDDS_ENVIRONMENT_FILE指定。环境文件的优先级高于系统环境变量，一个例外是，ROS_DISCOVERY_SERVER可以在运行时通过环境变量动态修改。3.1. FASTDDS_DEFAULT_PROFILES_FILE、SKIP_DEFAULT_XMLFASTDDS_DEFAULT_PROFILES_FILE、SKIP_DEFAULT_XML用于定义默认XML配置文件路径的环境变量。比如：export FASTDDS_DEFAULT_PROFILES_FILE=/home/user/profiles.xml加载profile文件的优先级为：优先加载环境变量指定的文件路径。如果环境变量未设置，则默认加载工作目录下的文件DEFAULT_FASTDDS_PROFILES.xml，这个默认文件名是写死在Fast DDS源码中的。当定义了SKIP_DEFAULT_XML环境变量并设置为1时，Fast DDS将跳过加载默认XML配置文件，以与Discovery Server的配置分开（CliDiscoveryManager::load_XML_file）。此时FASTDDS_DEFAULT_PROFILES_FILE、SKIP_DEFAULT_XML环境变量仍然起作用，即SKIP_DEFAULT_XML只决定是否跳过默认XML配置文件（工作目录下的DEFAULT_FASTDDS_PROFILES.xml）的加载。3.2. FASTDDS_BUILTIN_TRANSPORTS环境变量FASTDDS_BUILTIN_TRANSPORTS用于指定DomainParticipant选用的内置传输类型，该环境变量直接影响DDS实体的通信方式。仅当use_builtin_transports为true时，才会使用内置传输类型。传输类型有以下几种： 传输类型 描述 NONE 不使用任何内置传输，用户需要手动添加传输实例，否则DomainParticipant创建失败。 DEFAULT 实例化UDPv4和SHM传输。在可行的情况下，参与者会优先选择SHM。 DEFAULTv6 实例化UDPv6和SHM传输。在可行的情况下，参与者会优先选择SHM。 SHM 仅实例化SHM，适用于同一主机上的进程间通信。 UDPv4 仅实例化UDPv4。 UDPv6 仅实例化UDPv6。 LARGE_DATA 实例化UDPv4、TCPv4和SHM。其中UDP用于参与者发现阶段多播，参与者的活跃度和应用程序数据交付通过TCP/SHM进行。 FASTDDS_BUILTIN_TRANSPORTS带有一些参数，主要的参数有（相关代码RTPSParticipantImpl::setup_transports）： 参数名称 描述 max_msg_size / uint32_t DDS传输层（RTPS）的最大消息大小，单位为字节。 sockets_size / uint32_t 操作系统收发缓冲区大小（socket的SO_SNDBUF / SO_RCVBUF选项），单位为字节。 non_blocking / bool 是否启用非阻塞模式。如果启用，在操作系统缓冲区已满的时候，消息会被丢弃。 tcp_negotiation_timeout / uint32_t TCP连接的协商超时时间，单位为毫秒。仅在LARGE_DATA传输类型下有效。 max_msg_size、sockets_size这两个参数应该设置为消息的最大大小，否则可能导致分片，从而降低传输效率。在使用LARGE_DATA传输类型时，将non_blocking设置为true应该是一个优选。另外，在LARGE_DATA传输类型下，才能将max_msg_size设置大于65500 KB，其他传输类型下将导致创建参与者失败。TCPv4传输使用以下配置初始化： calculate_crc、check_crc和apply_security设置为false。 enable_tcp_nodelay设置为true。 keep_alive_thread和accept_thread使用默认配置。一个示例配置如下：export FASTDDS_BUILTIN_TRANSPORTS=LARGE_DATA?max_msg_size=200KB\u0026amp;sockets_size=1MB\u0026amp;non_blocking=true\u0026amp;tcp_negotiation_timeout=503.3. 发现协议相关环境变量：ROS_DISCOVERY_SERVER、ROS2_EASY_MODE、ROS_SUPER_CLIENTFast DDS发现机制分为两种：基于多播的Simple Discovery Protocol (SDP)（去中心化），和基于客户端-服务器的Discovery Server。当参与者DomainParticipant实例的发现协议为SIMPLE、SERVER或者BACKUP时，环境变量ROS_DISCOVERY_SERVER起作用（其他发现协议时无效）。如果定义了ROS_DISCOVERY_SERVER环境变量，会对DomainParticipant实例初始化的发现协议作一些更新： 当发现协议是SIMPLE时，更新为CLIENT； 当发现协议是SERVER/BACKUP时，保持不变。相关代码：RTPSDomainImpl::client_server_environment_attributes_override。ROS_DISCOVERY_SERVER格式举例： 192.168.2.23，默认使用UDPv4协议，UDP协议的默认端口11811。 TCPv4:[127.0.0.1]、TCPv4:[127.0.0.1]:42100，TCP协议的默认端口42100。TCP协议用于与服务器通信。 域名（使用DNS解析）：eprosima.com、eprosima.com:4567，即不指定协议时，被解析为UDP协议。TCPv4:[eprosima.com]、TCPv4:[eprosima.com]:42100（UDP格式类似）。如上的格式中，也可以指定TCPv6/UDPv6协议，格式类似于TCPv4/UDPv4。一个包含多个地址的示例如下：export ROS_DISCOVERY_SERVER=\"84.22.259.329:8888;localhost:1234\"3.3.1. ROS_SUPER_CLIENT环境变量ROS_SUPER_CLIENT用于启用Discovery Server的超级客户端模式。SUPER_CLIENT可以从Server收到所有DomainParticipant实例的订阅信息（等同于Server间同步），即可以监控所有client的状态。一般用作telemetry或监控工具。启用方法如下：export ROS_SUPER_CLIENT=true # TRUE/FALSE, True/False, 1/0，true/false环境变量ROS_SUPER_CLIENT是配合ROS_DISCOVERY_SERVER使用的。只有当ROS_DISCOVERY_SERVER环境变量被定义，且DomainParticipant实例的发现协议为SIMPLE时（默认为SIMPLE，其他发现方式为显式定义，此时不应该改变），环境变量ROS_SUPER_CLIENT才会生效。相关代码：RTPSDomainImpl::run_easy_mode_discovery_server。3.3.2. ROS2_EASY_MODE环境变量ROS2_EASY_MODE也要求DomainParticipant实例的发现协议为SIMPLE时（默认为SIMPLE），并更改发现机制为Discovery Server。同时，将DomainParticipant变更为SUPER_CLIENT，以保持跟SIMPLE发现协议的时候一样可以收到所有其他参与者的信息。当定义了ROS2_EASY_MODE环境变量并设置为1时，会在本地创建一个Discovery Server，这个本地server连接到远程Discovery Server，启动之后作为守护进程驻守后台。本地的DomainParticipant实例将连接到本地的Discovery Server，以避免直接与远程server通信（减小连接overhead、传输开销）。3.4. FASTDDS_STATISTICS环境变量FASTDDS_STATISTICS用于启用使能DomainParticipant实例的统计功能：启用统计信息DataWriter，即对外发布指定的统计信息。使用Fast DDS的监控工具Fast DDS Monitor，或者其他参与者可以订阅这些统计信息。Fast DDS Monitor可视化展示延迟、丢包、吞吐量、发现事件等指标。如下示例启用节点的延迟测量统计：export FASTDDS_STATISTICS=\"HISTORY_LATENCY_TOPIC;NETWORK_LATENCY_TOPIC\"实际的统计信息，是环境变量FASTDDS_STATISTICS中指定的统计主题（topic）与环境文件中指定的统计信息配置的合集。 编译时，需要启用编译选项FASTDDS_STATISTICS。A. 资料 Fast DDS 入门 一口气看完 FastDDS ，耗时666分钟制作 Fast DDS 官方文档" },

  { "title": "NVIDIA GPU 架构：SP、SM 与 LSU 工作原理详解", "url": "/blog/2026/NVIDIA-GPU-SP-SM-LSU/", "categories": "CUDA",
  "tags": "CUDA, GPU", "date": "2026-04-09", "content":
  "本文整理自 NVIDIA 开发者论坛的两个讨论帖，重点探讨 SM 内部功能执行单元（FMA / LSU 等）的调度模型，以及 Ampere 架构中 LSU（Load/Store Unit）的详细工作机制。以下内容主要来自 NVIDIA 工程师 Greg 以及资深版主 Robert Crovella 的回答。功能执行单元与 Warp-wide 调度GPU 中所有指令的调度都以 warp（32 个 thread） 为粒度。SM 内的每一类功能执行单元（functional unit）——包括 FMA 单元（即 CUDA Core）、LSU（LD/ST Unit）、Tensor Core 等——每个单元每周期只处理 1 条指令、1 个 thread。因此，要在单个时钟周期内完成一条 warp-wide 指令（涵盖 32 个 thread），就需要 32 个同类功能单元并行工作。以浮点乘加指令（FMUL/FMA）为例：如果希望一条 FMUL 在 1 个周期内完成整个 warp 的处理，就需要 32 个 FMA 单元同时工作，每个 FMA 单元负责 1 个 thread 的那一条指令。如果某个 SM sub-partition 只有 16 个 FMA 单元，那么一条 warp-wide 的 FMUL 就需要 2 个周期才能完全 issue。LSU 遵循完全相同的逻辑。一条 warp-wide 的 Load/Store 指令（如 LDG、STG）同样需要 32 个 LSU 单元来支撑单周期 issue。如果 SM sub-partition 中 LSU 数量不足 32 个，那么该指令的 issue 就需要跨越多个周期。这里的 issue 是指指令进入 LSU pipeline 的 front end，后续还有地址 coalescing、cache 访问等流水线阶段。 Tensor Core 是一个例外——它的调度粒度与上述逻辑不同，单个 Tensor Core 在单周期内可以处理更大的矩阵片段，不遵循 “1 指令/1 thread/1 cycle” 的模型。SM 子分区与 Warp Scheduler在 Ampere 等现代架构中，一个 SM 被划分为若干 sub-partition（processing block）。每个 sub-partition 拥有自己的 Warp Scheduler、register file、以及一组 functional unit。以 Ampere 的 A100 为例，一个 SM 包含 4 个 sub-partition，因此有 4 个 Warp Scheduler，每个 sub-partition 每周期可以 issue 1 条指令，整个 SM 的算术指令吞吐峰值为 4 IPC。然而，内存相关指令走的是另一条路径。所有 sub-partition issue 的 Load/Store 指令都汇集到 SM 级别的 MIO（Memory I/O）单元，而 MIO 每周期只能向 LSU pipeline 发送 1 条指令。这意味着即便 4 个 Warp Scheduler 都在 issue 内存指令，最终能进入 LSU pipeline 的也只有 1 条/周期——这是整个 SM 的内存指令 IPC 瓶颈所在，从 4 IPC 降至 1 IPC。LSU 管线的工作流程从 Warp Scheduler 到 LSU Pipeline当一个 warp 准备执行 Load/Store 指令时，Warp Scheduler 将该指令送入 MIO instruction queue（一个 shallow queue）。在指令从 MIOC（MIO Controller）dispatch 之前，必须等到所有 register operand 都已从 register file 中读取完毕。读取完成后，指令进入 ready 状态，随后被 MIO 以每周期 1 条的速率 dispatch 到 LSU pipeline。整个过程是异步的：指令一旦被 dispatch 到 LSU pipeline，该 warp 可以继续执行后续不依赖本次 Load/Store 结果的指令。只有当后续指令需要使用 Load 返回的数据时，该 warp 才会因 register 未就绪而 stall。LSU Front End 与 Back Pressure 机制Robert Crovella 强调，LSU 的 front end 并没有一个固定大小的 request queue。在理想情况下，LSU 可以持续以每周期 1 条指令（1 op/thread/clock）的速率接受新请求，不存在队列积压。然而 LSU 只是 memory pipe 的 front end，它后面连接着 L1 cache、L2 cache、DRAM controller 等一系列硬件。当下游出现拥塞（例如大量 non-coalesced 访问导致 transaction 数量暴增），back pressure 信号会沿 pipeline 反向传播，最终以一种 on/off signal 的形式关闭 LSU 的入口——此时 LSU 立即停止接受新请求，后续的内存指令会引发 stall。这种 stall 既不是典型的 data dependency stall，也不一定是 register reservation stall，其具体分类在 NVIDIA 的公开文档中并没有详细说明。Scoreboard 与寄存器复用指令进入 LSU pipeline 后，编译器可以选择两种 scoreboard 更新策略。一种是在 MIO 将指令 dispatch 到 LSU 后立即释放 scoreboard，允许该 register 被后续指令写入新值——此时 register 中的地址数据已经被 LSU 读走，不再需要保留。另一种是等到 Load 指令的数据真正返回（retire）后才释放 scoreboard，确保更强的顺序保证。据 Scott Gray 在 maxas wiki 中的观察，Read Dependency Barrier 的延迟大约为 20 个 clock——这是指令从 issue 到进入 in-flight 状态所需的时间，而完整的 Load 指令可能需要更多的 clock 才能返回数据。如果一个 warp 执行到某条指令时，所依赖的 register 还未就绪（scoreboard 未释放），该 warp 就会进入 long scoreboard stall。此时 Warp Scheduler 会切换到其他 eligible warp 继续执行，直到该 warp 的 scoreboard 解除为止。Wavefront 与地址合并地址合并（Coalescing）LSU pipeline 接收到一条 warp-wide 的 Load/Store 指令后，会将其 dispatch 到 shared memory pipe 或 tagged pipe（L1TEX）。在 tagged pipe 中，LSU 为 warp 中每个 thread 计算 cache tag（即地址所在的 32-byte sector），然后将落在同一 sector 内的 thread 合并为一个 wavefront。每个 wavefront 代表一次对 L1 cache 的访问请求。在 GV100 及后续架构上，L1TEX 的 tag stage 每周期可以 resolve 4 sets × 4 sectors。如果单个 wavefront 无法覆盖所有 thread（即地址分散在超过 4 个 set 的 sector 中），LSU 会继续在 tag stage 生成新的 wavefront，每周期生成一个，直到该指令的所有 thread 都被处理完毕。合并效率对吞吐的影响Coalescing 的程度直接决定了 LSU pipeline 的实际吞吐。以 32-bit Load 为例：如果 warp 中 32 个 thread 访问连续的 4-byte 地址（总计 128 字节，覆盖 4 个 sector），那么仅需 1 个 t-stage wavefront 加上 4 个 miss-stage wavefront 就能完成——这是最理想的情况。反过来，如果一条 Store 指令的 32 个 thread 分别写入 32 个不同的 sector（完全离散的地址），就会生成 32 个 wavefront，每个 wavefront 处理 1 个 sector。由于 wavefront 每周期只能生成 1 个，这条指令将独占 tag stage 长达 32 个周期，严重降低有效吞吐。这也正是 coalesced 访问与 non-coalesced 访问之间的本质差异：前者用少量 transaction 完成请求（高效率），后者产生大量 transaction（低效率），并可能触发前述的下游 back pressure 机制。Load 与 Store 的异同Load 和 Store 指令在 LSU pipeline 中的 coalescing 流程是一致的，区别在于后续阶段。Load 指令在 L1 cache 响应后，LSU 从每个 sector 中提取对应数据写入各 thread 的目标 register，然后更新 scoreboard 通知 warp 数据就绪。Store 指令则是 fire-and-forget 操作：数据和地址被送入 write pipeline 后，source register 很快就可以被释放，不必等到写操作真正完成。术语表 术语 说明 SM (Streaming Multiprocessor) GPU 核心执行单元，包含 Warp Scheduler、functional unit、register file、shared memory / L1 cache 等 SP (Streaming Processor) / CUDA Core SM 内部的算术 functional unit，每个单元每周期处理 1 条指令、1 个 thread sub-partition (processing block) SM 的子分区，每个拥有独立的 Warp Scheduler、register file 和一组 functional unit Warp Scheduler 每个 sub-partition 中的指令调度器，每周期从 eligible warp 中选择一个并 issue 1 条指令 warp 32 个 thread 的集合，是 GPU 调度和执行的最小粒度 issue Warp Scheduler 将指令发送到对应 functional unit 的动作 dispatch MIO Controller 将指令从 instruction queue 发送到 LSU pipeline 的动作 IPC (Instructions Per Cycle) 每周期指令吞吐，SM 算术峰值为 4 IPC，内存指令受 MIO 限制为 1 IPC functional unit SM 内的执行单元统称，包括 FMA unit、INT unit、LSU、Tensor Core、SFU 等 FMA (Fused Multiply-Add) 浮点乘加单元，即通常所说的 CUDA Core LSU (Load/Store Unit) 执行内存访问指令（LDG/STG/LDS/STS 等）的 functional unit MIO (Memory I/O) SM 级别的内存指令控制单元，汇集所有 sub-partition 的 Load/Store 指令，每周期向 LSU pipeline dispatch 1 条 MIOC (MIO Controller) MIO 的控制器，负责从 instruction queue 中取出 ready 指令并 dispatch register file 寄存器文件，每个 sub-partition 拥有独立的 register file，存放 thread 的局部数据 register operand 指令的寄存器操作数，dispatch 前必须从 register file 读取完毕 scoreboard 记录 register 依赖关系的硬件结构，防止 RAW（Read After Write）冒险 long scoreboard stall warp 因等待 register 数据就绪（如 Load 未返回）而进入的长延迟 stall Read Dependency Barrier 读依赖屏障，约 20 clock 延迟，指令从 issue 到进入 in-flight 状态的时间 eligible warp 所有操作数就绪、无 stall 的 warp，可被 Warp Scheduler 选中执行 pipeline 流水线，指令按阶段依次通过的硬件通路 front end pipeline 的前端/入口阶段 back pressure 下游硬件拥塞时，向上游发送的反压信号，导致 front end 停止接受新请求 stall 流水线停顿，warp 因某种原因无法继续执行，类似 CPU pipeline 中的 bubble bubble CPU pipeline 中因 hazard 插入的空操作周期，GPU 中对应 stall，通过 warp 切换实现 latency hiding memory pipe 内存访问硬件通路，包括 LSU（front end）、L1/L2 cache、DRAM controller 等 L1TEX L1 Texture/Data cache，处理 global memory 和 texture 访问的缓存子系统 tag stage (t-stage) L1TEX pipeline 中计算 cache tag、进行地址分组的阶段 miss stage L1TEX pipeline 中处理 cache miss 的阶段 coalescing 地址合并，将 warp 中多个 thread 的内存请求按 sector 合并以减少 transaction 数 sector 32-byte 的 cache 数据单元，是 L1 cache 访问的最小粒度 wavefront LSU 在 tag stage 中将同一 sector 的 thread 合并后生成的一次 cache 访问请求，每周期生成 1 个 transaction 一次实际的 cache/memory 数据传输操作 coalesced / non-coalesced 合并/非合并访问，前者 transaction 少（高效），后者 transaction 多（低效）可触发 back pressure in-flight 指令已 issue 但尚未 retire（完成），正在 pipeline 中执行 retire 指令执行完毕，结果写回 register file 或 cache fire-and-forget Store 指令的执行模式：数据送入 write pipeline 后 source register 即可释放，不等待写完成 latency hiding 延迟隐藏，GPU 通过大量 warp 并发，在某 warp stall 时切换到其他 eligible warp 继续执行 参考资料 How does the LSU execute Load/Store instructions in the Ampere architecture? - NVIDIA Developer Forums Questions about SP and SM - NVIDIA Developer Forums Nsight Compute Profiling Guide - Memory Tables L1" },

  { "title": "Ubuntu 26.04 安装 Docker 和 Docker Compose", "url": "/blog/2026/%E5%AE%89%E8%A3%85Docker-Docker-Compose/", "categories": "Tools, Blog",
  "tags": "jekyll, al-folio, ubuntu, docker", "date": "2026-02-26", "content":
  "1. 依赖安装先检查是否已经安装了Docker和Docker Compose，如果安装了，则先卸载旧版本：sudo apt remove -y docker docker-engine docker.io containerd runcsudo rm -rf /var/lib/docker /var/lib/containerd安装依赖项：sudo apt update \u0026amp;\u0026amp; sudo apt upgrade -y# 安装证书、curl、gnupg 等基础依赖sudo apt install -y ca-certificates curl gnupg lsb-release2. 添加国内源首先添加GPP密钥，另外再添加阿里云的Docker源：# 创建密钥存储目录sudo mkdir -p /etc/apt/trusted.gpg.d# 导入阿里云 Docker GPG 密钥（避免签名验证失败）curl -fsSL https://mirrors.aliyun.com/docker-ce/linux/ubuntu/gpg | sudo gpg --dearmor -o /etc/apt/trusted.gpg.d/docker.gpg# 添加适配 Ubuntu 24.04（noble）的阿里云 Docker 源echo \"deb [arch=$(dpkg --print-architecture) signed-by=/etc/apt/trusted.gpg.d/docker.gpg] https://mirrors.aliyun.com/docker-ce/linux/ubuntu $(lsb_release -cs) stable\" | sudo tee /etc/apt/sources.list.d/docker.list \u0026gt; /dev/null# 再次更新包索引（加载新添加的 Docker 源）sudo apt update3. 安装 Docker 和 Docker Compose，并配置Docker服务# 安装 Docker CE、Containerd、Docker Compose 插件sudo apt install -y docker-ce docker-ce-cli containerd.io docker-buildx-plugin docker-compose-plugin启动并设置开机自启（启动Docker，设置开机自启，验证Docker状态）：sudo systemctl start dockersudo systemctl enable dockersudo systemctl status docker将当前用户加入 docker 组（允许非 root 用户运行 Docker 命令）：# 将当前用户加入 docker 组sudo usermod -aG docker $USER# 生效组配置（无需重启，重新登录终端即可）sudo apt install util-linux-extra # 安装 newgrp 命令newgrp docker# 验证权限（无需 sudo 能执行即成功）docker ps4. 配置 Docker 镜像加速首先创建 Docker 配置文件目录（如果不存在）：sudo mkdir -p /etc/dockersudo touch /etc/docker/daemon.json配置文件内容如下：{ \"builder\": { \"gc\": { \"defaultKeepStorage\": \"100GB\", \"enabled\": true } }, \"data-root\": \"/home/docker/MyDocker\", \"experimental\": false, \"registry-mirrors\": [ \"https://dockerproxy.net\", \"http://mirrors.ustc.edu.cn\", \"http://mirror.azure.cn\", \"https://hub.rat.dev\" ]}验证配置是否生效：# 重启 Docker 使加速配置生效sudo systemctl daemon-reloadsudo systemctl restart docker# 验证加速器是否生效docker info | grep -i mirror输出中能看到配置的镜像地址，说明加速生效。5. 验证安装是否成功–运行一个测试容器# 运行 hello-world 测试容器验证安装成功docker run hello-world若输出 “Hello from Docker!” 相关内容，说明安装+配置全部成功。A. 参考资料 Ubuntu 24.04.1 LTS 安装 Docker（国内源版，全程可访问） Ubuntu 24.04 LTS Docker 和 Docker Compose 安装指南" },

  { "title": "C++ Traits", "url": "/blog/2026/C++-Traits/", "categories": "Cpp",
  "tags": "Cpp", "date": "2026-02-06", "content":
  "C++ Traits 是一种元编程技术，将类型信息封装在一个类中，供算法或其他模板使用，以期达到从算法/逻辑中分离类型信息的目的。比如元素类型、一些常量定义、指令选择、对齐要求等等。 问题 Traits 如何解决 内置类型（int, float）不能添加成员 Traits 外挂信息 不同类型需要不同的算法策略 模板特化选择策略 想要零开销的编译期多态 全部在编译期决议 接口与实现解耦 算法只依赖 Traits 接口 1. 自定义 Traits 常见模式1.1 类型信息萃取目的：从不同类型中统一提取需要的信息。// 主模板 (可以留空或给默认值)template \u0026lt;typename T\u0026gt;struct NumericTraits;// 对 float 特化template \u0026lt;\u0026gt;struct NumericTraits\u0026lt;float\u0026gt; { using type = float; using compute_type = float; // 计算时用的类型 using accum_type = float; // 累加器类型 static constexpr int bits = 32; static constexpr float epsilon = 1e-7f; static constexpr float max_val = 3.4028235e+38f;};// 对 half 特化template \u0026lt;\u0026gt;struct NumericTraits\u0026lt;__half\u0026gt; { using type = __half; using compute_type = float; // half 计算时提升到 float using accum_type = float; static constexpr int bits = 16;};// 对 int8_t 特化template \u0026lt;\u0026gt;struct NumericTraits\u0026lt;int8_t\u0026gt; { using type = int8_t; using compute_type = int32_t; // int8 用 int32 累加 using accum_type = int32_t; static constexpr int bits = 8;};// --- 使用 ---template \u0026lt;typename T\u0026gt;void gemm_kernel() { using Traits = NumericTraits\u0026lt;T\u0026gt;; using AccumType = typename Traits::accum_type; AccumType accumulator = 0; // 自动选择合适的累加器类型 // ...}1.2. 策略选择目的：根据类型选择不同的算法实现。// 内存拷贝策略 Traitstemplate \u0026lt;typename T, int Size\u0026gt;struct CopyTraits;// 小数据: 逐元素拷贝template \u0026lt;typename T\u0026gt;struct CopyTraits\u0026lt;T, 1\u0026gt; { static void copy(T* dst, const T* src) { *dst = *src; }};// 4字节对齐: 用 uint32_t 一次拷贝template \u0026lt;\u0026gt;struct CopyTraits\u0026lt;float, 4\u0026gt; { static void copy(float* dst, const float* src) { // 使用向量化加载 *reinterpret_cast\u0026lt;float4*\u0026gt;(dst) = *reinterpret_cast\u0026lt;const float4*\u0026gt;(src); }};// 128位对齐: 用 LDG.128 指令template \u0026lt;\u0026gt;struct CopyTraits\u0026lt;__half, 8\u0026gt; { static void copy(__half* dst, const __half* src) { // 8 个 half = 128 bit, 用 128-bit load uint4 tmp = *reinterpret_cast\u0026lt;const uint4*\u0026gt;(src); *reinterpret_cast\u0026lt;uint4*\u0026gt;(dst) = tmp; }};1.3. 特征探测目的：检测类型是否具有某些特征，以启用/禁用特定代码路径。// 检测类型是否有 .size() 方法template \u0026lt;typename T, typename = void\u0026gt;struct has_size_method : std::false_type {};template \u0026lt;typename T\u0026gt;struct has_size_method\u0026lt;T, std::void_t\u0026lt;decltype(std::declval\u0026lt;T\u0026gt;().size())\u0026gt;\u0026gt; : std::true_type {};// --- 使用 ---static_assert(has_size_method\u0026lt;std::vector\u0026lt;int\u0026gt;\u0026gt;::value, \"\"); // truestatic_assert(!has_size_method\u0026lt;int\u0026gt;::value, \"\"); // true// 检测类型是否有嵌套类型 value_typetemplate \u0026lt;typename T, typename = void\u0026gt;struct has_value_type : std::false_type {};template \u0026lt;typename T\u0026gt;struct has_value_type\u0026lt;T, std::void_t\u0026lt;typename T::value_type\u0026gt;\u0026gt; : std::true_type {};1.4. 标签分发目的：通过标签类型分发不同的实现。// Tagsstruct RowMajorTag {};struct ColMajorTag {};// Layout Traitstemplate \u0026lt;typename Layout\u0026gt;struct LayoutTraits;template \u0026lt;\u0026gt;struct LayoutTraits\u0026lt;RowMajorTag\u0026gt; { static int index(int row, int col, int ldim) { return row * ldim + col; } static constexpr char name[] = \"RowMajor\";};template \u0026lt;\u0026gt;struct LayoutTraits\u0026lt;ColMajorTag\u0026gt; { static int index(int row, int col, int ldim) { return col * ldim + row; } static constexpr char name[] = \"ColMajor\";};// 算法根据 tag 自动适配template \u0026lt;typename LayoutTag\u0026gt;void fill_matrix(float* data, int M, int N, int ld) { using LTraits = LayoutTraits\u0026lt;LayoutTag\u0026gt;; for (int i = 0; i \u0026lt; M; ++i) for (int j = 0; j \u0026lt; N; ++j) data[LTraits::index(i, j, ld)] = static_cast\u0026lt;float\u0026gt;(i * N + j);}2. Traits 核心技术手段2.1. 特化/偏特化 完全特化（Template Specialization）：为特定类型提供完整实现。 偏特化（Partial Specialization）：为满足某些条件的类型提供实现（如指针类型、数组类型等）。模板特化例子：template\u0026lt;typename T\u0026gt;struct my_is_void { static const bool value = false;};template\u0026lt;\u0026gt;struct my_is_void\u0026lt;void\u0026gt; { static const bool value = true;};偏特化例子：template\u0026lt;typename T\u0026gt;struct is_pointer { static const bool value = false;};template\u0026lt;typename T\u0026gt;struct is_pointer\u0026lt;T*\u0026gt; { static const bool value = true;};2.2. typename在 Traits 中经常需要定义嵌套类型（如 type、compute_type 等），使用 typename 来显式告诉编译器该名字是一个类型。template \u0026lt;typename T\u0026gt;struct NumericTraits { using type = T; // 定义一个嵌套类型};template \u0026lt;typename T\u0026gt;void foo() { using MyType = typename NumericTraits\u0026lt;T\u0026gt;::type; // 使用嵌套类型 // ...}2.3. SFINAE（Substitution Failure Is Not An Error）在 SFINAE 中，如果某个表达式无效（如访问了不存在的成员），编译器会将该特化视为无效，而不是报错，从而允许其他特化继续匹配。在 SFINAE 中，编译器优先特化/偏特化版本，C++ 偏特化规则：更特化（更具体）的版本优先。最后才是主模板版本。2.3.1. void_t 和 SFINAE 结合实现特征探测// void_t 是 SFINAE 的瑞士军刀template \u0026lt;typename... Ts\u0026gt;using void_t = void; // C++17 标准已有// 主模板：利用 void_t 探测嵌套类型template \u0026lt;typename T, typename = void\u0026gt;struct element_type_of { using type = T; // 默认: 类型本身};// 偏特化：如果 T 有 element_type 嵌套类型，就用它template \u0026lt;typename T\u0026gt;struct element_type_of\u0026lt;T, void_t\u0026lt;typename T::element_type\u0026gt;\u0026gt; { using type = typename T::element_type; // 如果有 element_type 就用它};// 使用// element_type_of\u0026lt;std::unique_ptr\u0026lt;int\u0026gt;\u0026gt;::type → 得到 int// element_type_of\u0026lt;float\u0026gt;::type → 得到 floatvoid_t是一个通用的的类型萃取容器，能将任意数量的表达式包裹为 void 类型。void_t\u0026lt;X\u0026gt; 本身永远是 void，但若 X 是非法类型表达式，模板替换就会失败。上述代码片段中，以 float 为例，偏特化得到：element_type_of\u0026lt; float, void_t\u0026lt;typename float::element_type\u0026gt; \u0026gt; ^^^^^^^^^^^^^^^^^^^^^^^^^^^ float 没有 ::element_type 成员 → 类型表达式非法 → 替换失败以 std::unique_ptr\u0026lt;int\u0026gt; 为例，偏特化得到：void_t\u0026lt;typename unique_ptr\u0026lt;int\u0026gt;::element_type\u0026gt; = void_t\u0026lt;int\u0026gt; ← ::element_type = int，合法！ = void ← void_t 永远返回 void最终得到：element_type_of\u0026lt; unique_ptr\u0026lt;int\u0026gt;, void \u0026gt;2.3.2. SFINAE 结合 enable_if 实现函数重载// enable_if 的原理:template \u0026lt;bool Cond, typename T = void\u0026gt;struct enable_if {}; // 默认: 没有 type 成员template \u0026lt;typename T\u0026gt;struct enable_if\u0026lt;true, T\u0026gt; { using type = T; // 条件为 true 时才有 type};// 使用: 条件不满足 → enable_if\u0026lt;false\u0026gt;::type 不存在 → SFINAEtemplate \u0026lt;typename T\u0026gt;typename std::enable_if\u0026lt;std::is_integral\u0026lt;T\u0026gt;::value, T\u0026gt;::typedouble_it(T x) { return x * 2;}template \u0026lt;typename T\u0026gt;typename std::enable_if\u0026lt;std::is_floating_point\u0026lt;T\u0026gt;::value, T\u0026gt;::typedouble_it(T x) { return x * 2.0;}double_it(42); // is_integral = true → 第一个 ✅double_it(3.14); // is_floating = true → 第二个 ✅2.4. constexpr 与 static 编译期常量template \u0026lt;typename T\u0026gt;struct AlignmentTraits { // 编译期常量 static constexpr int alignment = alignof(T); static constexpr int size_bytes = sizeof(T); // 编译期函数 (C++14) static constexpr int max_vector_width() { if constexpr (sizeof(T) == 1) return 16; // 128-bit / 8-bit if constexpr (sizeof(T) == 2) return 8; // 128-bit / 16-bit if constexpr (sizeof(T) == 4) return 4; // 128-bit / 32-bit return 1; }};2.5. 嵌套类型别名template \u0026lt;typename T\u0026gt;struct AlignmentTraits { // 编译期常量 static constexpr int alignment = alignof(T); static constexpr int size_bytes = sizeof(T); // 编译期函数 (C++14) static constexpr int max_vector_width() { if constexpr (sizeof(T) == 1) return 16; // 128-bit / 8-bit if constexpr (sizeof(T) == 2) return 8; // 128-bit / 16-bit if constexpr (sizeof(T) == 4) return 4; // 128-bit / 32-bit return 1; }};3. CuTe 中使用 Traits 的例子以SM80_CP_ASYNC_CACHEALWAYS为例，其 Traits 定义如下：// CuTe 风格的 Copy_Traits (简化示意)// 对应 cp.async 指令的 Traitstemplate \u0026lt;typename CopyInstr\u0026gt;struct Copy_Traits;// 针对 SM80_CP_ASYNC_CACHEALWAYS\u0026lt;cute::uint128_t\u0026gt; 的特化template \u0026lt;\u0026gt;struct Copy_Traits\u0026lt;SM80_CP_ASYNC_CACHEALWAYS\u0026lt;cute::uint128_t\u0026gt;\u0026gt; { // ---- 关联类型 ---- // 指令操作的寄存器类型 // 源: 1个 gmem 指针 (uint128_t) // 目标: 1个 smem 指针 (uint128_t) using SRegisters = cute::tuple\u0026lt;uint128_t const*\u0026gt;; // source (gmem ptr) using DRegisters = cute::tuple\u0026lt;uint128_t*\u0026gt;; // dest (smem ptr) // Layout 描述: 一个线程搬运多少数据、如何排布 // ThrID: 哪些线程参与 // ValLayoutSrc / ValLayoutDst: 值的布局 using ThrID = Layout\u0026lt;_1\u0026gt;; // 1个线程 using ValLayoutSrc = Layout\u0026lt;_1\u0026gt;; // 搬运1个128-bit值 using ValLayoutDst = Layout\u0026lt;_1\u0026gt;; // ---- 核心操作 ---- template \u0026lt;typename TS, typename TD\u0026gt; CUTE_HOST_DEVICE static void copy(TS const\u0026amp; src, TD\u0026amp; dst) { // 调用 PTX 内联汇编 // cp.async.ca.shared.global [dst], [src], 16; cute::cp_async\u0026lt;16\u0026gt;(dst, src); }};Copy_Atom 如何使用 Traits// 简化版 Copy_Atomtemplate \u0026lt;typename Traits, typename Element\u0026gt;struct Copy_Atom { using TraitsType = Traits; using ElementType = Element; // 从 Traits 中提取信息 using SrcRegisters = typename Traits::SRegisters; using DstRegisters = typename Traits::DRegisters; // 执行拷贝 — 委托给 Traits template \u0026lt;typename SrcEngine, typename SrcLayout, typename DstEngine, typename DstLayout\u0026gt; CUTE_HOST_DEVICE void copy(Tensor\u0026lt;SrcEngine, SrcLayout\u0026gt; const\u0026amp; src, Tensor\u0026lt;DstEngine, DstLayout\u0026gt; \u0026amp; dst) const { // 实际调用 Traits::copy Traits::copy(src.data(), dst.data()); }};// ===== 用户代码 =====// 选择 Traits → 决定用什么指令using CopyTraits = Copy_Traits\u0026lt;SM80_CP_ASYNC_CACHEALWAYS\u0026lt;uint128_t\u0026gt;\u0026gt;;// 创建 Atomusing CopyAtom = Copy_Atom\u0026lt;CopyTraits, cute::half_t\u0026gt;;// 构建 TiledCopy (多线程协作版本)using TiledCopy = decltype( make_tiled_copy( CopyAtom{}, Layout\u0026lt;Shape\u0026lt;_32, _4\u0026gt;\u0026gt;{}, // 线程布局: 32x4 = 128线程 Layout\u0026lt;Shape\u0026lt;_1, _8\u0026gt;\u0026gt;{} // 每线程搬运的值布局 ));5. C++ 标准库中 Traits 相关的常用工具5.1. 类型判断5.1.1. 基本类型判断 返回 ::value，bool 类型，C++17 起可用 _v 后缀简写。td::is_void\u0026lt;T\u0026gt; // T 是 void？std::is_null_pointer\u0026lt;T\u0026gt; // T 是 std::nullptr_t？ (C++14)std::is_integral\u0026lt;T\u0026gt; // T 是整数类型？(bool/char/int/long...)std::is_floating_point\u0026lt;T\u0026gt; // T 是浮点？(float/double/long double)std::is_arithmetic\u0026lt;T\u0026gt; // T 是算术类型？(整数 或 浮点)std::is_enum\u0026lt;T\u0026gt; // T 是枚举？std::is_class\u0026lt;T\u0026gt; // T 是 class/struct？std::is_function\u0026lt;T\u0026gt; // T 是函数类型？std::is_array\u0026lt;T\u0026gt; // T 是数组类型？(T[] 或 T[N])5.1.2. 复合类型判断std::is_pointer\u0026lt;T\u0026gt; // T 是指针？(int*, 但不含成员指针)std::is_reference\u0026lt;T\u0026gt; // T 是引用？(左值引用 或 右值引用)std::is_lvalue_reference\u0026lt;T\u0026gt; // T 是左值引用？std::is_rvalue_reference\u0026lt;T\u0026gt; // T 是右值引用？std::is_member_pointer\u0026lt;T\u0026gt; // T 是成员指针？std::is_const\u0026lt;T\u0026gt; // T 有顶层 const？std::is_volatile\u0026lt;T\u0026gt; // T 有顶层 volatile？std::is_signed\u0026lt;T\u0026gt; // T 是有符号类型？std::is_unsigned\u0026lt;T\u0026gt; // T 是无符号类型？5.1.3. 类型关系判断// ===== 两个类型的关系 =====std::is_same\u0026lt;T, U\u0026gt; // T 和 U 是完全相同的类型？std::is_base_of\u0026lt;Base, Derived\u0026gt; // Base 是 Derived 的基类？std::is_convertible\u0026lt;From, To\u0026gt; // From 能隐式转换为 To？std::is_assignable\u0026lt;T, U\u0026gt; // T = U 赋值合法？std::is_constructible\u0026lt;T, Args...\u0026gt; // T(Args...) 构造合法？std::is_invocable\u0026lt;F, Args...\u0026gt; // F(Args...) 可调用？ (C++17)std::is_invocable_r\u0026lt;R, F, Args...\u0026gt; // F(Args...) 可调用且返回 R？ (C++17)5.2. 类型变换5.2.1. 修饰符增删// ===== 去除修饰 =====std::remove_const\u0026lt;T\u0026gt; // const int → intstd::remove_volatile\u0026lt;T\u0026gt; // volatile int → intstd::remove_cv\u0026lt;T\u0026gt; // const volatile int → intstd::remove_reference\u0026lt;T\u0026gt; // int\u0026amp;/int\u0026amp;\u0026amp; → intstd::remove_pointer\u0026lt;T\u0026gt; // int* → intstd::remove_extent\u0026lt;T\u0026gt; // int[10] → intstd::remove_all_extents\u0026lt;T\u0026gt; // int[3][4] → int// ===== 添加修饰 =====std::add_const\u0026lt;T\u0026gt; // int → const intstd::add_volatile\u0026lt;T\u0026gt; // int → volatile intstd::add_cv\u0026lt;T\u0026gt; // int → const volatile intstd::add_lvalue_reference\u0026lt;T\u0026gt; // int → int\u0026amp;std::add_rvalue_reference\u0026lt;T\u0026gt; // int → int\u0026amp;\u0026amp;std::add_pointer\u0026lt;T\u0026gt; // int → int*5.2.2. decay// decay 模拟\"按值传参\"时的类型退化:// - 去掉引用// - 去掉顶层 cv// - 数组 → 指针// - 函数 → 函数指针std::decay\u0026lt;const int\u0026amp;\u0026gt; // → intstd::decay\u0026lt;int[10]\u0026gt; // → int*std::decay\u0026lt;int(double)\u0026gt; // → int(*)(double)std::decay\u0026lt;const int\u0026amp;\u0026amp;\u0026gt; // → int// --- 实用示例: 存储任意传入的值 ---template \u0026lt;typename T\u0026gt;struct Storage { using StoredType = std::decay_t\u0026lt;T\u0026gt;; StoredType value; Storage(T\u0026amp;\u0026amp; v) : value(std::forward\u0026lt;T\u0026gt;(v)) {}};5.2.3. std::conditional// ===== conditional: 编译期三目运算符 =====std::conditional\u0026lt;true, int, double\u0026gt;::type // → intstd::conditional\u0026lt;false, int, double\u0026gt;::type // → double// --- 实用示例: 根据大小选计算类型 ---template \u0026lt;typename T\u0026gt;struct ComputeTypeSelector { using type = std::conditional_t\u0026lt; (sizeof(T) \u0026lt;= 2), // 半精度/int8 等小类型 float, // → 提升到 float 计算 T // → 原类型计算 \u0026gt;;};// ComputeTypeSelector\u0026lt;half\u0026gt;::type → float// ComputeTypeSelector\u0026lt;float\u0026gt;::type → float// ComputeTypeSelector\u0026lt;double\u0026gt;::type → double5.2.4. common_type// 求多个类型的\"公共类型\" (类似三目运算符的推导)std::common_type\u0026lt;int, double\u0026gt;::type // → doublestd::common_type\u0026lt;int, long, float\u0026gt;::type // → floatstd::common_type\u0026lt;int, unsigned int\u0026gt;::type // → unsigned int// --- 实用示例: 通用 max ---template \u0026lt;typename T, typename U\u0026gt;std::common_type_t\u0026lt;T, U\u0026gt; generic_max(T a, U b) { return a \u0026gt; b ? a : b;}generic_max(1, 2.5); // 返回 doublegeneric_max(1L, 2); // 返回 long5.3. SFINAE 相关工具5.3.1. std::enable_if// 原理:// enable_if\u0026lt;true, T\u0026gt;::type = T// enable_if\u0026lt;false, T\u0026gt;::type = 不存在 → SFINAE// ---- 用法1: 放在返回类型 ----template \u0026lt;typename T\u0026gt;std::enable_if_t\u0026lt;std::is_integral_v\u0026lt;T\u0026gt;, T\u0026gt;bit_count(T val) { return __builtin_popcount(val);}// ---- 用法2: 放在模板参数默认值 (更简洁) ----template \u0026lt;typename T, std::enable_if_t\u0026lt;std::is_integral_v\u0026lt;T\u0026gt;, int\u0026gt; = 0\u0026gt;T bit_count_v2(T val) { return __builtin_popcount(val);}// ---- 用法3: 放在函数参数 ----template \u0026lt;typename T\u0026gt;T bit_count_v3(T val, std::enable_if_t\u0026lt;std::is_integral_v\u0026lt;T\u0026gt;\u0026gt;* = nullptr) { return __builtin_popcount(val);}5.3.2. std::void_t// void_t\u0026lt;Ts...\u0026gt; = void (无论 Ts 是什么)// 但如果 Ts 中有无效类型 → 替换失败 → SFINAE// 检测是否有 .size() 方法template \u0026lt;typename T, typename = void\u0026gt;struct has_size : std::false_type {};template \u0026lt;typename T\u0026gt;struct has_size\u0026lt;T, std::void_t\u0026lt;decltype(std::declval\u0026lt;T\u0026gt;().size())\u0026gt;\u0026gt; : std::true_type {};// 检测是否有嵌套 value_typetemplate \u0026lt;typename T, typename = void\u0026gt;struct has_value_type : std::false_type {};template \u0026lt;typename T\u0026gt;struct has_value_type\u0026lt;T, std::void_t\u0026lt;typename T::value_type\u0026gt;\u0026gt; : std::true_type {};// 检测是否支持 \u0026lt;\u0026lt; 输出template \u0026lt;typename T, typename = void\u0026gt;struct is_printable : std::false_type {};template \u0026lt;typename T\u0026gt;struct is_printable\u0026lt;T, std::void_t\u0026lt; decltype(std::declval\u0026lt;std::ostream\u0026amp;\u0026gt;() \u0026lt;\u0026lt; std::declval\u0026lt;T\u0026gt;())\u0026gt;\u0026gt; : std::true_type {};5.3.3. std::declval// std::declval\u0026lt;T\u0026gt;() 在不构造对象的情况下，\"假装\"有一个 T 类型的值// 只能在 decltype / sizeof 等不求值上下文中使用// 用途: 探测表达式是否合法template \u0026lt;typename T, typename U\u0026gt;using add_result_t = decltype(std::declval\u0026lt;T\u0026gt;() + std::declval\u0026lt;U\u0026gt;());// add_result_t\u0026lt;int, double\u0026gt; → double// add_result_t\u0026lt;string, string\u0026gt; → string// add_result_t\u0026lt;int, string\u0026gt; → 替换失败// 也可以检测成员函数返回类型template \u0026lt;typename Container\u0026gt;using iterator_t = decltype(std::declval\u0026lt;Container\u0026gt;().begin());// iterator_t\u0026lt;std::vector\u0026lt;int\u0026gt;\u0026gt; → std::vector\u0026lt;int\u0026gt;::iterator5.4. 编译期计算工具5.4.1. std::integral_constant// integral_constant: 将编译期常量包装成类型std::integral_constant\u0026lt;int, 42\u0026gt; // ::value = 42, ::type = 自身// 最常见的两个特化:std::true_type // = integral_constant\u0026lt;bool, true\u0026gt;std::false_type // = integral_constant\u0026lt;bool, false\u0026gt;// --- 实用: 自定义 trait 继承它们 ---template \u0026lt;typename T\u0026gt;struct is_gpu_type : std::false_type {}; // 默认 falsetemplate \u0026lt;\u0026gt;struct is_gpu_type\u0026lt;__half\u0026gt; : std::true_type {}; // 特化为 truetemplate \u0026lt;\u0026gt;struct is_gpu_type\u0026lt;__nv_bfloat16\u0026gt; : std::true_type {};// 使用:static_assert(is_gpu_type\u0026lt;__half\u0026gt;::value);if constexpr (is_gpu_type\u0026lt;T\u0026gt;::value) { /* ... */ }5.4.2. std::integer_sequence// 生成编译期整数序列std::integer_sequence\u0026lt;int, 0, 1, 2, 3, 4\u0026gt;std::index_sequence\u0026lt;0, 1, 2, 3, 4\u0026gt; // 简写 (size_t)std::make_index_sequence\u0026lt;5\u0026gt; // 自动生成 0,1,2,3,4// --- 实用: 展开 tuple ---template \u0026lt;typename Tuple, size_t... Is\u0026gt;void print_tuple_impl(const Tuple\u0026amp; t, std::index_sequence\u0026lt;Is...\u0026gt;) { ((std::cout \u0026lt;\u0026lt; std::get\u0026lt;Is\u0026gt;(t) \u0026lt;\u0026lt; \" \"), ...); // fold expression}template \u0026lt;typename... Ts\u0026gt;void print_tuple(const std::tuple\u0026lt;Ts...\u0026gt;\u0026amp; t) { print_tuple_impl(t, std::make_index_sequence\u0026lt;sizeof...(Ts)\u0026gt;{});}print_tuple(std::make_tuple(1, 3.14, \"hello\"));// 输出: 1 3.14 hello// --- 实用: 编译期循环 ---template \u0026lt;size_t... Is\u0026gt;void unrolled_copy(float* dst, const float* src, std::index_sequence\u0026lt;Is...\u0026gt;) { ((dst[Is] = src[Is]), ...); // 展开成 dst[0]=src[0]; dst[1]=src[1]; ...}5.4.3. if constexpr (C++17)// 编译期 if — 不满足的分支完全不编译 — 替代大量 enable_iftemplate \u0026lt;typename T\u0026gt;auto convert(T val) { if constexpr (std::is_same_v\u0026lt;T, std::string\u0026gt;) { return std::stoi(val); // string → int } else if constexpr (std::is_floating_point_v\u0026lt;T\u0026gt;) { return static_cast\u0026lt;int\u0026gt;(val); // float → int } else if constexpr (std::is_integral_v\u0026lt;T\u0026gt;) { return val; // int → int } else { static_assert(always_false\u0026lt;T\u0026gt;, \"Unsupported type\"); }}// 辅助: 永远为 false 的依赖模板 (防止 static_assert 直接触发)template \u0026lt;typename\u0026gt; constexpr bool always_false = false;5.6. 常用场景速查表 我想做什么 用什么 判断是不是整数/浮点 is_integral_v, is_floating_point_v 判断两个类型是否相同 is_same_v\u0026lt;T, U\u0026gt; 去掉 const/引用 remove_cv_t\u0026lt;remove_reference_t\u0026lt;T\u0026gt;\u0026gt; 或 decay_t 根据条件选类型 conditional_t\u0026lt;cond, A, B\u0026gt; 求公共类型 common_type_t\u0026lt;T, U\u0026gt; 按条件启用/禁用重载 enable_if_t (C++11) / if constexpr (C++17) / concept (C++20) 检测类型有没有某成员 void_t + 偏特化 / is_detected 检测表达式是否合法 decltype(expr) + declval 判断能不能用 memcpy is_trivially_copyable_v 判断是不是空基类(EBO) is_empty_v 判断能不能调用 is_invocable_v\u0026lt;F, Args...\u0026gt; 编译期循环/展开 index_sequence + fold expression 自定义 bool trait 继承 true_type / false_type 有符号↔无符号 make_signed_t / make_unsigned_t 组合多个条件 conjunction_v / disjunction_v (C++17) A. 资料 C++ Traits Classes How do we use void_t for SFINAE?" },

  { "title": "道格拉斯-普克算法(Douglas–Peucker algorithm)", "url": "/blog/2025/%E9%81%93%E6%A0%BC%E6%8B%89%E6%96%AF-%E6%99%AE%E5%85%8B%E7%AE%97%E6%B3%95/", "categories": "Algorithm",
  "tags": "Algorithm", "date": "2025-12-28", "content":
  "道格拉斯-普克算法（Douglas–Peucker algorithm），又称为Ramer-Douglas-Peucker算法，是一种用于简化曲线的算法。它通过减少曲线上的点数来近似表示原始曲线，同时尽量保持其形状和特征。可以应用于计算机图形学、地理信息系统（GIS）、甚至CAD领域。算法的具体实现逻辑如下： 在轨迹曲线在曲线首尾两点A，B之间连接一条直线AB，该直线为曲线的弦； 遍历曲线上其他所有点，求每个点到直线AB的距离，找到最大距离的点C，最大距离记为maxDistance 比较该距离maxDistance与预先定义的阈值epsilon大小，如果maxDistance \u0026lt; epsilon，则将该直线AB作为曲线段的近似，舍去AB之间的所有点，曲线段处理完毕； 若maxDistance \u0026gt;= epsilon，则使C点将曲线AB分为AC和CB两段，并分别对这两段进行（1）~（3）步处理； 当所有曲线都处理完毕时，依次连接各个分割点形成的折线，即为原始曲线的路径。" },

  { "title": "CMake支持库收集", "url": "/blog/2025/CMake%E6%94%AF%E6%8C%81%E5%BA%93/", "categories": "CMake",
  "tags": "CMake", "date": "2025-12-23", "content":
  "CMake Toolchain files for Windows CMake project_options" },

  { "title": "QGC代码架构解析：飞行前检查（起飞条件）", "url": "/blog/2025/QGC%E9%A3%9E%E8%A1%8C%E5%89%8D%E6%A3%80%E6%9F%A5/", "categories": "QGC",
  "tags": "QGC", "date": "2025-12-18", "content":
  "QGC中，飞行状态指示器（Flight Status Indicator）显示的状态列表： 状态名称 说明 准备好飞速 (_绿色背景) 载具已准备就绪，可以起飞 准备好飞行 (_黄色背景) 载具已准备好在当前飞行模式下飞行，但有些警告可能造成问题 尚未准备好 载具没有准备好飞行，也不会起飞 解锁 载具已解锁并准备起飞。 飞行 载具正在飞行 着陆 载具正在着陆 通信丢失 QGC已失去与载具的通信 更多信息，参考飞行视图工具栏。1. 解锁检查项 气压计（Barometer） 指南针（Compass） GPS 锁定（GPS lock） 惯性导航系统（INS） 参数（Parameters） 遥控通道（RC Channels） 电路板电压（Board voltage） 电池电量（Battery Level） 空速（Airspeed） 日志记录可用（Logging Available） 硬件安全开关（Hardware safety switch） GPS 配置（GPS Configuration） 系统（System） ：Safety Setup (ArduPilot)检查实现函数：HealthAndArmingCheckReport::update。QGC提供了可选的飞行前检查清单功能：Pre Flight Checklist。备注： 通常情况下，不需要手动解锁飞机。简单地起飞或开始执行任务时，飞行器会自动解锁。 不同的飞行器类型（多旋翼、固定翼、地面车、潜水器）可能有略微不同的解锁检查项，但核心检查项是相似的。2. MAV_MODEMAV_MODE枚举定义了飞行器的标准飞行模式。不同的飞行模式决定了飞行器的行为和控制方式，设置命令是MAV_CMD_DO_SET_MODE。// common/common.htypedef enum MAV_MODE{ MAV_MODE_PREFLIGHT=0, /* System is not ready to fly, booting, calibrating, etc. No flag is set. | */ MAV_MODE_MANUAL_DISARMED=64, /* System is allowed to be active, under manual (RC) control, no stabilization | */ MAV_MODE_TEST_DISARMED=66, /* UNDEFINED mode. This solely depends on the autopilot - use with caution, intended for developers only. | */ MAV_MODE_STABILIZE_DISARMED=80, /* System is allowed to be active, under assisted RC control. | */ MAV_MODE_GUIDED_DISARMED=88, /* System is allowed to be active, under autonomous control, manual setpoint | */ MAV_MODE_AUTO_DISARMED=92, /* System is allowed to be active, under autonomous control and navigation (the trajectory is decided onboard and not pre-programmed by waypoints) | */ MAV_MODE_MANUAL_ARMED=192, /* System is allowed to be active, under manual (RC) control, no stabilization | */ MAV_MODE_TEST_ARMED=194, /* UNDEFINED mode. This solely depends on the autopilot - use with caution, intended for developers only. | */ MAV_MODE_STABILIZE_ARMED=208, /* System is allowed to be active, under assisted RC control. | */ MAV_MODE_GUIDED_ARMED=216, /* System is allowed to be active, under autonomous control, manual setpoint | */ MAV_MODE_AUTO_ARMED=220, /* System is allowed to be active, under autonomous control and navigation (the trajectory is decided onboard and not pre-programmed by waypoints) | */ MAV_MODE_ENUM_END=221, /* | */} MAV_MODE;结合文章QGC代码架构解析：MAVLink Mission Protocol，以及 QGC 航点管理查看MAV_CMD相关命令。" },

  { "title": "QGC代码架构解析：MAVLink Mission Protocol，以及 QGC 航点管理", "url": "/blog/2025/qgc_mission_manager/", "categories": "QGC",
  "tags": "QGC", "date": "2025-12-15", "content":
  "本文厘清如下几个问题： 航点协议（Mission Protocol）有哪些命令/消息； 微服务的处理流程，以及流程中通信双方使用到的命令ID/请求ID-响应ID； 带参数的命令/消息，其参数格式，以及单位； 航点文件格式。1. MAVLink Mission Protocol航点协议主要实现航点集合的上传、下载、清除。另外有一些其他消息，以及附带的参数/枚举。航点上传/下载使用到如下几个消息定义： 消息名称 说明 发送方 MISSION_REQUEST_LIST 启动航点集下载动作 地面站 MISSION_COUNT 航点数量 地面站/飞控 MISSION_REQUEST_INT 请求航点 地面站/飞控 MISSION_ITEM_INT 请求航点 飞控/地面站 MISSION_ACK 航点响应 飞控/地面站 1.1. 航点的上传/下载流程航点上传/下载流程图（左边为QGC请求下载，右边为QGC请求上传）：对比发现规律： QGC请求下载时，使用MISSION_REQUEST_LIST来启动，而请求上传时，使用MISSION_COUNT来启动（注意：不论上传/下载，都是QGC发起的）； MISSION_COUNT既可以作为飞控响应QGC的下载请求，也可以作为QGC请求飞控上传航点的启动消息； 启动之后，中间航点的请求，双方使用MISSION_REQUEST_INT \u0026lt;–\u0026gt; MISSION_ITEM_INT配对，即一个完整的单个航点传输流程。航点发送方使用MISSION_REQUEST_INT请求，接收方使用MISSION_ITEM_INT响应。 最后使用MISSION_ACK作为结束标示。在这里，ACK消息的含义是结束，即整个流程结束了，这样理解更合理。 从文档看，MISSION_ACK只用在上传/下载流程里面。其他消息里面没有使用到。上传/下载的流程，不具有对称性，给理解带来了一定的混乱。流程步骤在语义上理解也不顺畅。1.1.1. 协议实现注意事项 在流程图的Start timeout处，需要实现超时重传机制： MISSION_REQUEST_LIST超时没有响应，重传该命令若干次； MISSION_COUNT超时没有响应，重传该命令若干次； MISSION_REQUEST_INT超时没有响应，重传该命令若干次； MISSION_REQUEST_INT请求的航点，需要按序号顺序请求以及应答。如果收到的MISSION_ITEM_INT中的航点顺序不对，需要丢弃该航点数据，并重新请求。 当上传/下载航点时，飞机会返回一个opaque_id（类似整个航点集合计算得到的哈希），用于避免不不必要的再次上传/下载。当QGC上传时，飞机在最后一个消息返回该值：MISSION_ACK.opaque_id。当QGC下载时，飞机在MISSION_COUNT中返回该值：MISSION_COUNT.opaque_id。 当上传/下载的过程中失败时（对方提前返回MISSION_ACK并包含错误消息），QGC或者飞机应该终止当前流程，并恢复使用上一次的航点集。 协议没有说明，上传过程中，最后一步，如果飞机没有返回MISSION_ACK，应该如何处理。QGC的实际处理方式是，认为上传成功并完成，但是opaque_id没有更新。1.2. 航点消息结构：消息 MISSION_ITEM_INT，以及 MAV_CMD航点不仅仅只有坐标等数据，还包含动作含义，即MAV_CMD其实是一个命令+参数数据。MISSION_ITEM_INT就是用于发送这些MAV_CMD子命令+参数的。MAV_CMD分为如下几类： MAV_CMD_NAV_*：导航类命令（起飞、降落、返回RTL、悬停、飞到指定航点位置），比如MAV_CMD_NAV_WAYPOINT表示普通航点，MAV_CMD_NAV_LOITER_UNLIM表示无限悬停等。 MAV_CMD_DO_*：动作类命令，比如MAV_CMD_DO_CHANGE_SPEED表示改变速度，MAV_CMD_DO_SET_RELAY表示设置继电器等； MAV_CMD_CONDITION_*：命令执行条件，比如MAV_CMD_CONDITION_DELAY表示等待一段时间之后，再执行下一个航点MAV_CMD。从Ardupilot文档看，MAV_CMD_CONDITION_*命令是作用于MAV_CMD_DO_*命令，参考Ardupilot – Mission Commands – Conditional commands。主要的MAV_CMD_NAV命令列表： Command ID Name Description 16 MAV_CMD_NAV_WAYPOINT Navigate to a specific waypoint 21 MAV_CMD_NAV_LAND Land at the specified location 22 MAV_CMD_NAV_TAKEOFF Take off and ascend to specified altitude 20 MAV_CMD_NAV_RETURN_TO_LAUNCH Return to launch/home location 80 MAV_CMD_NAV_ROI Sets region of interest for camera 82 MAV_CMD_NAV_SPLINE_WAYPOINT Navigate using a spline path 所有MAV_CMD命令的完整列表，以及参数，可以参考MAVLink协议文档：Commands (MAV_CMD)。1.3. 航点命令的参数：MISSION_ITEM_INT的参数：Frame（坐标系）在使用MISSION_ITEM_INT消息发送航点命令(Mission Item)时（包括MAV_CMD_NAV_*命令，以及MAV_CMD_DO_*命令），需要指定坐标系frame，比如WGS84坐标系，NED坐标系，或者在WGS84坐标系的修改，如高度改为相对HOME点高度，或者地形高度。坐标系枚举定义见：MAV_FRAME。官方文档整理下来，有关Frame的描述，没有完全讲清楚，整个MAVLink协议中，有若干个命令使用到Frame作为参数。常见的需要使用到Frame的命令： Mission Protocol中少量MAV_CMD_DO。 COMMAND_INT需要使用Frame作为命令参数：Command Protocol。根据文档描述，针对MISSION_ITEM_INT，目前APM以及PX4仅支持global类型的Frame，见文档描述Mission Items (MAVLink Commands)。另外，有少量的MAV_CMD_DO命令，里面也带有frame参数，很奇怪，有些混乱。更多信息： MAVLink – Frames \u0026amp; Positional Information Ardupilot – Navigation commands MAVLink – Commands (MAV_CMD) MAVLink – MISSION_ITEM_INT (73)1.4. 航点命令的参数：MISSION_ITEM_INT的参数：param1 ~ param7前四个参数param1 ~ param4，具体含义，以及单位，由具体的MAV_CMD命令决定，且数据类型就是float，即如果是其他类型，需要static_cast转换为float。针对MAV_CMD_NAV命令，则是坐标信息，其中param5、param6（经纬度）是全局坐标，即1e7。高度参数param7，其含义由frame指定（但应该全部都是global类型的），global坐标系列表： MAV_FRAME_GLOBAL_INT， MAV_FRAME_GLOBAL_RELATIVE_ALT_INT， MAV_FRAME_GLOBAL_TERRAIN_ALT， MAV_FRAME_GLOBAL_TERRAIN_ALT_INT， MAV_FRAME_GLOBAL。另外，如果frame是MAV_FRAME_MISSION，表示这param5 ~ param7不是坐标，所以实现发送/接收的时候，不需要将param5 ~ param7乘以1e7。根据协议，如果是frame的类型是local的，则发送/接收的时候param5、param6的精度应该是1e4，单位是米（参考协议文档Frames \u0026amp; Positional Information）。总结：协议实现的时候： 如果frame是MAV_FRAME_MISSION：则param5、param6按原样发送值，可能需要static_cast\u0026lt;int\u0026gt;转换。 如果frame是global类型的：则param5、param6需要乘以1e7进行发送，接收时需要除以1e7。 如果frame是local类型的：则param5、param6需要乘以1e4进行发送，接收时需要除以1e4。 其他param1 ~ param4，以及param7，按原样发送/接收，可能需要static_cast\u0026lt;float\u0026gt;转换。另外，这个规则，应该也适用于COMMAND_INT命令COMMAND_INT。且： 使用COMMAND_INT命令时，有些参数没有使用到，部分样例参考：ardupilot – Commands supported by Copter 这些浮点类型的参数，值nan也有有意义的，比如维持原值，具体搜索QGC代码中qQNaN()的使用地方。1.5. 航点管理中其他命令/消息如下两个消息，用来监控航点执行进度及状态： MISSION_CURRENT：当前航点改变通知消息，由飞控广播，其他信息：序号seq，一起当前飞机航点状态，比如是否是暂停等。 MISSION_ITEM_REACHED：与MISSION_CURRENT类似，QGC中没有处理该消息。清除航点使用消息：MISSION_CLEAR_ALL。2. QGC 航点管理实现PlanManager实现Mission Protocol的协议。由于MAVLink v2中将航点(flight plans)、地理围栏(geofences)、降落点(rally/safe points)都放在Mission Protocol中，在请求上传(MISSION_COUNT)/下载(MISSION_REQUEST_INT)，以及传输(MISSION_ITEM_INT)时，带有MAV_MISSION_TYPE字段，确定是哪种Mission Type，参考协议文档：Mission Protocol – Mission Types。 MissionManager：拓展航点的协议MAV_MISSION_TYPE_MISSION的固件实现：主要实现Ardupilot相关的实现，以及一些功能。 GeoFenceManager：实现MAV_MISSION_TYPE_FENCE。 RallyPointManager：实现MAV_MISSION_TYPE_RALLY。2.1. 航点协议上传/下载功能的实现PlanManager使用状态机实现上传/下载流程。定义一个TransactionType_t枚举，表示当前的事务类型，以及一个AckType_t表示期望的ACK类型：typedef enum { AckNone, ///\u0026lt; State machine is idle AckMissionCount, ///\u0026lt; MISSION_COUNT message expected AckMissionItem, ///\u0026lt; MISSION_ITEM expected AckMissionRequest, ///\u0026lt; MISSION_REQUEST is expected, or MISSION_ACK to end sequence AckMissionClearAll, ///\u0026lt; MISSION_CLEAR_ALL sent, MISSION_ACK is expected AckGuidedItem, ///\u0026lt; MISSION_ACK expected in response to ArduPilot guided mode single item send} AckType_t;typedef enum { TransactionNone, TransactionRead, TransactionWrite, TransactionRemoveAll} TransactionType_t;请求下载入口函数：PlanManager::loadFromVehicle，请求上传入口函数：PlanManager::writeMissionItems(const QList\u0026lt;MissionItem*\u0026gt;\u0026amp; missionItems)。以及一个handle函数，用于处理收到的消息： PlanManager::_handleMissionCount：处理MISSION_COUNT消息（请求下载）； PlanManager::_handleMissionItem：处理MISSION_ITEM_INT消息（请求下载）； PlanManager::_handleMissionRequestInt：处理MISSION_REQUEST_INT消息（请求上传）；PlanManager中，有两个函数处理ACK消息：_handleMissionAck，_ackTimeout，这两个函数驱动状态机的流转：判断ACK与请求步骤是否匹配，发送一下一个航点数据/请求下一个航点数据，以及结束流程。2.2. 其他模块 MissionItem：表示单个航点数据结构，航点管理模块的基础数据类。 MissionController：提供QML接口访问航点数据。 PlanMasterController：航点管理模块的顶层控制类（入口），提供QML接口访问航点（MissionController）、电子围栏（GeoFenceController）、降落点（RallyPointController）。以及从文件中加载/保存（包含kml文件格式）。 MissionCommandTree：提供MAV_CMD命令树结构，供QML界面使用。2.3. 航点文件格式QGC中，航点文件格式使用JSON格式，文件扩展名为.plan，保存目录样例：C:\\Users\\Administrator\\Documents\\QGroundControl Daily\\Missions保存时，将航点数据、电子围栏数据、降落点数据，保存到同一个文件中。保存及加载函数入口：PlanMasterController::saveToJson();MissionController::save(QJsonObject\u0026amp; json);// 加载实现bool _loadJsonMissionFileV2(const QJsonObject\u0026amp; json, QmlObjectListModel* visualItems, QString\u0026amp; errorString);另外还可以导出/导入kml格式的航点文件，记录的字段格式有所不同，另外测试发现加载有些BUG。航点有SimpleItem，以及ComplexItem，一个简单的航点文件样例：{ \"fileType\": \"Plan\", \"version\": 1, \"groundStation\": \"QGroundControl\", \"mission\": { \"cruiseSpeed\": 5, \"hoverSpeed\": 3, \"items\": [ { \"AMSLAltAboveTerrain\": null, \"autoContinue\": true, \"command\": 16, \"frame\": 3, \"params\": [0, 0, 0, null, 47.397742, 8.545594, 488], \"type\": \"SimpleItem\" }, { \"AMSLAltAboveTerrain\": null, \"autoContinue\": true, \"command\": 16, \"frame\": 3, \"params\": [0, 0, 0, null, 47.397825, 8.545632, 488], \"type\": \"SimpleItem\" } ] }}有关航点文件格式的更多信息，参考QGC代码仓库文档： Mission Command Tree Plan File Format3. 参考文档 MAVLink Mission Protocol Ardupilot – Mission Commands" },

  { "title": "AI工具收集及使用笔记，持续更新", "url": "/blog/2025/AI%E5%B7%A5%E5%85%B7%E6%94%B6%E9%9B%86%E5%8F%8A%E4%BD%BF%E7%94%A8%E7%AC%94%E8%AE%B0/", "categories": "AI",
  "tags": "AI", "date": "2025-12-14", "content":
  "为了让内容更条理清晰，本文对原有的 AI 工具和笔记进行了重新分类和整理。1. AI 编程辅助与 Agent 工具Pi Agentpi agent：第三方开源agent，使用CLI/TUI界面交互，可以接入市面上LLM供应商。github – pi。使用/login选择厂商并设置API Key，使用/scoped-models选进或者剔除模型列表，使用/model选择正在使用的模型，/resume恢复上次的会话，命令列表：Slash Commands。有Packages以及Models扩展安装。Pi Agent 增加 GLM-5.2 模型GLM-5.2是一个多模态大模型，上下文长度1M token。创建并编辑~/.pi/agent/models.json，内容如下：{ \"providers\": { \"zai\": { \"api\": \"openai-completions\", \"apiKey\": \"\u0026lt;YOUR_ZAI_API_KEY\u0026gt;\", \"baseUrl\": \"https://api.z.ai/api/coding/paas/v4\", \"compat\": { \"supportsDeveloperRole\": false, \"supportsReasoningEffort\": false }, \"models\": [ { \"id\": \"glm-5.2\", \"name\": \"GLM-5.2\", \"contextWindow\": 200000, \"maxTokens\": 131072, \"reasoning\": true }, { \"id\": \"glm-4.6\", \"name\": \"GLM-4.6\", \"contextWindow\": 200000, \"maxTokens\": 131072, \"reasoning\": true } ] } }} 可以在~/.bashrc添加环境变量ZAI_API_KEY以设置API Key，并在~/.pi/agent/models.json中使用${ZAI_API_KEY}替代\u0026lt;YOUR_ZAI_API_KEY\u0026gt;。查看模型列表：pi --list-models | grep -i zai可以在~/.pi/agent/models.json设置GLM-5.2为默认模型:\"defaultProvider\": \"zai\",\"defaultModel\": \"glm-5.2\" zai_pi.md：参考资料，给 Pi Agent 添加 GLM 模型。Pi Agent 增加本地部署的 gemma 4 模型参见另外一篇记录的博客NVIDIA Jetson Orin AGX 安装，以及配套的docker compose仓库：Gemma-4 Local Servers on Jetson Orin。另外，也可以在Pi agent中添加局域网部署的Gemma 4 模型，即在~/.pi/agent/models.json中进行配置，具体请参考代码仓库Gemma-4 Local Servers on Jetson Orin。Antigravity agy CLI（原Gemini CLI）Gemini CLI 是基于 Node.js 的，Antigravity agy CLI 改为使用 Go 实现，速度快很多。安装（安装完成之后可能需要重新启动终端）：# macOS/Linuxcurl -fsSL https://antigravity.google/cli/install.sh | bash# Windows PowerShellirm https://antigravity.google/cli/install.ps1 | iex# 测试是否可用agy --version# 启动agy进入交互式界面agy# 去掉权限提示，直接操作沙箱外的文件系统agy --dangerously-skip-permissions为了方便记忆，可以在~/.bashrc中添加alias：# AGY 别名 (免确认权限运行)alias dagy='agy --dangerously-skip-permissions'alias agy-skip='agy --dangerously-skip-permissions'认证及 VPN 代理设置首次启动，需要登录并认证。每次启动agy都会进行区域检测，针对国内区域有限制，认证阶段agy会提示Google OAuth失败，需要打开科学上网工具并设置代理，或者在~/.bashrc中添加代理配置：# Linux / macOSexport http_proxy=\"http://192.168.11.139:7890\"export https_proxy=\"http://192.168.11.139:7890\"# Windows (PowerShell)$env:http_proxy=\"http://192.168.11.139:7890\"$env:https_proxy=\"http://192.168.11.139:7890\"# 验证代理是否生效curl -I https://oauth2.googleapis.com 设置代理之后，认证完成之后，进入agy交互，还是会提示”Eligibility check failed: Your current account is not eligible for Antigravity, because it is not currently available in your location.”。不过似乎不影响使用。可能使用VPN的TUN模式（增强模式/虚拟网卡模式）可以完全解决，不太熟悉，尝试没有成功。agy卡死之后，使用Ctrl+D强制退出，而不是Ctrl+C。可以免费试用的 LLM 有：Gemini 3.5 Flash (High / Medium)。其他几个也可以使用，的额度比较有限，比较容易出现额度已满，包括：Claude Sonnet 4.6 (Thinking)、Claude Opus 4.6 (Thinking)、GPT-OSS 120B (Medium)。其中 Claude Opus 4.6 适合复杂架构设计及深度推理。在进入 agy 之后，还是会提示”Eligibility check failed: Your current account is not eligible for Antigravity, because it is not currently available in your location.”，其原因为 Google 账号归属地限制，但是似乎不影响使用（使用免费LLM测试可用，高端LLM可能会有限制）。此时需要修改 Google Terms of Service，此时显示的”Country version”是 china 或者 Hong Kong。需要将地区修改美国/特拉华州(Delaware)，修改时，原因选择”Other reason”，并填写原因，比如：I have a valid US mailing address and intend to use Google services primarily from the United States for both personal and professional purposes. Changing my region to the US will align my account with my current and anticipated usage patterns.如果后期需要绑定支付方式，填写如下地址（咨询DeepSeek，特拉华州税比较干净）：147 S College AveNewark, DE 19711解释如下：纽瓦克市（Newark）街道地址 (Street Address)：147 S College Ave城市 (City)：Newark州 (State)：DE邮编 (ZIP Code)：19711 一招解决antigravity登录账号区域限制问题 随笔】招商银行visa金卡申请和绑定海外谷歌账户支付：绑定VISA卡获取更多免费LLM额度据说可以通过在Google AI Studio上申请API Key，获取更多额度。Google AI Studio也是一个网页版的免费平台，聊天入口：https://aistudio.google.com/prompts/new_chat，API Key申请入口在Dashboard菜单栏中。export ANTIGRAVITY_API_KEY=\"你的_AI_Studio_API_Key\"其他的获取免费LLM的方式，根据Google AI回复，有Google One Pro / Ultra 订阅。解决使用过程中 Region Not Supported 错误在使用过程中，依旧会检测用户的区域是否合法，不合法的报错”Agent execution terminated due to error”，就是因为模型其检测到区域不在美国。解决办法为在hosts文件中添加 DNS：# linux 路由文件路径为 /etc/hosts# windows 路由文件路径为 C:\\Windows\\System32\\drivers\\etc\\hosts# 添加类似如下 DNS 地址123.123.123.123 api.agy.io添加 DNS 之后，刷新 DNS 缓存：# linux 清理缓存命令sudo resolvectl flush-cachessudo resolvectl statistics # 验证缓存是否清零# 查询并测试 DNS 解析resolvectl query google.com# windows 更新 DNS 缓存信息ipconfig /flushdns参考： AGY CLI 地区限制破解实测：3 种可行方案，无需科学上网agy 命令列表常用/有用的命令：# 查看各个模型用量情况/usage# 恢复以前的会话（从列表中选择恢复的会话）/resume# 配置项目设置/config# 查看已经安装的skills/skills# 选择正在使用的模型/model# 针对复杂的任务，使用plan模式/planning# fast模式，使用完成之后，记得恢复成planning模式/fast# 查看操作键?# 退出/exit # 或者 /quit 或者 Ctrl+D Antigravity CLI Cheatsheet：比较完整的命令列表，以及操作键列表（一个简单的github仓库） Best practices for Antigravity CLI：官方资料一些第三方资源 Agent Skills：这是一个通用的 agent skills 仓库，其中就包括导入到Antigravity CLI 的技能。 Antigravity-Proxy：代理注入，操作简单，专门为国内用户设计。仅适用于Windows。资料： 亲身体验 Antigravity CLI：Google 官方的使用教程 Antigravity CLI (agy): 安装、从 Gemini CLI 迁移和解决桌面应用冲突 (2026) 谷歌 Antigravity CLI 踩坑与破局全指南（Gemini CLI 替代品） Antigravity 登录失败怎么办？从 Account not eligible 到 oauth-success 的完整排查 Gemini 网页版CodeX安装 CodeX CLI 版本：# set proxy: linux \u0026amp; macOSexport http_proxy=\"http://192.168.11.139:7890\"export https_proxy=\"http://192.168.11.139:7890\"# set proxy: windows$env:HTTP_PROXY=\"http://127.0.0.1:7890\"$env:HTTPS_PROXY=\"http://127.0.0.1:7890\"# 推荐使用本地安装的方式，即不使用npm安装，不需要nodejs环境# 同时更新也使用下面同一条命令curl -fsSL https://chatgpt.com/codex/install.sh | sh# 验证：查看版本codex -V# 定义alias，禁用沙箱模式alias dcodex='codex --yolo'安装过程中，需要验证手机号，使用接码平台：SMS-Activation短信接收服务，该平台是收费的，且要求开启的VPN代理区域与号码所在区域一致，例如都在美国。 CodeX 桌面版下载地址：Codex app。不论CodeX CLI，还是桌面版，其都需要稳定的VPN。设置 CodeX 的代理：要分别在两个文件中添加设置，针对Windows，分别是：C:\\Users\\Administrator\\.codex\\.env，C:\\Users\\Administrator\\.codex\\config.toml，其中config.toml是 CodeX 的配置文件。增加的配置如下：添加 API Key：OpenAI有两种计费方式：API Key，以及按类似订阅登录的方式，参考Codex / GPT-5 国内怎么用?ChatGPT Plus 一键登录、额度与限制说明(程序员向,2026)。另一个讲解搭建 CodeX 比较清楚的博客：OpenAI API Key 获取完整教程：国内开发者AI编程与GPT-5.4-Codex集成实战。主要有：访问官网并注册/登录 -\u0026gt; 申请 API Keys页面 -\u0026gt; 开通计费（Billing）。另外，可以使用合规中转/聚合API平台（如UIUI API）。Open AI 官方申请 API Key 以及绑定支付 VISA 卡网页：https://platform.openai.com/home。可以购买虚拟 VISA 卡，搜索虚拟visa卡申请查找购买网站。或者使用第三方网站购买 Open AI，参见上述的Codex / GPT-5 国内怎么用?ChatGPT Plus 一键登录、额度与限制说明(程序员向,2026)，其中一个 ChatGPT代购网站：Glouth。获取到 API Key 之后，添加系统环境变量OPENAI_API_KEY。 经过实践，需要完成billing，即绑定VISA卡（也可能需要付费），才可以使用 CodeX。ChatGPT的Credit额度分为API Key额度，以及登录订阅的额度。登录订阅有免费的额度（20260807，GPT-5.6 luna），但是需要先执行codex login的方式 -\u0026gt; 打开链接 -\u0026gt; 账号登录认证。对于远程登录，则使用codex login --device-auth进行登录认证。认证完成之后，后面就可以使用codex登录并使用了。一些搜集的资料： github – OpenAI Codex: The Complete Guide：OpenAI Codex: The Complete Guide · 橙皮书系列 · GPT-5.5 时代的 AI 编程实战手册 Codex PPT SkillOpenCode 以及 GLM-5.2OpenCode CLI 安装命令：# 设置国内镜像，以及全局安装npm config set registry https://registry.npmjs.org/npm install -g opencode-ai# 检查安装是否成功opencode --versionOpenCode 常用命令 分类 命令 说明 初始化 /init 项目初始化，自动扫描项目目录、识别技术栈、生成项目配置文件 文件 /open 打开本地文件进行编辑分析   /save 保存当前修改的代码   /list 查看当前项目所有文件   /run 直接运行当前代码文件 代码优化 /debug 自动排查代码报错、修复BUG   /refactor 一键重构代码、精简冗余逻辑、规范代码格式   /test 自动生成对应单元测试用例 会话管理 /new 新建代码文件、初始化项目模板   /clear 清空当前对话记录   /help 查看全部内置命令   /exit 退出 OpenCode 模型 /connect [provider] 连接模型提供商   /model [model-name] 切换使用的模型   /models 查看可用模型列表   /mode plan 切换到规划模式   /mode build 切换到构建模式   /compact 压缩上下文，减少 token 使用 OpenCode 使用技巧# 引用单个文件# e.g. 帮我分析这段代码@file src/main.py# 引用多个文件# e.g. 对比这两个文件的差异@file src/app.js src/utils.js# 引用特定行# e.g. 优化这段代码的性能@file src/app.js:10-30# 引用函数# e.g. 重构这个验证函数@function UserService.validate# 引用目录# e.g. 分析这些组件的设计模式@dir src/componentsOpenCode 资料 OpenCode download：OpenCode 还提供桌面版APP（测试版） OpenCode 中文Prime agent：自进化 agent# 安装 Prime agentcurl -fsSL https://app.primeintellect.ai/prime-agent/install.sh | sh# 启动prime-agentCopilot 资源及使用笔记Copilot 中的 Instructions、Prompts、SkillsInstructions、Prompts、Skills 是 Copilot 中的三个重要概念，分别用于指导模型行为、提供输入提示和定义特定功能，他们之间的区别在于： Instructions = 「你必須這樣做」（規則） Prompts = 「幫我做這件事」（範本） Skills = 「遇到這類問題時請這樣處理」（專家知識）Copilot 添加Qt的 Instructions、Prompts、SkillsQt官方在github上的一个仓库，提供了相关的资源：Qt AI Skills，可以放在项目中使用、也可以放在用户目录中作为全局配置使用。另外，这套Skills可以接入几乎所有工具中，比如Claude Code、CodeX等等十几种。在安装的时候，可以复选创建哪些工具的配置。在Windows系统上，使用方法如下：# 首先下载并安装 Copilot CLI，在 Copilot CLI 中使用如下# Copilot CLI地址：https://github.com/github/copilot-cligh skill install TheQtCompanyRnD/agent-skills qt-qml-review在Linux系统上，使用方法如下：git clone https://github.com/TheQtCompanyRnD/agent-skills.gitcp -r skills/qt-qml-review ~/.copilot/skills/qt-qml-reviewChrome 中解除 Gemini 地区限制关闭掉所有 Chrome 进程之后，备份之后，使用文本编辑器打开并修改文件C:\\Users\\Administrator\\AppData\\Local\\Google\\Chrome\\User Data\\Local State：is_glic_eligible 由 false 改为 truevariations_country 由 cn 改为 us# variations_permanent_consistency_country 可能会出现在两处，注意全部替换variations_permanent_consistency_country 由 cn 改为 us保存至，再重新打开 Chrome，打开chrome://flags，搜索glic，将glic设置为Enabled。如果还不行，则再将 Chrome 语言改为英语。参考资料：如何在 Google Chrome 中强制开启 Gemini AI 侧边栏。2. AI 平台接入与模型使用使用NVIDIA免费模型通过使用claude-nvidia-proxy，将claude客户端连接到NVIDIA的API，可以使用NVIDIA提供的免费模型。其流程为：claude 客户端 → claude-nvidia-proxy → NVIDIA API → 模型推理结果返回给claude客户端。** 安装claude客户端# 标准安装方式（linux）curl -fsSL https://claude.ai/install.sh | bash# 标准安装方式（windows）irm https://claude.ai/install.ps1 | iex# 国内安装，使用淘宝镜像sudo npm install -g @anthropic-ai/claude-code --registry=https://registry.npmmirror.com# NPM卸载sudo npm uninstall -g @anthropic-ai/claude-code新增并编辑claude的配置文件，添加字段跳过官方登陆引导：{ \"skipOnboarding\": true} Windows: C:\\Users\\你的用户名\\.claude.json macOS/Linux: ~/.claude.json测试claude：# 进入claude交互式界面claude# 选择模型/model**获取NVIDIA API密钥登陆https://build.nvidia.com/explore/discover，注册账号并获取API密钥。密钥期限为半年，另外，可以多次获取密钥，重复申请，之前的密钥将失效，最后一个有效。**编译安装及配置claude-nvidia-proxygit clone https://github.com/zhangrr/claude-nvidia-proxy.gitcd claude-nvidia-proxy# 编译安装GOOS=linux GOARCH=amd64 go install -trimpath -ldflags \"-s -w\" .在claude-nvidia-proxy的安装目录创建config.json，内容如下：{ \"nvidia_url\": \"https://integrate.api.nvidia.com/v1/chat/completions\", \"nvidia_key\": \"你的NVIDIA API密钥\"}**启动claude-nvidia-proxy及claude客户端创建一个脚本start-claude.sh，用于设置claude的环境变量并启动claude：#!/bin/zshSUPPORTED_MODELS=( \"deepseek-ai/deepseek-v3.2\" \"moonshotai/kimi-k2-thinking\" \"minimaxai/minimax-m2.7\" \"z-ai/glm4.7\")DEFAULT_MODEL=\"deepseek-ai/deepseek-v3.2\"print_models() { echo \"支持的模型列表：\" for m in \"${SUPPORTED_MODELS[@]}\"; do if [[ \"$m\" == \"$DEFAULT_MODEL\" ]]; then echo \" $m (默认)\" else echo \" $m\" fi done}if [[ \"$1\" == \"-h\" || \"$1\" == \"--help\" ]]; then echo \"用法: $0 [模型名称]\" echo \"\" print_models exit 0fiMODEL=${1:-$DEFAULT_MODEL}# 校验模型是否在支持列表中valid=0for m in \"${SUPPORTED_MODELS[@]}\"; do if [[ \"$m\" == \"$MODEL\" ]]; then valid=1 break fidoneif [[ $valid -eq 0 ]]; then echo \"错误: 不支持的模型 \\\"$MODEL\\\"\" echo \"\" print_models exit 1fiexport ANTHROPIC_BASE_URL=http://localhost:3001export ANTHROPIC_AUTH_TOKEN=\"nvapi-你的真实key\"export ANTHROPIC_DEFAULT_HAIKU_MODEL=${MODEL}export ANTHROPIC_DEFAULT_SONNET_MODEL=${MODEL}export ANTHROPIC_DEFAULT_OPUS_MODEL=${MODEL}echo \"使用模型: ${MODEL}\"claude随后，运行该脚本，以及启动claude-nvidia-proxy，即可使用NVIDIA的免费模型。如果claude客户端是运行在Windows系统上，或者另外一台电脑上，需要修改脚本中的base url为claude-nvidia-proxy所在电脑的IP地址。运行之前，确保claude-nvidia-proxy监听的端口可用，使用命令sudo ufw allow 3001放行该端口。**更新 NIM 免费模型列表如上脚本中的模型列表，可以通过软件工具freellm更新，代码仓库：freellm。本地安装以及使用命令如下：# 需要创建uv虚拟环境，并激活uv pip install -e .# 保存到 JSON 文件freellm list --provider nvidia --output out/nvidia.json# 保存筛选结果filter-models -i out/nvidia.json -o out/top_models.json**资料 NVIDIA Explore claude-nvidia-proxy 黄大善人免费的的nvidia glm和minimax模型应用于Claude Code：博客介绍了如何使用claude-nvidia-proxy连接NVIDIA的免费模型，并在claude客户端中使用这些模型进行交互。 零成本上手AI：英伟达nvidia免费大模型全攻略：介绍一些用法，比如使用其他工具连接NVIDIA的免费模型。在 VSCode 中使用 NVIDIA NIM 服务及模型首先在VSCode中安装插件NVIDIA NIM Provider。安装完成之后，在VSCode的Chat边栏的模型选择弹出菜单中，点击管理语言模型打开模型管理界面，点击添加新的语言模型，在模型管理弹窗中，选择NVIDIA NIM，并配置：名称（默认名称为NVIDIA NIM），以及API_KEY。配置完成之后，NVIDIA NIM栏目下面会显示很多可用的LLM模型，同时在退出弹窗之后，在Chat的模型选择中也可选择并使用这些模型。NIM 中的模型推荐列表通用LLM：Llama 4 Maverick（17B 128E）：Meta最新一代LLM模型，128k token上下文窗口，适合需要处理长文本的应用场景。对中文的支持也相当不错。DeepSeek V4：DeepSeek最新旗舰LLM模型，支持1M token上下文窗口，适合需要处理超长文本的应用场景。应该是编写代码的首选。Qwen 3.5 系列：应该也是最新的LLM模型。GLM 5.1：智谱AI+清华，中文语言模型的领头羊。MiniMax M2.7：最新的 M2.7 是一個 2300 億參數的 MoE（混合專家）模型，特別強在推理和多模態任務，2026年4月上架NVIDIA NIM。NVidia 自家的 LLM：Nemotron 3 Super 120B：旗艦版，適合複雜推理和結構化工具使用。Nemotron 3 Nano Omni：多模態版本，可以同時處理圖片、影片、語音和文字。特殊用途 LLM（语音、图像、嵌入）：Riva（TTS）：NVIDIA 自家的文字/语音互转模型。Llama 3.2 Vision：支持图像输入的多模态LLM。Qwen Image Edit：可以使用文字指令直接编辑图片。安装 NVIDIA NemoClawTODONemoClaw 资料 NVIDIA NemoClaw：github托管的资料，安装及使用NemoClaw。 把 NemoClaw 跑起来：NVIDIA官方的安装指南。 Running NemoClaw on macOS：一个gist，介绍了如何在macOS上安装和运行NVIDIA NemoClaw。 NemoClaw：NVIDIA NemoClaw的GitHub仓库。OpenRoute 资源及使用笔记OpenRouter 上面有很多LLM，其中有一些免费的LLM可以使用，列表地址：OpenRouter Models，要接入OpenRouter，也需要配置API Key，地址：OpenRouter API Keys，例如设置如下：export OPENROUTER_API_KEY=\"你的OpenRouter API Key\"3. AI 垂类工具收集代碼与知识图谱生成 BabelDOC: 一款基于AI的文档生成工具，能够自动从代码库中提取信息并生成详细的技术文档，支持多种编程语言和格式。 graphify：基于AI的知识库生成工具，会搜集目录中的代码、文档、网页等内容，生成知识图谱。代价是需要消耗LLM资源。 codebase-memory-mcp：代码仓库知识图谱，对C++不友好，或者说不支持 graphify：另外一个代码知识图谱工具，支持C++ code-graph：又一个代码知识图谱，支持C++ Mimir：使用Go实现的代码知识图谱视听媒体处理 Pixelle-Video：阿里云团队出品：只需输入一个 主题，AI 全自动短视频引擎–撰写视频文案、生成 AI 配图/视频、合成语音解说、添加背景音乐、一键合成视频。 OpenMontage：开源视频生成系统 MOSS-TTS：MOSS-TTS主打高保真，相对QWen TTS，Voice Clone时间比较长（QWen TTS似乎限制3秒），但QWen TTS时延小（\u0026lt;96ms）。另外，MOSS-TTS nano版本占用资源很小，可以在CPU上运行。图像与辅助设计 ComfyUI：AI绘图工具，提供本地APP，以及在线方式。 DrawnixDesktop：github 开源白板工具，一体化白板，包含思维导图、流程图、自由画等 design.md：Google 出品，UI设计skills以及工具 archify：架构图/流程图等画图skills及工具计算机视觉处理 supervision：Roboflow 团队开源的 CV 后处理工具库GPU相关 agent-gpu-skills：包含 cutlass、CuTe、CuTeDSL、triton ptx-isa-markdown：PTX指令集的Markdown版本文档处理与格式转换 MinerU：github项目，将pdf/docx等文档转换为markdown。 Tocify – 快速给PDF添加书签 Quarkdown：github项目，提供专业的mardown渲染，如论文等。4. 日常实用工具与资源媒体下载与管理 yt-dlp: 一个命令行视频下载工具，支持从YouTube及其他多个网站下载视频，功能强大且持续更新。 Bilibili-Mass-Unfollower: 一个用于批量取消关注Bilibili用户的工具。初中物理/数学/化学虚拟实验室 云学习实验室 PhET 免费在线 STEM 虚拟实验室（物理·化学·生物·数学） 云鱼的物理实验室 ChinaTextBook通用开发资源通用Skills agent-engineer：Google工程师addyosmani写的教程 agent-skills：Google工程师addyosmani写的agent skills" },

  { "title": "AI工具：使用docling转换文档", "url": "/blog/2025/AI-%E4%BD%BF%E7%94%A8docling%E8%BD%AC%E6%8D%A2%E6%96%87%E4%BB%B6/", "categories": "AI",
  "tags": "AI", "date": "2025-12-14", "content":
  "1. 环境搭建环境搭建参考：Python venv 环境搭建及 VSCode 环境配置。pip install docling2. 使用 docling转换文档\"\"\"docx 转 Markdown 脚本用法： 转换单个文件：python docx2md.py 文件.docx 转换整个目录：python docx2md.py 目录路径/ 指定输出目录：python docx2md.py 文件.docx -o 输出目录/\"\"\"import argparseimport sysfrom pathlib import Pathfrom docling.document_converter import DocumentConverterdef convert_file(input_path: Path, output_dir: Path) -\u0026gt; None: \"\"\"将单个 docx 文件转换为 Markdown。\"\"\" print(f\"正在转换：{input_path}\") converter = DocumentConverter() result = converter.convert(str(input_path)) md_content = result.document.export_to_markdown() output_dir.mkdir(parents=True, exist_ok=True) output_path = output_dir / (input_path.stem + \".md\") output_path.write_text(md_content, encoding=\"utf-8\") print(f\"已保存：{output_path}\")def main() -\u0026gt; None: parser = argparse.ArgumentParser(description=\"将 docx 文件转换为 Markdown\") parser.add_argument(\"input\", help=\"输入的 .docx 文件或包含 .docx 文件的目录\") parser.add_argument( \"-o\", \"--output\", default=None, help=\"输出目录（默认与输入文件同目录）\" ) args = parser.parse_args() input_path = Path(args.input) if not input_path.exists(): print(f\"错误：路径不存在：{input_path}\", file=sys.stderr) sys.exit(1) # 收集待转换文件 if input_path.is_dir(): files = list(input_path.rglob(\"*.docx\")) if not files: print(f\"目录中未找到 .docx 文件：{input_path}\", file=sys.stderr) sys.exit(1) else: if input_path.suffix.lower() != \".docx\": print(f\"错误：不是 .docx 文件：{input_path}\", file=sys.stderr) sys.exit(1) files = [input_path] for file in files: output_dir = Path(args.output) if args.output else file.parent convert_file(file, output_dir) print(f\"\\n完成，共转换 {len(files)} 个文件。\")if __name__ == \"__main__\": main()A. 资料 docling docling文档–使用" },

  { "title": "QGC代码架构解析：QGC初始加载及状态机", "url": "/blog/2025/qgc_init_load_states/", "categories": "QGC",
  "tags": "QGC", "date": "2025-12-13", "content":
  "在收到飞机发来的心跳包后，消息发送给MultiVehicleManager，在MultiVehicleManager中，检查组件ID是否是MAV_COMP_ID_AUTOPILOT1，以及vehicleType不是GCS等之后，会创建一个Vehicle对象，并进入初始化流程。 mavlink_message_t中已经包含了sysid、compid信息。而心跳包mavlink_heartbeat_t中则包含了vehicleType、firmwareType等信息。 需要注意的是，与飞控通信中，组件ID是固定的：MAV_COMP_ID_AUTOPILOT1，即每个飞控的组件ID都是1。飞机发现、创建及初始化流程如下图所示：主要初始流程入口在InitialConnectStateMachine中以状态机实现，且部分子流程也是以状态机实现（至多嵌套了三层状态机）： 由于获取信息需要使用同步方式，在InitialConnectStateMachine状态机中，使用回调方式处理应答Ack，在回调中进入下一个处理阶段。参见：1. 请求的实现，以及模拟同步请求。static constexpr const StateMachine::StateFn _rgStates[] = { _stateRequestAutopilotVersion, _stateRequestStandardModes, _stateRequestCompInfo, _stateRequestParameters, _stateRequestMission, _stateRequestGeoFence, _stateRequestRallyPoints, _stateSignalInitialConnectComplete};1. 请求的实现，以及模拟同步请求请求飞机信息，使用MAV_CMD_REQUEST_MESSAGE命令字，请求对应的消息ID（即子命令，比如请求飞机版本信息MAVLINK_MSG_ID_AUTOPILOT_VERSION），以及子命令的参数。另外，使用命令MAV_CMD_SET_MESSAGE_INTERVAL让飞机定期周期响应子命令消息。飞机端收到MESSAGE消息之后，先返回一个响应Ack（Ack中包含msgid，以及响应码，比如MAV_RESULT_ACCEPTED）。QGC收到该消息，继续处理之前发送的请求，实现代码主要有两个函数入口：Vehicle::requestMessage，Vehicle::_handleCommandAck(mavlink_message_t\u0026amp; message)，以及一个主要的数据成员QMap\u0026lt;int, QMap\u0026lt;int, RequestMessageInfo_t*\u0026gt;\u0026gt; _requestMessageInfoMap。MAV_CMD_REQUEST_MESSAGE消息的文档：How to Request \u0026amp; Stream Messages。处理流程示例：你的程序 飞控 | | | 1. 发送 MAV_CMD_SET_MESSAGE_INTERVAL |----------------------------------------→ (请求：以1Hz发送BATTERY_STATUS) | | | 处理请求 | 2. 接收 COMMAND_ACK |←---------------------------------------- (确认已接受) | | 3. 等待 BATTERY_STATUS | | | 自动发送电池信息（周期: 1秒） | ← 接收 BATTERY_STATUS #1 | | ← 接收 BATTERY_STATUS #2 | （自动循环，无需请求） | ← 接收 BATTERY_STATUS #3 | | ← 接收 BATTERY_STATUS #4 |2. 请求飞机版本信息(MAVLINK_MSG_ID_AUTOPILOT_VERSION)主要获取飞机的编号、固件的vender_id、product_id，固件版本信息，以及capabilities（64位bitmask），capabilities相关枚举定义在MAVLink协议的MAV_PROTOCOL_CAPABILITY中。3. 请求飞机标准模式(MAVLINK_MSG_ID_AVAILABLE_MODES)主要获取飞机支持的标准模式。获取的模式列表用于设置给FirmwarePlugin，并在Vehicle中使用。4. 请求组件元数据（META）信息由于这一步请求处理多个类型META数据文件，整个处理放在单独的模块（源码文件）ComponentInformationManager中，且也使用状态机来实现：请求General元数据、Param元数据、Events元数据、Actuator元数据。General是指组件的信息（主要是飞控自身），而Events，Actuator不一定每个组件都有。请求每个子分类的META数据，分为几个步骤（还是用状态机实现），在RequestMetaDataTypeStateMachine中实现：请求数据文件的地址URI（返回URI，以及CRC），根据URI请求META数据文件内容（具有缓存功能，先比较CRC，不相等再请求远程META文件），即数据类型描述文件。比如下一个步骤请求飞机的参数信息，就需要将请求到的参数生成Fact，而Fact的类型信息就来自于参数的META数据文件。 在请求META数据文件的过程中，定义了两个数据类：1. struct CompInfo::Uris：存放META文件的URI，CRC，以及其他一些信息（如Fallback请求信息）。2. CompInfo，以及继承自CompInfo的上述几种META文件对应的子类：CompInfoGeneral，CompInfoParam，CompInfoEvents，CompInfoActuators。其中最重要也是首先需要请求的META文件是CompInfoGeneral，因为这个文件里面包含了设备支持的META文件类型列表（数据成员QMap\u0026lt;COMP_METADATA_TYPE, Uris\u0026gt; _supportedTypes;）。后续几个请求，要先检查设备是否支持该类型的META文件。 在ComponentInformationManager中，维护了一个QMap\u0026lt;uint8_t /* compId */, QMap\u0026lt;COMP_METADATA_TYPE, CompInfo*\u0026gt;\u0026gt; _compInfoMap;，用于存放每个组件的各类META文件对象。这与上面所述的逻辑连接起来。 Events是MAVLink协议中的一个系统事件和诊断机制，用于飞机端向地面站实时报告系统事件、警告和错误。相关文档：Events Interface (WIP)。比如可能有如下事件：飞机端事件流： ├─ 电池低电量事件 ├─ GPS 信号丢失事件 ├─ IMU 过热警告 ├─ 传感器校准失败 ├─ 电机故障检测 └─ 健康检查失败 (Health \u0026amp; Arming Checks)4.1. META数据使用流程请求到的META数据文件，主要用于创建FactMetaData对象，进而创建Fact对象。以请求参数的META数据文件为例，流程如下所示：┌───────────────────────────────────────────────────┐│ 飞机端 (Autopilot) │└───────────────────────────────────────────────────┘ │ │ MAVLink │ ┌─────────────▼──────────────┐ │ ComponentInformation │ │ ┌──────────────────────┐ │ │ │ COMP_METADATA_TYPE │ │ │ │ _PARAMETER │ │ │ │ (JSON 文件 URI) │ │ │ └──────────────────────┘ │ └─────────────┬──────────────┘ │ 下载 JSON │ ┌─────────────▼──────────────┐ │ CompInfoParam.setJson() │ │ (解析 JSON 文件) │ └─────────────┬──────────────┘ │ ┌─────────────▼──────────────────────────────┐ │ FactMetaData::createFromJsonObject() │ │ (将 JSON 转换为 FactMetaData 对象) │ └─────────────┬──────────────────────────────┘ │ ┌─────────────▼──────────────────────────────┐ │ ParameterManager │ │ _nameToMetaDataMap[paramName] │ │ (存储所有参数的元数据) │ └────────────────────────────────────────────┘代码执行流程：// 1️⃣ 初始化连接时，请求参数的元数据_stateRequestCompInfoEvents() └─► _requestTypeStateMachine.request( _compInfoMap[MAV_COMP_ID_AUTOPILOT1][COMP_METADATA_TYPE_PARAMETER] );// 2️⃣ 下载 JSON 文件后，调用 setJson()CompInfoParam::setJson(const QString\u0026amp; metadataJsonFileName){ // 3️⃣ 解析 JSON 文件 QJsonDocument jsonDoc = // 从文件读取 QJsonArray rgParameters = jsonDoc[\"QGC_PARAMETERS\"].toArray(); // 4️⃣ 为每个参数创建 FactMetaData 对象 for (QJsonValue parameterValue : rgParameters) { FactMetaData* newMetaData = FactMetaData::createFromJsonObject(parameterValue.toObject(), ...); // 5️⃣ 存储到 map 中 _nameToMetaDataMap[newMetaData-\u0026gt;name()] = newMetaData; }}// 6️⃣ 后续使用时FactMetaData* meta = _compInfoParam-\u0026gt;factMetaDataForName(\"PARAM_NAME\");// 使用 meta 来验证、转换参数值4.2. 对应的 MAVLink 服务请求META数据使用微服务Component Metadata Protocol (WIP)，命令字：MAVLINK_MSG_ID_COMPONENT_METADATA。针对各个META数据类型，提供了枚举定义COMP_METADATA_TYPE。请求流程图如下所示：sequenceDiagram; participant Client participant Server Note over Server, Client: Client: Request component information. Client-\u0026gt;\u0026gt;Server: MAV_CMD_REQUEST_MESSAGE(param1=397) Client--\u0026gt;\u0026gt;Client: Start ACK receive timeout Server-\u0026gt;\u0026gt;Client: CMD_ACK Server-\u0026gt;\u0026gt;Client: COMPONENT_METADATA( uri, file_crc) Note over Server, Client: Client check file at uri has changed (using CRC in file_crc). Note over Server, Client: Client download file at uri using MAVFTP and parse. Note over Server, Client: Client download other metadata types referenced in general metadata\u0026lt;br\u0026gt; (from device or Internet).5. 请求系统参数列表这个步骤请求飞机的所有参数，使用MAVLink的微服务Parameter Protocol，在QGC的ParameterManager模块中实现，见上一篇QGC代码架构解析：MAVLink参数服务及QGC参数管理模块。请求参数，依赖于上一步骤，即请求的参数META数据文件，用于创建参数对应的FactMetaData对象。6. 请求任务列表（航点列表）这个步骤使用Mission Protocol，直接调用PlanManager::loadFromVehicle下载飞机航点信息，进行初始化同步。具体参考下一篇QGC代码架构解析：MAVLink Mission Protocol，以及 QGC 航点管理。由于MAVLink v2中，Mission Protocol不仅仅包含航点，还包含地理围栏(GeoFence)、降落点(Rally Points)等信息。所有这些部分都实现在PlanManager以及其继承子类中。7. 请求地理围栏列表参考6。8. 请求降落点列表参考6。9. 初始化完成完成初始化，发送signal，通知QML界面。" },

  { "title": "QGC代码架构解析：MAVLink参数服务及QGC参数管理模块", "url": "/blog/2025/qgc_parameter_module/", "categories": "QGC",
  "tags": "QGC", "date": "2025-12-12", "content":
  "MAVLink参数服务网页：Parameter Protocol1. 微服务：Parameter Protocol基本流程为请求-\u0026gt;响应返回。请求/消息列表： PARAM_REQUEST_LIST：请求所有参数。随后远端会周期性发送所有参数，直到发送完毕。请求数据结构： target_system（uint8_t）：System ID target_component（uint8_t）：Component ID PARAM_REQUEST_READ：请求单个参数。请求数据结构： target_system（uint8_t）：System ID target_component（uint8_t）：Component ID param_id（char[16]）：参数名称，与param_index二选一 param_index（int16_t）：参数索引（-1表示忽略） PARAM_SET：设置单个参数。请求数据结构： target_system（uint8_t）：System ID target_component（uint8_t）：Component ID param_id（char[16]）：参数名称 param_value（float）：参数值 param_type（uint8_t）：参数数据类型枚举 PARAM_VALUE：参数值响应。响应数据结构（系统ID以及组件ID在mavlink_message_t中）： param_id（char[16]）：参数名称 param_value（float）：参数值 param_type（uint8_t）：参数数据类型枚举 param_count（uint16_t）：参数总数 param_index（uint16_t）：当前参数索引 知识点总结： 获取/设置所有子系统的参数：target_component设置为 MAV_COMP_ID_ALL（QGC就是采取这种方式）。 参数响应消息中，带有参数总数和当前索引，QGC可以判断是否接收完毕。 param_value有两种格式存储参数值：转换为float类型；或者直接按照原始数据复制到param_value域。Ardupilot直接使用memcpy的方式（小端）。MAVLink参数协议中，定义了相应的标志位，表示使用哪种方式存储。 当QGC发送设置命令PARAM_SET之后，飞机端会返回一个PARAM_VALUE消息作为确认：Parameter Protocol – Write Parameters。 PX4固件支持缓存机制，即通过返回名称为PARAM_HASH，值为哈希值的PARAM_VALUE消息，确认参数没有变化：Parameter Protocol – PX4。 设置参数非法等原因导致飞机拒绝设置时，飞机端返回STATUS_TEXT消息。1.1. 参数存储代码示例// 定义如下结构体，适用于原始字节流存储参数值（Bytewise）// https://github.com/mavlink/c_library_v2/blob/master/mavlink_types.hMAVPACKED(typedef struct param_union { union { float param_float; int32_t param_int32; uint32_t param_uint32; int16_t param_int16; uint16_t param_uint16; int8_t param_int8; uint8_t param_uint8; uint8_t bytes[4]; }; uint8_t type;}) mavlink_param_union_t;mavlink_param_union_t param;int32_t integer = 20000;param.param_int32 = integer;param.type = MAV_PARAM_TYPE_INT32;// Then send the param by providing the float bytes to the send functionmavlink_msg_param_set_send(xxx, xxx, param.param_float, param.type, xxx);1.2. 协议限制及缺陷 没有同步保证一致性机制，即发送修改命令，飞机端的参数可能已经更新了，跟QGC不一样了。安全的方式是：设置时，QGC将修改前的值带入，让飞机端进行比较。 如果有多个关联参数需要一起设置的，没有原子操作机制。新增的服务Extended Parameter Protocol，数据类型可以支持uint64_t以及字节流：MAVPACKED(typedef struct { union { float param_float; double param_double; int64_t param_int64; uint64_t param_uint64; int32_t param_int32; uint32_t param_uint32; int16_t param_int16; uint16_t param_uint16; int8_t param_int8; uint8_t param_uint8; uint8_t bytes[MAVLINK_MSG_PARAM_EXT_SET_FIELD_PARAM_VALUE_LEN]; }; uint8_t type;}) param_ext_union_t; Extended Parameter Protocol github – mavlink-devguide: Extended Parameter Protocol1.3. 资料 github – PX4 Drone Autopilot2. QGC参数管理模块参数管理模块ParameterManager主要提供参数的请求，缓存功能，以及参数的设置功能。以及对外提供参数访问接口（参数管理模块不直接对UI层提供参数列表，由各个AutoPilotPlugin中各个部件模块提供比较合理），故参数模块应该作为一个公共模块，提供给plugin使用。另外，还提供了离线参数加载功能，离线参数文件定义在FirmwarePlugin中。下载的参数，存储在成员变量_mapCompId2FactMap中：QMap\u0026lt;int /* comp id */, QMap\u0026lt;QString /* parameter name */, Fact*\u0026gt;\u0026gt; _mapCompId2FactMap;缓存的文件名格式为\u0026lt;系统ID\u0026gt;_\u0026lt;组件ID\u0026gt;.v2，存储目录定义在ParameterManager::parameterCacheDir()中。参数存储格式为：typedef QPair\u0026lt;int /* FactMetaData::ValueType_t */, QVariant /* Fact::rawValue */\u0026gt; ParamTypeVal;typedef QMap\u0026lt;QString /* parameter name */, ParamTypeVal\u0026gt; CacheMapName2ParamTypeVal;// 展开之后为// QMap\u0026lt;QString, std::pair\u0026lt;int, QVariant\u0026gt;\u0026gt;// 即 QMap\u0026lt;参数名称, \u0026lt;参数类型枚举, 参数值\u0026gt;\u0026gt;2.1. 参数请求在QGC初始化时，请求所有参数： 从网站（不是飞控）下载参数文件，仅针对APM固件，入口：APMAirframeComponentController::loadParameters； 使用FTP从飞控下载参数文件，入口函数refreshAllParameters； 通过参数服务请求参数，入口函数refreshAllParameters。请求完成之后，都需要调用_checkInitialLoadComplete。在ParameterManager中，跟请求相关的主要的数据结构定义：QMap\u0026lt;int, int\u0026gt; _paramCountMap; ///\u0026lt; Key: Component id, Value: count of parameters in this componentQMap\u0026lt;int, QMap\u0026lt;int, int\u0026gt;\u0026gt; _waitingReadParamIndexMap; ///\u0026lt; Key: Component id, Value: Map { Key: parameter index still waiting for, Value: retry count }QMap\u0026lt;int, QList\u0026lt;int\u0026gt;\u0026gt; _failedReadParamIndexMap; ///\u0026lt; Key: Component id, Value: failed parameter index2.2. 参数设置由于生成的参数列表，存在_mapCompId2FactMap中，及每个参数以Fact表示，使用Fact值变更的信号，设置单个参数。" },

  { "title": "QGC代码架构解析：FirmwarePlugin与AutopilotPlugin", "url": "/blog/2025/qgc_firmwareplugin_autopilotplugin/", "categories": "QGC",
  "tags": "QGC", "date": "2025-12-11", "content":
  "概念： FirmwarePlugin：固件插件，表示某种飞控固件（如APM、PX4等）。 AutoPilotPlugin：表示某种飞控固件实现的不同飞机类型，比如固定翼、旋翼等。1. FirmwarePlugin 与 AutoPilotPlugin 的关系从逻辑关系看（主要从分类角度），AutoPilotPlugin需要从FirmwarePlugin创建，例如APM的固件，创建APM相关的AutoPilotPlugin。从实现看，FirmwarePlugin最多也有三层继承，以APM固件为例： FirmwarePlugin（基类） APMFirmwarePlugin（表示APM固件） ArduPlaneFirmwarePlugin（表示APM的固定翼飞机） ArduCopterFirmwarePlugin（表示APM的多旋翼飞机） ArduRoverFirmwarePlugin（表示APM的地面车） 是在APMFirmwarePlugin这一层创建AutoPilotPlugin，即针对APM固件，实际只有一种AutoPilotPlugin实现。1.1. FirmwarePlugin 实例的创建MAVLink协议心跳包中，包含了固件类型和飞行器类型两个字段：typedef struct __mavlink_heartbeat_t { uint32_t custom_mode; /*\u0026lt; A bitfield for use for autopilot-specific flags*/ uint8_t type; /*\u0026lt; Vehicle or component type. For a flight controller component the vehicle type (quadrotor, helicopter, etc.). For other components the component type (e.g. camera, gimbal, etc.). This should be used in preference to component id for identifying the component type.*/ uint8_t autopilot; /*\u0026lt; Autopilot type / class. Use MAV_AUTOPILOT_INVALID for components that are not flight controllers.*/ uint8_t base_mode; /*\u0026lt; System mode bitmap.*/ uint8_t system_status; /*\u0026lt; System status flag.*/ uint8_t mavlink_version; /*\u0026lt; MAVLink version, not writable by user, gets added by protocol because of magic data type: uint8_t_mavlink_version*/} mavlink_heartbeat_t;通信层在收到心跳包之后，最终传递给MultiVehicleManager，然后在创建Vehicle中，创建对应的FirmwarePlugin实例：// void Vehicle::_commonInit(LinkInterface* link)_firmwarePlugin = FirmwarePluginManager::instance()-\u0026gt;firmwarePluginForAutopilot(_firmwareType, _vehicleType);2. FirmwarePlugin 的职责获取对应的AutoPilotPlugin：virtual AutoPilotPlugin *autopilotPlugin(Vehicle *vehicle) const;主要功能是，定义支持的飞行模式，并通过接口提供给外部。以及，一些飞行模式的控制实现：比如TakeOff模式中，起飞前检查（解锁等），发送起飞命令；再比如，发送Mission命令，以及发送之前的检查。一些接口定义：virtual QList\u0026lt;MAV_CMD\u0026gt; supportedMissionCommands(QGCMAVLink::VehicleClass_t /*vehicleClass*/) const;virtual QStringList flightModes(Vehicle* /*vehicle*/) const;virtual void startTakeoff(Vehicle *vehicle) const;virtual void startMission(Vehicle *vehicle) const;另一个重要的接口是，定义飞控的离线参数文件，例如APM固定翼飞机：QString offlineEditingParamFile(Vehicle *vehicle) const override { return QStringLiteral(\":/FirmwarePlugin/APM/Plane.OfflineEditing.params\"); }3. AutoPilotPlugin 的职责FirmwarePlugin实现通信协议及逻辑，而AutoPilotPlugin实现界面呈现。AutoPilotPlugin目录下，定义了AutoPilotPlugin模块，以及VehicleComponent模块。APM，PX4两个子目录实现各自的AutoPilotPlugin子类，以及各自包含的诸多组件（继承自VehicleComponent）。实现的组件包括：Vehicle Configuration面板中的各个功能模块，比如遥控器、飞行模式、传感器、参数等页面（参看QGC运行页面）。" },

  { "title": "APM/Pixhawk常用飞行模式", "url": "/blog/2025/ardupilot-flight-modes/", "categories": "Ardupilot",
  "tags": "Ardupilot, QGC", "date": "2025-12-10", "content":
  "1. 多旋翼：手动飞行模式如下几种飞行模式是手动控制模式，即受遥控器控制： Stabilize（稳定模式/姿态模式） Altitude（定高模式） Position（定点模式） Offboard（板外模式/指令模式）1.1. Stabilize：稳定模式（姿态模式） 俯仰、横滚摇杆控制飞机对应的角度（注意是角度）； 油门控制飞机的上升/下降速度，以及其他轴的移动速度（注意是速度，类似汽车油门）； 偏航控制飞机的旋转速率（水平方向）。当摇杆回到中立位置时，飞机会自动保持当前的姿态（俯仰角0，横滚角0，偏航角0）和高度。但不会自动保持位置：可能会朝着风力的方向漂移，此时需要控制油门以保持高度。资料： PX4 Guide – 位置模式（多旋翼）1.2. Altitude：定高模式定高模式与稳定模式类似：俯仰、横滚摇杆控制飞机的角度。但是油门控制逻辑是：油门摇杆以预定的最大速率（和其他轴上的移动速度）控制上升速度。 当摇杆归中之后，如果飞机在水平方向飞行，则持续运动，直到被风的阻力减速停下。如果刮风，飞机会朝着风的方向漂移。 高度保持依赖传感器（气压计或激光测距仪等）来维持高度。资料： PX4 Guide – 定高模式（多旋翼）1.3. Position：定点模式 横滚、俯仰摇杆分别控制飞机在左右和前后方向上的地面水平加速度。 油门摇杆控制飞机的上升/下降速度。 偏航摇杆控制飞机的旋转速率（水平面方向）。当摇杆回到中立位置时，飞机会自动保持当前位置和高度，但不保持当前方向（当有外力改变水平朝向之后，会保持新的角度）。依赖GPS获取绝对位置，以及磁罗盘获取航向。如果这两个组件有失效，则不能进入该模式；如果在该模式下失效，则进入失效处理。资料： PX4 Guide – 位置模式（多旋翼）1.4. Offboard：板外模式/指令模式指令控制模式，即通过地面站发送切换指令： 切换指令需要带位置、偏航角等参数； 需要按指定周期发送指令，否则会触发失效处理。2. 多旋翼：自动飞行模式TBD3. 摇杆遥控器摇杆控制：俯仰(pitch)、横滚(roll)、偏航(yaw)、油门(throttle)：对应的飞机运动： 俯仰 =\u0026gt; 上升/下降。 横滚 =\u0026gt; 向左/右倾斜并转弯。 偏航 =\u0026gt; 机尾向左/右转动并转弯。 油门 =\u0026gt; 改变前进速度。以上指固定翼飞机，其他类型飞机类似。资料： PX4 Guide – 固定翼飞机基础飞行指南参考资料 PX4 Guide – PX4多旋翼无人机飞行模式（Flight Mode） APM/Pixhawk常用飞行模式讲解：http://www.nufeichuiyun.com/?p=1128 【无人机】多旋翼无人机控制器架构，PX4控制器，PID控制 开放航空航天仿真工具集锦" },

  { "title": "地面站 Helios 以及 Dart/Flutter 环境搭建", "url": "/blog/2025/%E5%9C%B0%E9%9D%A2%E7%AB%99-helios-%E4%BB%A5%E5%8F%8A-dart-flutter-%E7%8E%AF%E5%A2%83%E6%90%AD%E5%BB%BA/", "categories": "Ardupilot",
  "tags": "Ardupilot, QGC, Helios, Dart, Flutter", "date": "2025-12-08", "content":
  "1. Flutter 环境搭建从Flutter 南京大学镜像下载并解压 Flutter SDK，当前最新版本3.41.7。随后，将其子目录bin添加到系统环境变量Path中。并添加如下两个环境变量（以Linux环境为例）：export FLUTTER_STORAGE_BASE_URL=\"https://mirrors.cernet.edu.cn/flutter\"export PUB_HOSTED_URL=\"https://mirrors.cernet.edu.cn/dart-pub\" # pub get2. Helios 地面站从Helios GitHub 仓库下载源码。A. 资源 Dart中文网站 Flutter 文档 Helios 官方网站" },

  { "title": "ArduPilot 笔记", "url": "/blog/2025/ardupilot-notes/", "categories": "Ardupilot",
  "tags": "Ardupilot, QGC", "date": "2025-12-08", "content":
  "1. ArduPilot SITL 编译设置windows/Cygwin环境下git选项：# 忽略权限位变化git config --global core.fileMode false# 信任当前目录git config --global --add safe.directory /cygdrive/e/work/flight/ardupilot# 关闭 CRLF 自动转换（避免与 Windows 工具冲突）git config --global core.autocrlf false下载 ArduPilot 代码：git clone https://github.com/ardupilot/ardupilot.gitcd ardupilotgit submodule update --init --recursivegit submodule foreach --recursive \"git reset --hard HEAD\"编译 ArduPilot SITL：# 编译及开发环境准备./Tools/environment_install/install-prereqs-ubuntu.sh -y./waf configure --board sitl./waf plane # ArduPlane# 生成 build/sitl/bin/arduplane1.1. 初步运行使用自定义python脚本运行 SITLgithub – ap-swarm-launcher：# 启动两个 ArduPlane SITL 实例，设置数据目录为 ~/tmp/arduplaneuv run ap-sitl-swarm --model plane -n 2 --data-dir ~/tmp/arduplane --no-multicast --tcp-base-port 5760 --home 31.8269,117.2280,30 ~/tmp/arduplane/arduplane2. Windows上使用Cygwin编译从cygwin 下载并安装setup-x86_64.exe。在cygwin环境中安装选择以下软件包：autoconf automake ccache gcc-g++ git libtool make gawk libexpat-devel libxml2-devel python39 python39-future python39-lxml python39-pip libxslt-devel python39-devel procps-ng zip gdb ddd xterm cmake另外，还需要安装如下软件包，以部分解决Cygwin中的终端启动不了的问题：xterm xorg-server xinit font-util unifont-fonts在cygwin环境中安装pip包：# 确保pip已经安装并升级到最新版本python3 -m ensurepip --default-pippython3 -m pip install --upgrade pippip install pymavlink pyserial empy==3.3.4 MAVProxy pexpect lxml在cygwin环境中编译（此时不需要也不能运行脚本install-prereqs-ubuntu.sh）：# 进入 ardupilot 根目录# ./waf configure --board sitl --debug# ./waf -j8 plane -v./waf configure --board sitl./waf plane2.1. SITL 运行运行sim_vehicle.py需要MAVProxy，首先需要从github上下载并安装MAVProxy。在Cygwin环境中，没有支持的console输出的terminal，需要作一些设置或修改，以显示console输出窗口。有两种方式：使用X server，或者mintty终端。首先，在cygwin中启动X Server（已弃用，可以不用启动）：startxwin \u0026amp;export DISPLAY=:0启动sim_vehicle.py前，需要修改Tools\\autotest\\run_in_terminal_window.sh脚本，添加mintty支持：elif [ -n \"$(which mintty 2\u0026gt;/dev/null)\" ]; then # Cygwin native terminal - no X11 fonts required mintty --hold always -T \"$name\" -e \"$@\" \u0026amp;然后，在cygwin环境中启动sim_vehicle.py：cd /cygdrive/e/work/flight/ardupilot/ArduPlane../Tools/autotest/sim_vehicle.py --map --console启动之后，观察LOG信息：SIM_VEHICLE: Run ArduPlaneSIM_VEHICLE: \"/cygdrive/e/work/flight/ardupilot/Tools/autotest/run_in_terminal_window.sh\" \"ArduPlane\" \"/cygdrive/e/work/flight/ardupilot/build/sitl/bin/arduplane\" \"--model\" \"plane\" \"--speedup\" \"1\" \"--slave\" \"0\" \"--sim-address=127.0.0.1\" \"-I0\"SIM_VEHICLE: Run MavProxySIM_VEHICLE: \"/usr/bin/cygstart\" \"-w\" \"mavproxy.exe\" \"--retries\" \"5\" \"--out\" \"127.0.0.1:14550\" \"--master\" \"tcp:127.0.0.1:5760\" \"--sitl\" \"127.0.0.1:5501\" \"--map\" \"--console\"RiTW: Starting ArduPlane : /cygdrive/e/work/flight/ardupilot/build/sitl/bin/arduplane --model plane --speedup 1 --slave 0 --sim-address=127.0.0.1 -I0可以看到SITL通过TCP:5760端口与MAVProxy通信，MAVProxy通过TCP:14550端口与QGroundControl通信。2.1.1. 更改SITL仿真的HOME坐标在Tools\\autotest\\locations.txt文件中定义了一些预设的HOME坐标，默认加载HOME点是CMAC。可以通过NAME选择其他地点，比如：../Tools/autotest/sim_vehicle.py -L Unalga --map --console也可以在locations.txt中添加NAME+坐标。例如添加合肥坐标：Hefei=31.8206,117.2272,27,0 #Hefei, Anhui, China (WGS84)此时启动命令为：../Tools/autotest/sim_vehicle.py -L Hefei --map --console2.1.2. 命令交互–固定翼 ArduPlane启动sim_vehicle.py之后，会启动一个console窗口、一个map窗口，以及一个terminal窗口。可以在terminal窗口中输入命令来控制仿真，比如：# 在 cygwin 环境中，由于 MAVProxy 是在 windows 环境下安装的（不是在 cygwin 环境安装的），# 所以 MAVProxy 命令接收的路径是 windows 的路径格式。# 另外，可以使用相对路径，但是测试这种方式发现不可靠。STABILIZE \u0026gt; wp load \"E:\\work\\flight\\ardupilot\\Tools\\autotest\\Generic_Missions\\CMAC-circuit.txt\"STABILIZE \u0026gt; mode guidedGUIDED \u0026gt; arm throttleGUIDED \u0026gt; takeoff 40GUIDED \u0026gt; mode autoGUIDED \u0026gt;GUIDED \u0026gt; mode rtl # 返航模式# 自动降落返回HOME点RTL \u0026gt; param set RTL_AUTOLAND 1RTL \u0026gt; mode autoland# 从 SITL / 飞控 下载航线文件，并更新 Mapmp list具体操作步骤参考官方教程文档：Plane SITL/MAVProxy Tutorial。另外参考知乎文章：ArduPilot 软件在环仿真SITL（SITL+MAVProxy）。 当飞机处于MISSION_RUNNING状态时（Armed + MISSION_RUNNING），不能接收航线修改命令。飞控接收新航线，会冲掉当前的航线。 其他wp命令：wp list、wp clear。 航线文件格式参见官方文档：Plan File Format。 txt格式文件格式：File Formats。 sim_vehicle.py的使用：Using SITL。 如果使用我修改过的HOME地点Hefei，此时对应的航线文件Hefei-circuit.txt内容如下：QGC WPL 1100 0 0 16 0.000000 0.000000 0.000000 0.000000 31.820600 117.227200 27.000000 11 0 3 22 15.000000 0.000000 0.000000 0.000000 31.823705 117.226818 41.029999 12 0 3 16 0.000000 0.000000 0.000000 0.000000 31.823953 117.223507 100.000000 13 1 3 16 0.000000 0.000000 0.000000 0.000000 31.817075 117.224346 100.000000 14 0 3 16 0.000000 0.000000 0.000000 0.000000 31.817407 117.226696 100.000000 15 0 3 16 0.000000 0.000000 0.000000 0.000000 31.824266 117.225872 100.000000 16 0 3 177 2.000000 -1.000000 0.000000 0.000000 0.000000 0.000000 0.000000 17 0 3 16 0.000000 0.000000 0.000000 0.000000 31.824266 117.225872 100.000000 1在mavProxy中，加载航线文件的命令为：wp load d:\\Hefei-circuit.txt2.1.3. 参数文件SITL启动时，都会加载默认的参数文件，比如plane的默认文件在代码仓库中的路径是Tools/autotest/models/plane.parm。运行sim_vehicle.py时，可以通过加载自定义参数文件来修改/添加SITL的仿真参数，比如使用如下命令加载用户自定义的参数文件：../Tools/autotest/sim_vehicle.py --map --console --param-file my_params.parm参数文件my_params.parm内容：# 查看当前参数param show# 修改特定参数param set BATTERY_CAPACITY 5200param set SIM_SPEEDUP 22.1.4. 有关SITL仿真的资源 Copter SITL/MAVProxy Tutorial：使用sim_vehicle.py的一个官方文档。 SITL setup on Windows using Cygwin (not recommended)2.1.5. 启动仿真以及连接QGroundControlQGC创建一个UDP:14550端口的连接，连接到MAVProxy的TCP:14550端口。MAVProxy会将SITL的数据转发给QGC。2.1.6. Cygwin终端美化编辑Cygwin的.bashrc文件（比如在windows中绝对路径为C:\\cygwin64\\home\\Administrator\\.bashrc）：##################### Git completion and promptsource git-completion.bashsource git-prompt.sh# Git prompt 状态标记GIT_PS1_SHOWDIRTYSTATE=1 # * 未暂存变更，+ 已暂存变更GIT_PS1_SHOWUNTRACKEDFILES=1 # ? 未追踪文件GIT_PS1_SHOWUPSTREAM=\"auto\" # \u0026lt; 落后，\u0026gt; 超前，\u0026lt;\u0026gt; 分叉# PS1: 蓝色目录名(basename) + 黄色 git 状态 + 普通 $PS1='\\[\\e[1;34m\\]\\W\\[\\e[0m\\]\\[\\e[1;33m\\]$(__git_ps1 \" (%s)\")\\[\\e[0m\\] \\$ '################# 目录导航别名alias ..='cd ..'alias ...='cd ../..'alias c='clear'alias h='history'alias ll='ls -alF'alias la='ls -A'alias l='ls -CF'alias grep='grep --color'其中，git-completetion.bash以及git-prompt.sh来自https://github.com/git/git/tree/master/contrib/completion。参考： Bash Prompt Generator：在线生成 PS1 Configuration CYGWIN - BASHRC：参考的Cygwin command prompt代码片段2.2. SITL 与 JSBSim在Cygwin环境中，编译并安装JSBSim：git clone git://github.com/JSBSim-Team/jsbsim.gitcd jsbsim \u0026amp;\u0026amp; mkdir build \u0026amp;\u0026amp; cd buildcmake -DSYSTEM_EXPAT=ON -DBUILD_SHARED_LIBS=ON -DINSTALL_JSBSIM_PYTHON_MODULE=ON -DCMAKE_CXX_FLAGS_RELEASE=\"-O3 -march=native -mtune=native\" -DCMAKE_C_FLAGS_RELEASE=\"-O3 -march=native -mtune=native\" -DCMAKE_BUILD_TYPE=Release ..make -j8 \u0026amp;\u0026amp; make installcmake --install . --component pymodules 编译命令参考：https://github.com/JSBSim-Team/jsbsim/blob/master/doc/DevelopersDocs.md。启动仿真命令：../Tools/autotest/sim_vehicle.py -v ArduPlane -f jsbsim --console --mapA. 资料 Using SITL with AirSim search: ardupilot airsim Ardupilot – SimulationA.1. 一些地面站收集 ADOSMissionControl：一个基于typescript的地面站 Qt AI Skills：QtCompany 发布的Qt/QML 开发 skills Manual client setup：Qt/QML开发 MCP 安装指南，适用于 CodeX、antigravity、Claude Code。" },

  { "title": "内存模型(Memory Model)：从多处理器到高级语言", "url": "/blog/2025/memory_model2/", "categories": "CPU",
  "tags": "CPU, Cpp", "date": "2025-12-07", "content":
  "复制自：github笔记 – GHScan大神 – Memory_Model前言谁需要关心这个主题？ 实现同步原语、Lockless算法的并行计算工程师 实现操作系统内核或驱动，和DMA设备打交道的系统工程师 实现高级语言Memory Model的编译器工程师 实现处理器Memory Model的CPU工程师如何阅读本文？阅读时注意本文的组织结构，第一遍阅读时可跳过斜体的技术细节。本文将回答的问题 x86为什么允许Store-Load乱序，而不是其他？见TSO的优点 std::memoryorder_seq_cst非常慢、会严重损害性能？见WO_ std::memoryorder_acquire和std::memory_order_consume的区别？后者换作std::memory_order_relaxed会怎样？见Dependent Loads_ std::memoryorder_acq_rel和std::memory_order_seq_cst的区别？见Store Atomicity和IRIW_基础知识共享存储器多处理器SMP是通过一个共享的地址空间进行通信和协调的多处理器系统，与之相对的概念是集群(Cluster)，后者每个节点都有独立的地址空间。SMP又包括： NUMA(Non-Uniform Memory Access)：大量的处理器通过互连网络连接在一起，每个节点访问本地存储器时延迟较短，访问其他节点有不同的延迟 UMA(Uniform Memory Access)：又叫SMP(Symmetric Multiprocessor)，所有的处理器通过总线连接在一起，访问一个共享的存储器，每个节点的存储器访问延迟相同。目前常见的桌面多核处理器就是这种结构现代处理器中的关键技术\\(^{[2]}\\) - 流水线(Pipeline)：每个指令的执行需要多个步骤，线代CPU通过流水线的方式允许同时执行多个指令，从而提高功能单元的利用率和系统总吞吐。支持流水线的CPU的IPC(Instructions Per Cycle) 可以达到1，哪怕每条指令实际上需要多个时钟周期才能完成。- `动态分支预测(Dynamic Branch Prediction)`：带流水线的CPU需要每个时钟发射1条指令，但只有分支指令执行结束后才能确定下条指令是什么，这就导致`流水线停顿(Pipeline Stall)`。为避免分支指令导致的流水线停顿，一种对策是分支预测，即在发射分支指令之后，马上预测下条指令的地址并发射，如果分支指令执行结束后发现预测错误，则撤销之前的操作取正确的指令重新发射。这里预测失败导致的撤销开销，叫`分支预测惩罚(Mispredict Penalty)`，由于现代系统的分支预测正确率很高，摊还后的惩罚开销往往可以接受。动态分支预测是基于分支指令历史进行的，现代CPU的预测正确率在大部分场合可以高达95%以上；相对的，静态分支预测是基于固定分支选择策略、源码中的Hint，或根据编译器的得到的Profile信息来完成的。- `动态多发射(超标量，Superscalar)`：为更好的利用富裕的功能单元，CPU希望IPC能够超过1，这就要求每个时钟发射多条指令。支持超标量的处理器，需要处理同时发射的多条指令间的数据依赖关系，这个复杂性限制了动态发射窗口的大小。与之相对的是静态多发射，即由编译器或程序员往一个`发射包(Issue Packet)`中填充多条无关指令，然后同时发射和执行，典型的例子是`超长指令字(Very Long Instruction Word)`体系结构。- `乱序执行(Out-of-Order Execution)`：当前面的指令由于特种类型的功能单元不足、存储器延迟或操作数没有计算出来，必须停顿时，CPU可以发射后续的无关指令，从而乱序执行。乱序指令有效的提高了CPU利用率，掩盖了各种停顿。 - `寄存器重命名(Register Renaming)`：CPU通过寄存器重命名的方式使得物理寄存器数目超过指令集中的逻辑寄存器，缓解了如x86指令集中的寄存器不足的问题。 - 寄存器重命名另一大作用是，避免由于`假数据依赖(False Data Dependence)`导致的停顿；具体来说，寄存器重命名解决了`WAW Hazard`和`WAR Hazard`。- `推断执行(Speculative Execution)`：支持动态分支预测和乱序执行的处理器，需要保留一个`重排序缓冲区(Reorder Buffer)`，用来对乱序执行的指令进行`顺序提交(In-Order Commit)`。重排序缓冲区为推断失败时的撤销提供了解决方案，只需要清空分支指令后的所有指令即可。另外，顺序提交也为`精确异常(Precise Exception)`提供了基础，这是操作系统中断和`缺页故障(Page Fault)`处理的关键。推断执行的指导思想是“加速大概率事件”。- `写缓冲区(Writer Buffer)`：CPU在写存储器时，不直接访问存储器或Cache，而是将要写的数据放入一个写缓冲区，然后继续执行后面的指令，这缓解了写存储器导致的停顿。- `硬件多线程(Hardware Multithreading)`：上面列举的优化策略都旨在改进`指令级并行(Instruction-Level Parallelism)`，另一种提高CPU利用率、掩盖停顿的做法是硬件多线程，即为每个CPU准备两套寄存器和指令计数器，同时从两个逻辑线程取指令并轮询发射，由于这两个线程的指令间没有依赖，并不会引入额外的停顿。支持推断执行的硬件多线程技术又叫`同时多线程(Simultaneously Multithreading)`。Cache Coherence\\[^{[2,4]}\\]多处理器系统中(这里讨论的处理器，也包括DMA设备)，为减少存储器访问延迟，会为每个处理器添加本地的Cache(比如Core i7-6700K每个Core各有32KB的Instruction Cache和32KB的Data Cache)，引入本地Cache会导致数据多副本问题：某处理器更新了存储器中的一个块，同一个块可能在其他处理器的Cache中还有过期副本。多处理器需要专门的硬件来实现Cache Coherence Protocol。本文仅仅讨论Write-Back \u0026amp; Write-Invalide \u0026amp; Snooping-Based的Cache，对Write-Through/Write-Broadcast/Directory-Based的Cache的讨论类似Cache Coherence Protocol MSI协议：每个Cache Line上有几个标记位用来标志其处于Modified/Shared/Invalid中的某个状态，当处理器读写该Cache Line时，会根据其状态进行状态迁移并发送相应的协议消息以保持多副本数据的一致性 Shared状态：处理器读操作引起的Read Miss会令该Cache Line以Shared状态从存储器读入本地Cache中；如果之前该Cache Line以Modified状态被某处理器持有，那监听到这个Read Miss的处理器用它持有的该Cache Line最新的副本响应源处理器，并更新存储器 Modified状态：处理器写操作引起的Write Miss会令Cache Line以Modified状态从存储器读入本地Cache中；如果之前该Cache Line被一个或多个处理器以Shared状态持有，则写操作处理器将向他们发送Invalidate Message令它们持有的Cache Line失效；如果之前该Cache Line被某处理器以Modified状态持有，则写操作处理器向它发送Read-Invalidate Message，目标处理器收到消息后将其持有的Cache Line标记为Invalid并回应以最新的数据副本同时更新存储器 Invalid状态：收到Invalidate或Read-Invalidate Message的处理器会将对应的Cache Line置为Invalid状态 MESI协议：在MSI协议的基础上，从Shared状态中分离出Exclusive状态来避免独占Cache Line的处理器第一次写Cache Line时发出的Coherence Message，从而减少总线流量 即，相比Shared到Modified状态的迁移，第一次写独占Cache Line时进行的是Exclusive到Modified的迁移，不必发送Invalidate Message了 Write Buffer 与 Invalidate Queue 写缓冲区(Write Buffer)：写缓冲区不仅可以在单处理器中被用来掩盖写延迟，也可以用来掩盖Cache Coherence Protocol的延迟。处理器为确保其他处理器看见自己的写操作，需要等待其他处理器在处理完自己发出的Invalidate Message后回应以Acknowledge Message，这个时间可能是几十上百个时钟周期，而利用Write Buffer，处理器可以在将一个写操作放入Write Buffer的同时就发出Invalidate Message，在收到Acknowledge Message后才从Write Buffer中移除并以Modified状态写入自己的本地Cache，而在收到来自其他处理器的回应前可以继续执行其他指令 失效队列(Invalidate Queue)：当处理器工作负载很高时，可能来不及处理、回应来自其他处理器的Coherence Message，从而将等待回应的其他处理器阻塞，而利用一个消息缓冲区缓存这些请求并立刻回应源处理器表示消息已收到，可以提高响应速度。处理器在收到Invalidate Message后可以先将其放入Invalidate Queue并立刻回应以Acknowledge Message，并在之后不忙或发生Cache Miss的时候再回来处理Invalidate Queue中缓存的请求 编译器优化技术编译器被允许在不改变单线程程序执行结果的前提下，进行各种优化。常见编译器优化 公共子表达式删除(Common Subexpression Elimination)\\(^{[17]}\\) a = b * c + g; //----------\u0026gt; tmp = b * c;d = b * c * e; // rewrite a = tmp + g; // d = tmp * e; 死代码删除(Dead Code Elimination) while (!flag); //-----------\u0026gt; loop: // jmp loop 寄存器分配(Register Allocation)：下例将对g的2次读取优化成了1次 a = g; //-----------\u0026gt; load %r1, 0($mem1)b += a; // rewrite add %r2, %r2, %r1a = g; //c += a; // add %r3, %r3, %r1 指令调度(Instruction Scheduling)：下例重排的后一条指令不必等待前一条的结果减少了停顿 load %r0, 0($mem0) // load %r0, 0($mem0)mul %r1, %r1, %r0 //-----------\u0026gt; load %r2, 0($mem2)store 0($mem1), %r1 // rewrite mul %r1, %r1, %r0load %r2, 0($mem2) // mul %r3, %r3, %r2mul %r3, %r3, %r2 // store 0($mem1), %r1store 0($mem3), %r3 // store 0($mem3), %r3 Memory Model什么是Memory ModelMemory Model(Memory Consistency/Memory Consistency Model)是系统和程序员之间的规范，它规定了在一个共享存储器的多线程程序中的存储器访问应该表现出怎样的行为。这个规范影响了系统的性能，因为它决定了多处理器/编译器能应用哪些优化；也影响了可编程性(Programmability)，因为多线程程序的正确性取决于Memory Model，从而约束了程序员的编程方式一些解释 运行在单处理器(且没有DMA设备等)上的单线程程序无需考虑Memory Model问题，因为处理器/编译器的优化都保证对程序员透明，即程序看起来像是按自然顺序(Program Order)执行的 一般讨论Memory Model中的读写乱序，是指对任意地址读写的乱序；但从实现层面讲，由于系统必须保证乱序对执行读写的线程透明，而相同地址的读写存在数据相关Data Dependence，所以系统实际上只能对不同地址的读写进行乱序 Memory Model的讨论，一般分为两层：汇编语言中的Memory Model由具体多处理器规定和实现，不可移植；高级语言中的Memory Model由高级语言标准来规定，由编译器和目标多处理器共同实现，可移植 为什么需要Memory Model在单线程程序中，编译器的各种优化如冗余代码消除、循环融合(Loop Fusion)、指令调度等技术，会重写代码造成存储器访问顺序变化，而寄存器分配会改变存储器访问次数；类似的，处理器的乱序执行机制也会通过让后一条指令先执行，来掩盖前一条指令导致的流水线停顿和存储器停顿。尽管会改变不同地址存储器的访问顺序，但系统的这些优化对程序员是透明的，看起来程序仍然是在顺序执行。然而在多线程程序中，编译器和多处理器并无手段自动发现多个线程间的协作关系，使得那些可能改变存储器访问顺序和次数的优化，同时对多个线程透明。没有程序员的帮助，要保持多线程程序的正确性，系统只能禁用这些作用在共享存储器上的优化，而这将严重损害性能。为最大限度保留编译器和多处理器的优化能力，同时使多线程程序的执行结果是可预测的，系统需要程序员的帮助。最后的方案是，系统提供所谓Memory Model的规范，程序员通过规范中同步设施(各种内存屏障(Memory Barrier)和Atomic指令)来标记多个线程间的协作关系，使得不仅是单线程，系统的优化对多线程程序也将透明。互斥锁(Exclusive Lock)是被最广泛支持的同步机制，编译器和多处理器会确保基于锁同步的多线程程序看起来就像是有多个同时执行的顺序线程。而一旦离开锁的庇护，程序员要么直面各种优化作用下的混乱世界，要么和实现同步原语(Synchronization Primitives)的系统工程师站在同一起跑线，捡起Memory Model这个更细粒度、更微妙的武器，在乱序优化的多线程世界中重建秩序。几个反直觉的例子当没有正确进行线程间同步时，以下例子都是可能的，它们反映了某些Relaxed Memory Model的行为。Store-Store乱序或Load-Load乱序{ data == 0, flag == 0 }// thread 0 thread 1//---------------------------------------- store 0($data), $1 loop: store 0($flag), $1 load %r0, 0($flag) beq %r0, $0, loop load %r1, 0($data)可能r1 == 0。本例在ARM上能重现Store-Load乱序{ x == 0, y == 0 }// thread 0 thread 1//---------------------------------------- store 0($x), $1 store 0($y), $1 load %r0, 0($y) load %r1, 0($x)可能r0 == 0 \u0026amp;\u0026amp; r1 == 0。本例在ARM/x86上能重现Dependent Loads乱序\\[^{[5]}\\]{ A == 1, B == 2, C == 3, P == \u0026amp;A, Q == \u0026amp;C }// thread 0 thread 1//-------------------------------------- store 0($B), $4 BARRIER store 0($P), $B load %r0, 0($P) load %r1 0(%r0)可能r0 == \u0026amp;B \u0026amp;\u0026amp; r1 == 2。本例在DEC Alpha上能重现Non-Causality/Non-Transitivity\\[^{[3]}\\]{ flag0 == 0, flag1 == 0 }// thread 0 thread 1 thread 2//----------------------------------------------------------------- store 0($flag0), $1 loop: load %r0, 0($flag0) beq %r0, $0, loop BARRIER store 0($flag1), $1 loop: load %r1, 0($flag1) beq %r1, $0, loop BARRIER load %r2, 0($flag0)可能r2 == 0。本例在不支持Causality的系统中能重现IRIW(Independent Read Independent Write)\\[^{[3]}\\]// thread 0 thread 1 thread 2 thread 3//--------------------------------------------------------------------------------- store 0($data1), 1 store 0($data2), 1 load %r1, 0($data1) load %r3, 0($data2) BARRIER BARRIER load %r2, 0($data2) load %r4, 0($data1)在r1 == 1 \u0026amp;\u0026amp; r3 == 1的前提下，可能r2 == 0 \u0026amp;\u0026amp; r4 == 0，即thread 2和3看见了不同的写顺序。本例在不支持Atomic Store的系统中能重现，比如某些NUMA和带SMT的UMA系统Memory Model的属性构成Memory Model主要规定不同地址上的读写操作在其他处理器看来会否乱序，以及，一个写操作是否同时被其他处理器观察到。Memory Ordering Load-Load Order：不同地址上的读操作会否乱序 Load-Store Order：读操作和后面另一个地址上的写操作会否乱序 Store-Load Order：写操作和后面的读操作会否乱序 Store-Store Order：不同地址上的写操作会否乱序 Dependent Loads Order：当第二条读操作的地址取决于前一条读操作的结果时，会否乱序写原子性(Store Atomicity)写原子性是指，处理器的写操作是否同时被所有处理器看到。根据写操作的同时性，从弱到强排序： Load Other’s Store Early \u0026amp;\u0026amp; Non-Causality：允许写操作被自己及个别其他处理器先看到，不支持Causality。写序列可能以不同顺序被多个处理器观察到 Load Other’s Store Early \u0026amp;\u0026amp; Causality：允许写操作被自己及个别其他处理器先看到，支持Causality Load Own Store Early：只允许写操作被自己先看到。写序列以相同顺序被多个处理器观察到 Atomic Store：所有处理器同时看到写操作 多处理器对Memory Model的影响\\[^{[4]}\\]现代处理器可能出于优化的考虑，放松下面这些限制的一条或多条。Memory Ordering的实现 Load-Load Order：当前一条Load指令因为操作数未就绪(RAW Hazard)或Cache Miss而必须等待时，乱序执行的CPU可能先执行后一条Load指令以掩盖停顿。如果Memory Model不允许这种乱序或程序员在两次Load之间插入了一条Barrier指令，那第一条Load执行完毕后，如果发现有Invalidate Message令后面要Load的Cache Line失效，就应该通过清空ROB撤销乱序执行的Load和后续指令然后重新执行 Load-Store Order：同上，当前一条Load因为各种原因等待时，后一条Store指令可能被先执行。如果Memory Model或Barrier指令限制这种乱序，在推断执行的CPU中往往不用做任何事，因为Store指令是在提交阶段才更新存储器 Store-Load Order：当CPU和Cache间有写缓冲(Write Buffer)时，写操作会被放入写缓冲而不是更新Cache，这可能导致Load指令执行过后，写操作仍然对其他处理器不可见。如果Memory Model或Barrier指令限制这种乱序，这要求在执行读操作之前，先Flush整个写缓冲，令结果写到Cache中，这里的开销非常大。由于开销非常大，现代处理器都允许这种乱序，且只有Full Barrier限制这种乱序 Store-Store Order：当CPU和Cache间有一个Non-FIFO或Coalescing的写缓冲时，如果前一个写操作导致Cache Miss或和其他写操作合并(因为写同一条Cache Line)，后一条结果可能先写入Cache。如果Memory Model或Barrier指令限制这种乱序，则要求采用FIFO的写缓冲，或在遇到Barrier指令时Flush整个写缓冲 Dependent Loads Order\\(^{[4,5]}\\)：即使处理器在语句”obj-\u0026gt;field = new_value”和”g_obj = obj;”间插入Barrier指令，强制Flush写缓冲区，另一个处理器的相关读操作”obj = g_obj; value = obj-\u0026gt;field”在取得新对象指针的同时仍然可能看到旧的字段值，因为新的字段值还在处理器的Invalidate Queue里面，要稍后才能处理。 写原子性的实现\\[^{[1,3]}\\] Load Other’s Store Early \u0026amp;\u0026amp; Non-Causality：在NUMA和支持硬件多线程的UMA中，Cache Coherence Message可能先抵达较近的逻辑处理器，从而导致写操作被部分处理器先看到。 Load Other’s Store Early \u0026amp;\u0026amp; Causality：一般通过类似Acquire/Release的Barrier指令可以获得Causality。 Load Own Store Early：这里要求除自己外的其他处理器同时看到写操作，开销很大，这就是不严格的所谓Atomic Store。 Atomic Store：该情形甚至不允许转发自己的写缓冲中的值，很少见，所以更多是理论价值 编译器对Memory Model的影响在不改变单线程执行结果的前提下，编译器倾向于进行各种代码变换来实施优化，这些优化会改变存储器访问的顺序和次数，故前面列举的Load-Load/Store-Load乱序等都可能发生。另外，一般来说，相同地址的内存访问受Data Dependence约束往往无法重排，但是一些激进优化也可能导致Dependent Loads乱序\\(^{[16]}\\)。Memory Consistency和Cache Coherence的区别Cache Coherence是多处理器的本地Cache导致多个数据副本在Cache和存储器间不一致的问题，Memory Model是多处理器和编译器优化导致存储器操作被多个处理器观察到的顺序不一致的问题。此外，Memory Consistency和Cache Coherence还有这些明显区别\\(^{[3]}\\)： 前者的研究对象是多个地址，后者的研究对象是单个Cache Line 就正确性而言，前者对程序员可见，后者对程序员透明，尽管后者影响性能 Memory Model可以实现在只有Incoherent Cache甚至没有Cache的多处理器系统上理想的Memory Model多处理器和编译器的诉求对多处理器而言，允许所有存储器操作的乱序可以提供最多的优化机会，哪怕其中一些顺序它能高效实现；允许部分处理器先观察到写结果，能够减少在Cache Coherence Protocol上阻塞的时间。类似的，编译器优化也期望能够自由地对代码做变换，只要这种变换不改变单线程执行的结果。故，多处理器和编译器都倾向于提供一种允许任意乱序、没有原子写要求的Memory Model。程序员的诉求程序员需要的Memory Model是，运行在多处理器系统上的共享内存多线程程序，看起来就像是并行或交错的多个顺序执行的线程，没有任何乱序、写操作也立即全局可见，同时每个线程和单线程程序运行得一样块，多线程的整个程序性能随处理器个数伸缩。调和二者的矛盾理想的Memory Model，能够尽量满足上面系统和程序员的诉求：允许各种乱序、非原子写，但只要程序员按一定的模式来组织程序、协调多个线程的通信，那这些混乱都是透明的。SC for DRF programs\\(^{[3]}\\)当两个线程同时访问一个地址，其中至少一个是写操作时，我们说发生了Data-Race。有一类良好组织Data-Race Free程序，运行在Sequential Consistency和Relaxed Consistency的Memory Model上都能得到一样的结果，这类程序的行为可以仅仅通过程序顺序(Programm Order)去推断，我们说这类程序叫SC for DRF programs(Sequential Consistency for Data-Race Free programs)。也就是说，SC for DRF的程序能用Sequential Consistency的Memory Model去推断正确性，却同时具备Relaxed Memory Model的性能(这往往意味着程序不依赖Atomic Store)基于存储器分类的Memory Model存储器分类： 私有数据(Private Data)：简称p_data，只被固定线程访问的数据。没有Data-Race、不需要Cache Coherence；其上的操作可以进行任何乱序，包括跨越s_data和sync_var的读写。 p_data对应的Cache Line一般常驻在特定处理器的本地Cache中，除非操作系统为负载均衡执行Migration。编译器能够识别局部变量但却无法很好识别thread local的堆数据；部分处理器如ARM允许将存储器标记为非共享的\\(^{[7]}\\)，从而禁用Cache Coherence并执行各种优化 共享数据(Shared Data)：简称s_data，允许多个线程访问，但任意时刻只被一个Owner持有。没有Data-Race，需要Cache Coherence；其上的读写可以进行任何乱序，但不能跨过sync_var。 只读的s_data会存在于多个处理器的本地Cache中，读写的s_data会随着Onwer的变化在多个处理器的本地Cache中迁移。在多处理器和高级语言中默认的存储器就是s_data 同步变量(Synchronization Variable)：简称sync_var，允许多个线程同时访问。有Data-Race，需要Cache Coherence；sync_var的读写之间不能进行乱序。 sync_var会被同时用于读写，在高争用(High Contention)的情况下对应的Cache Line会频繁出现在多个本地Cache中，然后因为特定处理器的写操作而失效。编译器通过高级语言提供的atomic或volatile声明识别出sync_var，进而生成Barrier指令来限制sync_var之间或sync_var与s_data之间的乱序，ARM也提供了Strongly-Ordered的声明以禁用sync_var之间的乱序 编程模式：利用上面的存储器分类，我们很容易实现锁操作，从而得到程序片段如”access p_data; lock(\u0026amp;sync_var); access s_data; unlock(\u0026amp;sync_var); access p_data”，可以看到，中间临界区(Critical Section)中对s_data的访问可以乱序，但不能跨过lock/unlock，两边对p_data的访问可以任意乱序包括调度进临界区。多处理器中的Memory Model强一致性模型(Strong Consistency Model)SC(Sequential Consistency)定义：将一个多线程程序各线程的操作按程序顺序(Program Order)交错在一起得到一个程序，如果这个单线程程序的执行结果和原来的多线程程序一样，我们就说执行程序的这个多处理器系统的Memory Model是Sequential Consistency的。 Leslie B. Lamport定义的SC\\(^{[18]}\\)：He first called a single processor (core) sequential if “the result of an execution is the same as if the operations had been executed in the order specified by the program.” He then called a multiprocessor sequentially consistent if “the result of any execution is the same as if the operations of all processors (cores) were executed in some sequential order, and the operations of each individual processor (core) appear in this sequence in the order specified by its program.”Memory Model的属性： LL/LS/SL/SS/DL乱序：不允许 Store Atomicity：Load Own Store Early弱一致性模型(Relaxed Consistency Model / Weak Ordering)在SC的基础上，我们逐步放松各个属性的限制，会依次得到各种弱一致性模型。允许SL乱序的模型TSO(Total Store Order)Memory Model的属性： SL乱序：允许 LL/LS/SS/DL乱序：不允许 Store Atomicity：Load Own Store Early优点：推断执行的CPU能高效实现TSO，具体来说： Load-Load Order：保持。前一个Load遇到Cache Miss的情况下，允许后一个Load先执行，但如果前一个Load返回发现Invalidate Queue当中的消息将使得后续的Load失效，则清空ROB撤销乱序执行重新发射指令 Load-Store Order：保持。推断执行CPU是顺序提交的，而Store指令的写存储器发生在提交阶段，此时Load指令已经Retire Store-Load Order：乱序。要使得Store指令先更新存储器，必须禁用写缓冲区或者在Load指令之前Flush写缓冲区 Store-Store Order：保持。利用FIFO、Non-Coalescing的写缓冲区，每条Store指令是顺序更新存储器的 锁实现的简化：Load/Store分别等效于Acquire/Release操作，实现锁时无需Barrier。Acquire指的是一个读操作，其后的Load/Store不允许乱序到前面，而Release是一个写操作，其前的Load/Store不允许乱序到后面。因为在TSO中，LL、LS乱序是不允许的，故Load指令直接可以用作Acquire；类似的，因为LS、SS乱序是不允许的，Store指令也可被用作Release。实现SpinLock等互斥锁时，Acquire操作可以被用于Lock，Release操作可以被用于Unlock，它们一起避免了临界区内的共享存储器的读写被乱序到临界区外，从而避免了Data-Race。由于在TSO中Load、Store分别等效于Acquire、Release，故在x86等系统中实现锁是不需要Barrier指令的。评价：这是一种相对较强的Memory Model，被x86采用\\(^{[21]}\\)，程序员可以简单地把它抽象成带写缓冲区(Write Buffer)的多处理器系统。PC(Processor Consistency)和TSO的区别： Store Atomicity：Load Other’s Store Early优点：相比TSO，放松了写一致性的要求，降低了Coherence Protocol上的开销。允许SS乱序的模型PST(Partial Store Order)和TSO的区别： SS乱序：允许优点：相比TSO，允许Non-FIFO和Coalescing的写缓冲区，允许多个Store指令合并或重叠的执行。允许LL/LS乱序的模型WO(Weak Ordering)定义：在WO中，共享存储器被划分为Data(对应前文s_data)和Synchronizing Variable(对应前文sync_var)。对Synchronizing Variable的读写要保持相对顺序，对Data的读写允许乱序，但不能跨过对Synchronizing Variable的操作。 Michel Dubols等人定义的WO\\(^{[19]}\\)：In a multiprocessor system, storage acceses are weakly ordered if 1) accesses to global synchronizing variables are strongly ordered and if 2) no access to a synchronizing variable is issued in a processor before all previous global data accesses have been performed and if 3) no access to global data is issued by a processor before a previous access to a synchronizing varible has been performedMemory Model的属性： LL/LS/SL/SS乱序：允许(Data读写)，不允许(Synchronizing Variable读写) DL乱序：不允许 Store Atomicity：Load Own Store Early评价：WO是一种粗粒度的Memory Model，程序员可以先把所有存在Data-Race的存储器标记为Synchronizing Variable来确保程序正确性，之后再优化，故易于编程。因为一般程序中Data的访问远多于Synchronizing Variable，其上的操作都可乱序，所以WO的性能也相当不错。WO是一种在性能和可编程性上有很好折中的Memory Model。RC(Release Consistency)定义：在RC中，共享存储器操作被划分为Ordinary(对应前文s_data)和Special，其中Special进步一划分为Sync(对应前文sync_var)和NSync，Sync又包括Acquire和Release。Ordinary操作的乱序不允许向前跨过Acquire、不允许向后跨过Release，根据Special操作之间是Sequential Consistency或Processor Consistency，RC又可以进一步细分为RCsc和RCpc。RCsc：Memory Model的属性： LL/LS/SL/SS乱序：允许(Ordinary操作) DL乱序：不允许 Store Atomicity：Load Own Store EarlyRCpc：Memory Model的属性： LL/LS/SL/SS乱序：允许(Ordinary操作) DL乱序：不允许 Store Atomicity：Load Other’s Store Early \u0026amp;\u0026amp; Causality Kourosh Gharachorloo等人定义的RCpc\\(^{[20]}\\)：(A) before an ordinary LOAD or STORE access is allowed to perform with respect to any other processor, all previous acquire access must be performed, and (B) before a release access is allowed to perform with respect to any other processor, all previous ordinary LOAD and STORE accesses must be performed, and (C) special accesses are processor consistent with respect to one another评价：当Ordinary操作(对应前文s_data)仅出现在临界区内时，用WO实现的锁和RCsc等价，比RCpc版本多了写原子性；在更复杂的Lockless算法中，RCpc提供了比WO更细粒度的重排序控制。RC是一种提供细粒度控制的常见Memory Model。ARMv7定义：ARMv7的存储器被划分为Strongly-Ordered(对应前文sync_var)、Device和Normal，其中Normal又分为Shareable(对应前文s_data)和Non-Shareable(对应前文p_data)。Strongly-Ordered和Device上的操作之间不允许乱序，Normal的操作可以任意乱序。要限制Normal操作的顺序需要显示的Barrier指令。Memory Model的属性： LL/LS/SL/SS乱序：允许(Normal操作)，不允许(Strongly-Ordered/Device操作) DL乱序：不允许 Store Atomicity：Load Own Store Early允许DL乱序的模型DEC AlphaMemory Model的属性： LL/LS/SL/SS/DL乱序：允许 Store Atomicity：Load Own Store Early评价：DEC Alpha的生命期实际上在2001年已经宣告结束，我们之所以讨论它，是因为它是唯一一款允许Dependent Loads乱序的多处理器，这使得它有几乎最Weak的Memory Model，而这又促成了Linux Kernel的Barrier设计(Data Dependency Barrier)以及C++标准的memory model设计(std::memory_order_consume)。 引自Paul E. McKenney\\(^{[4]}\\)：Alpha is interesting because, with the weakest memory ordering model, it reorders memory operations the most aggressively. It therefore has defined the Linux-kernel memory-ordering primitives, which must work on all CPUs, including Alpha. Understanding Alpha is therefore surprisingly important to the Linux kernel hacker高级语言中的Memory ModelLinux Kernel和高级语言标准都定义了自己的Memory Model，其中有专门的同步操作用于线程间协作。编译器负责将这个抽象的Memory Model映射到目标多处理器上，如果多处理器自己的Memory Model相对更强，那么上层Memory Model的同步操作可能退化成普通的存储器访问；如果目标多处理器的Memory Model相对更弱，则部分上层同步操作可能生成处理器提供的Barrier或Atomic指令来强制顺序和写原子性。Linux Kernel\\[^{[5]}\\]Linux Kernel的代码早于Memory Model的概念被引进C语言(2011)，所以必须自己面对多处理器下存储器操作乱序的问题。为更好的可移植性，Kernel假设多处理器有一个最弱的Memory Model(不强于DEC Alpha)、编译器也会进行任意乱序。为在这种环境下正确的编写多线程程序和设备驱动，Linux使用了如下Barrier：编译器Barrier READ_ONCE(x) 和 WRITE_ONCE(x)：用于存储器的读写，避免这些读写彼此被编译器乱序或优化掉处理器Barrier Store Barrier：前后的Store指令不能越过Barrier乱序，搭配Load Barrier或Data Dependency Barrier Load Barrier：前后的Load指令不能越过Barrier乱序，搭配Store Barrier Data Dependency Barrier：依赖于前一条Load指令结果的第二条Load指令不能越过前一条Load乱序，搭配Store Barrier。这是Load Barrier的优化版本，用于仅需要避免读相关乱序的场合，例如Producer更新对象字段后将对象指针传给Consumer使用 General Barrier：相当于Store+Load Barrier，搭配Load/Store/Data Dependency Barrier Acquire操作：之后的读写不能向前乱序，搭配Release操作 Release操作：之前的读写不能向后乱序，搭配Acquire操作 其中编译器Barrier是通过C语言的volatile实现的(比如将目标地址转化为用volatile修饰的指针后再读写)，而处理器Barrier根据目标多处理器的Memory Model生成读写指令和必要的Barrier指令，保证了Causality。C++语言标准\\[^{[14]}\\]std::memory_order：std::atomic的变量上的操作，可以用以下参数来强制某种顺序或写原子性： std::memory_order_relaxed：仅保持变量自身读写的相对顺序 std::memory_order_consume：依赖于该读操作的后续读写，不能往前乱序；另一个线程上std::memory_order_release之前的相关写序列，在std::memory_order_consume同步之后对当前线程可见 std::memory_order_acquire：之后的读写不能往前乱序；另一个线程上std::memory_order_release之前的写序列，在std::memory_order_acquire同步之后对当前线程可见 std::memory_order_release：之前的读写不能往后乱序；之前的写序列，对使用std::memory_order_acquire/std::memory_order_consume同步的线程可见 std::memory_order_acq_rel：两边的读写不能跨过该操作乱序；写序列仅在同步线程之间可见 std::memory_order_seq_cst：两边的读写不能跨过该操作乱序；写序列的顺序对所有线程相同 实现Memory Model\\[^{[13]}\\]C++中的volatile/std::atomic有以下用法，分别对应不同的Memory Model，编译器可以沿用实现这些Memory Model时的手段支持这些用法： volatile：只应用于单线程，等同Linux Kernel中的编译器Barrier。(MSVC中为向后兼容赋予了volatile等同Java的volatile的职责，即读是Acquire写是Release) std::memory_order_relaxed + Read-Modify-Write：原子操作，保持变量自身读写的相对顺序 std::memory_order_seq_cst：相当于WO。因为WO良好的可编程性\\(^{[1]}\\)，该标志是标准库的默认值，一般用它快速实现算法，进一步优化才考虑下面的RCpc std::memory_order_acq_rel：相当于弱化的WO(写原子性换成了Load Other’s Store Early \u0026amp;\u0026amp; Causality，故和std::memory_order_seq_cst相比通不过IRIW测试) std::memory_order_acquire + std::memory_order_release：相当于RCpc std::memory_order_consume + std::memory_order_release：相当于弱化的RCpc。同Linux Kernel中的Data Dependency Barrier + Store Barrier 几个简单的例子这里用先快速实现、再优化(先WO，再RCpc)，两步走的方式，为几个简单算法加上std::memory_order标志。自旋锁(SpinLock)Version 1：基于WO(std::memory_order_seq_cst)struct SpinLock { void lock() { for (;;) { while (lock_.load()); if (!lock_.exchange(true)) break; } } void unlock() { lock_.store(false); } std::atomic\u0026lt;bool\u0026gt; lock_ = {false};};因为WO本来就是最容易编程的Memory Model，同步变量读写彼此不会乱序，临界区内外的操作也不能跨过同步变量读写乱序，故只要以默认标志编写算法即可。Version 2：基于RCpc(std::memory_order_acquire + std::memory_order_release)struct SpinLock2 { void lock() { for (;;) { while (lock_.load(std::memory_order_relaxed)); if (!lock_.exchange(true, std::memory_order_acquire)) break; } } void unlock() { lock_.store(false, std::memory_order_release); } std::atomic\u0026lt;bool\u0026gt; lock_ = { false };};RCpc是细粒度的Memory Model，程序正确性不容易保证，需要我们逐项自检： 线程安全吗？——线程安全性由改写前的Version 1保证 三种存储器操作的顺序 禁止临界区内s_data读写乱序到外面吗(Data-Race)？——lock的最后一个操作是std::memory_order_acquire，避免了向前乱序；unlock的最后一条操作是std::memory_order_release，避免了向后乱序 允许临界区外p_data读写乱序到里面吗(优化)？——lock中只有std::memory_order_acquire，故允许前面的代码向后乱序；unlock中只有std::memory_order_release，故允许后面的代码向前乱序 禁止sync_var读写间乱序吗(正确性)？——单个std::atomic变量的读写不被乱序 写原子性的要求？——Load Other’s Store Early \u0026amp;\u0026amp; Non-Causality 另一种分析正确性的方法，是根据Happens-Before关系推导读写锁(Readers-Writer Lock)Version 1：基于WOstruct RWLock { void rlock() { for (;;) { while (writer_.load()); readers_.fetch_add(1); if (!writer_.load()) break; readers_.fetch_sub(1); } } void runlock() { readers_.fetch_sub(1); } void wlock() { while (writer_.exchange(true)); while (readers_.load() \u0026gt; 0); } void wunlock() { writer_.store(false); } std::atomic\u0026lt;bool\u0026gt; writer_ = {false}; std::atomic\u0026lt;int\u0026gt; readers_ = {0};};Version 2：基于RCpcstruct RWLock2 { void rlock() { for (;;) { while (writer_.load(std::memory_order_acquire)); readers_.fetch_add(1, std::memory_order_acquire); if (!writer_.load(std::memory_order_acquire)) break; readers_.fetch_sub(1, std::memory_order_acquire); } } void runlock() { readers_.fetch_sub(1, std::memory_order_release); } void wlock() { while (writer_.exchange(true, std::memory_order_acquire)); while (readers_.load(std::memory_order_acquire) \u0026gt; 0); } void wunlock() { writer_.store(false, std::memory_order_release); } std::atomic\u0026lt;bool\u0026gt; writer_ = {false}; std::atomic\u0026lt;int\u0026gt; readers_ = {0};};自检： 线程安全吗？——由Version 1保证 三种存储器操作的顺序 禁止临界区内s_data读写乱序到外面吗(Data-Race)？——rlock/wlock的最后一个操作是std::memory_order_acquire，runlock/wunlock的最后一条操作是std::memory_order_release 允许临界区外p_data读写乱序到里面吗(优化)？——rlock/wlock中只有std::memory_order_acquire，runlock/wunlock中只有std::memory_order_release 禁止sync_var读写间乱序吗(正确性)？——rlock/wlock内全是std::memory_order_acquire，避免了两个不同地址读写间的乱序 写原子性的要求？——Load Other’s Store Early \u0026amp;\u0026amp; Non-Causality双重检查的单例模式(Double-Checked Locking in Singleton)C++11标准的6.7节确保了传统的static局部变量用作Singleton已经线程安全没必要用Double-Checked，这里仅作演示目的。Version 1：基于WOtemplate\u0026lt;typename T\u0026gt;struct Singleton { static T* get() { T *p = instance_s.load(); if (p == nullptr) { std::lock_guard\u0026lt;std::mutex\u0026gt; guard(mutex_s); p = instance_s.load(); if (p == nullptr) { p = new T(); instance_s.store(p); } } return p; } static std::atomic\u0026lt;T*\u0026gt; instance_s; static std::mutex mutex_s;};Version 2：基于弱化的RCpc(std::memory_order_consume + std::memory_order_release)template\u0026lt;typename T\u0026gt;struct Singleton2 { static T* get() { T *p = instance_s.load(std::memory_order_consume); if (p == nullptr) { std::lock_guard\u0026lt;std::mutex\u0026gt; guard(mutex_s); p = instance_s.load(std::memory_order_relaxed); if (p == nullptr) { p = new T(); instance_s.store(p, std::memory_order_release); } } return p; } static std::atomic\u0026lt;T*\u0026gt; instance_s; static std::mutex mutex_s;};自检： 线程安全吗？——由Version 1保证 存储器操作顺序 创建对象是在写入instance_s指针前完成的吗？——std::memory_order_release保证了new T()中的操作不会被乱序到后面 访问对象是在读取instance_s指针后完成的吗？(这是典型的Dependent Loads场景) 第1次load：std::memory_order_consume避免了get()返回后的字段访问被乱序到前面(避免编译器激进优化\\(^{[16]}\\)和多处理器的Dependent Loads\\(^{[4]}\\)) 第2次load：当p非nullptr返回时，使用std::memory_order_relaxed无法避免get()后的字段访问被乱序到前面(从CPU实际执行的指令序列来看)，但这种乱序是安全的，因为当前线程的mutex_s上的lock已经与对象创建线程上的unlock同步过，对象字段和instance_s不会再修改 写原子性的要求？——Load Other’s Store Early \u0026amp;\u0026amp; Non-Causality 推荐阅读数字标号是推荐顺序，其他资料作为参考。主要资源 Shared Memory Consistency Models: A Tutorial 补充阅读： Computer Architecture: A Quantitative Approach A Primer on Memory Consistency and Cache Coherence Is Parallel Programming Hard, And, If So, What Can You Do About It? The C++11 Memory Model and GCC 补充阅读： std::memory_order What Every Systems Programmer Should Know About Lockless Concurrency References [1] Shared Memory Consistency Models: A Tutorial [2] Computer Architecture: A Quantitative Approach [3] A Primer on Memory Consistency and Cache Coherence [4] Is Parallel Programming Hard, And, If So, What Can You Do About It? [5] Linux Kernel Memory Barriers [6] A Tutorial Introduction to the ARM and POWER Relaxed Memory Models [7] ARM Cortex-A Series. Programmer’s Guide [8] JSR 133 (Java Memory Model) FAQ [9] The JSR-133 Cookbook for Compiler Writers [10] The Java Language Specification, Java SE 9 Edition [11] The C# Memory Model in Theory and Practice [12] C# Language Specification 5.0 [13] The C++11 Memory Model and GCC [14] std::memory_order [15] What Every Systems Programmer Should Know About Lockless Concurrency [16] C++ Data-Dependency Ordering: Atomics and Memory Model [17] Wiki: Compiler optimizations [18] How to Make a Multiprocessor Computer that Correctly Executes Multiprocess Programs [19] Memory access buffering in multiprocessors [20] Memory Consistency and Event Ordering in Scalable Shared-Memory [21] x86-TSO: A Rigorous and Usable Programmer’s Model for x86 Multiprocessors [22] Time, Clocks and the Ordering of Events in a Distributed System" },

  { "title": "再理解 std::condition_variable 条件变量", "url": "/blog/2025/conditonal_variable/", "categories": "Cpp",
  "tags": "Cpp", "date": "2025-12-06", "content":
  "有如下几个问题需要厘清： 工作原理，即流程； 虚假唤醒与唤醒丢失； notify_one 与 notify_all。 lock_guard 与 unique_lock。#include \u0026lt;atomic\u0026gt;#include \u0026lt;condition_variable\u0026gt;#include \u0026lt;iostream\u0026gt;#include \u0026lt;thread\u0026gt;std::mutex mutex_;std::condition_variable condVar;std::atomic\u0026lt;bool\u0026gt; dataReady{false}; // (0)void waitingForWork() { std::cout \u0026lt;\u0026lt; \"Waiting \" \u0026lt;\u0026lt; std::endl; std::unique_lock\u0026lt;std::mutex\u0026gt; lck(mutex_); condVar.wait(lck, []{ return dataReady.load(); }); // (1) std::cout \u0026lt;\u0026lt; \"Running \" \u0026lt;\u0026lt; std::endl; // do the work}void setDataReady() { { std::lock_guard\u0026lt;std::mutex\u0026gt; lck(mutex_); // (2) // prepare the data dataReady = true; } std::cout \u0026lt;\u0026lt; \"Data prepared\" \u0026lt;\u0026lt; std::endl; condVar.notify_one(); // (3)}int main(){ std::thread t1(waitingForWork); std::thread t2(setDataReady); t1.join(); t2.join(); return 0;}其中，带谓词(predicate)的wait等效于：std::unique_lock\u0026lt;std::mutex\u0026gt; lck{mutex_}; // (a)while(![]{return dataReady.load();}) { // (b) //time window(1) condVar.wait(lck); // (c)}consumer获取到锁之后： 如果谓词返回false，则通过wait操作释放锁（调用lck.unlock()），并进入等待状态，直到notify唤醒。 被唤醒之后，重新获取锁（调用lck.lock()），进入时间窗口(1)，再次检查谓词。 如果谓词返回true，则继续持锁继续后续的工作，直到释放锁。参见std::condition_variable::wait – 中文版。1. 工作原理及流程即，对producer，consumer两者需要确保互锁，并且producer进行notify的时候，consumer已经进入等待状态。如果不使用带谓词的wait，则会出现唤醒丢失的问题，即producer发送notify的时候，consumer还没有进入等待状态，原因就是需要等待lock：std::unique_lock\u0026lt;std::mutex\u0026gt; lck{mutex_};condVar.wait(lck);并且consumer进入等待状态之后，producer、consumer会进入死锁状态。另一方面，如果不使用原子变量dataReady，则会出现虚假唤醒的问题。即，如果dataReady是一个普通变量，可能由于缓存不一致性，导致consumer读取到的值不正确，如果读取到的值为true，则出现虚假唤醒。2. notify_one 与 notify_all当使用notify_all的时候，后续被唤醒的consumer需要等第一个consumer释放锁之后，才能重新获取锁，然后再经过一次谓词检查，此时，可能谓词返回false，则继续进入等待状态。3. lock_guard 与 unique_lock对于producer，使用lock_guard即可，因为只需要持锁一次操作完成之后随即释放，没有再次持锁的需求。也可以使用``unique_lock，但需要保证在notify`之前释放锁。对于consumer，必须使用unique_lock，因为wait需要传入一个unique_lock对象，并且在wait过程中会释放锁，等待被唤醒之后重新获取锁继续工作。参考 cppreference: std::condition_variable C++ Core Guidelines: Be Aware of the Traps of Condition Variables 中文翻译：条件变量condition_variable的使用及陷阱" },

  { "title": "使用 libevent 实现时间戳调度", "url": "/blog/2025/timestamp_schedule_using_libevent/", "categories": "Cpp",
  "tags": "Libevent, Cpp", "date": "2025-12-05", "content":
  "1. 介绍libevent底层使用不同的事件通知机制：linux 使用epoll，Windows使用IOCP。它提供了一个统一的接口来处理网络事件和定时事件。使用libevent的几个基本操作步骤：1.1. 创建 event_base 对象，作为事件循环的核心struct event_base *base = event_base_new(); event_base只能在一个线程中使用。1.2. 创建事件处理器struct event *ev = event_new(base, fd, EV_READ | EV_PERSIST, my_read_callback, NULL);fd是文件描述符，EV_READ | EV_PERSIST指定了事件类型（这里是读事件并且是持久的，即事件被触发后不会自动删除），my_read_callback是事件触发时调用的函数，最后一个参数是传递给回调函数的用户数据。回调函数在同一个线程中执行。1.3. 添加事件到事件循环中event_add(ev, NULL);1.4. 启动事件循环event_base_dispatch(base);判断事件循环是否应该退出：const auto got_break = event_base_got_break(base);const auto got_exit = event_base_got_exit(base);return got_break == 0 \u0026amp;\u0026amp; got_exit == 0;1.5. 清理资源event_free(ev);event_base_free(base);2. 使用 libevent 实现时间戳调度创建event_base对象：void PcapngPlayer::start() { // ... if (base_ = event_base_new(); !base_) { spdlog::error(\"Could not initialize libevent!\"); return; } // ... scheduleNextPacket(); event_base_dispatch(base_);}创建event，实现本次帧时间计算及调度启动：void PcapngPlayer::scheduleNextPacket() { if (curr_idx_ \u0026gt; end_idx_ || curr_idx_ \u0026gt;= udp_packets_.size()) { event_base_loopbreak(base_); return; } const int64_t diff = timestamp_u64(udp_packets_[curr_idx_]) - curr_ts_; timeval wait = diff \u0026gt; 0 ? timeval_from_us(diff) : timeval_from_us(0); timer_event_ = evtimer_new(base_, timerCallback, this); evtimer_add(timer_event_, \u0026amp;wait);}实现定时器回调函数，销毁本次定时器event，以及下一帧启动：void PcapngPlayer::timerCallback(evutil_socket_t fd, short what, void* arg) { PcapngPlayer* player = static_cast\u0026lt;PcapngPlayer*\u0026gt;(arg); // ... player-\u0026gt;curr_ts_ = timestamp_u64(udp_info); player-\u0026gt;curr_idx_++; event_free(player-\u0026gt;timer_event_); player-\u0026gt;timer_event_ = nullptr; player-\u0026gt;scheduleNextPacket();}最后，需要在启动之后，轮询等待事件循环退出：while (true) { // 检查libevent事件循环是否退出 // ... std::this_thread::sleep_for(std::chrono::milliseconds(100));}3. libevent事件循环int event_base_loop(struct event_base *base, int flags) { // ... 初始化与锁 ... while (!done) { // 1. 检查是否需要立即停止 if (base-\u0026gt;event_gotterm) { break; } // 2. 计算后端等待时间 (Timeout) // 从最小堆中取出最近的一个定时器时间 tv_p = \u0026amp;tv; if (!base-\u0026gt;event_count_active \u0026amp;\u0026amp; !(flags \u0026amp; EVLOOP_NONBLOCK)) { timeout_next(base, \u0026amp;tv_p); } else { // 如果有激活事件，或者是非阻塞模式，则不等待 evutil_timerclear(\u0026amp;tv); } // 3. 调用后端 (Backend Dispatch) // 例如调用 epoll_wait，将就绪的 fd 对应的 event 插入激活队列 res = base-\u0026gt;evsel-\u0026gt;dispatch(base, tv_p); // 4. 处理超时事件 timeout_process(base); // 5. 处理激活队列 (Process Active Events) if (base-\u0026gt;event_count_active) { event_process_active(base); } else if (flags \u0026amp; EVLOOP_ONCE) { // 如果是 EVLOOP_ONCE 模式且没有事件被激活，则退出 done = 1; } } // ... 清理 ... return 0;} Event Base 与 Event Loop" },

  { "title": "QML基础语法积累", "url": "/blog/2025/QML%E7%AC%94%E8%AE%B0/", "categories": "Qt",
  "tags": "QML, Qt", "date": "2025-12-03", "content":
  "1. 信号与 Connections信号一般用于组件封装，在外部声明的组件内声明信号处理器。Connections则提供另外一种方式：即在被调用组件外部使用信号并使用信号处理器。1.1. 信号的语法及使用信号语法：signal \u0026lt;signalName\u0026gt;([\u0026lt;parameterName\u0026gt;:\u0026lt;parameterType\u0026gt;[,...]])信号处理器语法：// on: 固定关键字// Signal: 信号名，首字母必须大写on\u0026lt;Signal\u0026gt;示例：Item { signal mySignal(param1: int, param2: string) MouseArea { anchors.fill: parent onClicked: { mySignal(42, \"Hello\") } } // Qt6 要求信号处理器显式声明参数 onMySignal: (param1, param2) =\u0026gt; { console.log(\"Signal received with param1:\", param1, \"and param2:\", param2) }}1.2. 属性值改变信号QML类型提供内建属性值改变信号，这个属性属性值改变就会自动发出信号。属性值改变信号的命名规则是 on\u0026lt;PropertyName\u0026gt;Changed。示例：Rectangle { width: 100 height: 100 color: \"red\" onWidthChanged: width =\u0026gt; { console.log(\"Width changed to:\", width); } MouseArea { anchors.fill: parent onClicked: { width += 10 } }}1.3. ConnectionsConnections 允许在组件外部处理信号。语法如下：import QtQuickimport QtQuick.ControlsWindow { width: 640 height: 480 visible: true title: qsTr(\"Hello World\") // 连接button1，信号被分发到这 Connections { target: button1 function onClicked() { console.log(\"cnnection func\") } } // 被连接对象 Button { id: button1 width: 200 height: 100 text: \"button1\" font.pixelSize: 30 onClicked: { // 直接定义信号处理器的使用方式 console.log(\"button1 func\") } }}1.4. 自定义C++/QML交互：C++模块属性及在QML中使用 在C++模块中，使用 Q_PROPERTY 宏定义属性，并指定 NOTIFY 信号。此时，QML信号处理器的语法与属性值改变信号相同：on\u0026lt;PropertyName\u0026gt;Changed。 在QML中，可以使用Connections连接C++对象的信号，也可以使用直接绑定到C++属性的方式。示例：class NuoQianMapController : public QObject { Q_OBJECT Q_PROPERTY(QGeoCoordinate center READ center WRITE setCenter NOTIFY centerChanged) Q_PROPERTY(double zoomLevel READ zoomLevel WRITE setZoomLevel NOTIFY zoomLevelChanged) // ...省略其他代码...};Connections使用示例：map { Connections { target: NuoQianMapController function onCenterChanged() { // 处理中心点变化 } function onZoomLevelChanged() { // 处理缩放级别变化 } }}直接绑定属性示例：map { center: NuoQianMapController.center zoomLevel: NuoQianMapController.zoomLevel // 也可以使用property绑定，如果是复杂类型，则在QML中property定义为var类型 property var mapCenter: NuoQianMapController.center} 注意：如果C++信号中包含不被QML支持的参数类型，则无法在QML中使用该信号处理器。另外，信号不能同名，即不能重载，QML无法区分。2. Model-View-DelegateQML中的Model-View-Delegate模式用于处理数据的显示和交互。Model表示数据源，View负责显示数据，而Delegate定义了每个数据项的外观和行为。模型数据一般定义成如下几种类型：QStringList，QVariantList，QObjectList，QAbstractItemModel。使用方式是：在view中指定model属性为数据模型，然后使用delegate属性定义每个数据项的外观和行为。示例：class NuoQianCoverageAreaManager : public QObject { Q_OBJECT Q_PROPERTY(QQmlListProperty\u0026lt;NuoQianCoverageArea\u0026gt; coverageAreas READ coverageAreas NOTIFY coverageAreasChanged) Q_PROPERTY(int count READ count NOTIFY coverageAreasChanged) // ...省略其他代码...};// view.qmlMapItemView { id: view model: NuoQianCoverageAreaManager.coverageAreas delegate: CoverageAreaView { areaData: modelData }}// delegate实现// CoverageAreaView.qmlMapItemGroup { id: coverageGroup property var areaData: null // 覆盖区域填充 MapPolygon { path: areaData ? areaData.borderPath : [] color: areaData ? areaData.color : \"#4444FF88\" border.color: areaData ? Qt.darker(areaData.color, 1.5) : \"#3333DD\" border.width: 2 visible: areaData \u0026amp;\u0026amp; areaData.visible \u0026amp;\u0026amp; path.length \u0026gt; 0 z: 40 } // ...省略其他代码...} 注意：有两种方式将model与view连接起来：deletegate 与 Repeater。比如Qt内置的ListView，GridView等，使用Repeater，会将所有数据都加载进来，使用delegate将根据视图需要进行加载。 注意：delegate会针对每一项数据生成一个delegate实例。个人理解：增加delegate这一个桥接层，可以在model与view之间进行自定义的处理，比如数据过滤，数据转换等，以及使用DelegateChooser动态切换不同的delegate。 Model/View Programming Writing QML Extensions with C++" },

  { "title": "交叉编译 Qt 5.15.2", "url": "/blog/2025/cross_build_qt_arm64_linaro_7.5/", "categories": "Qt",
  "tags": "Qt", "date": "2025-09-20", "content":
  "1. 下载交叉编译器下载地址目录：aarch64-linux-gnu gcc-linaro-7.5.0-2019.12-x86_64_aarch64-linux-gnus sysroot-glibc-linaro-2.25-2019.12-aarch64-linux-gnu下载之后，将编译器及sysroot解压到/opt目录下：/opt/gcc-linaro-7.5.0-2019.12-x86_64_aarch64-linux-gnu/opt/sysroot-glibc-linaro-2.25-2019.12-aarch64-linux-gnuQt 5.15.2下载目录：Qt 5.15.2 。选择下载 qt-everywhere-src-5.15.2.tar.xz。2. 修改Qt源码修改头文件：qtbase/src/corelib/global/qglobal.h，include添加：// 第44行开始的地方添加 \u0026lt;limits\u0026gt;，解决configure的时候报错找不到\u0026lt;limits\u0026gt;以及limits相关错误# include \u0026lt;limits\u0026gt;// ....拷贝（不拷贝，configure 的时候报错：Invalid target platform ‘aarch64-linux-gnu-g++’）：cp -r qtbase/mkspecs/linux-aarch64-gnu-g++ qtbase/mkspecs/aarch64-linux-gnu-g++修改文件qtbase/mkspecs/aarch64-linux-gnu-g++/qmake.conf：# modifications to g++.confQMAKE_CC = /opt/gcc-linaro-7.5.0-2019.12-x86_64_aarch64-linux-gnu/bin/aarch64-linux-gnu-gccQMAKE_CXX = /opt/gcc-linaro-7.5.0-2019.12-x86_64_aarch64-linux-gnu/bin/aarch64-linux-gnu-g++QMAKE_LINK = /opt/gcc-linaro-7.5.0-2019.12-x86_64_aarch64-linux-gnu/bin/aarch64-linux-gnu-g++QMAKE_LINK_SHLIB = /opt/gcc-linaro-7.5.0-2019.12-x86_64_aarch64-linux-gnu/bin/aarch64-linux-gnu-g++# modifications to linux.confQMAKE_AR = /opt/gcc-linaro-7.5.0-2019.12-x86_64_aarch64-linux-gnu/bin/aarch64-linux-gnu-ar cqsQMAKE_OBJCOPY = /opt/gcc-linaro-7.5.0-2019.12-x86_64_aarch64-linux-gnu/bin/aarch64-linux-gnu-objcopyQMAKE_NM = /opt/gcc-linaro-7.5.0-2019.12-x86_64_aarch64-linux-gnu/bin/aarch64-linux-gnu-nm -PQMAKE_STRIP = /opt/gcc-linaro-7.5.0-2019.12-x86_64_aarch64-linux-gnu/bin/aarch64-linux-gnu-strip3. 编译 Qt 源码如下为编译一个比较精简的Qt库，去掉了很多模块。#!/bin/shcd build-qt5-qtbaseyes | rm -rf *export SYSROOT=/opt/sysroot-glibc-linaro-2.25-2019.12-aarch64-linux-gnuexport CROSS_PREFIX=aarch64-linux-gnu-export DEST_PREFIX=/opt/qt5-aarch64-gcc-linaro-7.5.0export TOOLCHAIN_PREFIX=/opt/gcc-linaro-7.5.0-2019.12-x86_64_aarch64-linux-gnu/binexport PATH=$TOOLCHAIN_PREFIX:$PATHexport PKG_CONFIG_LIBDIR=${SYSROOT}/lib/pkgconfig:${SYSROOT}/share/pkgconfigexport PKG_CONFIG_PATH=${SYSROOT}/lib/pkgconfig:${SYSROOT}/share/pkgconfig../configure -pkg-config -xplatform aarch64-linux-gnu-g++ -prefix $DEST_PREFIX -release -opensource -confirm-license -sysroot $SYSROOT -nomake tests -nomake examples -skip qtwebengine -skip qt3d -skip qtwebview -skip qtnetworkauth -skip qtserialport -skip qtsensors -skip qtmultimedia -skip qtdoc -skip qtmacextras -skip qtandroidextras -no-opengl# 添加skip qtlocation，qtlocation编译报错# ../configure -pkg-config -xplatform aarch64-linux-gnu-g++ -prefix $DEST_PREFIX -release -opensource -confirm-license -sysroot $SYSROOT -nomake tests -nomake examples -skip qtwebengine -skip qt3d -skip qtwebview -skip qtnetworkauth -skip qtserialport -skip qtsensors -skip qtmultimedia -skip qtdoc -skip qtmacextras -skip qtandroidextras -skip qtlocation -no-opengl# 去掉几乎所有界面相关模块，只保留QtGui，QtWidget界面组件# ../configure -pkg-config -xplatform aarch64-linux-gnu-g++ -prefix $DEST_PREFIX -release -opensource -confirm-license -sysroot $SYSROOT -nomake tests -nomake examples -skip qtwebengine -skip qt3d -skip qtwebview -skip qtnetworkauth -skip qtserialport -skip qtsensors -skip qtmultimedia -skip qtdoc -skip qtmacextras -skip qtandroidextras -skip qtlocation -skip qtdeclarative -skip qtquickcontrols -skip qtquickcontrols2 -skip qtgraphicaleffects -no-openglmake -j8sudo make install4. 资料 qt5.15.2 交叉编译 arm64" },

  { "title": "使用 C++ template 构建一个通信用序列化反序列化模块", "url": "/blog/2025/%E6%9E%84%E5%BB%BA%E4%B8%80%E4%B8%AA%E9%80%9A%E4%BF%A1%E7%94%A8%E5%BA%8F%E5%88%97%E5%8C%96%E5%8F%8D%E5%BA%8F%E5%88%97%E5%8C%96%E6%A8%A1%E5%9D%97/", "categories": "Cpp",
  "tags": "Cpp", "date": "2025-08-26", "content":
  "1. 序列化 PacketSerializer分为以下几个步骤： 需要一个 buffer 来存储序列化过程中的数据，使用 std::vector 作为底层存储。 提供一个模板实现的基础函数 writeToBuffer，将基本类型数据写入 buffer。 提供一个模板函数 pack，支持将基础类型数据，以及其数组，以及 vector 类型进行序列化。 在 pack 函数的基础上，提供一个模板函数 packMultiple，支持将多个数据打包进行序列化。 提供接口函数 finalize，添加包头、包尾，组装成完整的协议，返回序列化后的数据 buffer。1.1. writeToBuffer template \u0026lt;typename T\u0026gt; void writeToBuffer(const T\u0026amp; value) { static_assert(std::is_integral_v\u0026lt;T\u0026gt; || std::is_floating_point_v\u0026lt;T\u0026gt;, \"Type must be arithmetic\"); const uint8_t* bytes = reinterpret_cast\u0026lt;const uint8_t*\u0026gt;(\u0026amp;value); for (size_t i = 0; i \u0026lt; sizeof(T); i++) { buffer_.push_back(bytes[i]); } } template \u0026lt;typename T\u0026gt; void writeArrayToBuffer(const T* array, size_t count) { static_assert(std::is_integral_v\u0026lt;T\u0026gt; || std::is_floating_point_v\u0026lt;T\u0026gt;, \"Array element type must be arithmetic\"); for (size_t i = 0; i \u0026lt; count; i++) { writeToBuffer(array[i]); } }1.2. pack template \u0026lt;typename T\u0026gt; typename std::enable_if_t\u0026lt;std::is_integral_v\u0026lt;T\u0026gt; || std::is_floating_point_v\u0026lt;T\u0026gt;, PacketSerializer\u0026amp;\u0026gt; // pack(const T\u0026amp; value) { writeToBuffer(value); return *this; } template \u0026lt;typename T, size_t N\u0026gt; typename std::enable_if_t\u0026lt;std::is_integral_v\u0026lt;T\u0026gt; || std::is_floating_point_v\u0026lt;T\u0026gt;, PacketSerializer\u0026amp;\u0026gt; // pack(const T (\u0026amp;array)[N]) { writeArrayToBuffer(array, N); return *this; } template \u0026lt;typename T\u0026gt; typename std::enable_if_t\u0026lt;std::is_integral_v\u0026lt;T\u0026gt; || std::is_floating_point_v\u0026lt;T\u0026gt;, PacketSerializer\u0026amp;\u0026gt; // pack(const std::vector\u0026lt;T\u0026gt;\u0026amp; vec) { if (!vec.empty()) { writeArrayToBuffer(vec.data(), vec.size()); } return *this; }1.3. packMultiple template \u0026lt;typename T, typename... Args\u0026gt; PacketSerializer\u0026amp; packMultiple(const T\u0026amp; first, const Args\u0026amp;... args) { pack(first); if constexpr (sizeof...(args) \u0026gt; 0) { packMultiple(args...); } return *this; }1.4. 打包原始结构体数据提供另外一种类型打包接口，用于直接序列化 1 字节对齐的结构体数据。 template \u0026lt;typename T\u0026gt; PacketSerializer\u0026amp; packRawData(const T\u0026amp; structData) { static_assert(std::is_trivially_copyable_v\u0026lt;T\u0026gt;, \"Type must be trivially copyable for direct packing\"); const uint8_t* bytes = reinterpret_cast\u0026lt;const uint8_t*\u0026gt;(\u0026amp;structData); return appendBytes(bytes, sizeof(T)); }1.5. finalize组装成完成的数据包，并返回序列化后的 packet。 std::vector\u0026lt;uint8_t\u0026gt; finalize() { std::vector\u0026lt;uint8_t\u0026gt; packet; uint16_t total_length = 2 + 2 + static_cast\u0026lt;uint16_t\u0026gt;(buffer_.size()) + 1; packet.push_back(static_cast\u0026lt;uint8_t\u0026gt;(frame_header_ \u0026amp; 0xFF)); packet.push_back(static_cast\u0026lt;uint8_t\u0026gt;((frame_header_ \u0026gt;\u0026gt; 8) \u0026amp; 0xFF)); packet.push_back(static_cast\u0026lt;uint8_t\u0026gt;(total_length \u0026amp; 0xFF)); packet.push_back(static_cast\u0026lt;uint8_t\u0026gt;((total_length \u0026gt;\u0026gt; 8) \u0026amp; 0xFF)); packet.insert(packet.end(), buffer_.begin(), buffer_.end()); const uint8_t checksum = calculateChecksum(packet.data(), packet.size()); packet.push_back(checksum); return packet; }1.6. 便利函数template \u0026lt;typename... Args\u0026gt;std::vector\u0026lt;uint8_t\u0026gt; createPacket(uint16_t frame_header, const Args\u0026amp;... args) { PacketSerializer serializer(frame_header); if constexpr (sizeof...(args) \u0026gt; 0) { serializer.packMultiple(args...); } return serializer.finalize();}template \u0026lt;typename T\u0026gt;std::vector\u0026lt;uint8_t\u0026gt; createRawDataPacket(uint16_t frame_header, const T\u0026amp; structData) { static_assert(std::is_trivially_copyable_v\u0026lt;T\u0026gt;, \"Type must be trivially copyable for direct packing\"); PacketSerializer serializer(frame_header); serializer.packRawData(structData); return serializer.finalize();}2. 反序列化 PacketDeserializer首先解出header。payload部分的解析，与序列化提供相同步骤的接口函数 unpack，unpackRawData。并提供便利接口：template \u0026lt;typename T, typename... Args\u0026gt;bool unpackMultiple(PacketDeserializer\u0026amp; deserializer, T\u0026amp; first, Args\u0026amp;... args) { if (deserializer.unpack(first) == 0) { return false; } if constexpr (sizeof...(args) \u0026gt; 0) { return unpackMultiple(deserializer, args...); } return true;} 实际应用中，由于存在一些运行时判断，实际在使用 PacketDeserializer 的时候，需要依据具体的协议格式，进行相应的逻辑处理，即针对这些特殊协议实现一个特化版本的解析封装。比如，某个字段的值决定后续字段的类型和数量等。3. 提取结构体中的字段进行序列化由于通信需要，通信中，通信双方需要设置以及接收参数结构体中的字段。梳理需求，提取出接口： 根据字段名称字符串（field name），提取结构体中的字段值：get_field_value_by_name，get_field_value_by_name_as。 根据结构体中的字段（field），获取字段的名称字符串（field name）：get_field_name_by_field。 由于C++ 17 支持不完善，不能提取字段名称字符串，此处使用手动定义字段名-字段偏移地址-字段类型映射表。 定义需要的数据结构。定义支持的字段数据类型枚举 field_type_e，枚举对应的数据类型大小映射表 field_type_sizes，以及字段描述结构体 FieldInfo。// clang-format offusing field_var_t =std::variant\u0026lt;int8_t, uint8_t, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t, float, double, std::string\u0026gt;;enum field_type_e : uint8_t { field_type_i8, field_type_ui8, field_type_i16, field_type_ui16, field_type_i32, field_type_ui32, field_type_i64, field_type_ui64, field_type_f32, field_type_f64, field_type_cstrarr,field_type_cstrptr, field_type_count};constexpr std::array\u0026lt;size_t, field_type_count\u0026gt; field_type_sizes = { sizeof(int8_t), sizeof(uint8_t), sizeof(int16_t), sizeof(uint16_t), sizeof(int32_t), sizeof(uint32_t), sizeof(int64_t), sizeof(uint64_t), sizeof(float), sizeof(double), 0, 0};// clang-format onstruct FieldInfo { std::string_view name; std::size_t offset; field_type_e type;};3.1. 根据输入的结构体实例，以及域字段（field），获取字段名称template \u0026lt;typename StructType, typename FieldType, size_t FieldInfoN\u0026gt;std::string get_field_name_by_field(const StructType\u0026amp; sinst, const FieldType\u0026amp; field, const std::array\u0026lt;FieldInfo, FieldInfoN\u0026gt;\u0026amp; fields) { const void* base = static_cast\u0026lt;const void*\u0026gt;(std::addressof(sinst)); const void* field_addr = static_cast\u0026lt;const void*\u0026gt;(std::addressof(field)); const ptrdiff_t offset = static_cast\u0026lt;const char*\u0026gt;(field_addr) - static_cast\u0026lt;const char*\u0026gt;(base); for (const auto\u0026amp; f : fields) { if (f.offset == offset) { return std::string(f.name); } } return {};} 实现：根据结构体实例的地址，以及字段的地址，计算字段的偏移地址 offset。然后在手动定义的字段描述信息数组 fields 中查找该偏移地址对应的字段名称，并返回。3.2. 根据字段名称（field name）获取结构体中该字段的值，返回一个 field_var_ttemplate \u0026lt;typename StructType, size_t FieldInfoN\u0026gt;std::tuple\u0026lt;field_var_t, field_type_e\u0026gt; get_field_value_by_name(const StructType\u0026amp; sinst, const std::string\u0026amp; name, const std::array\u0026lt;FieldInfo, FieldInfoN\u0026gt;\u0026amp; fields) { std::string_view sv{name}; // clang-format off using extractor_t = field_var_t (*)(const char* base, std::size_t off); auto extract_i8 = [](const char* base, std::size_t off) -\u0026gt; field_var_t { return *reinterpret_cast\u0026lt;const int8_t*\u0026gt;(base + off); }; auto extract_ui8 = [](const char* base, std::size_t off) -\u0026gt; field_var_t { return *reinterpret_cast\u0026lt;const uint8_t*\u0026gt;(base + off); }; auto extract_i16 = [](const char* base, std::size_t off) -\u0026gt; field_var_t { return *reinterpret_cast\u0026lt;const int16_t*\u0026gt;(base + off); }; auto extract_ui16 = [](const char* base, std::size_t off) -\u0026gt; field_var_t { return *reinterpret_cast\u0026lt;const uint16_t*\u0026gt;(base + off); }; auto extract_i32 = [](const char* base, std::size_t off) -\u0026gt; field_var_t { return *reinterpret_cast\u0026lt;const int32_t*\u0026gt;(base + off); }; auto extract_i64 = [](const char* base, std::size_t off) -\u0026gt; field_var_t { return *reinterpret_cast\u0026lt;const int64_t*\u0026gt;(base + off); }; auto extract_ui64 = [](const char* base, std::size_t off) -\u0026gt; field_var_t { return *reinterpret_cast\u0026lt;const uint64_t*\u0026gt;(base + off); }; auto extract_ui32 = [](const char* base, std::size_t off) -\u0026gt; field_var_t { return *reinterpret_cast\u0026lt;const uint32_t*\u0026gt;(base + off); }; auto extract_f32 = [](const char* base, std::size_t off) -\u0026gt; field_var_t { return *reinterpret_cast\u0026lt;const float*\u0026gt;(base + off); }; auto extract_f64 = [](const char* base, std::size_t off) -\u0026gt; field_var_t { return *reinterpret_cast\u0026lt;const double*\u0026gt;(base + off); }; auto extract_cstrarr = [](const char* base, std::size_t off) -\u0026gt; field_var_t { const char* p = reinterpret_cast\u0026lt;const char*\u0026gt;(base + off); return std::string(p); }; auto extract_cstrptr = [](const char* base, std::size_t off) -\u0026gt; field_var_t { const char* p = *reinterpret_cast\u0026lt;const char* const*\u0026gt;(base + off); return p ? std::string(p) : std::string(); }; static constexpr std::array\u0026lt;extractor_t, static_cast\u0026lt;size_t\u0026gt;(field_type_count)\u0026gt; extractor_table = { extractor_t(+extract_i8), extractor_t(+extract_ui8), extractor_t(+extract_i16), extractor_t(+extract_ui16), extractor_t(+extract_i32), extractor_t(+extract_i64), extractor_t(+extract_ui64), extractor_t(+extract_ui32), extractor_t(+extract_f32), extractor_t(+extract_f64), extractor_t(+extract_cstrarr), extractor_t(+extract_cstrptr)}; auto pred = [\u0026amp;sv](const FieldInfo\u0026amp; field) { return field.name == sv; }; const auto it = std::find_if(fields.begin(), fields.end(), pred); if (it != fields.end()) { const char* base = reinterpret_cast\u0026lt;const char*\u0026gt;(\u0026amp;sinst); const std::size_t off = it-\u0026gt;offset; const auto idx = static_cast\u0026lt;size_t\u0026gt;(it-\u0026gt;type); if (idx \u0026lt; extractor_table.size()) { const extractor_t ext = extractor_table[idx]; return {ext(base, off), it-\u0026gt;type}; } } // clang-format on return {field_var_t{}, field_type_e{}};} 实现关键点：根据定义的可解析数据类型，定义对应的提取函数 extractor_t，并建立提取函数表 extractor_table。根据字段名称查找字段信息，然后根据字段类型，调用对应的提取函数，获取字段值。最后，封装为 field_var_t 返回。字段类型信息，来源于手动定义的字段描述信息数组 fields。根据字段名称查找得到该字段的 field_type_e。3.3. 另外一个版本：根据字段名称（field name）获取结构体中该字段的值，直接返回类型Ttemplate \u0026lt;typename T\u0026gt;std::optional\u0026lt;T\u0026gt; make_copy(const void* src, std::size_t sz) { static_assert(std::is_trivially_copyable_v\u0026lt;T\u0026gt;, \"T must be trivially copyable\"); if (sizeof(T) != sz) return std::nullopt; T v; std::memcpy(\u0026amp;v, src, sz); return v;}template \u0026lt;typename StructType, typename FieldType, size_t FieldInfoN\u0026gt;std::optional\u0026lt;FieldType\u0026gt; get_field_value_by_name_as(const StructType\u0026amp; sinst, const std::string\u0026amp; name, const std::array\u0026lt;FieldInfo, FieldInfoN\u0026gt;\u0026amp; fields) { std::string_view sv{name}; auto it = std::find_if(fields.begin(), fields.end(), [\u0026amp;](const FieldInfo\u0026amp; f) { return f.name == sv; }); if (it == fields.end()) return std::nullopt; const char* base = reinterpret_cast\u0026lt;const char*\u0026gt;(\u0026amp;sinst); const std::size_t off = it-\u0026gt;offset; const field_type_e ft = it-\u0026gt;type; if constexpr (std::is_trivially_copyable_v\u0026lt;FieldType\u0026gt;) { using extractor_t = std::optional\u0026lt;FieldType\u0026gt; (*)(const char* base, std::size_t off); // clang-format off // lambdas must be non-capturing to convert to function pointer auto ex_i8 = [](const char* base, std::size_t off) -\u0026gt; std::optional\u0026lt;FieldType\u0026gt; { return make_copy\u0026lt;FieldType\u0026gt;(base + off, sizeof(int8_t)); }; auto ex_ui8 = [](const char* base, std::size_t off) -\u0026gt; std::optional\u0026lt;FieldType\u0026gt; { return make_copy\u0026lt;FieldType\u0026gt;(base + off, sizeof(uint8_t)); }; auto ex_i16 = [](const char* base, std::size_t off) -\u0026gt; std::optional\u0026lt;FieldType\u0026gt; { return make_copy\u0026lt;FieldType\u0026gt;(base + off, sizeof(int16_t)); }; auto ex_ui16 = [](const char* base, std::size_t off) -\u0026gt; std::optional\u0026lt;FieldType\u0026gt; { return make_copy\u0026lt;FieldType\u0026gt;(base + off, sizeof(uint16_t)); }; auto ex_i32 = [](const char* base, std::size_t off) -\u0026gt; std::optional\u0026lt;FieldType\u0026gt; { return make_copy\u0026lt;FieldType\u0026gt;(base + off, sizeof(int32_t)); }; auto ex_ui32 = [](const char* base, std::size_t off) -\u0026gt; std::optional\u0026lt;FieldType\u0026gt; { return make_copy\u0026lt;FieldType\u0026gt;(base + off, sizeof(uint32_t)); }; auto ex_i64 = [](const char* base, std::size_t off) -\u0026gt; std::optional\u0026lt;FieldType\u0026gt; { return make_copy\u0026lt;FieldType\u0026gt;(base + off, sizeof(int64_t)); }; auto ex_ui64 = [](const char* base, std::size_t off) -\u0026gt; std::optional\u0026lt;FieldType\u0026gt; { return make_copy\u0026lt;FieldType\u0026gt;(base + off, sizeof(uint64_t)); }; auto ex_f32 = [](const char* base, std::size_t off) -\u0026gt; std::optional\u0026lt;FieldType\u0026gt; { return make_copy\u0026lt;FieldType\u0026gt;(base + off, sizeof(float)); }; auto ex_f64 = [](const char* base, std::size_t off) -\u0026gt; std::optional\u0026lt;FieldType\u0026gt; { return make_copy\u0026lt;FieldType\u0026gt;(base + off, sizeof(double)); }; auto ex_invalid = [](const char* /*base*/, std::size_t /*off*/) -\u0026gt; std::optional\u0026lt;FieldType\u0026gt; { return std::nullopt; }; static const std::array\u0026lt;extractor_t, static_cast\u0026lt;size_t\u0026gt;(field_type_count)\u0026gt; extractors = { extractor_t(+ex_i8), extractor_t(+ex_ui8), extractor_t(+ex_i16), extractor_t(+ex_ui16), extractor_t(+ex_i32), extractor_t(+ex_i64), extractor_t(+ex_ui64), extractor_t(+ex_ui32), extractor_t(+ex_f32), extractor_t(+ex_f64), extractor_t(+ex_invalid), extractor_t(+ex_invalid)}; // clang-format on const auto idx = static_cast\u0026lt;size_t\u0026gt;(ft); if (idx \u0026lt; extractors.size()) { return extractors[idx](base, off); } return std::nullopt; } // 非 trivially copyable if constexpr (std::is_same_v\u0026lt;FieldType, std::string\u0026gt;) { if (ft == field_type_cstrarr) { const char* p = reinterpret_cast\u0026lt;const char*\u0026gt;(base + off); return std::string(p); } if (ft == field_type_cstrptr) { const char* const* pp = reinterpret_cast\u0026lt;const char* const*\u0026gt;(base + off); return pp \u0026amp;\u0026amp; *pp ? std::string(*pp) : std::string{}; } return std::nullopt; } return std::nullopt;} 实现关键点：同样根据定义的可解析数据类型，定义对应的提取函数 extractor_t，并建立提取函数表 extractors。根据字段名称查找字段信息，然后根据字段类型，调用对应的提取函数，获取字段值。最后，封装为 std::optional 返回。字段类型信息的获取，与函数 get_field_value_by_name 相同。3.4. 根据字段名称（field name），设置结构体中该字段的值template \u0026lt;typename StructType, typename FieldType, size_t FieldInfoN\u0026gt;inline ptrdiff_t set_field_by_name(StructType\u0026amp; ui, const std::string\u0026amp; name, const FieldType\u0026amp; value, const std::array\u0026lt;FieldInfo, FieldInfoN\u0026gt;\u0026amp; fields, size_t cstrarr_size = 0) { std::string_view sv{name}; auto it = std::find_if(fields.begin(), fields.end(), [\u0026amp;](const FieldInfo\u0026amp; f) { return f.name == sv; }); if (it == fields.end()) return -1; char* base = reinterpret_cast\u0026lt;char*\u0026gt;(\u0026amp;ui); const auto offset = static_cast\u0026lt;ptrdiff_t\u0026gt;(it-\u0026gt;offset); if constexpr (std::is_trivially_copyable_v\u0026lt;FieldType\u0026gt;) { if (field_type_sizes[it-\u0026gt;type] == sizeof(FieldType)) { std::memcpy(base + offset, \u0026amp;value, sizeof(FieldType)); return offset; } return -2; // 类型不匹配 } if constexpr (std::is_convertible_v\u0026lt;FieldType, const char*\u0026gt;) { // char*, char[], or string literal if (it-\u0026gt;type == field_type_cstrarr) { if (cstrarr_size == 0) return -3; // 需指定数组长度 char* arr = base + offset; std::strncpy(arr, static_cast\u0026lt;const char*\u0026gt;(value), cstrarr_size - 1); arr[cstrarr_size - 1] = '\\0'; return offset; } if (it-\u0026gt;type == field_type_cstrptr) { char** pptr = reinterpret_cast\u0026lt;char**\u0026gt;(base + offset); if (*pptr) { std::strcpy(*pptr, static_cast\u0026lt;const char*\u0026gt;(value)); return offset; } return -4; // 指针为空 } return -2; } // 支持 std::string if constexpr (std::is_same_v\u0026lt;FieldType, std::string\u0026gt;) { if (it-\u0026gt;type == field_type_cstrarr) { if (cstrarr_size == 0) return -3; char* arr = base + offset; std::strncpy(arr, value.c_str(), cstrarr_size - 1); arr[cstrarr_size - 1] = '\\0'; return offset; } if (it-\u0026gt;type == field_type_cstrptr) { char** pptr = reinterpret_cast\u0026lt;char**\u0026gt;(base + offset); if (*pptr) { std::strcpy(*pptr, value.c_str()); return offset; } return -4; } return -2; } return -5; // 不支持的类型}3.5. 调用实例定义映射表如下：static_assert(std::is_standard_layout_v\u0026lt;ps_formation_param_t\u0026gt;, \"userinfo not standard layout\");constexpr size_t kFormationParamNamesCount = boost::pfr::tuple_size\u0026lt;ps_formation_param_t\u0026gt;::value + ps_formation_param_t::kVTableSize - 1;using namespace struct_reflect;// clang-format offconstexpr std::array\u0026lt;struct_reflect::FieldInfo, kFormationParamNamesCount\u0026gt; kFormationParamNames = { FieldInfo{std::string_view(\"Formation_Desired_Position_Offset_N\"),offsetof(ps_formation_param_t, Formation_Desired_Position_Offset_N),field_type_e::field_type_f64}, FieldInfo{std::string_view(\"Formation_Desired_Position_Offset_E\"),offsetof(ps_formation_param_t, Formation_Desired_Position_Offset_E),field_type_e::field_type_f64}, FieldInfo{std::string_view(\"Formation_Desired_Position_Offset_D\"),offsetof(ps_formation_param_t, Formation_Desired_Position_Offset_D),field_type_e::field_type_f64}, FieldInfo{std::string_view(\"swarmNum\"),offsetof(ps_formation_param_t, swarmNum),field_type_e::field_type_ui8}, FieldInfo{std::string_view(\"hController_k\"),offsetof(ps_formation_param_t, hController_k),field_type_e::field_type_f64}, FieldInfo{std::string_view(\"hController_f\"),offsetof(ps_formation_param_t, hController_f),field_type_e::field_type_f64}, FieldInfo{std::string_view(\"hController_D\"),offsetof(ps_formation_param_t, hController_D),field_type_e::field_type_f64}, FieldInfo{std::string_view(\"coefs_k\"),offsetof(ps_formation_param_t, coefs_k),field_type_e::field_type_f64}, FieldInfo{std::string_view(\"coefs_f\"),offsetof(ps_formation_param_t, coefs_f),field_type_e::field_type_f64}, FieldInfo{std::string_view(\"coefs_u2a\"),offsetof(ps_formation_param_t, coefs_u2a),field_type_e::field_type_f64}, FieldInfo{std::string_view(\"coefs_uff\"),offsetof(ps_formation_param_t, coefs_uff),field_type_e::field_type_f64}, FieldInfo{std::string_view(\"vTable_0\"),offsetof(ps_formation_param_t, vTable) + sizeof(double) * 0,field_type_e::field_type_f64}, FieldInfo{std::string_view(\"vTable_1\"),offsetof(ps_formation_param_t, vTable) + sizeof(double) * 1,field_type_e::field_type_f64}, FieldInfo{std::string_view(\"vTable_2\"),offsetof(ps_formation_param_t, vTable) + sizeof(double) * 2,field_type_e::field_type_f64}, FieldInfo{std::string_view(\"vTable_3\"),offsetof(ps_formation_param_t, vTable) + sizeof(double) * 3,field_type_e::field_type_f64}, FieldInfo{std::string_view(\"vTable_4\"),offsetof(ps_formation_param_t, vTable) + sizeof(double) * 4,field_type_e::field_type_f64}, FieldInfo{std::string_view(\"vTable_5\"),offsetof(ps_formation_param_t, vTable) + sizeof(double) * 5,field_type_e::field_type_f64}, FieldInfo{std::string_view(\"vTable_6\"),offsetof(ps_formation_param_t, vTable) + sizeof(double) * 6,field_type_e::field_type_f64}, FieldInfo{std::string_view(\"vTable_7\"),offsetof(ps_formation_param_t, vTable) + sizeof(double) * 7,field_type_e::field_type_f64}, FieldInfo{std::string_view(\"hSafe\"),offsetof(ps_formation_param_t, hSafe),field_type_e::field_type_f64}, FieldInfo{std::string_view(\"asb_offset_n\"),offsetof(ps_formation_param_t, asb_offset_n),field_type_e::field_type_f64}, FieldInfo{std::string_view(\"asb_offset_e\"),offsetof(ps_formation_param_t, asb_offset_e),field_type_e::field_type_f64}, FieldInfo{std::string_view(\"L1\"),offsetof(ps_formation_param_t, L1),field_type_e::field_type_f64}, FieldInfo{std::string_view(\"kVel\"),offsetof(ps_formation_param_t, kVel),field_type_e::field_type_f64}, FieldInfo{std::string_view(\"robustLat\"),offsetof(ps_formation_param_t, robustLat),field_type_e::field_type_f64}, FieldInfo{std::string_view(\"asbController_kPos\"),offsetof(ps_formation_param_t, asbController_kPos),field_type_e::field_type_f64}, FieldInfo{std::string_view(\"asbController_kAsb\"),offsetof(ps_formation_param_t, asbController_kAsb),field_type_e::field_type_f64}, FieldInfo{std::string_view(\"asbController_tol1\"),offsetof(ps_formation_param_t, asbController_tol1),field_type_e::field_type_f64}, FieldInfo{std::string_view(\"asbController_tol2\"),offsetof(ps_formation_param_t, asbController_tol2),field_type_e::field_type_f64}, FieldInfo{std::string_view(\"asbController_r1\"),offsetof(ps_formation_param_t, asbController_r1),field_type_e::field_type_f64}, FieldInfo{std::string_view(\"asbController_r2\"),offsetof(ps_formation_param_t, asbController_r2),field_type_e::field_type_f64}, FieldInfo{std::string_view(\"fmtController_kLat\"),offsetof(ps_formation_param_t, fmtController_kLat),field_type_e::field_type_f64}, FieldInfo{std::string_view(\"fmtController_kProj\"),offsetof(ps_formation_param_t, fmtController_kProj),field_type_e::field_type_f64}, FieldInfo{std::string_view(\"fmtWpStartIdx\"),offsetof(ps_formation_param_t, fmtWpStartIdx),field_type_e::field_type_ui16}, FieldInfo{std::string_view(\"takeIdx\"),offsetof(ps_formation_param_t, takeIdx),field_type_e::field_type_ui16}, FieldInfo{std::string_view(\"returnIdx\"),offsetof(ps_formation_param_t, returnIdx),field_type_e::field_type_ui16}, FieldInfo{std::string_view(\"Vn\"),offsetof(ps_formation_param_t, Vn),field_type_e::field_type_f64}, FieldInfo{std::string_view(\"rm\"),offsetof(ps_formation_param_t, rm),field_type_e::field_type_f64},};// clang-format on 在此，对结构体作限制：要求结构体以及其字段均为标准布局类型（standard layout type），以确保字段偏移地址的正确性。其他代码省略。" },

  { "title": "static_cast 与 reinterpret_cast", "url": "/blog/2025/static_cast_vs_reinterpret_cast/", "categories": "Cpp",
  "tags": "Cpp", "date": "2025-08-25", "content":
  "1. static_cast 编译时类型检查 只允许安全的、有意义的类型转换 会进行必要的类型转换计算static_cast不能用于在不同类型的指针之间互相转换，也不能用于整型和指针之间的互相转换。示例：// 数值类型转换int i = 42;double d = static_cast\u0026lt;double\u0026gt;(i);// 基类与派生类之间的转换（向上转换总是安全的）Derived* derived = new Derived();Base* base = static_cast\u0026lt;Base*\u0026gt;(derived);// void* 转换为具体类型指针void* ptr = malloc(sizeof(int));int* intPtr = static_cast\u0026lt;int*\u0026gt;(ptr);2. reinterpret_cast 几乎不进行类型检查 直接重新解释内存中的位模式 非常危险，需要程序员确保安全性示例：// 指针与整数之间的转换int* ptr = new int(42);uintptr_t addr = reinterpret_cast\u0026lt;uintptr_t\u0026gt;(ptr);// 不相关类型指针之间的转换char* charPtr = reinterpret_cast\u0026lt;char*\u0026gt;(ptr);// 函数指针转换void (*funcPtr)() = reinterpret_cast\u0026lt;void(*)()\u0026gt;(some_address);3. 对比 特性 static_cast reinterpret_cast 安全性 高，编译时检查 低，几乎无检查 性能 可能有运行时开销 无运行时开销 用途 合理的类型转换 位模式重解释 可移植性 好 依赖平台 示例场景 数值转换、继承关系 指针转换、底层操作 4. ref reinterpret_cast 和 static_cast" },

  { "title": "C++ template 学习资料", "url": "/blog/2025/cpp_template_study_res/", "categories": "Cpp",
  "tags": "Cpp, Template", "date": "2025-08-13", "content":
  "现代C++模板教程 github – Modern-Cpp-templates-tutorial 现代C++模板教程 C++ Templates 2ed github – Cpp-Templates-2ed Cpp-Templates-2ed github – CppTemplateTutorial boost – Simple C++11 metaprogramming" },

  { "title": "使用 CMake CPack NSIS 打包 Qt 应用", "url": "/blog/2025/qt_cpack_nsis/", "categories": "CMake",
  "tags": "CMake, Qt", "date": "2025-07-18", "content":
  "if(WIN32 AND NOT UNIX) get_target_property(qmake_executable Qt6::qmake IMPORTED_LOCATION) get_filename_component(_qt_bin_dir \"${qmake_executable}\" DIRECTORY) find_program(WINDEPLOYQT_EXECUTABLE windeployqt HINTS \"${_qt_bin_dir}\") message(STATUS \"Using windeployqt: ${WINDEPLOYQT_EXECUTABLE}\") message(STATUS \"qt bin dir: ${_qt_bin_dir}\") add_custom_command(TARGET ${target_name} POST_BUILD COMMAND \"${WINDEPLOYQT_EXECUTABLE}\" $\u0026lt;$\u0026lt;CONFIG:Debug\u0026gt;:--debug\u0026gt; $\u0026lt;$\u0026lt;NOT:$\u0026lt;CONFIG:Debug\u0026gt;\u0026gt;:--release\u0026gt; --libdir $\u0026lt;TARGET_FILE_DIR:${target_name}\u0026gt; --verbose 0 --compiler-runtime --no-opengl-sw --force --plugindir \"$\u0026lt;TARGET_FILE_DIR:${target_name}\u0026gt;/plugins\" $\u0026lt;TARGET_FILE:${target_name}\u0026gt; COMMENT \"Deploying Qt...\" ) if(QT_VERSION_MAJOR EQUAL 6) qt_finalize_executable(${target_name}) endif() include(GNUInstallDirs) install(PROGRAMS $\u0026lt;TARGET_FILE:${target_name}\u0026gt; DESTINATION . # TYPE BIN ) install(DIRECTORY $\u0026lt;TARGET_FILE_DIR:${target_name}\u0026gt;/ DESTINATION . FILES_MATCHING PATTERN \"*.dll\" ) # CPack # InstallRequiredSystemLibraries 以及选项 # https://cmake.org/cmake/help/latest/module/InstallRequiredSystemLibraries.html set(CMAKE_INSTALL_SYSTEM_RUNTIME_DESTINATION TRUE) set(CMAKE_INSTALL_UCRT_LIBRARIES TRUE) set(CMAKE_INSTALL_SYSTEM_RUNTIME_DESTINATION \".\") include(InstallRequiredSystemLibraries) set(CPACK_GENERATOR \"IFW\") set(CPACK_IFW_ROOT \"D:/dev_libs/Qt/Tools/QtInstallerFramework/4.10\") set(CPACK_PACKAGE_VERSION_MAJOR \"1\") set(CPACK_PACKAGE_VERSION_MINOR \"0\") # set human names to exetuables set(CPACK_PACKAGE_EXECUTABLES \"${PROJECT_NAME}\" \"MAVLink Monitor installer\") set(CPACK_CREATE_DESKTOP_LINKS \"${PROJECT_NAME}\")endif()include(CPACK)include(CPackIFW)编译完成之后，生成CPackConfig.cmake，在build目录下执行：cpack.exe ./CPackConfig.cmake参考 Revisiting the Qt Installer Framework with CMake cmake-qt-packaging-example 基于 CMake 的打包 AwesomeWorld-Cmake How to use CMake/CPack to create beautiful GUI-installers gist gist – Configuring Qt 6.x and CMake for windows 10 cpack_setup_nsis.cmake [CMake CPack使用NSIS打包安装程序](https://blog.bug-maker.com/archives/511.html)" },

  { "title": "CMake export 命令以及 install 命令", "url": "/blog/2025/cmake_export_and_install/", "categories": "CMake",
  "tags": "CMake", "date": "2025-06-11", "content":
  "1. export 导出一个库export命令用于导出一个库，导出的库可以被其他项目使用。如下cmake命令生成库gcFactSystem，并导出库以及库的头文件（使用PUBLIC）。set(target_name \"gcFactSystem\")# compile library gcFactSystem# export library gcFactSystem to gcFactSystemConfig.cmaketarget_include_directories(${target_name} PUBLIC ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/validator)export(TARGETS ${target_name} FILE \"${CMAKE_CURRENT_BINARY_DIR}/${target_name}Targets.cmake\")生成gcFactSystemTargets.cmake，里面包含了目标库，以及头文件包含路径：set_target_properties(gcFactSystem PROPERTIES INTERFACE_INCLUDE_DIRECTORIES \"D:/work/ground_station/dev/jhatcgcs/src/FactSystem;D:/work/ground_station/dev/jhatcgcs/src/FactSystem/validator\" INTERFACE_SOURCES \"\\$\u0026lt;\\$\u0026lt;BOOL:\\$\u0026lt;TARGET_PROPERTY:QT_CONSUMES_METATYPES\u0026gt;\u0026gt;:D:/work/ground_station/dev/jhatcgcs/build/src/FactSystem/meta_types/qt6gcfactsystem_debug_metatypes.json\u0026gt;\")set_target_properties(gcFactSystem PROPERTIES IMPORTED_IMPLIB_DEBUG \"D:/work/ground_station/dev/jhatcgcs/bin/gcFactSystem.lib\" IMPORTED_LINK_DEPENDENT_LIBRARIES_DEBUG \"Qt6::Core;Qt6::Qml;gcLogging;gcMAVLink\" IMPORTED_LOCATION_DEBUG \"D:/work/ground_station/dev/jhatcgcs/bin/gcFactSystem.dll\")工程中的其他模块使用被export出来的库：target_link_libraries(${target_name} PRIVATE gcFactSystem)2. install 安装库# 引入要用到的 CMake 模块include(CMakePackageConfigHelpers)include(GNUInstallDirs)# 基本安装及 Targets 文件的生成install(TARGETS ${PROJECT_NAME} EXPORT ${PROJECT_NAME}-targets RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR} LIBRARY DESTINATION ${CMAKE_INSTALL_LIBDIR} ARCHIVE DESTINATION ${CMAKE_INSTALL_LIBDIR} INCLUDES DESTINATION ${CMAKE_INSTALL_INCLUDEDIR})# Targets 文件的安装install(EXPORT ${PROJECT_NAME}-targets FILE ${PROJECT_NAME}-targets.cmake NAMESPACE ${PROJECT_NAME}:: DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/${PROJECT_NAME})# Config 文件的导出（与安装？）configure_package_config_file( ${PROJECT_SOURCE_DIR}/cmake/config.cmake.in ${PROJECT_BINARY_DIR}/${PROJECT_NAME}-config.cmake INSTALL_DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/${PROJECT_NAME})# ConfigVersion 文件的导出write_basic_package_version_file( ${PROJECT_NAME}-config-version.cmake VERSION ${PACKAGE_VERSION} COMPATIBILITY AnyNewerVersion)# Config 和 ConfigVersion 文件的安装install(FILES ${PROJECT_BINARY_DIR}/${PROJECT_NAME}-config.cmake ${PROJECT_BINARY_DIR}/${PROJECT_NAME}-config-version.cmake DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/${PROJECT_NAME})# 头文件的安装install(DIRECTORY ${PROJECT_SOURCE_DIR}/include/${PROJECT_NAME} DESTINATION ${CMAKE_INSTALL_INCLUDEDIR})2.1. config.cmake 的生成及安装生成config.cmake文件：include(CMakePackageConfigHelpers)configure_package_config_file( ${PROJECT_SOURCE_DIR}/cmake/config.cmake.in ${PROJECT_BINARY_DIR}/${PROJECT_NAME}-config.cmake INSTALL_DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/${PROJECT_NAME})安装config.cmake文件：install(FILES ${PROJECT_BINARY_DIR}/${PROJECT_NAME}-config.cmake ${PROJECT_BINARY_DIR}/${PROJECT_NAME}-config-version.cmake DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/${PROJECT_NAME})2.2. 头文件的安装设置需要导出的头文件路径：target_include_directories(${PROJECT_NAME} PUBLIC ${${PROJECT_NAME}_INCLUDES})设置导出文件更好的方式，区分build和install，分别引用不同路径下的include目录：target_include_directories(${PROJECT_NAME} PUBLIC $\u0026lt;BUILD_INTERFACE:${${PROJECT_NAME}_INCLUDES}\u0026gt; # 编译项目时引用项目下的头文件目录 $\u0026lt;INSTALL_INTERFACE:include\u0026gt; # 安装项目时引用 `\u0026lt;PREFIX\u0026gt;/include` ，比如 `/usr/include`)安装头文件：install(DIRECTORY ${PROJECT_SOURCE_DIR}/include/${PROJECT_NAME} DESTINATION ${CMAKE_INSTALL_INCLUDEDIR})3. 参考 CMake ——库的安装与导出" },

  { "title": "MAVLink消息的打包和解包", "url": "/blog/2025/MAVLink_msg_pack_unpack/", "categories": "QGC",
  "tags": "QGC, MAVLink", "date": "2025-06-06", "content":
  "以心跳包为例，消息格式定义在common.xml中。1. 打包心跳包的打包函数为mavlink_msg_heartbeat_encode，将heartbeat作为msg的payload，并计算MAVLink消息的其余域，即完成填充所有内容到msg中。static inline uint16_t mavlink_msg_heartbeat_encode(uint8_t system_id, uint8_t component_id, mavlink_message_t* msg, const mavlink_heartbeat_t* heartbeat) { return mavlink_msg_heartbeat_pack(system_id, component_id, msg, heartbeat-\u0026gt;type, heartbeat-\u0026gt;autopilot, heartbeat-\u0026gt;base_mode, heartbeat-\u0026gt;custom_mode, heartbeat-\u0026gt;system_status);}其他msgid消息类似，调用mavlink_msg_xxx_encode，并调用mavlink_msg_xxx_pack函数进行打包。其中mavlink_msg_heartbeat_pack函数首先将mavlink_heartbeat_t中的内容填充到msg的payload中，然后调用mavlink_finalize_message计算MAVLink消息的其余域，即完成填充所有内容到msg中。static inline uint16_t mavlink_msg_heartbeat_pack(uint8_t system_id, uint8_t component_id, mavlink_message_t* msg, uint8_t type, uint8_t autopilot, uint8_t base_mode, uint32_t custom_mode, uint8_t system_status){#if MAVLINK_NEED_BYTE_SWAP || !MAVLINK_ALIGNED_FIELDS char buf[MAVLINK_MSG_ID_HEARTBEAT_LEN]; _mav_put_uint32_t(buf, 0, custom_mode); _mav_put_uint8_t(buf, 4, type); _mav_put_uint8_t(buf, 5, autopilot); _mav_put_uint8_t(buf, 6, base_mode); _mav_put_uint8_t(buf, 7, system_status); _mav_put_uint8_t(buf, 8, 3); memcpy(_MAV_PAYLOAD_NON_CONST(msg), buf, MAVLINK_MSG_ID_HEARTBEAT_LEN);#else mavlink_heartbeat_t packet; packet.custom_mode = custom_mode; packet.type = type; packet.autopilot = autopilot; packet.base_mode = base_mode; packet.system_status = system_status; packet.mavlink_version = 3; memcpy(_MAV_PAYLOAD_NON_CONST(msg), \u0026amp;packet, MAVLINK_MSG_ID_HEARTBEAT_LEN);#endif msg-\u0026gt;msgid = MAVLINK_MSG_ID_HEARTBEAT; return mavlink_finalize_message(msg, system_id, component_id, MAVLINK_MSG_ID_HEARTBEAT_MIN_LEN, MAVLINK_MSG_ID_HEARTBEAT_LEN, MAVLINK_MSG_ID_HEARTBEAT_CRC);}其中，包序号seq存储在内部数组m_mavlink_status[]对应channel中，每次pack操作之后，该channel对应的seq自增1。获取通道函数为：mavlink_status_t* mavlink_get_channel_status(uint8_t chan);最终，mavlink_finalize_message完成填充msg结构体除payload之外的所有域，包括计算得到的crc。在mavlink_finalize_message的调用子函数中，还有两个操作涉及到v1和v2协议的处理： 在v2版本中，payload末尾进行0填充，通过调用_mav_trim_payload函数实现。 在v2版本中，如果需要进行signing操作，通过调用mavlink_sign_packet生成signature。1.1. 消息的发送在封包完成之后，调用mavlink_msg_to_send_buffer函数将包转换为字节流，然后调用方就可以通过TCP/UDP等传输层发送了。uint16_t mavlink_msg_to_send_buffer(uint8_t *buf, const mavlink_message_t *msg);2. 解包每接收到一个字节，就需要调用mavlink_frame_char解包，并更新对应channel的mavlink_status_t结构体中的状态机（parse_state成员域），直到解析完最后一个字节（16位CRC的最后一个字节）。uint8_t mavlink_frame_char(uint8_t chan, uint8_t c, mavlink_message_t* r_message, mavlink_status_t* r_mavlink_status){ return mavlink_frame_char_buffer(mavlink_get_channel_buffer(chan), mavlink_get_channel_status(chan), c, r_message, r_mavlink_status);}同样的，获取mavlink_status_t结构体的函数为mavlink_get_channel_status。在mavlink_frame_char_buffer函数中，就是判断mavlink_status_t结构体中的状态（parse_state成员域）进行解析，并更细parse_state。在解析完一个完成的mavlink_message_t消息体之后，返回MAVLINK_FRAMING_OK表示完成一个数据包的解析，并填充外面给的参数r_message，以及r_mavlink_status。2.1. 状态机流转图3. channel 与 system idMAVLink协议中，channel和system id是两个不同的概念。关于channel概念，见官方文档Multiple Streams (“channels”)。 channel是指MAVLink消息的传输通道，每新建一个TCP/UDP/串口连接，即创建一个channel。可以理解为建立一个UDP Address + port，或一个TCP Address + port，就会创建一个channel。 system id是标示不同的设备，比如地面站、飞控设备、以及其他设备。 一个channel可以对应多个system id，比如一个channel可以对应多个飞控设备。同时，一个system id也可以对应多个channel，比如一个地面站可以同时连接多个飞控设备。在QGC中，LinkInterface::_allocateMavlinkChannel用于给一个传输层连接分配一个channel，根据MAVLink中的宏定义，一般容量为4个或16个。4. 参考及资料 mavlink（五）C库源码分析 Using C MAVLink Libraries (mavgen)" },

  { "title": "MAVLink协议", "url": "/blog/2025/MAVLink_intro/", "categories": "QGC",
  "tags": "QGC, MAVLink", "date": "2025-05-14", "content":
  "1. MAVLink v2 协议MAVLink协议格式文档：Packet Serialization1.1. MAVLink协议格式重要字段MAVPACKED(typedef struct __mavlink_message { uint16_t checksum; ///\u0026lt; sent at end of packet uint8_t magic; ///\u0026lt; protocol magic marker uint8_t len; ///\u0026lt; Length of payload uint8_t incompat_flags; ///\u0026lt; flags that must be understood uint8_t compat_flags; ///\u0026lt; flags that can be ignored if not understood uint8_t seq; ///\u0026lt; Sequence of packet uint8_t sysid; ///\u0026lt; ID of message sender system/aircraft uint8_t compid; ///\u0026lt; ID of the message sender component uint32_t msgid:24; ///\u0026lt; ID of message in payload uint64_t payload64[(MAVLINK_MAX_PAYLOAD_LEN+MAVLINK_NUM_CHECKSUM_BYTES+7)/8]; uint8_t ck[2]; ///\u0026lt; incoming checksum bytes uint8_t signature[MAVLINK_SIGNATURE_BLOCK_LEN];}) mavlink_message_t; System ID：系统ID(sysid)，用于区分不同的无人机，标示发送者的sisid； Component ID：组件ID(compid)，用于区分无人机的不同模块，标示发送者的compid； Message ID：消息ID(msgid)，用于区分不同的消息类型，不同类型消息(GPS数据、遥测信息、飞行控制指令等)，需要再次解析payload得到有意义的消息体； sequence number：序列号(seq)，用于统计数据包丢失；1.1.1. LinkIdLinkId位于signature的第一个字节，在多链路通信系统中（如wifi，4G，卫星通信）用于区分不同的物理链路，可用于该链路信号质量检测及调整。1.1.2. 时间同步MAVLink要求所有节点设备的时间同步，以保证消息的顺序性。mavlink（一）帧格式和特性1.2. MAVLink协议解析参照ardupilot仓库代码相关解析，handle_xxx等成员方法，例如GCS_MAVLINK_Copter::handle_message(const mavlink_message_t \u0026amp;msg)。void GCS_MAVLINK_Copter::handle_message(const mavlink_message_t \u0026amp;msg){ switch (msg.msgid) {#if MODE_GUIDED_ENABLED case MAVLINK_MSG_ID_SET_ATTITUDE_TARGET: handle_message_set_attitude_target(msg); break; case MAVLINK_MSG_ID_SET_POSITION_TARGET_LOCAL_NED: handle_message_set_position_target_local_ned(msg); break; case MAVLINK_MSG_ID_SET_POSITION_TARGET_GLOBAL_INT: handle_message_set_position_target_global_int(msg); break;#endif#if AP_TERRAIN_AVAILABLE case MAVLINK_MSG_ID_TERRAIN_DATA: case MAVLINK_MSG_ID_TERRAIN_CHECK: copter.terrain.handle_data(chan, msg); break;#endif#if TOY_MODE_ENABLED case MAVLINK_MSG_ID_NAMED_VALUE_INT: copter.g2.toy_mode.handle_message(msg); break;#endif default: GCS_MAVLINK::handle_message(msg); break; }}2. MAVLink 消息2.1. MAVLink 消息类型MAVLink不同消息类型，分配有不同的msgid，与msgid相应的，带有不同的子数据结构，以进行二次解析获取该类型消息的内容。常见消息类型有： 心跳消息（Heartbeat）：用于定期确认系统是否在线，并报告系统状态(如飞行模式、硬件状态等)。 遥测数据（Telemetry）：包括飞行数据、传感器数据、GPS位置、飞行状态等。 控制指令：如起飞、降落、飞行模式切换等控制命令。 传感器数据：包括加速度、陀螺仪、磁力计、气压计、GPS等传感器采集的数据。 状态报告：如电池电量、飞行器姿态、航向等系统状态信息。 日志和错误报告：用于记录系统日志信息，如飞行轨迹、错误码、事件日志等。2.2. MAVLink消息的广播与路由有些类型的消息，包(packet)中包含target_system、target_component。如果target_system和target_component不为0，则该消息是定向发送给特定系统或组件的。target_system如果为0，则为广播消息，所有在线的系统都会收到这个消息。target_component如果为0，则为广播消息，同一个系统内的所有在线的组件都会收到这个消息。飞控器的消息的路由，由类MAVLink_routing处理。在内部，在接收到消息之后，根据发送方的sysid、compid，会存储一个channel（IP/PORT，串口号等）与(target_system, target_component)的映射。具体路由的规则见ArduPilot Development Site – MAVLink Routing in ArduPilot，以及MAVLink Developer Guide – Routing。2.3. MAVLink 消息的确认与重传MAVLink协议支持消息的确认与重传机制。发送方在发送消息后，会等待接收方的确认响应。如果在一定时间内没有收到确认响应，发送方会重新发送该消息，直到收到确认或达到最大重试次数。3. 资料 MAVLink Wireshark抓包插件" },

  { "title": "QGC 笔记以及资料", "url": "/blog/2025/QGC%E7%AC%94%E8%AE%B0%E5%8F%8A%E8%B5%84%E6%96%99/", "categories": "QGC",
  "tags": "QGC", "date": "2025-04-24", "content":
  "1. 资料1.1. QGC资料收集 官方user guide 官方dev guide PR: QML: Various easy optimizations Simple GCS – imGUI Gazebo Sensors：介绍了Gazebo中各种传感器的使用，包括IMU、GPS、相机等，这些传感器的数据可以通过MAVLink协议发送给QGC。 知乎 – 非线性MC：自动驾驶仿真、路径规划相关的文章。 QOpenHD：一个基于QGC的开源FPV图传系统，支持高清视频传输和遥控功能。 QGeoView：基于QGraphicsView的地图显示，无QML。1.2. QGC(master) 编译需要同时按照runtime版本和devel版本，安装完成之后，设置环境变量GSTREAMER_1_0_ROOT_MSVC_X86_64：GSTREAMER_1_0_ROOT_MSVC_X86_64 = d:\\dev_libs\\gstreamer\\1.0\\msvc_x86_64\\ Download GStreamer Gstreamer安装和使用 pkg-config for windows gstreamer 中文教程 常见问题：GStreamer的使用 PX4 – Simulation2. QGC 核心架构图示UML 核心类图（下载本地放大查看）：核心系统分析图（下载本地放大查看）：引用： CSDN – QGC(GGroundControl) 系统核心架构图3. QGC / MAVLink 通信流程flowchart TD A[QGC Application] --\u0026gt; B[VehicleLinkManager] B --\u0026gt; C1[UDP Link] B --\u0026gt; C2[Serial Link] B --\u0026gt; C3[TCP Link] C1 \u0026amp; C2 \u0026amp; C3 --\u0026gt; D[MAVLinkProtocol] D --\u0026gt; E[MAVLink Messages] E --\u0026gt; F[Vehicle Manager] F --\u0026gt; G1[Vehicle] G1 -- Telemetry/Status/Command --\u0026gt; E %% Additional flows for message routing and plugins E --\u0026gt; H[QGC Plugins and UI Components] H --\u0026gt; I[User Interaction / Display] %% Optional: GCS communication B --\u0026gt; J[Other GCS] J --\u0026gt; E %% Note style D fill:#f9f,stroke:#333,stroke-width:2px %% Legend subgraph Legend [Legend] direction LR L1[Boxes: QGC Components] L2[Arrows: Data/Message Flow] end通信流程说明 QGroundControl 启动后，由VehicleLinkManager管理与飞行器的所有链接（如UDP、串口、TCP等）。 各类物理/网络连接（UDP、串口、TCP）通过VehicleLinkManager建立。 所有数据流进入MAVLinkProtocol模块进行MAVLink消息的解析与封装。 MAVLink消息被分发到Vehicle Manager，并进一步与车辆（如无人机）交互。 车辆端会通过回传消息（如遥测、状态、命令响应等）同样经过上述流程返回QGC。 QGC插件和UI组件可订阅和处理这些消息，实现用户交互和数据显示。 多地面站（GCS）通信时，消息也可在多个QGC实例间流转。 以上流程图和说明基于 QGC dev guide – Communication Flow。 LinkManager 建立一个UDP监听端口(14550)，等待飞行器心跳包； 字节流经过MAVLinkProtocol转换为MAVLink消息包； 如果是心跳包，MAVLinkProtocol将消息包转发给MultiVehicleManager； MultiVehicleManager接收到心跳包，创建一个新的Vehicle对象，并将其添加到_vehicles列表中； Vehicle根据对应的_firmwareType、_vehicleType创建对应的FirmwarePlugin； FirmwarePlugin根据_vehicleType创建对应的AutoPilotPlugin； Vehicle实例创建完成，其持有的ParameterLoader实例发送PARAM_REQUEST_LIST给飞行器，请求飞行器参数列表（使用pramater protocol）； 飞行器参数列表加载完成之后，Vehicle对应的MissionManager实例持有的MissionManager请求飞行任务items（使用mission protol）； 心跳数据流转：UDPWorker::dataReceived -\u0026gt; UDPLink::_onDataReceived / LinkInterface::bytesReceived -\u0026gt; MAVLinkProtocol::receiveBytes / MAVLinkProtocol::vehicleHeartbeatInfo -\u0026gt; MultiVehicleManager::_vehicleHeartbeatInfo / new Vehicle。3.1. 资料 QGC通信流程 Communication Flow QML与C++交互" },

  { "title": "C++对象模型--多继承", "url": "/blog/2025/C++%E5%AF%B9%E8%B1%A1%E6%A8%A1%E5%9E%8B-%E5%A4%9A%E7%BB%A7%E6%89%BF/", "categories": "Cpp",
  "tags": "Cpp", "date": "2025-03-24", "content":
  "1. 多继承–无虚拟继承class Base1 {public: int a; int b;};class Base2 {public: int c; int d;};class Derive : public Base1 , public Base2 {public: int e; int f;};内存布局顺序为：Base1的成员变量 -\u0026gt; Base2的成员变量 -\u0026gt; Derive的成员变量。如下图所示（图中一格表示4字节）：1.1 指针调整当进行Derive / Base 指针赋值或比较时，编译器对Base / Derive 指针进行偏移调整。int main() { Derive* d_ptr = new Derive(); Base2* b2_ptr = d_ptr; Base1* b1_ptr = d_ptr; printf(\"address of d_ptr = %p\\n\", d_ptr); printf(\"address of b1_ptr = %p\\n\", b1_ptr); printf(\"address of b2_ptr = %p\\n\", b2_ptr); // 指针调整 if (d_ptr == b2_ptr) { // 明明两个指针在数值上不相同，但从C++的语义上看，两个指针指向同一个对象，所以编译器还是进行指针调整。 printf(\"d_ptr == b2_ptr\\n\"); } if (d_ptr == b1_ptr) { printf(\"d_ptr == b1_ptr\\n\"); }}address of d_ptr = 0x560982eebe70address of b1_ptr = 0x560982eebe70address of b2_ptr = 0x560982eebe78d_ptr == b2_ptrd_ptr == b1_ptr参考 CPP对象模型 深度探索C++对象模型_读书笔记.pdf" },

  { "title": "Windows编译安装VTK, TCL/TK, OCC", "url": "/blog/2025/windows-build-tcl-tk-occ/", "categories": "OCCT",
  "tags": "OCCT, VTK", "date": "2025-03-11", "content":
  "1. 预编译及安装的三方库 zlib freeType FreeImage1.1. freeTypefreeType在Windows下编译成动态库，需要做些修改：根目录下CMakeLists.txt：# add_library(freetype ....add_library(freetype SHARED...) freeType编译之后，cmake配置OCC时找不到freeType相关库，使用官方提供的编译好的三方库替代OCC Release。1.2. FreeImageFreeImage选择下载编译好的文件（没有Debug版本）。或选择第三方修改的仓库FreeImage-Cmake2. VTKTODO3. TCL/TK下载及编译TCL/TK 8.6.16源码： TCL 8.6.16 http://prdownloads.sourceforge.net/tcl/tcl8616-src.zip TK 8.6.16 http://prdownloads.sourceforge.net/tcl/tk8616-src.zip分别修改TCL及TK子目录win下的rules.vc文件：# SUFX = tsgx# 修改为以下内容SUFX = sgx编译及安装TCL：nmake -f makefile.vc INSTALLDIR=d:\\dev_libs\\occnmake -f makefile.vc install INSTALLDIR=d:\\dev_libs\\occ编译及安装TK：nmake -f makefile.vc INSTALLDIR=d:\\dev_libs\\occ TCLDIR=D:\\work\\3rd\\occ_packages\\tcl8.6.16nmake -f makefile.vc install INSTALLDIR=d:\\dev_libs\\occ TCLDIR=D:\\work\\3rd\\occ_packages\\tcl8.6.164. OCC OCC Release参考 tcl/tk编译 Build 3rd-parties" },

  { "title": "系统监控工具套件sysstat的使用", "url": "/blog/2025/%E7%B3%BB%E7%BB%9F%E7%9B%91%E6%8E%A7%E5%B7%A5%E5%85%B7sysstat/", "categories": "Performance",
  "tags": "Performance, Linux", "date": "2025-03-06", "content":
  "# https://sysstat.github.io/sudo apt-get install sysstatsudo dpkg-reconfigure sysstat # select \"Yes\"sar（System Activity Reporter 系统活动情况报告）是目前Linux上最为全面的系统性能分析工具之一，可以从多方面对系统的活动进行报告，包括：文件的读写情况、系统调用的使用情况、磁盘I/O、CPU效率、内存使用状况、进程活动及IPC有关的活动等。Linux内核维护着一些内部计数器，这些计数器包含了所有的请求及其完成时间和I/O块数等信息，sar命令从所有的这些信息中计算出请求的利用率和比例，以便找出瓶颈所在。1. sar 命令的使用用法语法：用法: sar [ 选项 ] [ \u0026lt;时间间隔\u0026gt; [ \u0026lt;次数\u0026gt; ] ]主选项和报告： -b I/O 和传输速率信息状况 -B 分页状况 -d 块设备状况 -I { \u0026lt;中断\u0026gt; | SUM | ALL | XALL } 中断信息状况 -m 电源管理信息状况 -n { \u0026lt;关键词\u0026gt; [,...] | ALL } 网络统计信息 关键词可以是： DEV 网卡 EDEV 网卡 (错误) NFS NFS 客户端 NFSD NFS 服务器 SOCK Sockets (套接字) (v4) IP IP 流 (v4) EIP IP 流 (v4) (错误) ICMP ICMP 流 (v4) EICMP ICMP 流 (v4) (错误) TCP TCP 流 (v4) ETCP TCP 流 (v4) (错误) UDP UDP 流 (v4) SOCK6 Sockets (套接字) (v6) IP6 IP 流 (v6) EIP6 IP 流 (v6) (错误) ICMP6 ICMP 流 (v6) EICMP6 ICMP 流 (v6) (错误) UDP6 UDP 流 (v6) -q 队列长度和平均负载 -r 内存利用率 -R 内存状况 -S 交换空间利用率 -u [ ALL ] CPU 利用率 -v Kernel table 状况 -w 任务创建与系统转换统计信息 -W 交换信息 -y TTY 设备状况 -o {\u0026lt;文件路径\u0026gt;} 将命令结果以二进制格式存放在指定文件中 怀疑CPU存在瓶颈，可用sar -u和sar -q等来查看 怀疑内存存在瓶颈，可用sar -B、sar -r和sar -W等来查看 怀疑I/O存在瓶颈，可用sar -b、sar -u和sar -d等来查看1.1. CPU利用率sar -u [ \u0026lt;时间间隔\u0026gt; [ \u0026lt;次数\u0026gt; ] ]使用示例：sar -u 1 3Linux 2.6.32-642.13.1.el6.x86_64 (upfor106) 2018年04月25日 _x86_64_ (1 CPU)10时33分08秒 CPU %user %nice %system %iowait %steal %idle10时33分09秒 all 0.00 0.00 0.00 0.00 0.00 100.0010时33分10秒 all 0.99 0.00 0.99 0.00 0.00 98.0210时33分11秒 all 0.00 0.00 0.00 0.00 0.00 100.00平均时间: all 0.33 0.00 0.33 0.00 0.00 99.33输出说明：CPU：all 表示统计信息为所有 CPU 的平均值。%user：显示在用户级别(application)运行使用 CPU 总时间的百分比%nice：显示在用户级别，用于nice操作，所占用 CPU 总时间的百分比%system：在核心级别(kernel)运行所使用 CPU 总时间的百分比%iowait：显示用于等待I/O操作占用 CPU 总时间的百分比%steal：管理程序(hypervisor)为另一个虚拟进程提供服务而等待虚拟 CPU 的百分比%idle：显示 CPU 空闲时间占用 CPU 总时间的百分比： 1. 若 %iowait 的值过高，表示硬盘存在I/O瓶颈 2. 若 %idle 的值高但系统响应慢时，有可能是 CPU 等待分配内存，此时应加大内存容量 3. 若 %idle 的值持续低于1，则系统的 CPU 处理能力相对较低，表明系统中最需要解决的资源是 CPU1.2. 队列长度和平均负载sar -q [ \u0026lt;时间间隔\u0026gt; [ \u0026lt;次数\u0026gt; ] ]示例：sar -q 1 3Linux 2.6.32-696.13.2.el6.x86_64 (upfor163) 2018年04月25日 _x86_64_ (2 CPU)11时00分35秒 runq-sz plist-sz ldavg-1 ldavg-5 ldavg-1511时00分36秒 0 268 0.00 0.00 0.0011时00分37秒 0 268 0.00 0.00 0.0011时00分38秒 0 268 0.00 0.00 0.00平均时间: 0 268 0.00 0.00 0.00输出说明：runq-sz：运行队列的长度（等待运行的进程数）plist-sz：进程列表中进程（processes）和线程（threads）的数量ldavg-1：最后1分钟的系统平均负载（System load average）ldavg-5：过去5分钟的系统平均负载ldavg-15：过去15分钟的系统平均负载1.3. 内存利用率sar -r [ \u0026lt;时间间隔\u0026gt; [ \u0026lt;次数\u0026gt; ] ]使用示例：sar -r 1 3Linux 2.6.32-696.13.2.el6.x86_64 (upfor163) 2018年04月25日 _x86_64_ (2 CPU)10时53分00秒 kbmemfree kbmemused %memused kbbuffers kbcached kbcommit %commit10时53分01秒 2027760 2028732 50.01 145492 1243820 1163900 28.6910时53分02秒 2027620 2028872 50.02 145492 1243820 1163896 28.6910时53分03秒 2028100 2028392 50.00 145492 1243820 1163900 28.69平均时间: 2027827 2028665 50.01 145492 1243820 1163899 28.69输出说明：kbmemfree：这个值和 free 命令中的 free 值基本一致，所以它不包括 buffer 和 cache 的空间kbmemused：这个值和 free 命令中的 used 值基本一致,所以它包括 buffer 和 cache 的空间%memused：这个值是 kbmemused 和内存总量(不包括 swap)的一个百分比kbbuffers 和 kbcached：这两个值就是 free 命令中的 buffer 和 cachekbcommit：保证当前系统所需要的内存，即为了确保不溢出而需要的内存(RAM + swap)%commit：这个值是 kbcommit 与内存总量(包括 swap)的一个百分比1.4. 内存分页sar -B [ \u0026lt;时间间隔\u0026gt; [ \u0026lt;次数\u0026gt; ] ]sar -B 1 3Linux 2.6.32-696.13.2.el6.x86_64 (upfor163) 2018年04月25日 _x86_64_ (2 CPU)10时55分41秒 pgpgin/s pgpgout/s fault/s majflt/s pgfree/s pgscank/s pgscand/s pgsteal/s %vmeff10时55分42秒 0.00 0.00 5723.76 0.00 3356.44 0.00 0.00 0.00 0.0010时55分43秒 0.00 0.00 1185.00 0.00 312.00 0.00 0.00 0.00 0.0010时55分44秒 0.00 0.00 27.00 0.00 56.00 0.00 0.00 0.00 0.00平均时间: 0.00 0.00 2323.26 0.00 1248.50 0.00 0.00 0.00 0.00输出项说明：pgpgin/s：表示每秒从磁盘或SWAP置换到内存的字节数(KB)pgpgout/s：表示每秒从内存置换到磁盘或SWAP的字节数(KB)fault/s：每秒钟系统产生的缺页数，即主缺页与次缺页之和(major + minor)majflt/s：每秒钟产生的主缺页数pgfree/s：每秒被放入空闲队列中的页个数pgscank/s：每秒被 kswapd 扫描的页个数pgscand/s：每秒直接被扫描的页个数pgsteal/s：每秒钟从 cache 中被清除来满足内存需要的页个数%vmeff：每秒清除的页(pgsteal)占总扫描页(pgscank + pgscand)的百分比1.5. 系统页交换sar -W [ \u0026lt;时间间隔\u0026gt; [ \u0026lt;次数\u0026gt; ] ]使用示例：sar -W 1 3Linux 2.6.32-696.13.2.el6.x86_64 (upfor163) 2018年04月25日 _x86_64_ (2 CPU)11时01分45秒 pswpin/s pswpout/s11时01分46秒 0.00 0.0011时01分47秒 0.00 0.0011时01分48秒 0.00 0.00平均时间: 0.00 0.00输出项说明：pswpin/s：每秒系统换入的交换页面（swap page）数量pswpout/s：每秒系统换出的交换页面（swap page）数量1.6. 块设备sar -d [ \u0026lt;时间间隔\u0026gt; [ \u0026lt;次数\u0026gt; ] ]使用示例：sar -d 1 3Linux 2.6.32-696.13.2.el6.x86_64 (upfor163) 2018年04月25日 _x86_64_ (2 CPU)11时02分46秒 DEV tps rd_sec/s wr_sec/s avgrq-sz avgqu-sz await svctm %util11时02分47秒 dev252-0 0.00 0.00 0.00 0.00 0.00 0.00 0.00 0.0011时02分47秒 DEV tps rd_sec/s wr_sec/s avgrq-sz avgqu-sz await svctm %util11时02分48秒 dev252-0 6.06 0.00 64.65 10.67 0.00 0.00 0.00 0.0011时02分48秒 DEV tps rd_sec/s wr_sec/s avgrq-sz avgqu-sz await svctm %util11时02分49秒 dev252-0 0.00 0.00 0.00 0.00 0.00 0.00 0.00 0.00平均时间: DEV tps rd_sec/s wr_sec/s avgrq-sz avgqu-sz await svctm %util平均时间: dev252-0 2.00 0.00 21.33 10.67 0.00 0.00 0.00 0.00输出说明：tps: 每秒从物理磁盘 I/O 的次数。多个逻辑请求会被合并为一个 I/O 磁盘请求，一次传输的大小是不确定的rd_sec/s: 每秒读扇区的次数wr_sec/s: 每秒写扇区的次数avgrq-sz: 平均每次设备 I/O 操作的数据大小(扇区)avgqu-sz: 磁盘请求队列的平均长度await: 从请求磁盘操作到系统完成处理，每次请求的平均消耗时间，包括请求队列等待时间，单位是毫秒(1秒=1000毫秒)svctm: 系统处理每次请求的平均时间,不包括在请求队列中消耗的时间.%util: I/O请求占CPU的百分比，比率越大，说明越饱和 1. avgqu-sz 的值较低时，设备的利用率较高 2. 当%util的值接近 1% 时，表示设备带宽已经占满2. 参考链接 sysstat——系统性能监控神器 sar —— Linux 上最为全面的系统性能分析工具之一" },

  { "title": "待翻译 -- Performance Hints", "url": "/blog/2025/Performance-Hints/", "categories": "Performance",
  "tags": "Performance", "date": "2025-03-06", "content":
  "Performance Hints – Jeff Dean 微信中文翻译 All About Rooflines Algorithms for Modern Hardware" },

  { "title": "OCC 待学习资料记录", "url": "/blog/2025/OCCT-study-notes/", "categories": "OCCT",
  "tags": "OCCT", "date": "2025-03-06", "content":
  "1. 笔记1.1. 求交对二次曲线与二次曲面求交，使用解析几何的方法，计算出交点，依赖底层数学库math_DirectPolynomialRoots。类math_DirectPolynomialRoots可以对最多4次方程进行求解。如下曲面可以使用二次曲线表示，并使用解析几何求交： I_gp_Pln：二次曲面特例； I_gp_Sphere：解析球面； I_gp_Cylinder：解析柱面； I_gp_Cone：解析锥面；这些曲面都可以使用二次曲线表示，即这个二次曲线方程的参数确定的上述这些曲面。OCC中相应的类为IntAna_Quaric。二次曲线与自由曲面求交，使用数值计算方法，即Newton-Raphson迭代逐次逼近 数学之美：牛顿-拉夫逊迭代法原理及其实现。 OpenCASCADE 线面求交 解析几何求交之直线与二次曲面 Modeling Algorithms1.2. OCC中的实体表示 – ModellingOCC表示B-Rep，其核心概念有两个：几何，拓扑。几何表示实体的形状，如二次曲线方程，Bezier曲线、NURBS曲线等。拓扑为存储结构，是一个树形结构，组成一个实体的各个部分，如顶点、边、面、体、轮廓等。 OpenCascade Overview2. Quaoar Workshop 的学习资料 youtube教学视频 开源项目 Analysis Situs OCC论坛 Analysis Situs Forums gitlab – AnalysisSitus3. 开源项目 Mayo Get started with Mayo github – Mayo4. FreeCAD github – Module developer’s guide to FreeCAD source code – FreeCAD overview and architecture github – CAD modules in FreeCAD – CAD modules in FreeCAD" },

  { "title": "NVIDIA Jetson Orin AGX 安装", "url": "/blog/2025/NVIDIA-Jetson-Orin-AGX-%E5%AE%89%E8%A3%85/", "categories": "NVIDIA, Jetson",
  "tags": "NVIDIA, Jetson, AI", "date": "2025-03-01", "content":
  "1. NVIDIA Jetson Orin AGX 安装使用 Nvidia Jetson SDK Manager 安装 Jetson Orin AGX 系列设备的系统镜像。截止目前最新版为 7.2， Ubuntu 24.04。安装完成之后，安装如下软件包：# 安装 Pythonsudo apt updatesudo apt install python3sudo apt install python3-pip# 安装 JetPack（含 CUDA、cuDNN、TensorRT 等组件），需在 jtop 之前装好sudo apt install nvidia-jetpack# 安装 jtop（依赖 JetPack 组件才能正常显示信息）sudo pip3 install -U pipsudo pip3 install --break-system-packages -U jetson-stats# https://github.com/rbonghi/jetson_stats查看已经安装的组件：git clone https://github.com/jetsonhacks/jetsonUtilities.gitcd jetsonUtilitiespython3 jetsonInfo.py安装 GTest：sudo apt install libgtest-dev libgmock-dev1.A. 参考及资源1.A.1. 参考 NVIDIA Jetson Linux 36.4：官方资料页面，包含组件以及驱动源码下载列表 Nvidia Jetson AGX Orin开发板配置与使用：安装更多软件包，比如 SLAM，深度学习框架等。 Nvidia Jetson AGX Orin系统刷写1.A.2. 资源 Jetson AI Lab – 支持的LLM模型 JetPack 7.2: Jetson Software Goes Agentic with Jetson Linux 39.2 jetson skills：官方github仓库 Jetson AI Lab – Supported Models：Jetson AI Lab 模型库2. CLI 条件下连接 WIFI刷写完系统后，若只有命令行环境，先配置网络，后续的 apt、git、docker pull 等操作都需要联网。# 查看网卡列表nmcli device status# 扫描附近的WiFi网络nmcli device wifi list# 连接到指定的WiFi网络（替换SSID_NAME和WIFI_PASSWORD）sudo nmcli device wifi connect \"SSID_NAME\" password \"WIFI_PASSWORD\"3. 在 Jetson Orin AGX 上挂载 SSD，挂载 Docker 目录及用户目录 这里记录我在 Jetson Orin AGX 上挂载 M.2 NVMe SSD（约 1TB）的过程。前提是 SSD 已正确插入，系统能看到设备 /dev/nvme0n1（目前未分区）。所有操作均需要 root（使用 sudo）。3.1. 创建分区 \u0026amp; 格式化# 确认设备lsblk -o NAME,SIZE,MODEL# 使用 fdisk 创建一个主分区（占用整个磁盘）sudo fdisk /dev/nvme0n1# 在交互界面依次输入：# n → 新建分区# p → 主分区# 1 → 分区号# （回车）默认起始扇区# （回车）默认结束扇区（使用剩余全部空间）# w → 写入分区表并退出完成后得到 /dev/nvme0n1p1。格式化为 ext4（可根据需要换成 xfs、btrfs 等）：sudo mkfs.ext4 -L nvme_ssd /dev/nvme0n1p1查看 UUID（后续挂载会用到）：sudo blkid /dev/nvme0n1p1# 示例输出：/dev/nvme0n1p1: UUID=\"a1b2c3d4-e5f6-7890-abcd-ef1234567890\" TYPE=\"ext4\" PARTLABEL=\"nvme_ssd\"记下此 UUID（例如 a1b2c3d4-e5f6-7890-abcd-ef1234567890）。3.2. 创建挂载点 \u0026amp; 加入开机自动挂载（/etc/fstab）# 统一挂载点（可自行修改路径）sudo mkdir -p /mnt/ssd# 临时挂载测试sudo mount /dev/nvme0n1p1 /mnt/ssddf -hT /mnt/ssd # 应看到约 1 TB 可用# 写入 fstab（使用 UUID 更稳妥)echo \"UUID=a1b2c3d4-e5f6-7890-abcd-ef1234567890 /mnt/ssd ext4 defaults,noatime 0 2\" | sudo tee -a /etc/fstab# 检查并生效sudo mount -a以后每次开机都会自动把 /dev/nvme0n1p1 挂载到 /mnt/ssd。3.3. 把 Docker 镜像/容器数据放到 SSD3.3.1. 推荐方案：修改 Docker 的 data-root# 停止 Dockersudo systemctl stop docker# 创建目标目录并迁移原有数据（如果有）sudo mkdir -p /mnt/ssd/dockersudo rsync -aXS /var/lib/docker/. /mnt/ssd/docker/# 配置 Docker 使用新目录sudo mkdir -p /etc/dockersudo tee /etc/docker/daemon.json \u0026gt; /dev/null \u0026lt;\u0026lt;'EOF'{ \"data-root\": \"/mnt/ssd/docker\"}EOF# 重启 Dockersudo systemctl start dockersudo systemctl status docker # 确保无错误# 验证docker info | grep \"Docker Root Dir\"# 应显示： Docker Root Dir: /mnt/ssd/docker替代方案： 使用符号链接（停止 Docker 后 mv /var/lib/docker /var/lib/docker.bak \u0026amp;\u0026amp; ln -s /mnt/ssd/docker /var/lib/docker），效果相同但不如 daemon.json 直观。 这里写入的 daemon.json 只包含 data-root。在第 4 节部署 LLM 时，会把它整体替换为包含 runtimes、registry-mirrors 的完整版本，到时直接覆盖即可。3.4. 把用户主目录（/home/hxf0223）也放到 SSD如果计划在用户目录下存放大量 Docker 镜像、离线 LLM 模型等，建议把 整个用户主目录 挂载到 SSD。 注意：下面 mv /home/hxf0223 时，若当前正登录在该家目录下会失败。请先 cd /（或切到 root 用户）再操作，并确认没有进程占用家目录（可 sudo lsof +D /home/hxf0223 检查）。# 在 SSD 上创建对应目录sudo mkdir -p /mnt/ssd/home/hxf0223# 复制现有数据（保持权限、属性）sudo rsync -aXS /home/hxf0223/. /mnt/ssd/home/hxf0223/# 备份原目录sudo mv /home/hxf0223 /home/hxf0223.bak# 绑定挂载使 /home/hxf0223 指向 SSDsudo mount --bind /mnt/ssd/home/hxf0223 /home/hxf0223# 加入 fstab 以实现开机自动绑定挂载echo \"/mnt/ssd/home/hxf0223 /home/hxf0223 none bind 0 0\" | sudo tee -a /etc/fstab# 验证df -hT /home/hxf0223 # 应看到和 /mnt/ssd 相同的可用空间mount | grep \"on /home/hxf0223\"# 应显示类似： /mnt/ssd/home/hxf0223 on /home/hxf0223 type none (rw,relatime,bind) 如果希望 所有用户的主目录 都放在 SSD，可把 /mnt/ssd/home 挂载到 /home，步骤类似。3.5. 检查一切是否正常# 检查挂载情况df -hT /mnt/ssd /home/hxf0223 /var/lib/docker# 检查 Docker 实际使用的目录docker info | grep \"Docker Root Dir\"# 检查用户目录是否真的指向 SSDmount | grep \"on /home/hxf0223\"如果输出均符合预期，说明已经成功： 将 M.2 SSD 挂载为 /mnt/ssd（额外存储）。 将 Docker 数据目录迁移到 SSD（/var/lib/docker → /mnt/ssd/docker）。 将个人工作目录 /home/hxf0223 挂载到 SSD，后续的镜像拉取、模型下载、编译等 I/O 密集操作都会走高速 SSD。3.6. 常见问题 \u0026amp; 小贴士 问题 解决方案 开机后挂载失败 检查 /etc/fstab 是否有语法错误；用 sudo mount -a 测试；根据报错修正 UUID 或挂载选项。 Docker 启动后仍在用旧目录 确认 /etc/docker/daemon.json JSON 合法（末尾无多余逗号）；重启 Docker 前先 docker info 查看 Docker Root Dir。 用户目录下出现权限问题 迁移时使用了 -aXS 保持权限；如仍有问题，可 sudo chown -R $USER:$USER /home/hxf0223。 3.7. 验证迁移并清理 eMMC 上的备份以释放空间完成上述迁移后，建议检查数据是否确实已迁移到 SSD，并删除 eMMC 上可能残留的备份副本，以腾出存储空间。3.7.1. 验证 Docker 是否真的在使用 SSD# 显示 Docker 的实际数据根目录sudo docker info | grep -i \"Docker Root Dir\"预期输出：Docker Root Dir: /mnt/ssd/docker。如果看到上述路径，说明 Docker 已经在使用挂载到 SSD 的目录；此时 /var/lib/docker 只是一个普通目录（挂载点），实际数据不再占用 eMMC 空间。3.7.2. 验证 /home/hxf0223 是否已指向 SSD# 查看挂载情况mount | grep \"on /home/hxf0223\"预期输出（示例）：/mnt/ssd/home/hxf0223 on /home/hxf0223 type none (rw,relatime,bind)或者检查是否是 bind‑mount：df -hT /home/hxf0223 应该和 /mnt/ssd 显示相同的文件系统（ext4）和可用空间。如果是这样，用户的实际数据已经在 SSD 上，eMMC 上只保存了 mount point 本身（几乎不占空间）。3.7.3. 检查是否还有明显的备份占用空间按照之前的步骤，迁移后通常会有以下两种备份形式： 位置 可能的内容 大小检查命令 /var/lib/docker （原始目录） 已迁移的 Docker 数据（若未删除） sudo du -sh /var/lib/docker /home/hxf0223.bak （用户家目录备份） 完整的 /home/hxf0223 副本 sudo du -sh /home/hxf0223.bak（若存在） 示例检查：sudo du -sh /var/lib/dockersudo du -sh /home/hxf0223.bak 2\u0026gt;/dev/null || echo \"没有找到 .bak 目录\" 我这边 /var/lib/docker 只有约 212 KB，说明几乎没有占用空间；如果看到类似的几百 MB 或几 GB，那就需要考虑删除。3.7.4. 安全清理（在确认无误后）3.7.4.1. 删除 Docker 原始目录的残留（如果仍然占用明显空间） 注意：只有在确认 docker info 的 Docker Root Dir 已经是 /mnt/ssd/docker 且 /var/lib/docker 中不含需要的镜像/容器时才删除。# 先再次确认 Docker 在使用 SSD（见 3.7.1）sudo docker info | grep -i \"Docker Root Dir\"# 如果确认无误，删除残留目录（保险起见先改名再删）sudo mv /var/lib/docker /var/lib/docker.bak_$(date +%F)# 等待一会儿，确认一切正常后再彻底删除sudo rm -rf /var/lib/docker.bak_$(date +%F)3.7.4.2. 删除用户家目录备份（如果存在）if [ -d /home/hxf0223.bak ]; then # 再次确认当前家目录已经是挂载到 SSD（见 3.7.2） mount | grep \"on /home/hxf0223\" \u0026amp;\u0026amp; echo \"家目录已正确挂载\" # 先改名再删，防止误删 sudo mv /home/hxf0223.bak /home/hxf0223.bak_$(date +%F) sudo rm -rf /home/hxf0223.bak_$(date +%F)fi3.7.4.3. 可选：清理 Docker 未使用的镜像/缓存（进一步释放 SSD 空间）sudo docker system prune -af # 移除所有停止的容器、未使用的镜像、网络等# 如果只想保留最近的，可改为：# sudo docker system prune -a --filter \"until=24h\"3.7.5. 再次确认系统状态# 检查 eMMC 剩余空间（应该会有所提升)df -hT / # eMMC 根分区df -hT /mnt/ssd # SSD 挂载点df -hT /home/hxf0223 # 应该和上面一致如果一切正常，说明大量数据已迁移到 SSD，并在 eMMC 上释放了可用空间。后续拉取大型离线 LLM 模型、构建 Docker 镜像或进行其它 I/O 密集型工作都将享受到 SSD 的带宽和低延迟。4. 在 Jetson Orin AGX 上部署与运行 LLM (Gemma 4) 设备: NVIDIA Jetson AGX Orin (64GB 统一内存) 系统: Ubuntu 24.04, JetPack 7.2-b187, CUDA 13.2 部署方式: 推荐使用 Docker Compose + ModelScope 进行一键化管理与部署；同时保留 Raw Docker + HuggingFace Mirror 作为手动分步运行的备选方案。4.1. 运行 LLM 完整组件图flowchart TB subgraph 硬件层[\"🖥️ 硬件层\"] GPU[\"Jetson AGX Orin\\n64GB 统一内存\"] SSD[\"NVMe SSD\\n/mnt/ssd\"] end subgraph 系统驱动层[\"⚙️ 系统驱动层\"] JetPack[\"JetPack 7.2\\n(含 L4T 内核)\"] CUDA[\"CUDA 13.2\"] NVIDIA-RT[\"nvidia-container-runtime\\n(Docker 访问 GPU 的桥梁)\"] end subgraph 容器运行时层[\"📦 容器运行时层\"] Docker[\"Docker\"] Containerd[\"containerd\\n(镜像/容器存储 → SSD)\"] end subgraph 推理引擎层[\"🧠 推理引擎层（支持二选一 / 共存）\"] vLLM[\"vLLM 容器镜像\\nghcr.io/.../vllm:gemma4-jetson-orin\\n(高性能，内存占用大)\"] llamaCPP[\"llama.cpp 容器镜像\\nghcr.io/.../llama_cpp:latest-jetson-orin\\n(内存效率高)\"] end subgraph 模型数据层[\"💾 模型数据层（存在 SSD）\"] MS-Cache[\"ModelScope / HF 缓存\\n/mnt/ssd/modelscope/ 或 /mnt/ssd/huggingface/\"] QAT-WT[\"31B QAT 权重\\n(量化感知训练版本)\\n→ 给 vLLM / llama.cpp\"] GGUF-WT[\"31B GGUF 权重\\n→ 给 llama.cpp\"] end subgraph 网络源[\"🌐 网络/镜像源\"] ModelScope-Site[\"ModelScope.cn\\n(魔搭社区 - 国内加速)\"] HF-Mirror[\"hf-mirror.com\\n(HuggingFace 国内镜像)\"] Docker-Mirror[\"docker.nju.edu.cn\\n(Docker 镜像加速)\"] end subgraph 配置文件[\"📝 配置文件\"] DaemonJSON[\"/etc/docker/daemon.json\\n· data-root → SSD\\n· registry-mirrors\"] ContainerdTOML[\"/etc/containerd/config.toml\\n· root → SSD\"] BashRC[\"~/.bashrc\\n· MODELSCOPE_CACHE\\n· HF_HOME\\n· HF_ENDPOINT\"] end JetPack --\u0026gt; CUDA CUDA --\u0026gt; NVIDIA-RT NVIDIA-RT --\u0026gt; Docker Docker --\u0026gt; Containerd Containerd --\u0026gt; SSD Docker --\u0026gt; vLLM Docker --\u0026gt; llamaCPP Docker-Mirror -.-\u0026gt; Docker ModelScope-Site -.-\u0026gt; MS-Cache HF-Mirror -.-\u0026gt; MS-Cache vLLM -.-\u0026gt; QAT-WT llamaCPP -.-\u0026gt; GGUF-WT QAT-WT --\u0026gt; MS-Cache GGUF-WT --\u0026gt; MS-Cache DaemonJSON -.-\u0026gt; Docker ContainerdTOML -.-\u0026gt; Containerd BashRC -.-\u0026gt; MS-Cache SSD --\u0026gt; MS-Cache GPU --- SSD style GPU fill:#4a9,color:#fff style SSD fill:#4a9,color:#fff style vLLM fill:#e8731a,color:#fff style llamaCPP fill:#e8731a,color:#fff style QAT-WT fill:#d33,color:#fff style GGUF-WT fill:#393,color:#fff4.2. 环境准备与推理引擎下载安装为了防止写满宿主机板载的 eMMC 存储，并确保推理引擎可以利用 GPU 硬件加速，需要依次进行容器存储优化以及推理引擎的下载/安装。4.2.1. Docker 运行时及镜像加速配置编辑 /etc/docker/daemon.json。这里在第 3.3.1 的基础上合并了 runtimes（NVIDIA GPU 运行时）与 registry-mirrors（国内镜像加速）字段，是最终版本，直接整体替换原文件即可：{ \"runtimes\": { \"nvidia\": { \"args\": [], \"path\": \"nvidia-container-runtime\" } }, \"data-root\": \"/mnt/ssd/docker\", \"max-concurrent-downloads\": 6, \"registry-mirrors\": [\"https://docker.nju.edu.cn\"]}4.2.2. Containerd 存储配置先把目录建好，并生成默认配置（文件默认可能不存在），再编辑 /etc/containerd/config.toml 重定向 containerd 目录至 SSD：sudo mkdir -p /mnt/ssd/containerdsudo containerd config default | sudo tee /etc/containerd/config.toml \u0026gt; /dev/null将其中的 root 与 state 字段改为：root = \"/mnt/ssd/containerd\"state = \"/mnt/ssd/containerd/state\"保存配置后，重启容器相关服务以使改动生效：sudo systemctl restart containerdsudo systemctl restart docker4.2.3. 系统环境变量配置编辑当前用户的 ~/.bashrc，将模型缓存路径和镜像加速端点重定向到 SSD 上：# ModelScope (推荐：国内稳定高速下载)export MODELSCOPE_CACHE=/mnt/ssd/modelscope# HuggingFace (备用：配置国内镜像站)export HF_HOME=/mnt/ssd/huggingfaceexport HF_HUB_CACHE=/mnt/ssd/huggingface/hubexport HF_ENDPOINT=https://hf-mirror.com运行 source ~/.bashrc 激活变量。4.2.4. 推理引擎的下载与安装方法在 Jetson Orin AGX 上部署推理引擎（llama.cpp 或 vLLM）时，通常有以下三种安装方式可选：选项 A：拉取预编译 Docker 镜像（推荐，极简无痛）如果使用容器化运行，下载安装引擎等同于拉取预编译镜像。镜像中已自动打包编译好的 CUDA 版本的推理引擎。在已配置国内镜像加速的基础上，直接拉取所需的引擎镜像： 拉取 llama.cpp 推理引擎镜像： sudo docker pull ghcr.io/nvidia-ai-iot/llama_cpp:latest-jetson-orin 拉取 vLLM 推理引擎镜像： sudo docker pull ghcr.io/nvidia-ai-iot/vllm:gemma4-jetson-orin 选项 B：使用 jetson-containers 本地编译容器镜像（进阶定制）如果需要针对特定 JetPack 固件版本或定制参数打包，可以使用官方的 jetson-containers 工具链在本地自动编译和构建容器：# 1. 克隆 NVIDIA 官方的 jetson-containers 仓库git clone https://github.com/dusty-nv/jetson-containers.gitcd jetson-containers# 2. 安装依赖pip install -r requirements.txt# 3. 本地编译构建 llama.cpp 推理引擎容器镜像（会自动调取本机 CUDA 编译器编译源码）./build.sh llama_cpp# 4. 或者构建 vllm 推理引擎镜像./build.sh vllm选项 C：在宿主机直接源码编译安装（裸机运行备选）如果不使用 Docker，想在宿主机操作系统中直接编译运行 llama.cpp 推理引擎，按以下步骤操作：# 1. 安装编译所需的系统依赖项sudo apt updatesudo apt install -y build-essential cmake git libcurl4-openssl-dev libssl-dev# 2. 克隆官方仓库并包含第三方依赖子模块git clone https://github.com/ggml-org/llama.cppcd llama.cppgit submodule update --init --recursive# 3. 创建编译目录并使用 CMake 进行配置mkdir build \u0026amp;\u0026amp; cd build# 关键：针对 Orin 显卡架构指定 -DCMAKE_CUDA_ARCHITECTURES=87 启用 GPU 加载cmake .. -DGGML_CUDA=ON -DCMAKE_CUDA_ARCHITECTURES=87 -DGGML_CUDA_FA_ALL_QUANTS=ON# 4. 编译构建make -j$(nproc)# 5. 安装至系统路径（可选）sudo make install4.2.5. Docker 模式下推理引擎的更新步骤当 llama.cpp 或 vLLM 上游发布更新、需要更换运行环境中的推理引擎时，按部署方式二选一。方式 A：使用 Docker Compose 管理服务（推荐）# 进入部署代码目录cd /path/to/gemma4-server# 从镜像仓库中拉取最新的推理引擎镜像，并重建服务sudo docker compose pullsudo docker compose up -d --build 说明： docker compose pull：读取当前目录下的 docker-compose.yml，把其中所有服务引用的镜像从仓库下载到本地，只下载不启动容器。 docker compose up -d --build：按配置文件创建并启动全部服务。-d 表示后台运行（detach，不占用终端）；--build 表示启动前重新构建那些用 build: 定义的本地镜像（对纯 image: 拉取的服务无影响）。 两条分开执行的好处是：pull 若因网络失败不会影响正在运行的旧服务，等镜像确认下载完成后再 up 切换到新版本，更新过程更可控。 方式 B：使用手动 docker run 命令启动# 拉取最新镜像（llama.cpp / vLLM 二选一）sudo docker pull ghcr.io/nvidia-ai-iot/llama_cpp:latest-jetson-orinsudo docker pull ghcr.io/nvidia-ai-iot/vllm:gemma4-jetson-orin# 停止并删除当前正在运行的老版本容器（假设名为 gemma4）sudo docker stop gemma4随后重新执行原有的 docker run 启动指令，新创建的容器即会自动应用并运行最新的推理引擎。4.3. 现代化部署方案：Docker Compose + ModelScope (推荐)此方案已整理并发布在代码仓库 Gemma-4 Local Servers on Jetson Orin。使用 Docker Compose 可以一键拉起推理服务，且 ModelScope 国内下载速度更有保障。4.3.1. 克隆代码仓库git clone https://github.com/aispace02/gemma4-server.gitcd gemma4-server4.3.2. 下载模型权重 (通过 ModelScope)使用 ModelScope 的 Python SDK 或 CLI 下载模型权重（推荐采用 量化感知训练 (QAT) 版本以在 Jetson 上获得更高精度与速度）：# 1. 安装 modelscope 库pip install modelscope --break-system-packages# 若不想加 --break-system-packages，也可先建虚拟环境：# python3 -m venv ~/venv \u0026amp;\u0026amp; source ~/venv/bin/activate \u0026amp;\u0026amp; pip install modelscope# 2. 运行 Python 脚本将模型下载至 SSD 对应的 ModelScope 缓存路径python3 -c \"from modelscope import snapshot_download# 下载 Gemma 4 31B QAT 版本模型snapshot_download('google/gemma-4-31B-it-QAT', cache_dir='/mnt/ssd/modelscope')\"4.3.3. 使用 Docker Compose 一键启动推理服务根据需要的推理引擎，运行以下命令（Docker Compose 会自动读取配置文件，挂载宿主机的 SSD 模型目录，并开启服务）： 运行 vLLM 服务（适合大显存，高吞吐性能优先）： docker compose up -d vllm 运行 llama.cpp 服务（适合轻量部署，内存开销较小）： docker compose up -d llama-cpp 4.4. 备选手动方案：Raw Docker + HuggingFace Mirror如果不希望克隆额外的 compose 仓库，也可以通过纯手动命令下载权重和运行 Docker 容器。4.4.1. 下载模型权重 (HuggingFace 命令行方式)使用 hf 下载工具将指定模型拉取到 SSD 上的缓存目录中：export HF_ENDPOINT=https://hf-mirror.comexport HF_HOME=/mnt/ssd/huggingface# 下载 Gemma 4 31B AWQ 模型（vLLM 用）hf download cyankiwi/gemma-4-31B-it-AWQ-4bit模型默认下载位置：/mnt/ssd/huggingface/hub/models--cyankiwi--gemma-4-31B-it-AWQ-4bit/。 如果要用 llama.cpp，则下载 GGUF 格式权重： hf download ggml-org/gemma-4-31B-it-GGUF --include \"*Q4_K_M*\" 下载后位于 /mnt/ssd/huggingface/hub/models--ggml-org--gemma-4-31B-it-GGUF/。4.4.2. 手动启动 vLLM 容器服务sudo docker run -it --rm --pull always --name gemma4 --runtime=nvidia --network host \\ -v /mnt/ssd/huggingface:/data/models/huggingface \\ -e HF_ENDPOINT=https://hf-mirror.com \\ ghcr.io/nvidia-ai-iot/vllm:gemma4-jetson-orin \\ vllm serve cyankiwi/gemma-4-31B-it-AWQ-4bit \\ --port 18000 \\ --gpu-memory-utilization 0.70 \\ --max-model-len 32768 \\ --enable-auto-tool-choice \\ --reasoning-parser gemma4 \\ --tool-call-parser gemma44.4.3. 手动启动 llama.cpp 容器服务sudo docker run -it --rm --name gemma4 \\ --runtime=nvidia --network host \\ -v /mnt/ssd/huggingface:/data/models/huggingface \\ -e HF_ENDPOINT=https://hf-mirror.com \\ ghcr.io/nvidia-ai-iot/llama_cpp:latest-jetson-orin \\ llama-server -hf ggml-org/gemma-4-31B-it-GGUF:Q4_K_M --port 8080 这里 -hf ggml-org/gemma-4-31B-it-GGUF:Q4_K_M 会从挂载的 HF 缓存中读取（即 4.4.1 中下载的 GGUF 权重）；若缓存里没有，容器会通过 HF_ENDPOINT 镜像站自动拉取。 关于 Docker 镜像的双标签问题 在 docker images 中可能会看到 ghcr.io/nvidia-ai-iot/vllm... 与 ghcr.nju.edu.cn/nvidia-ai-iot/vllm...，两者具有完全相同的 Image ID。这代表南大镜像站的加速标签与官方标签指向同一份物理镜像，不会占用额外的磁盘空间。两个标签都请予以保留，不要删除。4.5. 推理引擎使用与选型对比 引擎特性 🟢 vLLM 🟣 llama.cpp 推荐模型格式 AWQ 量化格式、QAT 量化版本 GGUF 格式、QAT 量化版本 内存/显存开销 较大（Orin 64G 统合显存下需精细限制系统参数） 极小（对 Unified Memory 友好，支持更长上下文） 核心启动参数 --gpu-memory-utilization 0.70 (防止 OOM)--max-model-len 32768 (KV Cache 限制) llama-server -hf \u0026lt;model\u0026gt;:\u0026lt;quant\u0026gt; --port 8080参数极简，易用性高 优势场景 并发吞吐量高，适合做生产 API 部署 适合本地单机轻量推理、极速测试，对内存不足的情况适应度高 4.6. 日常运维操作 测试模型服务可用性（以 vLLM 18000 端口为例）： curl -sN http://127.0.0.1:18000/v1/chat/completions \\ -H 'Content-Type: application/json' \\ -d '{ \"model\": \"cyankiwi/gemma-4-31B-it-AWQ-4bit\", \"messages\": [{\"role\": \"user\", \"content\": \"你好\"}], \"chat_template_kwargs\": {\"enable_thinking\": true}, \"stream\": true }' 手动释放系统内存（当多次推理后内存紧张时）： sudo sysctl -w vm.drop_caches=3 监控容器与 GPU 运行状态： # 查看容器 CPU / 内存占用sudo docker stats# 查看 Jetson 硬件负载（使用 jtop 工具）jtop 4.7. 常见报错及排查方案1. 网络连接超时 ([Errno 101] Network is unreachable) 原因：容器启动时未能成功通过国内镜像站连接 HuggingFace/ModelScope。 解决：确保在 docker run 命令中正确传入了 -e HF_ENDPOINT=https://hf-mirror.com，或确认 ModelScope 缓存路径正确，并已在启动前通过脚本完成权重下载。2. 挂载路径寻找失败 (LocalEntryNotFoundError: Cannot find...) 原因：宿主机与容器内部的模型挂载映射路径不一致。 解决：vLLM 官方 Jetson 镜像预设其 HF_HOME 为 /data/models/huggingface，挂载时务必使用 -v /mnt/ssd/huggingface:/data/models/huggingface，不能使用默认的 /root/.cache。3. 设备内存不足 OOM 报错 (ValueError: Free memory... is less than desired) 原因：Jetson Orin 64GB 属于共享内存架构，除系统固定的 14GB 占用外，可用空间约为 47GB。vLLM 默认会尝试锁定 90% 物理内存导致溢出。 解决：将启动参数设为 --gpu-memory-utilization 0.70，或配合 --max-model-len 16384 缩减 KV Cache。如果仍然报错，建议改用内存开销更低的 llama.cpp 引擎。4.8. 资料链接参考 Jetson AI Labs – Supported Models Gemma 4 on Jetson 官方教程 ModelScope – Gemma 4 31B 官方页面 Gemma-4 Local Servers on Jetson Orin (配套 Compose 仓库) 齐思头条：Gemma 4 MTP/QAT 合入 llama.cpp 实现低内存推理" },

  { "title": "Ampere GPU 新特性", "url": "/blog/2025/CUDA-Ampere-Feature/", "categories": "CUDA",
  "tags": "CUDA", "date": "2025-02-28", "content":
  "1. Async Copy异步拷贝cp.async（即指令LDGSTS）支持4/8/16字节单位的拷贝，其中： 4/8字节单位的拷贝：L2 -\u0026gt; L1 -\u0026gt; SMEM； 16字节单位的拷贝(Bypass L1)：L2 -\u0026gt; SMEM。明显的，16字节单位的拷贝性能最高。另外，cp.async需要使用commit/wait指令来配合使用。A. 资料 4.11. Asynchronous Data Copies：官方文档，介绍cp.async CUDA 11 NEW FEATURES(pdf)：Ampere 新特性介绍 CUDA on NVIDIA GPU AMPERE MICROARCHITECTURE Taking your algorithms to the next level of performance(pdf)" },

  { "title": "使用 CuTe Tiled Copy、Tiled MMA 以及 Multi-Stage 实现高性能 GEMM", "url": "/blog/2025/CuTe-GEMM-TiledCopy-TiledMMA-Pipeline/", "categories": "CUDA",
  "tags": "CUDA", "date": "2025-02-27", "content":
  "代码： https://github.com/HPC02/cuda_perf/blob/master/src/cute_gemm_sm80/gemm_sm80.cu https://github.com/HPC02/cuda_perf/blob/master/src/cute_gemm_sm80/kernel_sm80.cuhTODO：GMEM -\u0026gt; SMEM 不会产生 bank conflicts？配置流程及约束概览： 定义 CTA tile 大小TODO 定义 GMEM -\u0026gt; SMEM 的 Tiled Copy 配置TODO 定义 Tiled MMA 配置（包含SMEM TiledCopy）TODO 定义 SMEM swizzle 配置，以及SMEM Layout（包含multi-stage）TODO 大部分内容已经在其他文章中记录： CUTLASS-Cute 初步(4.1)：MMA Swizzle – MMA、ldmatrix、smem swizzle； CUTLASS-Cute 初步(6)：CUDA GEMM 计算优化、Multi-Stage 与软流水(Software Pipelining)。 CUDA 笔记集合：其中第一个章节：cutlass/CuTe GEMM 中矩阵的存储方式 NT / TN / NN / TT。 本文主要记录一些第三方资料，见末尾附录。1. 定义 block tile 大小配置 CTA 大小为 MNK = 128 * 128 * 32，数据类型为FP16： constexpr auto bM = cute::Int\u0026lt;128 * 2 / sizeof(TA)\u0026gt;{}; constexpr auto bN = cute::Int\u0026lt;128 * 2 / sizeof(TB)\u0026gt;{}; constexpr auto bK = cute::Int\u0026lt;32\u0026gt;{}; constexpr auto cta_tiler = cute::make_shape(bM, bN, bK); // (bM, bN, bK) constexpr auto bP = cute::Int\u0026lt;3\u0026gt;{}; // pipeline1.1. Roofline 计算RTX 3060 Tensor Core FP16 理论峰值为51TFLOPS，内存带宽为 360GB/s。Roofline临界点为：51 * 1000 / 360 = 141.67 FLOPs/Byte。对于分块矩阵计算，loop over k的过程中，包含一次乘法、一次加法。每个 CTA tile 的计算量与 GMEM 搬运量之比（算术强度AI）：\\[\\text{AI}_{tile} = \\frac{2 \\times bM \\times bN \\times bK}{(bM \\times bK + bN \\times bK) \\times sizeof(FP16)} = \\frac{2 \\times 128 \\times 128 \\times 32}{((128 \\times 32) + (128 \\times 32)) \\times 2} = 64 \\text{FLOP/Byte}\\]TFLOPS 为：\\[TFLOPS = \\text{AI} \\times \\text{Bandwidth}_{\\text{GMEM}} = 64 \\times 360 / 1000 = 23.04 \\text{TFLOPS}\\] 明显的，增大bM和bN可以提升算术强度，从而提升性能。如上计算公式没有考虑到L2 Cache，如果考虑到L2 Cache，即CTA之间数据共享，理论计算公式为（以M*N*K=4096*4096*4096为例）：\\[\\text{AI}_{global} = \\frac{2 \\times M \\times N \\times K}{(M \\times b + N \\times K + M \\times N) \\times sizeof(FP16)} = \\frac{2 \\times 4096 \\times 4096 \\times 4096}{(4096 \\times 4096 + 4096 \\times 4096 + 4096 \\times 4096) \\times 2} = 1365 \\text{FLOP/Byte}\\]实测代码如下：double flops = 2.0 * M * N * K;double tflops = flops / (elapsed_ms * 1e-3) / 1e12;printf(\"%.2f TFLOPS\\n\", tflops);实测结果：cuBLAS: 5.24442 ms, 26.2067 TFLOPSCustom: 3.64926 ms, 37.6622 TFLOPS达到理论峰值的37.66 / 51 = 73.84%。2. Tiled MMA 配置TiledMMA 使用SM80_16x8x16_F16F16F16F16_TN，对应 PTX 指令m16n8k16，使用一个warp（32个线程协作）完成子块的MMA计算。打印的MMA Atom配置信息如下：MMA_Atom ThrID: _32:_1 Shape_MNK: (_16,_8,_16) LayoutA_TV: ((_4,_8),(_2,_2,_2)):((_32,_1),(_16,_8,_128)) LayoutB_TV: ((_4,_8),(_2,_2)):((_16,_1),(_8,_64)) LayoutC_TV: ((_4,_8),(_2,_2)):((_32,_1),(_16,_8))在 SMEM -\u0026gt; REG 的过程中，使用ldmatrix拷贝（具体为使用 CopyTraits：SM75_U16x8_LDSM_T）。ldmatrix以8 x FP16=128-bit为单位进行拷贝（可理解为：每个线程指向的SMEM要求8 x FP16连续）。SM75_U16x8_LDSM_T使用ldmatrix.x4指令，使用一个warp（32个线程），一次拷贝四个8x8 FP16=(32, 8)。SM75_U16x8_LDSM_T的打印信息如下：Copy_Atom ThrID: _32:_1 ValLayoutSrc: (_32,_8):(_8,_1) ValLayoutDst: ((_4,_8),(_1,_2,_4)):((_16,_1),(_1,_8,_64)) ValLayoutRef: ((_4,_8),(_1,_2,_4)):((_16,_1),(_1,_8,_64)) ValueType: 16b 参考https://zhuanlan.zhihu.com/p/696231622，有如下表述：“矩阵中连续的两行无需在shared memory中连续，但1行是连续的128-bit。也就是说，ldmatrix读取shared memory的单元是128-bit。”。Tiled MMA受上述SMEM的Tiled Copy约束，要求每个线程处理 8 x FP16数据，这个约束作用于A sub-tile和B sub-tile。A sub-tile已经满足要求，针对B sub-tile，SM80_16x8x16_F16F16F16F16_TN只给每个线程分配四个FP16，因此需要使用permutation参数（即mma_layout）使其满足SMEM TiledCopy要求： using MMATraits = cute::MMA_Traits\u0026lt;cute::SM80_16x8x16_F16F16F16F16_TN\u0026gt;; using MMAAtomShape = MMATraits::Shape_MNK; constexpr auto mma_atom = cute::MMA_Atom\u0026lt;MMATraits\u0026gt;{}; constexpr auto mma_atom_shape = MMAAtomShape{}; // (16, 8, 16) constexpr auto MMA_LAYOUT_M = 2, MMA_LAYOUT_N = 2, MMA_LAYOUT_K = 1; // 线程数扩充 constexpr auto NUM_MMA_TILE_M = 1, NUM_MMA_TILE_N = 2, NUM_MMA_TILE_K = 1; // 每个线程Atom数量扩充 constexpr auto MMA_TILE_M = cute::get\u0026lt;0\u0026gt;(mma_atom_shape) * NUM_MMA_TILE_M * MMA_LAYOUT_M; // 32 constexpr auto MMA_TILE_N = cute::get\u0026lt;1\u0026gt;(mma_atom_shape) * NUM_MMA_TILE_N * MMA_LAYOUT_N; // 32 constexpr auto MMA_TILE_K = cute::get\u0026lt;2\u0026gt;(mma_atom_shape) * NUM_MMA_TILE_K * MMA_LAYOUT_K; // 16 constexpr auto mma_layout = cute::make_layout(cute::make_shape(cute::Int\u0026lt;MMA_LAYOUT_M\u0026gt;{}, cute::Int\u0026lt;MMA_LAYOUT_N\u0026gt;{}, cute::Int\u0026lt;MMA_LAYOUT_K\u0026gt;{})); constexpr auto mma_tile = cute::make_tile(cute::Int\u0026lt;MMA_TILE_M\u0026gt;{}, cute::Int\u0026lt;MMA_TILE_N\u0026gt;{}, cute::Int\u0026lt;MMA_TILE_K\u0026gt;{}); constexpr auto tiled_mma = cute::make_tiled_mma(mma_atom, mma_layout, mma_tile);打印的TiledMMA配置信息如下：TiledMMA ThrLayoutVMNK: (_32,_2,_2,_1):(_1,_32,_64,_0) PermutationMNK: (_32,_32,_16)MMA_Atom ThrID: _32:_1 Shape_MNK: (_16,_8,_16) LayoutA_TV: ((_4,_8),(_2,_2,_2)):((_32,_1),(_16,_8,_128)) LayoutB_TV: ((_4,_8),(_2,_2)):((_16,_1),(_8,_64)) LayoutC_TV: ((_4,_8),(_2,_2)):((_32,_1),(_16,_8)) (2,2,1):(_1,2,4)：描述Atom的线程扩展配置，即在M、N、K三个维度上分别扩展2倍、2倍、1倍线程。 (32,32,16):(_1,32,1024)：描述Tile的大小配置，即解决的MNK规模。2.1. 约束这个定义的TiledMMA针对线程做了配置：在M方向及N方向均使用MMAAtom的两倍线程，在K方向上保持不变。即使用 2x2=4个warp，共128个线程协作完成一个MMA Tile的计算。TiledMMA的配置，对 GMEM =\u0026gt; SMEM拷贝过程中的线程划分形成约束，即在A/B子块的GMEM =\u0026gt; SMEM过程中，配置的线程数量也是 128 个线程。同时，TildMMA的配置，对输入A/B矩阵的tiler也形成约束，即要求分配给CTA的tile大小在M、N、K三个维度上分别是MMA_TILE_M=32、MMA_TILE_N=32、MMA_TILE_K=16的整数倍。A. 资料 cute 之 GEMM流水线 cute 之 高效GEMM实现 CUDA SGEMM优化笔记 从 GEMM 实践 CUDA 优化 https://github.com/NVIDIA/cutlass/blob/main/examples/cute/tutorial/sgemm_sm80.cu：官方源码，包含multi-stage pipeline实现3. Multi-Stage Pipeline3.1. SMEM 资源分配GMEM =\u0026gt; SMEM采用cp.async指令，重叠内存拷贝与GEMM计算以隐藏GMEM访问延迟。其需要的SMEM计算公式为：\\[\\text{SMEM} = \\text{numStages} \\times \\left( bM \\times bK \\times \\text{sizeof}(\\text{ElementA}) + bK \\times bN \\times \\text{sizeof}(\\text{ElementB}) \\right)\\]其中： $bM \\times bK$：加载自矩阵 $A$ 的分块大小 $bK \\times bN$：加载自矩阵 $B$ 的分块大小 $\\text{numStages}$：流水线级数（即 SMEM 中同时维护的分块副本数）配置部分代码如下： constexpr auto smem_shape_A = cute::make_shape(bM, bK, bP); // (bM, bK, bP) constexpr auto smem_shape_B = cute::make_shape(bN, bK, bP); // (bN, bK, bP) constexpr auto smem_layout_A = cute::tile_to_shape(smem_atom_layout_A_swizzled, smem_shape_A); constexpr auto smem_layout_B = cute::tile_to_shape(smem_atom_layout_B_swizzled, smem_shape_B); // GMEM -\u0026gt; SMEM 的 Tiled Copy 配置 constexpr auto gmem_tiled_copy_A = cute::make_tiled_copy(cute::Copy_Atom\u0026lt;CopyOperationA, TA\u0026gt;{}, thread_layout_A, vector_layout_A); constexpr auto gmem_tiled_copy_B = cute::make_tiled_copy(cute::Copy_Atom\u0026lt;CopyOperationB, TB\u0026gt;{}, thread_layout_B, vector_layout_B); // Tiled MMA 配置 // 见上面章节，此处省略...... // Tiled MMA 配置对 GMEM -\u0026gt; SMEM 过程的约束 CUTE_STATIC_ASSERT(cute::size(tiled_mma) == cute::size(thread_shape_C)); CUTE_STATIC_ASSERT(std::is_same_v\u0026lt;TA, cute::half_t\u0026gt; \u0026amp;\u0026amp; std::is_same_v\u0026lt;TB, cute::half_t\u0026gt; \u0026amp;\u0026amp; std::is_same_v\u0026lt;TC, cute::half_t\u0026gt;); // configure tiled copy from smem to register via tiled MMA using Copy_Atom_A = cute::Copy_Atom\u0026lt;cute::SM75_U16x8_LDSM_T, TA\u0026gt;; using Copy_Atom_B = cute::Copy_Atom\u0026lt;cute::SM75_U16x8_LDSM_T, TB\u0026gt;; constexpr auto smem_tiled_copy_A = cute::make_tiled_copy_A(Copy_Atom_A{}, tiled_mma); constexpr auto smem_tiled_copy_B = cute::make_tiled_copy_B(Copy_Atom_B{}, tiled_mma);A.1. 全流程优化参考资料 How to Optimize a CUDA Matmul Kernel for cuBLAS-like Performance: a Worklog CUTLASS Tutorial: Efficient GEMM kernel designs with Pipelining Nvidia Tensor Core-CUDA HGEMM Advanced Optimization：待阅读 Advanced Matrix Multiplication Optimization on NVIDIA GPUs：待阅读 0x_gemm_tutorial.md：cutlass/CuTe仓库教程 CUDA Techniques to Maximize Compute and Instruction Throughput(pdf) Developing CUDA Kernels to Push Tensor Cores to the Absolute Limit on NVIDIA A100 (pdf) Performance Analysis of CUDA-based General Matrix Multiplication through Memory Coalescing and Grid-Level Parallelization CUTLASS MMA Pipelined Header https://github.com/NVIDIA/cutlass/blob/main/test/unit/gemm/device/default_gemm_configuration.hpp CuTe– CUDA Tensors：Layout \u0026amp; Tensor基础A.2. Learn CUTLASS the hard way Learn CUTLASS the hard way! Learn CUTLASS the hard way - part 2!A.3. Triton Triton Linear Layout: Concept" },

  { "title": "GEMM 版本1：使用 CuTe 实现一个 naive GEMM", "url": "/blog/2025/GEMM1-Cute-naive-GEMM/", "categories": "CUDA",
  "tags": "CUDA", "date": "2025-02-26", "content":
  "1. navie tile GEMM 代码文件：navie tile GEMM基于分块矩阵乘法的简单实现，按照 Thread Block 将矩阵划分为多个tile进行计算，在 Thread Block内，再次将 tile 划分为多个子块，由每个线程负责计算子块。使用 Shared Memory 来缓存 tile 数据，减少全局内存访问次数。每个线程负责从全局内存中复制 tile 内的一小块内存到 Shared Memory。 循环展开#pragma unroll优化加载和计算部分的循环，提高指令级并行性（消耗更多寄存器资源），本代码测试整体运行时间提升15%左右。2. CuTe 版本 naive tile GEMM 代码文件：CuTe naive tile GEMM使用 CuTe 库重写的分块矩阵乘法，使用 slice-k 方法，即分块（tile）沿着 K 维度累加所有结果子矩阵。 使用 NVIDIA CuTe 库重写的分块矩阵乘法实现，采用 cute::gemm 期望的标准布局。2.1. 矩阵布局约定采用 BLAS/Fortran 风格的列主序 (Column-major)： Tensor Shape Stride 说明 A (M, K) (1, M) 列主序，M方向连续 B (N, K) (1, N) 列主序，存储 B^T C (M, N) (1, M) 列主序，M方向连续 矩阵参数，以及划分参数：M=1024，N=1024，K=1024*8，BM=64，BN=64，BK=16，TM=8，TN=8。关键点：B 矩阵以 (N, K) 形式存储，实际上是原始 B(K, N) 的转置。这是 cute::gemm 的标准输入格式。创建的矩阵 A、B、C 的 tensor 视图如下：// A: (M, K) 列主序, stride = (1, M)// B: (N, K) 列主序, stride = (1, N) -- 注意这里存的是B的转置// C: (M, N) 列主序, stride = (1, M)Tensor mA = make_tensor(make_gmem_ptr(Aptr), make_shape(M, K), make_stride(Int\u0026lt;1\u0026gt;{}, M));Tensor mB = make_tensor(make_gmem_ptr(Bptr), make_shape(N, K), make_stride(Int\u0026lt;1\u0026gt;{}, N));Tensor mC = make_tensor(make_gmem_ptr(Cptr), make_shape(M, N), make_stride(Int\u0026lt;1\u0026gt;{}, M)); kernel 的 dimensions 配置为：dim3 gridDim(N / BN, M / BM); // (16, 16)dim3 blockDim(BN / TN, BM / TM);// (8, 8)2.2. CuTe 实现：分块分割、线程分区、拷贝及计算以矩阵 A 为例，矩阵 A 在 M 纬度上，每个 Thread Block 负责处理 BM 行（复制 + GEMM）；在 K 维度上，Thread Block 负责处理 BK 列。Thread Block 以二维的方式划分，需要的 Thread Block 数量为 (M/BM, N/BN)，即划分的 tile 数量；每个 Thread Block 内的线程数量为 (BM/TM, BN/TN)，即同样以二维的方式将 tile 再次划分给 Thread Block 内的线程。使用 slice-k 方法，Thread Block 需要循环遍历 K 维度，将矩阵 A 分块加载到 SMEM 中。所以创建的 tile tensor 视图是一个三维 tensor： (BM, BK, K/BK) 大小。矩阵 B 同理，得到的三维 tensor：(BN, BK, K/BK)。由于分块计算得到的 C 矩阵 tile 大小为 (BM, BN)，所以需要取得矩阵 C 的 tile tensor。 矩阵 A、B 分块 GEMM 需要的 SMEM。其本质是缓存，大小为一个 tile 大小。即针对矩阵 A，Thread Block 每次从 GMEM 中依次取得 (BM, BK) 分块大小复制到 SMEM，然后沿着 K 维度循环 for(product)。定义如下：__shared__ T smemA[BM * BK];__shared__ T smemB[BN * BK];// sA: (BM, BK) 列主序, stride = (1, BM)// sB: (BN, BK) 列主序, stride = (1, BN)Tensor sA = make_tensor(make_smem_ptr(smemA), make_shape(Int\u0026lt;BM\u0026gt;{}, Int\u0026lt;BK\u0026gt;{}), make_stride(Int\u0026lt;1\u0026gt;{}, Int\u0026lt;BM\u0026gt;{}));Tensor sB = make_tensor(make_smem_ptr(smemB), make_shape(Int\u0026lt;BN\u0026gt;{}, Int\u0026lt;BK\u0026gt;{}), make_stride(Int\u0026lt;1\u0026gt;{}, Int\u0026lt;BN\u0026gt;{}));2.2.1. 分块操作// 创建全局内存 tensor 视图Tensor gA = local_tile(mA, make_tile(Int\u0026lt;BM\u0026gt;{}, Int\u0026lt;BK\u0026gt;{}), make_coord(blockIdx.y, _)); // (BM, BK, k)Tensor gB = local_tile(mB, make_tile(Int\u0026lt;BN\u0026gt;{}, Int\u0026lt;BK\u0026gt;{}), make_coord(blockIdx.x, _)); // (BN, BK, k)Tensor gC = local_tile(mC, make_tile(Int\u0026lt;BM\u0026gt;{}, Int\u0026lt;BN\u0026gt;{}), make_coord(blockIdx.y, blockIdx.x)); // (BM, BN)按照语义理解，A 矩阵使用 shape(BM, BK) 沿着 M 维度以及 K 维度进行切分，即将其划分为若干 tile，每个 tile 的大小为 (BM, BK)。而后，使用 make_coord(blockIdx.y, _) 取得第 blockIdx.y 个 tile，在 K 维度上使用 cute::_ 表示取所有 tile，所以生成了一个三维的 Tensor。由于 A 矩阵在 K 维度上被切分为多个 tile，所以最终生成的 tensor 维度为 (BM, BK, K/BK)，即表示有 K/BK 个二维 tile(BM, BK)。作为对比，矩阵 C 在 M 维度以及 N 维度上进行切分，得到若干个 tile，每个 tile 的大小为 (BM, BN)。而后使用完整的二维坐标 make_coord(blockIdx.y, blockIdx.x) 取得唯一的 tile，故其生成的 tensor 维度为 (BM, BN)。分块之后，每个 Thread Block 分到的 tile shape 如下： gA(64, 16, 512)：其中，1024*8 / 16 = 512，即这是一个 tile group gB(64, 16, 512)：其中，1024*8 / 16 = 512，即这是一个 tile group gC(64, 64)2.2.2. 线程划分：将 tile 划分给 Thread Block 内的线程分区复制 GMEM -\u0026gt; SMEM：前面已经按照 Thread Block 分块得到 tile，现在需要继续将划分细化到 Thread Block 内的每个线程。首先计算线程数量：// 线程块配置: (BN/TN, BM/TM) = (8, 8) = 64 线程constexpr int num_threads = (BM / TM) * (BN / TN);// 线程划分方法：按照 M 维度划分 A 矩阵，N 维度划分 B 矩阵Layout tA_copy = make_layout(make_shape(Int\u0026lt;num_threads\u0026gt;{}, Int\u0026lt;1\u0026gt;{}));Layout tB_copy = make_layout(make_shape(Int\u0026lt;num_threads\u0026gt;{}, Int\u0026lt;1\u0026gt;{}));// 得到本线程划分的 sub-tileTensor tAgA = local_partition(gA, tA_copy, tid); // 每个线程负责的gA部分Tensor tAsA = local_partition(sA, tA_copy, tid); // 每个线程负责的sA部分Tensor tBgB = local_partition(gB, tB_copy, tid); // 每个线程负责的gB部分Tensor tBsB = local_partition(sB, tB_copy, tid); // 每个线程负责的sB部分针对 tile A，只在 M 维度上进行线程划分，每个线程负责复制 BM / num_threads 行数据，在 K 维度上负责复制全部 K 列数据，维度信息保持，tAgA、tBgB 还是一个三维 tensor。每个线程在执行复制时，在 K 维度上每次复制 BK 列，一共需要 K / BK 次循环，才能完成 sub-tile 的加载以及 GEMM 得到最终的结果。矩阵 B 同理，只在 N 维度上进行线程划分。得到的线程分区 tensor layout 如下： tAgA(1, 16)：shape: (_1,_16,512), stride: (_0,1024,16384) tAsA(1, 16)：shape: (_1,_16), stride: (_0,64) tBgB(1, 16)：shape: (_1,_16,512), stride: (_0,1024,16384) tBsB(1, 16)：shape: (_1,_16), stride: (_0,64)计算分区：每个线程复制划分，是针对矩阵 A、B 进行的。计算 C 矩阵时，需要取得 A(TM, TK) 分块、B(TN, TK) 分块、C(TM, TN)，进行线程分块的 GEMM 计算。复制时的线程分块与计算时的线程分块可以分开，因为在 GMEM -\u0026gt; SMEM 复制之后，使用 __syncthreads 保证 Thread Block 的 A-tile、B-tile 都完成复制（更严谨的说，是在 GMEM -\u0026gt; SMEM 与 sub-tile GEMM 两个步骤之间同步）。// 计算时的线程划分方法: (BM/TM, BN/TN) = (8, 8)Layout tC = make_layout(make_shape(Int\u0026lt;BM / TM\u0026gt;{}, Int\u0026lt;BN / TN\u0026gt;{}));// local_partition: 按线程布局分配工作// Step\u0026lt;_1, X\u0026gt; 表示第 0 维参与分区，第 1 维不参与Tensor tCsA = local_partition(sA, tC, tid, Step\u0026lt;_1, X\u0026gt;{}); // (TM, BK)Tensor tCsB = local_partition(sB, tC, tid, Step\u0026lt;X, _1\u0026gt;{}); // (TN, BK)Tensor tCgC = local_partition(gC, tC, tid, Step\u0026lt;_1, _1\u0026gt;{}); // (TM, TN)线程分区之后，得到的每个线程的 tensor layout 如下：tCsA shape: (_8,_16), stride: (_8,_64)tCsB shape: (_8,_16), stride: (_8,_64)tCgC shape: (_8,_8), stride: (_8,8192)2.2.3. slice-k GEMM// 遍历K维度const int num_tile_k = K / BK;for (int k = 0; k \u0026lt; num_tile_k; k++) { // 从全局内存复制到共享内存 copy(tAgA(_, _, k), tAsA); copy(tBgB(_, _, k), tBsB); __syncthreads(); // 等待所有线程完成复制 // 使用 cute::gemm 执行矩阵乘法 // gemm 期望: A(M,K), B(N,K), C(M,N) - B是(N,K)形式 gemm(tCsA, tCsB, tCrC); // tCrC += tCsA * tCsB^T __syncthreads(); // 等待所有线程完成计算}在 tile 分区的时候，已经将 A-tile、B-tile 划分为 K/BK 个子 tile，即沿着 K 维度进行了分块。计算时，针对每一个 sub-tile 进行 GEMM，一共迭代 K / BK 次。 当矩阵是列主序时（比如矩阵 A、B 是 K-major），如果 Thread Block 内的线程任务划分也按照 K-major 进行，这样得到的访问矩阵内的元素的编号也是连续的，即访存合并。2.2.4. Bank Conflict 计算依据 outter-partition 划分的方式：Layout tC = make_layout(make_shape(Int\u0026lt;BM / TM\u0026gt;{}, Int\u0026lt;BN / TN\u0026gt;{}));Tensor tCsA = local_partition(sA, tC, tid, Step\u0026lt;_1, X\u0026gt;{}); // (TM, BK)Tensor tCsB = local_partition(sB, tC, tid, Step\u0026lt;X, _1\u0026gt;{}); // (TN, BK)线程在 tC layout 中的二维坐标：tC_row = tid % 8 (M 维度方向)tC_col = tid / 8 (N 维度方向)2.2.4.1. tCsA 访问方式及 bank conflict 分析Step\u0026lt;_1, X\u0026gt; 表示使用 M 维度参与分区。线程的 tCsA 起始行号为：\\[\\text{row_offset}_{A} = \\text{tC_row} \\times \\text{TM} = (\\text{tid} \\% 8) \\times 8\\]线程访问 sA 的地址计算公式：\\[\\begin{aligned}\u0026amp; \\text{addr}_{sA}[m, k] = (\\text{row_offset}_{A} + m) + k \\times 64 \\\\\u0026amp; m \\in [0, 7], k \\in [0, 15]\\end{aligned}\\]stride=8，可以理解为每个线程占据 8 个float 类型数据，则 4 个线程之后即产生 bank conflict。此时，每隔 4 个线程，访问的地址会落在同一个 bank 上，并产生 bank conflict。一个 warp 内 （tid - bankId） 对应表如下：[0 - 0, 1 - 8, 2 - 16, 3 - 24][4 - 0, 5 - 8, 6 - 16, 7 - 24][8 - 0, 9 - 8, 10 - 16, 11 - 24][12 - 0, 13 - 8, 14 - 16, 15 - 24] 其中，由于 m = (m+1) % 8，第 8 个线程地址跳跃 64 个 float。即一个 warp 产生 4 个 4-way bank conflict。2.2.4.2. tCsB 访问方式及 bank conflict 分析Step\u0026lt;X, _1\u0026gt; 表示使用 N 维度参与分区。每个线程的 tCsB 起始行号为：\\[\\text{row_offset}_{B} = \\text{tC_col} * \\text{TN} = (\\text{tid} / 8) \\times 8\\]线程访问 sB 的地址计算公式：\\[\\begin{aligned}\u0026amp; \\text{addr}_{sB}[n, k] = (\\text{row_offset}_{B} + n) + k \\times 64 \\\\\u0026amp; n \\in [0, 7], k \\in [0, 15]\\end{aligned}\\]此时，每 8 个线程一组，其访问地址都是同一个地址。下一组 8 个线程，其在smemB[]中的编号 +8。即每相邻的 8 个线程，产生一个 broadcast。对应表格如下：[0 - 0, 1 - 0, 2 - 0, 3 - 0, 4 - 0, 5 - 0, 6 - 0, 7 - 0][8 - 8, 9 - 8, 10 -8, 11 -8, 12 -8, 13 -8, 14 -8, 15 -8] 一个 warp 内，从第 8 个线程开始，访问编号跳转了 64 个 float。2.2.4.3. 总结发现，在当前情况下（sA 与 sB 布局相同，且划分大小相同），他们之前不一样的地方，来自于划分时，选择的维度不同： 首先，线程被划分为两个维度，且使用这两个维度分别去划分 sA 与 sB。 其次，由于使用了这两个维度进行划分，导致 M 维度是使用取余，N 维度是使用整除。这才是导致访问模式不同的根本原因。性能影响： tCsA 访问产生 bank conflict，影响性能。 tCsB 访问产生 broadcast，带宽利用率低。3. Stride 理解make_stride(s0, s1) 定义了沿各维度移动时的内存跳跃距离： stride(1, M) → 第0维步长=1（连续），第1维步长=M → 列主序 stride(M, 1) → 第0维步长=M，第1维步长=1（连续） → 行主序 简单记忆：Stride 为 1 的维度在内存中连续。4. CuTe 命名约定CuTe 采用三段式命名规则：t[分区者][存储空间][矩阵]：t C s A│ │ │ └─ 矩阵：A / B / C│ │ └─── 存储空间：g(global) / s(smem) / r(register)│ └───── 分区者（Partitioner）：A / B / C└─────── 前缀：t = thread-partitioned tensor4.1. 前缀 t表示该 Tensor 已经过线程分区（thread partitioning），即通过 localpartition 或 TiledMMA::partition* 切分后，当前线程所”看到”的子视图。4.2. 分区者标识：A / B / C这是命名中最容易混淆的部分。它不是指矩阵本身，而是用于分区的线程 Layout 对象： 标识 对应线程 Layout 作用 tA AThreadLayout tA 用于 Copy 阶段对 A 矩阵的线程分区 tB BThreadLayout tB 用于 Copy 阶段对 B 矩阵的线程分区 tC CThreadLayout tC / TiledMMA 用于 Math（MMA）阶段对 C/A/B 的线程分区 tC 的本质：C 的线程覆盖（overlay）布局，即 MMA 计算输出的线程分布模式。官方 Issue 原话：\"tC refers to the layout of the overlay threads of MMA.\"4.3. 存储空间标识：g / s / r 前缀 英文 含义 m matrix 完整矩阵的 tensor 视图 g global 全局内存中的 tile s shared 共享内存中的 tensor r register 寄存器中的 tensor 4.4. 矩阵标识：A / B / C指被分区的目标矩阵。4.5. 示例以 官方示例 https://github.com/NVIDIA/cutlass/blob/main/examples/cute/tutorial/sgemm_1.cu 为例，Copy 阶段与 Math 阶段对 sA 使用了两套不同的分区方式：// Copy 阶段：用 tA（32x8 线程布局）对 sA 做分区Tensor tAsA = local_partition(sA, tA, threadIdx.x); // (THR_M, THR_K)// ↑↑↑// tA 分区者 + s共享内存 + A矩阵// Math 阶段：用 tC（16x16 线程布局）对 sA 做分区Tensor tCsA = local_partition(sA, tC, threadIdx.x, Step\u0026lt;_1, X\u0026gt;{}); // (THR_M, BLK_K)// ↑↑↑// tC 分区者 + s共享内存 + A矩阵 💡 关键点：同一块 sA 共享内存，被 tA 分区用于 cute::copy，又被 tC 分区用于 cute::gemm。通过前缀就能从变量名直接识别”用哪套线程布局在操作”。 cutlass issue 中，命名提问：https://github.com/NVIDIA/cutlass/issues/1052#issuecomment-1680718745 官方文档有相关讲解：https://github.com/NVIDIA/cutlass/blob/main/media/docs/cpp/cute/0x_gemm_tutorial.md#copy-partitioning5. Q/A 为什么同一个矩阵有不同的分区？ 因为复制和计算时的线程分工不同。例如 sA：复制时：64个线程平均分配 BM×BK 元素 → tAsA计算时：每个线程取 TM×BK 子矩阵 → tCsA6. 补充知识：partitioning 方式CuTe 中，总共有三种 partitioning 方式：inner-partition、outer-partition 和 TV-layout-partition。其中local_tile和local_partition，分别用于不同层次的划分：   Inner-Partition Outer-Partition 接口 local_tile local_partition 用途 CTA 级别的 tile 划分 Thread 级别的数据分配 结果形状 (tile形状, rest) (thread形状, rest) “谁”是主角 tile 内部坐标 每个 thread 拥有的坐标 6.1. local_tile对于如下代码： // Define CTA tile sizes (static) auto cta_tiler = make_shape(bM, bN, bK); // (BLK_M, BLK_N, BLK_K) // Get the appropriate blocks for this threadblock auto cta_coord = make_coord(blockIdx.x, blockIdx.y, _); // (m,n,k) Tensor gA = local_tile(mA, cta_tiler, cta_coord, Step\u0026lt;_1, X,_1\u0026gt;{}); // (BLK_M,BLK_K,k) Tensor gB = local_tile(mB, cta_tiler, cta_coord, Step\u0026lt; X,_1,_1\u0026gt;{}); // (BLK_N,BLK_K,k) Tensor gC = local_tile(mC, cta_tiler, cta_coord, Step\u0026lt;_1,_1, X\u0026gt;{}); // (BLK_M,BLK_N)其中，A是在M方向切分得到tile，K方向上得到一个slice集合。B是按照N方向切分得到tile，K方向上得到一个slice集合。使用Step\u0026lt;_1, X, _1\u0026gt;表示选取A的M维度及K维度，等价于： // Use select\u0026lt;0,2\u0026gt; to use only the M- and K-modes of the tiler and coord Tensor gA = local_tile(mA, select\u0026lt;0,2\u0026gt;(cta_tiler), select\u0026lt;0,2\u0026gt;(cta_coord));local_tile内部拆分成两个步骤。首先使用tiler对输入mA进行mode拆分： // ((BLK_M,BLK_K),(m,k)) Tensor gA_mk = zipped_divide(mA, select\u0026lt;0,2\u0026gt;(cta_tiler));拆分的结果是，保留了tiler mode，得到inner mode (BLK_M, BLK_K)，以及rest mode (m, k)。接着使用coord从rest mode中选择对应的切片： // (BLK_M,BLK_K,k) Tensor gA = gA_mk(make_coord(_,_), select\u0026lt;0,2\u0026gt;(cta_coord)); 总结为两个步骤：inner-partition + slice selection。 zipped_divide底层使用的是composition。 官方相关文档： 0x_gemm_tutorial:cta-partitioning 03_tensor:inner-and-outer-partitioning 02_layout_algebra:composition-tilers资料 Matrix Multiplication Background User’s Guide。如何计算 GEMM 的性能指标 CuTe dense matrix-matrix multiply tutorial。CuTe GEMM 官方 Document。 cute 之 简单GEMM实现：reed 知乎文章 cute 之 MMA抽象：reed 知乎文章 CuTe Local Partition：Mao Lei博客 CUDA Matrix Multiplication Optimization：Mao Lei博客，GEMM优化步骤全解析 Colfax Research Cute Tutorial：Colfax Research Cute Tutorial参考代码 sgemm_sm80 官方示例：使用 CuTe 实现的 naive tile GEMM Awesome-CUDA-and-HPC Awesome CUDA 0x_gemm_tutorial 官方文档：CuTe GEMM 教程" },

  { "title": "CUTLASS-Cute 初步(6)：CUDA GEMM 计算优化、Multi-Stage 与软流水(Software Pipelining)", "url": "/blog/2025/Cute%E5%88%9D%E6%AD%A56-Pipeline-And-Multi-Stage/", "categories": "CUDA",
  "tags": "CUDA", "date": "2025-02-26", "content":
  "整个GEMM可用如下公式表示：\\[C[i, j] = \\sum_{k=0}^{nK-1} A[i, k] \\times B[k, j]\\] 层级 说明 Thread Block Tile 每个 CUDA 线程块（thread block）负责计算输出矩阵 C 的一个子块（tile） Warp Tile 在线程块内部，每个 warp（32个线程）负责计算 thread block tile 的一个子区域 Thread Tile 在 warp 内部，每个线程负责计算 warp tile 的一个更小的子区域 1. GEMM 计算步骤–分层 GEMM 结构依照硬件架构层次划分（也即 CUDA 编程模型），GEMM 计算可以分为多个层次：Thread Block Tile -\u0026gt; Warp Tile -\u0026gt; Thread Tile。即将一个大矩阵的算术运算，依次分解，直到最小的线程级别，一个线程计算一小部分的 tile。数据搬运过程分为几步：GMEM -\u0026gt; Shared Memory -\u0026gt; Register File -\u0026gt; CUDA Core。如下图所示： 完整的 GEMM 分层结构把数据从较慢的存储器搬运到较快的存储器，并在许多算术运算中对其进行重复利用，以提高计算强度。2. Thread Block Tile每个 Thread Block 负责计算输出矩阵 C 的一个 Tile（即几个分块）。对于几个矩阵计算C += A * B，每个 Thread Block 需要反复从输入矩阵中加载一个个的 Tile 并计算一个累加的矩阵乘积。如下图所示： 把一个 GEMM 问题分解为由单个线程块执行的计算。绿色所示的 C 的子矩阵，是由 A 的一个 tile 与 B 的一个子矩阵做矩阵乘积得到的。为此，沿 K 维（已被切分成多个 tile）循环，并把每个 tile 的矩阵乘积结果累加起来完成计算。在逻辑上，Thread Block Tile 又被分为若干个 Warp Tile，每个 Warp Tile 由一个 Warp 负责计算。线程块的输出 tile按空间被划分给各个 warp（如下图所示）。我们把用于存放这个输出 tile 的存储称为累加器（accumulators），因为它保存的是累加起来的矩阵乘积结果。每进行一次算术运算就会更新一次累加器，因此它必须驻留在 SM 中最快的存储器里：寄存器文件。 线程块的结构把 C 的一个 tile 划分给多个 warp，每个 warp 负责存储一个互不重叠的二维子块。每个 warp 都将其累加器元素存放在寄存器中。矩阵 A 和 B 的 tile 则存放在共享内存中，线程块内所有 warp 都可访问。3. Warp Tile当数据已存入共享内存后，每个 warp 通过在其线程块 tile 的 K 维上迭代、从共享内存中加载子矩阵（或称 fragment），并计算累加的外积。如下图所示： 单个 warp 通过在循环中把 A 和 B 的片段（fragments）从各自对应的共享内存（SMEM）tile 加载到寄存器（RF），并计算外积，从而形成累加的矩阵乘积。 上图也展示了多个 warp 之间如何从共享内存进行数据共享：同一线程块行（row）中的 warp 会加载相同的 A 片段；同一线程块列（column）中的 warp 会加载相同的 B 片段。4. Thread TileCUDA 编程模型是以线程块和单个线程来定义的。因此，warp 的结构需要映射到各个线程实际执行的操作上。由于线程之间不能互相访问寄存器，我们必须选取一种组织方式，使得保存在寄存器中的值能够被同一线程执行的多条算术指令反复复用。由此，在单个线程内部形成了一个二维分块（2D tiled）的结构，如下图的细节所示。每个线程向 CUDA 核心发出一串彼此独立的算术指令，并计算一个累加的外积。 单个线程（右）通过对寄存器中保存的 A 的片段（fragment）与 B 的片段（fragment）做外积，来参与 warp 级的矩阵乘法（左）。用绿色标示的 warp 累加器被分配给该 warp 内的各个线程，通常组织成一组二维小 tile。在上图中，warp 的左上象限以灰色标示。那 32 个小格分别对应一个 warp 内的 32 条线程。这种安排会使同一行中的多条线程各自去取 A 片段中的相同元素，而同一列中的多条线程各自去取 B 片段中的相同元素。为最大化计算强度，可以复制这一基础结构来构成完整的 warp 级累加器 tile，从而得到一个由 8×1 与 1×8 片段做外积得到的 8×8 的整体“线程 tile”。这在图中用绿色显示的四个累加器 tile予以说明。5. Multi-Stage 与占用率（Occupancy）Tiled GEMM会大量使用RF来存放A fragment/B fragment以及C accumulator，同时也需要分配较大的SMEM。对片上存储的相对高需求会限制占用率（occupancy），也就是单个SM上能并发运行的线程块数量上限。因此，GEMM的实现通常在每个SM中能容纳的warp与线程块数量远少于典型的GPU计算工作负载。GEMM为了把C子块常驻寄存器、A/B子块常驻共享内存，会吃掉大量片上资源导致占用率下降，用多开更多线程块来掩盖延迟（overlap，latency hiding）的常规手段不再有效。于是改为在同一线程块内重叠多个阶段：边算当前K-tile，边预取下一个K-tile，尽量让计算单元与带宽都忙起来。 在 CUTLASS 的 GEMM 主循环中交错执行的三条并发指令流。黑色箭头表示数据依赖关系（Math 依赖 S -\u0026gt; R）。当内存系统在从全局内存加载数据、且 SM 正在为下一轮线程 tile 加载片段时，线程通过为当前 tile 执行算术指令来让 SM 持续忙碌。在实践中，CUDA程序员通过在源码中交错编写各阶段的CUDA语句，来实现这些管线阶段之间的指令级并发，并依赖 CUDA 编译器在生成的机器码里安排合适的指令调度。广泛使用#pragma unroll与编译期常量可以让编译器完全展开循环并把数组元素映射到寄存器，这两点对实现可调优且高效的内核至关重要。可隐藏的实际时延量，取决于 线程块 / warp / 线程 三级tile的大小，以及SM内活动数学功能单元（FMA/WMMA 等）的吞吐。更大的tile通常带来更多数据复用与时延隐藏机会。5.1. Ampere 架构下的 Multi-Stage在Hopper架构中（sm90），可以基于warp-specialization对寄存器进行warp-specific分配，在此基础上在warp之间创建生产-消费模型（流水线），即Multi-Stage。在Ampere中，没有Hopper上的这些高级特性，只能将生产-消费结构的编码逻辑在同一线程之内实现，并使用cp.async实现latency-hiding。另外，由于使用软流水，导致A/B需要在原有基础上进一步分块，进而需要增加一级循环，共需要两层循环：一层用于loop-over BK，另外一层用于循环(CPY, CPY_M, CPY_K)中的CPY_K。5.2. Ampere 架构 Multi-Stage 编程中 A/B 的分块作为对比，使用单一TiledCopy进行分块搬运前，将tensor A分块为(BM, BK)，并使用ThrCopy获取其(CPY, CPY_M, CPY_K, k)，并使用loop-over BK沿着K方向逐次搬运并计算。流程代码如下： Tensor sA = make_tensor(make_smem_ptr(smemA), make_shape(Int\u0026lt;BM\u0026gt;{}, Int\u0026lt;BK\u0026gt;{}), make_stride(Int\u0026lt;1\u0026gt;{}, Int\u0026lt;BM\u0026gt;{})); Tensor sB = make_tensor(make_smem_ptr(smemB), make_shape(Int\u0026lt;BN\u0026gt;{}, Int\u0026lt;BK\u0026gt;{}), make_stride(Int\u0026lt;1\u0026gt;{}, Int\u0026lt;BN\u0026gt;{})); auto thr_copy_A = tiled_copy_A.get_thread_slice(tid); Tensor tAgA = thr_copy_A.partition_S(gA); // 源：全局内存 (CPY, CPY_M, CPY_K, k) Tensor tAsA = thr_copy_A.partition_D(sA); // 目标：共享内存 (CPY, CPY_M, CPY_K) auto thr_copy_B = tiled_copy_B.get_thread_slice(tid); Tensor tBgB = thr_copy_B.partition_S(gB); // 源：全局内存 Tensor tBsB = thr_copy_B.partition_D(sB); // 目标：共享内存 const int num_tile_k = K / BK; for (int k = 0; k \u0026lt; num_tile_k; k++) { copy(tiled_copy_A, tAgA(_, _, _, k), tAsA); copy(tiled_copy_B, tBgB(_, _, _, k), tBsB); }在Multi-Stage架构+软流水架构下，由于需要分离数据搬运与数据计算，导致在使用Multi-Stage进行GMEM =\u0026gt; SMEM搬运过程，以及SMEM =\u0026gt; REG搬运过程中，有两点不同。首先，在GMEM =\u0026gt; SMEM搬运过程中，需要为SMEM创建nStage个缓存，即缓存是以前的nStage倍。代码如下： // 配置代码 constexpr auto smem_shape_A = cute::make_shape(bM, bK, bP); // (bM, bK, bP) constexpr auto smem_shape_B = cute::make_shape(bN, bK, bP); // (bN, bK, bP) constexpr auto smem_layout_A = cute::tile_to_shape(smem_atom_layout_A_swizzled, smem_shape_A); constexpr auto smem_layout_B = cute::tile_to_shape(smem_atom_layout_B_swizzled, smem_shape_B); // kernel // partition gmem -\u0026gt; smem via gmem tiled copy auto gmem_thr_tiled_copy_A{gmem_tiled_copy_A.get_slice(threadIdx.x)}; auto gmem_tAgA = gmem_thr_tiled_copy_A.partition_S(gmem_block_tensor_A); // (CPY, CPY_M, CPY_K, k) auto gmem_tAsA = gmem_thr_tiled_copy_A.partition_D(smem_tensor_A); // (CPY, CPY_M, CPY_K, bP) auto gmem_thr_tiled_copy_B{gmem_tiled_copy_B.get_slice(threadIdx.x)}; auto gmem_tAgB = gmem_thr_tiled_copy_B.partition_S(gmem_block_tensor_B); // (CPY, CPY_N, CPY_K, k) auto gmem_tAsB = gmem_thr_tiled_copy_B.partition_D(smem_tensor_B); // (CPY, CPY_N, CPY_K, bP)可以看到，最终线程的gmem fragment跟以前一样，但是smem fragment layout多出一个轴bP。其次，在使用软流水优化SMEM =\u0026gt; REG搬运的过程中，以A为例，每次搬运的数据由(CPY, CPY_M, CPY_K)改为(CPY, CPY_M)，同样出于分离数据搬运/数据计算的需要；同时，计算量也由(MMA, MMA_M, MMA_K)改为(MMA, MMA_M)，即取原来的一个k'-slice。5.3. Ampere 架构实现 Multi-Stage 的流程在启动Multi-Stage的主循环之前，需要先对GMEM =\u0026gt; SMEM进行预取操作一次。代码实现如下： // prologue: 预填充 SMEM [0, PIPELINE-2]， 共(PIPELINE-1)个 tile 到 smem。 // 保证主循环有 pipeline slot 进行预加载，并保证 pipeline slot 不溢出， // 另外还需要 wait 能保证最新发出的 copy 仍在飞 for (auto pipeline_idx = 0; pipeline_idx \u0026lt; (NUM_SMEM_PIPELINE - 1); ++pipeline_idx) { cute::copy(gmem_tiled_copy_A, gmem_tAgA(cute::_, cute::_, cute::_, tile_idx_next), // gmem_tAsA(cute::_, cute::_, cute::_, pipeline_idx)); cute::copy(gmem_tiled_copy_B, gmem_tAgB(cute::_, cute::_, cute::_, tile_idx_next), // gmem_tAsB(cute::_, cute::_, cute::_, pipeline_idx)); cute::cp_async_fence(); --num_tiles_remain; if (num_tiles_remain \u0026gt; 0) { ++tile_idx_next; } } cute::cp_async_wait\u0026lt;NUM_SMEM_PIPELINE - 2\u0026gt;(); __syncthreads(); // prefetch first k' slice from SMEM =\u0026gt; REG // ...... prologue 总结： cp_async_fence()表示创建一个commit groups，这里一个commit groups包含一次A fragment的GMEM =\u0026gt; SMEM操作，以及一次B fragment的GMEM =\u0026gt; SMEM操作。 cp_async_wait\u0026lt;N\u0026gt;()表示等待直到有N个commit groups在飞（in flight）。语句cute::cp_async_wait\u0026lt;NUM_SMEM_PIPELINE - 2\u0026gt;()表示等待最先一批A fragment/B fragment完成。 首次预取循环代码中pipeline_idx \u0026lt; (NUM_SMEM_PIPELINE - 1)之所以循环次数取为NUM_SMEM_PIPELINE - 1，是因为如果预取所有的SMEM buffer，将导致主循环中再也写不了预取代码，即代码的逻辑结构需要。另外，如果取为NUM_SMEM_PIPELINE - 2，将导致cp_async_wait变为cute::cp_async_wait\u0026lt;NUM_SMEM_PIPELINE - 3\u0026gt;()，即当NUM_SMEM_PIPELINE为3的时候，需要等待所有拷贝操作完成，即Multi-Stage失效，无法并行数据搬运与数据计算。 NUM_SMEM_PIPELINE最小为3。 在prologue预取GMEM =\u0026gt; SMEM之后且预取SMEM =\u0026gt; REG之前，可以进行一些其他初始化工作，比如创建TiledMMA、初始化C fragment累加器，应该可以略微压缩一些时间。如下为Multi-Stage主循环的总结。在初次进入主循环的时候，loop-over BK的第一个分块已经在SMEM中就绪，剩余的NUM_SMEM_PIPELINE-2批数据在飞。要模拟实现数据搬运与数据计算的分离，就要保证每个循环中，发起的数据搬运领先数据计算一个位置，loop-over BK与k'循环都需要按照这个时序。在主循环中，需要继续安排这几个逻辑： 在内层k'循环的开始位置（k'==0），issue一次GMEM =\u0026gt; SMEM异步搬运，即保持流水逻辑（生产端）； 在内层k'循环的结束位置（k'==NUM_MMA_K_LOOP-1），等待k+1位置的fragment准备好，即到第k+1个循环的时候需要的计算数据。由于是异步搬运，所以需要两个指针分别标示： GMEM =\u0026gt; SMEM操作时smem(k+1 fragment)的位置：smem_pipeline_g2s_idx（对应生产端数据）； SMEM =\u0026gt; REG操作时smem(k fragment)的位置：smem_pipeline_s2r_idx（对应消费端数据）； 生产端fragment的位置始终领先消费端fragment一个位置。另外，在prologue之后，生产端的位置已经到了NUM_SMEM_PIPELINE-1，即在主循环中，从该位置开始issue下一个fragment。 int smem_pipeline_s2r_idx = 0; // 指示当前 SMEM -\u0026gt; REG pipeline index int smem_pipeline_g2s_idx = (NUM_SMEM_PIPELINE - 1); // 指示当前 GMEM -\u0026gt; SMEM pipeline index主循环代码如下： while (num_tiles_remain \u0026gt; -(NUM_SMEM_PIPELINE - 1)) { // loop-over K for (int mma_idx_k = 0; mma_idx_k \u0026lt; NUM_MMA_K_LOOP; ++mma_idx_k) { // loop-over BK within the tile // 计算下一个 tile 之前，wait 等待确保下一个 tile 已经 GMEM -\u0026gt; SMEM 完成 // 且 current_pipeline 指向下一个 tile 的数据 if (mma_idx_k == (NUM_MMA_K_LOOP - 1)) { smem_tCsA_current_pipeline = smem_tCsA(cute::_, cute::_, cute::_, smem_pipeline_s2r_idx); // (CPY, CPY_M, CPY_K) smem_tCsB_current_pipeline = smem_tCsB(cute::_, cute::_, cute::_, smem_pipeline_s2r_idx); // (CPY, CPY_N, CPY_K) // 等待 SMEM -\u0026gt; REG 的 prefetch 完成，保证最新的 SMEM slot 在飞 cute::cp_async_wait\u0026lt;NUM_SMEM_PIPELINE - 2\u0026gt;(); __syncthreads(); } // 加载下一个 MMA_K 的数据到 REG // ...... // perform prefetch next tile GMEM -\u0026gt; SMEM before this tile's MMA if (mma_idx_k == 0) { cute::copy(gmem_tiled_copy_A, gmem_tAgA(cute::_, cute::_, cute::_, tile_idx_next), // (CPY, CPY_M, CPY_K) gmem_tAsA(cute::_, cute::_, cute::_, smem_pipeline_g2s_idx)); cute::copy(gmem_tiled_copy_B, gmem_tAgB(cute::_, cute::_, cute::_, tile_idx_next), // (CPY, CPY_N, CPY_K) gmem_tAsB(cute::_, cute::_, cute::_, smem_pipeline_g2s_idx)); cute::cp_async_fence(); num_tiles_remain--; if (num_tiles_remain \u0026gt; 0) { ++tile_idx_next; } smem_pipeline_g2s_idx = smem_pipeline_s2r_idx; smem_pipeline_s2r_idx = (smem_pipeline_s2r_idx + 1) % NUM_SMEM_PIPELINE; } // perform MMA on the current MMA_K // ...... } }5.3.1. 一些实现细节分析TODO：最后阶段存在一些dummy copy。6. 软流水（Software Pipelining）在Ampere架构中，在GEMM kernel的最内层循环中，执行SMEM =\u0026gt;REG的搬运指令，以及FMA指令。可以通过类似Multi-Stage的方式，先后issue两条指令，达到掩盖（overlap）部分数据搬运指令的延时。6.1. 软流水（Software Pipelining）的原理指令能否并行执行，取决于一些因素，比如：能否分配到硬件资源、执行单元是否可用、数据是否存在依赖（REG/SMEM）等等。如果指令没有资源可用，则不会被issue（另外，资源分配以及issue，都是以warp为单位的）。SMEM =\u0026gt; REG指令由LSU（Load/Store Unit）处理。当一个warp需要执行SMEM =\u0026gt; REG时，如果SM可以分配出32个LSU资源，并且REG/SMEM也是可用的，则此时可以立即执行，否则就需要等待、或者查找其他可用的warp并执行可以执行的指令。执行单元LSU与Tensor Core/CUDA Core相互独立，他们之间存在的依赖就是REG/SMEM的读写依赖。使用double buffer，每次预取k+1位置的tile，并且计算第k位置的tile，可以实现软流水，从而掩盖部分SMEM =\u0026gt; REG数据搬运的延时。 有关硬件资源的调度，见另外一篇博客”NVIDIA GPU 架构：SP、SM 与 LSU 工作原理详解”。6.2. 软流水的代码示例SMEM =\u0026gt; REG使用数据预取，解除了共享内存访问（R/W）的依赖，可以节省一次 __syncthreads()。原始的GEMM主循环示意代码如下：for k: copy(GMEM[k] → SMEM) ← 写 SMEM __syncthreads() ← 1 等待所有线程写完 SMEM，才能安全读（read after write） gemm(SMEM, REG) ← 读 SMEM __syncthreads() ← 2 等待所有线程读完 SMEM，才能下一轮覆写（write after read）使用双缓冲后，循环示意代码如下：prologue: copy(GMEM[0] → smem[0])__syncthreads()for k: copy(GMEM[k+1] → smem[(k+1)%2]) ← 写 smem[1-cur] gemm(smem[k%2], REG) ← 读 smem[cur] __syncthreads() ← 1 只剩1次7. Multi-Stage 中 buffer 深度的计算RTX 3060算力如下： GMEM 带宽：360 GB/s Tensor Core 峰值性能（FP16）：12.74 TFLOPS记CTA处理的分块的大小为$BM \\times BN \\times BK$，峰值带宽为$P_m$，算力为$P_c$： 每个分块的计算量为：$2 \\times BM \\times BK \\times BN$，需要的时间为：$T_c = \\frac{2 \\times BM \\times BK \\times BN}{P_c}$。 每次分块需要搬运的数据量为$2 \\times (BM \\times BK + BK \\times BN)$，需要的时间为：$T_m = \\frac{2 \\times (BM \\times BK + BK \\times BN)}{P_m}$。得到访存比为：\\[\\text{ratio} = \\frac{T_c}{T_m} = \\frac{BM \\times BN \\times BK / P_c}{(BM \\times BK + BK \\times BN) / P_m} = \\frac{BM \\times BN \\times BK}{BM \\times BK + BK \\times BN} \\times \\frac{P_m}{P_c} = \\frac{BM \\times BN}{BM + BN} \\times \\frac{P_m}{P_c}\\]峰值算力/峰值带宽为：\\[\\frac{P_c}{P_m} = \\frac{12.74 \\text{ TFLOPS}}{360 \\text{ GB/s}} \\approx 35.4 \\text{ FLOP/byte}\\]以$BM=BN=128$为例，得到：\\[\\text{ratio} = \\frac{128 \\times 128}{128 + 128} \\times \\frac{1}{35.4} \\approx 18.3 \\text{ FLOP/byte}\\] 这个计算得到的访存比，其含义是访存的时间，是计算时间的18.3倍。因此，理想情况下，Multi-Stage的buffer深度应该大于18.3，才能完全掩盖访存的延时。7.1. Multi-Stage 为什么可以提高性能以理论的计算/访存比=3为例，没有pipeline的时候，有2/3的空隙未被填满，即算力只有理论值的1/3；当开辟两个pipeline的时候，还有1/3的空隙未被填满；当开辟三个pipeline的时候，可以使得计算填满空隙，消除data harzard导致的stall，达到理论峰值算力。从图示看，增加pipeline似乎是提高了内存带宽。实际情况是，由于L2的存在，使得有些SM可以从L2中直接获取数据，即不用每个SMEM每次都从GMEM中获取数据，消除了GMEM的访问时间。A. 资料 Hierarchical GEMM Structure Efficient GEMM in CUDA CUTLASS Tutorial: Efficient GEMM kernel designs with Pipelining：参考章节：Appendix: Pipelining for an Ampere GEMM 长文介绍矩阵乘法——从自己手搓到CUTLASS实现：待阅读，动图介绍GEMM的分层结构和数据流。 浅析GEMM优化multistage数怎么算：郑思泽（北京大学 计算机系统结构博士）的知乎文章A.1. 工具 Revezone APP：画板工具" },

  { "title": "CUTLASS-Cute 初步(5)：TV Layout", "url": "/blog/2025/Cute%E5%88%9D%E6%AD%A55-TV-Layout/", "categories": "CUDA",
  "tags": "CUDA", "date": "2025-02-26", "content":
  "TV-Layout 描述 CTA 中线程的 layout，以及每个线程可以访问到哪些数据。TV-Layout 的第一个 mode 定义线程在 CTA 中的分布，第二个 mode 定义每个线程处理的数据布局。见下面例子中的LayoutA_TV。数学表述形式为：$\\text{TV-Layout}:(t,v) -\u0026gt; \\text{linear index in tile}$。以 LayoutA_TV: ((_4,_8),(_2,_2,_2)):((_32,_1),(_16,_8,_128)) 为例：(t0,t1, v0,v1,v2) → t0×32 + t1×1 + v0×16 + v1×8 + v2×128Thread 5 (t0=1, t1=1), Value 0 (v0=0,v1=0,v2=0):→ 1×32 + 1×1 = 33 → A矩阵线性坐标 33Inverse TV-Layout 描述的是数据的逻辑坐标 coord 到线程ID的映射关系。比如给定 layout 的逻辑坐标 (m, n)，经过 inverse TV-Layout 得到 (threadID, valueID)，即： threadID，表示该逻辑坐标 (m, n) 属于 warp 中的哪个线程处理 valueID，表示该线程处理的第几个数据，即线程内数据的偏移，比如 reg[2] (M, K) -\u0026gt; (T, V) 分为两个步骤： 逻辑坐标(m, n) 计算得到一维坐标 Index = Thread_ID + (Value_ID * Thread_Group_Size) 转换为人可理解的二维坐标 (threadID, valueID)，比如 threadID = Index % Thread_Group_Size，valueID = Index / Thread_Group_Size Inverse TV-Layout 的数学表述形式为：$\\text{Inverse TV-Layout}: \\text{linear index} -\u0026gt; (t,v)$。 在 CopyOperation、MMAOperation 中，使用 print_latex、print_svg 打印的 layout，实际上是 Inverse TV-Layout。1. TV-Layout 例子以 cute::SM80_16x8x16_F16F16F16F16_TN 为例，其 TV-Layout 如下：MMA_Atom ThrID: _32:_1 Shape_MNK: (_16,_8,_16) LayoutA_TV: ((_4,_8),(_2,_2,_2)):((_32,_1),(_16,_8,_128)) LayoutB_TV: ((_4,_8),(_2,_2)):((_16,_1),(_8,_64)) LayoutC_TV: ((_4,_8),(_2,_2)):((_32,_1),(_16,_8)) 线程布局方式为 4x8。 一个 16x8x16(K维度)的矩阵，A = MxK = 16x16(T)，B = KxN = 16x8(N)。warp 中的每个线程需要从 A 中拿到 2x2x2=8 个值（16x16/32=8），从 B 拿到 2x2=4 个值（16x8/32=4）。MMAOperation 以及 MMA_Traits 定义如下：// MMA 16x8x16 TNstruct SM80_16x8x16_F16F16F16F16_TN{ using DRegisters = uint32_t[2]; using ARegisters = uint32_t[4]; using BRegisters = uint32_t[2]; using CRegisters = uint32_t[2]; CUTE_HOST_DEVICE static void fma(uint32_t \u0026amp; d0, uint32_t \u0026amp; d1, uint32_t const\u0026amp; a0, uint32_t const\u0026amp; a1, uint32_t const\u0026amp; a2, uint32_t const\u0026amp; a3, uint32_t const\u0026amp; b0, uint32_t const\u0026amp; b1, uint32_t const\u0026amp; c0, uint32_t const\u0026amp; c1) {#if defined(CUTE_ARCH_MMA_SM80_ENABLED) asm volatile( \"mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 \" \"{%0, %1},\" \"{%2, %3, %4, %5},\" \"{%6, %7},\" \"{%8, %9};\\n\" : \"=r\"(d0), \"=r\"(d1) : \"r\"(a0), \"r\"(a1), \"r\"(a2), \"r\"(a3), \"r\"(b0), \"r\"(b1), \"r\"(c0), \"r\"(c1));#else CUTE_INVALID_CONTROL_PATH(\"Attempting to use SM80_16x8x16_F16F16F16F16_TN without CUTE_ARCH_MMA_SM80_ENABLED\");#endif }};template \u0026lt;\u0026gt;struct MMA_Traits\u0026lt;SM80_16x8x16_F16F16F16F16_TN\u0026gt;{ using ValTypeD = half_t; using ValTypeA = half_t; using ValTypeB = half_t; using ValTypeC = half_t; using Shape_MNK = Shape\u0026lt;_16,_8,_16\u0026gt;; using ThrID = Layout\u0026lt;_32\u0026gt;; using ALayout = Layout\u0026lt;Shape \u0026lt;Shape \u0026lt; _4,_8\u0026gt;,Shape \u0026lt; _2,_2, _2\u0026gt;\u0026gt;, Stride\u0026lt;Stride\u0026lt;_32,_1\u0026gt;,Stride\u0026lt;_16,_8,_128\u0026gt;\u0026gt;\u0026gt;; using BLayout = Layout\u0026lt;Shape \u0026lt;Shape \u0026lt; _4,_8\u0026gt;,Shape \u0026lt;_2, _2\u0026gt;\u0026gt;, Stride\u0026lt;Stride\u0026lt;_16,_1\u0026gt;,Stride\u0026lt;_8,_64\u0026gt;\u0026gt;\u0026gt;; using CLayout = SM80_16x8_Row;};打印的 inverse TV-Layout 如下：2. 使用 TV-Layout 切分数据使用 TV-Layout 切分 CTA 的 tile 数据到线程得到线程的 sub-tile，使用 TV-Layout 比使用 local_tile、local_partition 简洁：只需要使用 TV-Layout 对输入的 layout 应用 composition 操作。// Construct a TV-layout that maps 8 thread indices and 4 value indices// to 1D coordinates within a 4x8 tensor// (T8,V4) -\u0026gt; (M4,N8)auto tv_layout = Layout\u0026lt;Shape \u0026lt;Shape \u0026lt;_2,_4\u0026gt;,Shape \u0026lt;_2, _2\u0026gt;\u0026gt;, Stride\u0026lt;Stride\u0026lt;_8,_1\u0026gt;,Stride\u0026lt;_4,_16\u0026gt;\u0026gt;\u0026gt;{}; // (8,4)// Construct a 4x8 tensor with any layoutTensor A = make_tensor\u0026lt;float\u0026gt;(Shape\u0026lt;_4,_8\u0026gt;{}, LayoutRight{}); // (4,8)// Compose A with the tv_layout to transform its shape and orderTensor tv = composition(A, tv_layout); // (8,4)// Slice so each thread has 4 values in the shape and order that the tv_layout prescribesTensor v = tv(threadIdx.x, _); // (4)2.1. 在 TiledMMA 中使用 TV-Layout 切分数据示例如下： using MMATraits = cute::MMA_Traits\u0026lt;cute::SM80_16x8x16_F16F16F16F16_TN\u0026gt;; using MMAAtomShape = MMATraits::Shape_MNK; constexpr auto mma_atom = cute::MMA_Atom\u0026lt;MMATraits\u0026gt;{}; constexpr auto mma_atom_shape = MMAAtomShape{}; // (16, 8, 16) constexpr auto MMA_LAYOUT_M = 2, MMA_LAYOUT_N = 2, MMA_LAYOUT_K = 1; // 线程数扩充 constexpr auto NUM_MMA_TILE_M = 1, NUM_MMA_TILE_N = 2, NUM_MMA_TILE_K = 1; // 每个线程Atom数量扩充 constexpr auto MMA_TILE_M = cute::get\u0026lt;0\u0026gt;(mma_atom_shape) * NUM_MMA_TILE_M * MMA_LAYOUT_M; // 32 constexpr auto MMA_TILE_N = cute::get\u0026lt;1\u0026gt;(mma_atom_shape) * NUM_MMA_TILE_N * MMA_LAYOUT_N; // 32 constexpr auto MMA_TILE_K = cute::get\u0026lt;2\u0026gt;(mma_atom_shape) * NUM_MMA_TILE_K * MMA_LAYOUT_K; // 16 constexpr auto mma_layout = cute::make_layout(cute::make_shape(cute::Int\u0026lt;MMA_LAYOUT_M\u0026gt;{}, cute::Int\u0026lt;MMA_LAYOUT_N\u0026gt;{}, cute::Int\u0026lt;MMA_LAYOUT_K\u0026gt;{})); constexpr auto mma_tile = cute::make_tile(cute::Int\u0026lt;MMA_TILE_M\u0026gt;{}, cute::Int\u0026lt;MMA_TILE_N\u0026gt;{}, cute::Int\u0026lt;MMA_TILE_K\u0026gt;{}); constexpr auto tiled_mma = cute::make_tiled_mma(mma_atom, mma_layout, mma_tile); auto thr_mma = tiled_mma.get_slice(threadIdx.x); auto mma_tCrA = thr_mma.partition_fragment_A(smem_tensor_A(cute::_, cute::_, cute::Int\u0026lt;0\u0026gt;{})); // (MMA, MMA_M, MMA_K) auto mma_tCrB = thr_mma.partition_fragment_B(smem_tensor_B(cute::_, cute::_, cute::Int\u0026lt;0\u0026gt;{})); // (MMA, MMA_N, MMA_K) // 切分 PIPELINE 维度，获取该维度的 slice(0)，从而 (MMA, MMA_M, MMA_K, bP) -\u0026gt; (MMA, MMA_M, MMA_K)另外，在 TiledCopy 中，使用partition_S、partition_D切分数据。资料 CuTe Thread-Value Layout：Mao Lei 的博客文章 CuTe Inverse Layout：Mao Lei 的博客文章 Tensor Core MMA Swizzle Layout Cute概念速通：待阅读" },

  { "title": "CUTLASS-Cute 初步(4.1)：MMA Swizzle -- MMA、ldmatrix、smem swizzle", "url": "/blog/2025/Cute%E5%88%9D%E6%AD%A54.1-mma-ldmatrix-swizzle/", "categories": "CUDA",
  "tags": "CUDA", "date": "2025-02-26", "content":
  "1. ldmatrix 指令 与 MMA 指令一些名词： LDS：LoaD Shared Memory，warp 指令，比如 LDS.32 表示加载 32 位数据到寄存器 LDSM：LoaD Shared Memory Matrix，Tensor Core 指令，ldmatrix的 SASS 表示ldmatrix指令为配合 Tensor Core 的 MMA 指令使用的，数据在RF中的布局与mma指令一致（准确的理解应该是：加载RF之后的Thread-Value布局）。ldmatrix指令格式为（以.x1、.x4为例）：ldmatrix.sync.aligned.m8n8.x1.shared.b16{.trans} { %0 }, [ %1 ]ldmatrix.sync.aligned.m8n8.x4.shared.b16{.trans} { %0, %1, %2, %3 }, [ %4 ]一个ldmatrix.x1指令加载一个8x8-BF16 = 128B矩阵，占用8个线程（比如0~7），并将从SMEM中加载的数据均分到一个warp 32个线程中，在warp线程中以32位寄存器存储。格式如下： ldmatrix.x1要求提供8 x (8-BF16 = 16B) SMEM地址（每个线程提供一个，共8个线程），且每个地址16B且连续。 REG均分规律：T0 SMEM（16B = 8-BF16 = 4-REG） =\u0026gt; T0、T1、T2、T3；T1 SMEM =\u0026gt; T4、T5、T6、T7；…；T7 SMEM =\u0026gt; T28、T29、T30、T31。ldmatrix.x4指令加载一个32x8-BF16 = 512B矩阵，占用一个warp（32个线程），依旧是每个线程提供一个16B SMEM地址。每个线程提供4-REG = 8-BF16，对应到上述给出的ldmatrix.x4指令格式。图示如下： ldmatrix.x4的加载分成4个phase，即四个阶段。每个phase加载一个8x8-BF16 = 128B矩阵，且每个phase占用8个线程（比如phase使用T0~T7）。 从图示可以看出来，ldmatrix.x4覆盖一个16x16-BF16矩阵，按线程将其划分为2x2的子块阵列，即这个16x16-BF16矩阵可以表示为((2, 2), (8, 8))。 在 PTX DOC 中，与上图对应的HMMA指令是MMA.m16n18k16，对应的MMA指令的layout为 Figure 79 of PTX doc，两个图示的layout一致。 原图修订来自知乎文章：cutlass swizzle机制解析（一） 1.1. ldmatrix 与 mma 指令布局关系以mma.sync.aligned.m16n8k8.row.col.f32.bf16.bf16.f32(封装为cute::SM80_16x8x16_F16F16F16F16_TN)为例： 该指令执行D = A * B + C。 A[M, K]为row-major，即对应到BLAS命名约定的T。 B[K, N]为col-major即对应到BLAS命名约定的N（即B[N, K] row-major）。以A fragment为例，该mma指令使用A[16,8]，每个线程使用4-BF16（两个连续）。其RF的thread-value布局如下，对应官方PTX文档9.7.14.5.7. Matrix Fragments for mma.m16n8k8中的图示：Figure 71 of PTX 9.0 doc：使用ldmatrix.m8n8.b16，需要两条ldmatrix.x1指令加载一个A fragment(16x8-BF16)，每条指令加载8x8-BF16。上图中的示意图（上半部分或下半部分）与官方PTX文档中给出的图示：Figure 104 of PTX doc 一致。 一条ldmatrix.x1指令加载8x8-BF16 = 128B，正好等于一个32-Bank宽度。另外，ldmatrix加载以16B为单位，在分析Bank Conflict的时候，可以将32-Bank简化为8-Bank。由指令mma.sync.aligned.m16n8k8.row.col.f32.bf16.bf16.f32可知，B fragment的布局为col-major。官方PTX文档9.7.14.5.7. Matrix Fragments for mma.m16n8k8中给出的图示：Figure 74 of PTX doc：对于B fragment的加载，需要使用一条ldmatrix.m8n8.b16.trans指令。1.2. MMA 指令的线程布局Taking m16n8k16 FP16 as an example, the calculation distribution of elements in each tile on the threads in the warp is shown in the figure below. It can be clearly found that the fragments calculated by each thread are discontinuous. 引用自：https://bruce-lee-ly.medium.com/nvidia-tensor-core-getting-started-with-mma-ptx-programming-508e44a6cb7d。从图中可以看出，MMA中，A fragment与C fragment都是row-major，B fragment是col-major。C fragment与A fragment的布局相同，这对连续进行MMA计算非常有利，比如Flash Attention中，三个矩阵相乘，前两个矩阵计算得到的中间结果，布局满足MMA要求，可以直接用于后续计算。另外，每个线程计算的元素在矩阵中是不连续的。1.3. ldmatrix 与 ldmatrix.trans 布局对比TEST(TiledCopy, Case02) { using namespace cute; TiledMMA mmaC = make_tiled_mma(MMA_Atom\u0026lt;SM80_16x8x16_F16F16F16F16_TN\u0026gt;{}, Layout\u0026lt;Shape\u0026lt;_1, _1\u0026gt;\u0026gt;{}); { Copy_Atom\u0026lt;SM75_U32x4_LDSM_N, cute::half_t\u0026gt; s2r_atom_A; TiledCopy s2r_copy_a = make_tiled_copy_A(s2r_atom_A, mmaC); std::cout \u0026lt;\u0026lt; \"TiledCopy s2r_copy_a: \" \u0026lt;\u0026lt; std::endl; print(s2r_copy_a); std::cout \u0026lt;\u0026lt; \"\\n-----------------------------\" \u0026lt;\u0026lt; std::endl; print_latex(s2r_copy_a); } { Copy_Atom\u0026lt;SM75_U16x4_LDSM_T, cute::half_t\u0026gt; s2r_atom_B; TiledCopy s2r_copy_b = make_tiled_copy_B(s2r_atom_B, mmaC); std::cout \u0026lt;\u0026lt; \"TiledCopy s2r_copy_b: \" \u0026lt;\u0026lt; std::endl; print(s2r_copy_b); std::cout \u0026lt;\u0026lt; \"\\n-----------------------------\" \u0026lt;\u0026lt; std::endl; print_latex(s2r_copy_b); }} 从上述代码，以及打印的Inverse TV-Layout中，不能看出ldmatrix与ldmatrix.trans的区别，仅作为参考。相关文章：https://zhuanlan.zhihu.com/p/1906775725278737888。 测试代码：https://github.com/HPC02/cuda_perf/blob/master/src/study_codes/study_tests/test_tiled_copy.cu。 MMA指令提供了几种形式，不同的指令参数，对A、B在SMEM中的布局要求如下：以cute::SM80_16x8x16_F16F16F16F16_TN为例，要求A的形式为：(M, K) row-major（即对应到BLAS命名约定的T），B的形式为(N, K) row-major（即(K, N) column-major，对应到BLAS命名约定的N）。在如下的完整代码示例中，使用的存储形式为A：(M, K) column-major，B：(N, K) column-major，即在我们的GEMM实现中，使用的是MN-major的存储形式。此时，需要使用ldmatrix.trans指令来加载A、B到寄存器中，即在SMEM =\u0026gt; REG的过程中，对A、B都需要进行转置加载。 总结：ldmatrix.trans用于矩阵A/B以列主序存储在共享内存时，硬件在加载过程中完成转置，使寄存器中的布局直接匹配 mma.sync对A/B操作数的期望格式。 如上参考自Mao Lei博客：CuTe ldmatrix 2. ldmatrix 指令的 Bank Conflict 分析以及 Swizzle2.1. 一个简单的 ldmatrix Bank Conflict 示例分析 测试代码：https://github.com/HPC02/cuda_perf/blob/master/src/study_codes/test_ldmatrix/test_ldmatrix.cu，参考自https://zhuanlan.zhihu.com/p/697228676。__global__ void ldmatrixBankConflicts() { // tile ((2,2), (8,8)) constexpr int TILE_SIZE = 4 * 8 * 8; __shared__ half aTile[TILE_SIZE]; const int tidx = threadIdx.x + blockDim.x * threadIdx.y; /* 按ldmatrix.x4.m8n8的访问顺序， tile(16, 16)拆分成4个(8, 8)子块，子块顺序为： 0 1 2 3 子块内内为行优先访问，stride为16。 */ const int sub_tid = tidx % 8, sub_tile_idx = tidx / 8; const int row_base = (sub_tile_idx % 2) * 8; const int col_base = (sub_tile_idx / 2) * 8; const int tile_offset = (row_base + sub_tid) * 16 + col_base; // const int aTile_index = tidx % 16 * 16 + tidx / 16 * 8; printf(\"tidx: %d, tile_offset: %d, bank_id:%d\\n\", tidx, tile_offset, tile_offset % 32); uint32_t regs[4]; uint32_t smem = __cvta_generic_to_shared(aTile + tile_offset); asm(\"ldmatrix.sync.aligned.m8n8.x4.shared.b16 { %0, %1, %2, %3 }, [ %4 ];\\n\" : \"=r\"(regs[0]), \"=r\"(regs[1]), \"=r\"(regs[2]), \"=r\"(regs[3]) : \"r\"(smem));}int main(void) { const uint3 block = {32, 1, 1}; const uint3 grid = {1, 1, 1}; ldmatrixBankConflicts\u0026lt;\u0026lt;\u0026lt;grid, block\u0026gt;\u0026gt;\u0026gt;(); cudaDeviceReset(); return 0;}按照ldmatrix.x4.m8n8的访问顺序，每个子块的layout为(8,8):(16,1)，导致一个ldmatrix.x1指令加载的8行，会出现一个4-way bank conflict（即需要分成两个transaction）。ldmatrix.x4.m8n8指令分成四个phase，每个phase出现一个4-way bank conflict，最终导致每个ldmatrix.x4.m8n8指令出现四个4-way bank conflict。以16B为单位，简化成8 bank。具体加载分析如下：T0 =\u0026gt; Bank 0 (8-BF16 = 16B) stride = 16-BF16T1 =\u0026gt; Bank 2 (8-BF16 = 16B) stride = 16-BF16T2 =\u0026gt; Bank 4 (8-BF16 = 16B) stride = 16-BF16T3 =\u0026gt; Bank 6 (8-BF16 = 16B) stride = 16-BF16T4 =\u0026gt; Bank 0 (8-BF16 = 16B) stride = 16-BF16T5 =\u0026gt; Bank 2 (8-BF16 = 16B) stride = 16-BF16T6 =\u0026gt; Bank 4 (8-BF16 = 16B) stride = 16-BF16T7 =\u0026gt; Bank 6 (8-BF16 = 16B) stride = 16-BF16即：产生了一个4-way bank conflict，且浪费了一半的SMEM带宽。2.2. 完整的 GMEM-\u0026gt;SMEM-\u0026gt;REG 拷贝排布分析以及 Bank Conflict 分析完整代码在第三部分，以下分析基于该代码示例。定义：A tile: (M=128, K=32) M-major，定义Tread-Value Layout：(16, 8)同样是M-major的。GMEM =\u0026gt; SMEM访问采用uint128_t向量化拷贝。SMEM =\u0026gt; REG使用ldmatrix.trans，每个线程提供8-half个元素的地址，单个线程的8-half地址沿着M方向连续，每个phase中，线程给出的地址沿着K方向排布。GMEM =\u0026gt; SMEM访问分析：每个线程拷贝一个uint128_t，按照简化8-bank分析，8个线程组成一个transaction，共128B，这128B地址连续，不会产生 bank conflict。另外，A tile在M方向上分为两次transction，之后从下一个K编号开始新的transaction，不会一次transaction跨越多个K编号。SMEM =\u0026gt; REG访问分析：每个线程提供一个8-half的地址，单个线程的8-half地址沿着M方向连续，每个phase中，线程给出的地址沿着K方向排布（参考上面Inverse TV-Layout中，B(N, K)的布局）。由于stride=128-half，即K方向地址不连续，故会产生4-way bank conflict。 结论： 在GMEM =\u0026gt; SMEM，以及SMEM =\u0026gt; REG的过程中，以每个线程拷贝uint128字长为单位进行的拷贝，每8个线程组成一个transaction，如果这8个线程拷贝的地址连续，则没有bank conflict；如果这8个线程拷贝的地址不连续，则会产生bank conflict。 ldmatrix指令中，由于每个线程提供的地址是沿着leading dimension排布的，并且SMEM的fast dimension方向长度大于uint128_t，导致前后线程地址不连续，从而产生bank conflict（即size(smem) == 128B，但是cosize(smem) \u0026gt; 128B）。 在GMEM =\u0026gt; SMEM过程中，线程排布沿着fast dimension的方向，能够保证每个transaction的地址连续，从而没有store bank conflict。如果调用kernel时block是二维的，则需要注意一个transaction内的线程是否是沿着fast dimension的方向。另外，如果取的是SMEM的sub-tile，则也需要注意是否会出现size(mem) == 128B，但是cosize(mem) \u0026gt; 128B的情况。 2.3. 使用 Swizzle 解决 ldmatrix bank conflict以Swizzle K-major 128B为例，定义的GMEM tile、SMEM tile布局如下：在切分CTA并创建GMEM tile（以及对应的SMEM tile）的时候，在K维度还可以有其他的切分方式，比如Swizzle K-major 64B、Swizzle M-major 32B等。使用128B K-major性能最好：此时连续的128B刚好填满一个L2 Cache Line。 GPU的L2 Cache Line大小为128B，分为4-Sectors，每个Sector为32B。同时，访问GMEM的时候，也是以Sector（32B）为单位，所以GMEM访问小于32B造成GMEM带宽浪费。如果是列主序的，相关的GMEM tile、SMEM tile见如下图示（Swizzle MN-major 128B）：3. GMEM -\u0026gt; SMEM -\u0026gt; REG 示例及分析 测试代码：https://github.com/HPC02/cuda_perf/blob/master/src/study_codes/test_gmem_smem_swizzle/test_async_cp_ldmatrix.cu 相关代码：https://github.com/HPC02/cuda_perf/blob/master/src/study_codes/test_ldmatrix/test_ldmatrix.cunamespace {constexpr int constexpr_log2(int n) { // return (n \u0026lt;= 1) ? 0 : 1 + constexpr_log2(n / 2);}} // namespace// GMEM -\u0026gt; SMEM (async copy + swizzle) -\u0026gt; REG (ldmatrix)// 数据流：全局内存 (列主序，M-major) -\u0026gt; 共享内存 (M-major + Swizzle\u0026lt;3,3,4\u0026gt;) -\u0026gt; 寄存器 (ldmatrix.x4.trans)template \u0026lt;typename T, typename AStride, typename ASmemLayout, typename GmemTiledCopyA, typename SmemTiledCopyA, typename TiledMMA\u0026gt;__global__ void kernel_load(T const* ptrA, int M, int K, AStride strideA, ASmemLayout smem_layout_A, GmemTiledCopyA gmem_tiled_copy_A, SmemTiledCopyA smem_tiled_copy_A, TiledMMA tiled_mma) { // 完整 GMEM 张量：shape=(M,K), stride=(1,M)，列主序（M-major） // A[m, k] 的地址偏移 = m*1 + k*M，M 方向连续，K 方向跨步 M（每列间隔 M*sizeof(T) 字节） auto gmem_tensor_A = cute::make_tensor(cute::make_gmem_ptr(ptrA), cute::make_shape(M, K), strideA); // (M,K):(1,M) // 共享内存张量：shape=(bM=128, bK=32)，M-major（M stride=1 连续存储） // 叠加 Swizzle\u0026lt;3,3,4\u0026gt; 后各行的列偏移被 XOR 打乱，用于消除 ldmatrix bank conflict __shared__ T smem_A[cute::cosize_v\u0026lt;ASmemLayout\u0026gt;]; auto tensor_smem_A = cute::make_tensor(cute::make_smem_ptr(smem_A), smem_layout_A); // (bM=128, bK=32) // Block tile：每个 CTA 负责 GMEM 中一块 (bM=128, bK=32) 的数据 // grid = (M/128, K/32)，blockIdx.x -\u0026gt; M 方向，blockIdx.y -\u0026gt; K 方向 auto gmem_block_tensor_A = cute::local_tile(gmem_tensor_A, cute::make_shape(cute::size\u0026lt;0\u0026gt;(smem_layout_A), cute::size\u0026lt;1\u0026gt;(smem_layout_A)), cute::make_coord(blockIdx.x, blockIdx.y)); // (bM=128, bK=32) // ---- Stage 1: GMEM -\u0026gt; SMEM via cp.async ---- // 线程布局 (16M, 8K)，stride=(1,16)，M-major（连续 thread ID 沿 M 方向排列） // 每线程每次 cp.async 加载 128bit=8 个 half，沿 M 方向连续读取 // 同一 K 列：16 线程 × 16B = 256B = 2 条 128B cache line，GMEM coalescing 100% // gmem_tAgA：当前线程负责的 GMEM 源分块，shape=(CPY=8, CPY_M, CPY_K) // gmem_tAsA：写入 SMEM 的目标分块（地址已按 swizzle XOR 重映射），shape 同上 auto gmem_thr_copy_A = gmem_tiled_copy_A.get_slice(threadIdx.x); auto gmem_tAgA = gmem_thr_copy_A.partition_S(gmem_block_tensor_A); // (CPY=8, CPY_M, CPY_K) auto gmem_tAsA = gmem_thr_copy_A.partition_D(tensor_smem_A); // (CPY=8, CPY_M, CPY_K) cute::copy(gmem_tiled_copy_A, gmem_tAgA, gmem_tAsA); cute::cp_async_fence(); cute::cp_async_wait\u0026lt;0\u0026gt;(); __syncthreads(); // ---- Stage 2: SMEM -\u0026gt; REG via ldmatrix ---- // SM75_U16x8_LDSM_T：ldmatrix.sync.aligned.x4.trans（转置加载） // 一次加载 4 个 8×8 half 矩阵 tiles = 64 字节，供 SM80_16x8x16 的 A operand 使用 // Swizzle 保证同一 warp 内 32 线程访问 SMEM 时落在不同 bank，消除 bank conflict auto thr_mma = tiled_mma.get_slice(threadIdx.x); auto mma_tCrA = thr_mma.partition_fragment_A(tensor_smem_A); // (MMA=(2,2), MMA_M, MMA_K) auto smem_thr_copy_A = smem_tiled_copy_A.get_slice(threadIdx.x); auto smem_tCsA = smem_thr_copy_A.partition_S(tensor_smem_A); // (CPY, CPY_M, CPY_K) SMEM 侧 auto smem_tCrA_view = smem_thr_copy_A.retile_D(mma_tCrA); // (CPY, CPY_M, CPY_K) REG 侧，与 mma_tCrA 共享存储 cute::copy(smem_tiled_copy_A, smem_tCsA, smem_tCrA_view);}void test_gmem_smem_ldmatrix() { using T = cute::half_t; using VectorType = cute::uint128_t; using CopyOp = cute::SM80_CP_ASYNC_CACHEALWAYS\u0026lt;VectorType\u0026gt;; constexpr int M = 4096, K = 4096; thrust::random::default_random_engine rng(42); thrust::random::uniform_real_distribution\u0026lt;float\u0026gt; dist(-1.0f, 1.0f); thrust::host_vector\u0026lt;T\u0026gt; h_A(M * K); thrust::generate(h_A.begin(), h_A.end(), [\u0026amp;]() { return T(__half(dist(rng))); }); thrust::device_vector\u0026lt;T\u0026gt; d_A_vec(h_A); T* d_A = thrust::raw_pointer_cast(d_A_vec.data()); // GMEM 列主序（M-major）：stride=(1, M)，A[m,k] = ptrA + m + k*M // M 方向 stride=1（连续），K 方向每列间隔 M*sizeof(T) 字节（非连续） auto const strideA = cute::make_stride(cute::Int\u0026lt;1\u0026gt;{}, M); // CTA tile 大小：bM=128（M 方向），bK=32（K 方向） // grid = (M/128, K/32) = (32, 128) 个 CTA constexpr auto bM = cute::Int\u0026lt;128\u0026gt;{}; constexpr auto bK = cute::Int\u0026lt;32\u0026gt;{}; // Swizzle\u0026lt;3,3,4\u0026gt;：对逻辑行号的 bit[6:4] XOR 到列地址的 bit[6:4]， // 保证同 warp 内 32 线程的 ldmatrix 访问落在 32 个不同的 SMEM bank constexpr auto MBase_A = constexpr_log2(sizeof(cute::uint128_t) / sizeof(T)); // 3 constexpr auto BBits_A = constexpr_log2(32 * 4 / sizeof(T)) - MBase_A; // 3 constexpr auto SShift_A = constexpr_log2(bM) - MBase_A; // 4 constexpr auto swizzle_A = cute::Swizzle\u0026lt;BBits_A, MBase_A, SShift_A\u0026gt;{}; // Swizzle\u0026lt;3,3,4\u0026gt; // smem_atom_layout_A：基础原子 layout，shape=(128M, 8K)，stride=(1,128)，M-major // 一列 8 个 uint128_t = 128 个 half，正好是 ldmatrix.x4 一个 phase 的加载量 // smem_atom_layout_A_swizzled：叠加 Swizzle，不改变 shape，只重映射地址 // smem_layout_A：tile_to_shape 将 K 扩展到 bK=32（沿 K 方向 tiling 4 次），shape=(128M, 32K) constexpr auto smem_atom_layout_A = cute::make_layout(cute::make_shape(bM, cute::Int\u0026lt;8\u0026gt;{})); // (128,8):(1,128) constexpr auto smem_atom_layout_A_swizzled = cute::composition(swizzle_A, smem_atom_layout_A); constexpr auto smem_layout_A = cute::tile_to_shape(smem_atom_layout_A_swizzled, cute::make_shape(bM, bK)); // (128,32), M-major // ---- GMEM -\u0026gt; SMEM tiled copy ---- // thread_layout_A：128 线程映射到 tile (bM=128, bK=32) 的方式 // shape=(16M, 8K)，stride=(1,16)，M-major（连续 thread ID 沿 M 方向增长） // thread t: M_pos = t%16，K_pos = t/16 // 与 GMEM M-major 对齐：相邻线程访问相邻地址 → 完美 coalescing constexpr auto thread_shape_A = cute::make_shape(cute::Int\u0026lt;16\u0026gt;{}, cute::Int\u0026lt;8\u0026gt;{}); // (16M, 8K) constexpr auto thread_stride_A = cute::make_stride(cute::Int\u0026lt;1\u0026gt;{}, cute::size\u0026lt;0\u0026gt;(thread_shape_A)); // (1, 16)，M-major constexpr auto thread_layout_A = cute::make_layout(thread_shape_A, thread_stride_A); // vector_layout_A：每次 cp.async 加载的向量形状 // uint128_t = 16 字节 = 8 个 half → shape=(8M, 1K)，沿 M 方向连续加载 // 16 线程 × 8 元素 = 128 half/列，恰好覆盖 bM=128 // 同一 K 列：16 线程 × 16B = 256B = 2 条 128B cache line，L2 利用率 100% constexpr auto NUM_ELEMENTS_A = sizeof(VectorType) / sizeof(T); // 8 fp16 per uint128_t constexpr auto vector_shape_A = cute::make_shape(cute::Int\u0026lt;NUM_ELEMENTS_A\u0026gt;{}, cute::Int\u0026lt;1\u0026gt;{}); // (8M, 1K) constexpr auto vector_stride_A = cute::make_stride(cute::Int\u0026lt;1\u0026gt;{}, cute::size\u0026lt;0\u0026gt;(vector_shape_A)); constexpr auto vector_layout_A = cute::make_layout(vector_shape_A, vector_stride_A); // gmem_tiled_copy_A：SM80_CP_ASYNC_CACHEALWAYS，128bit 粒度异步拷贝 // 整个 tile (128×32) = 8192 half = 16KB，128 线程 × 8 元素/次 × 4 轮 = 4096 元素/轮 constexpr auto gmem_tiled_copy_A = cute::make_tiled_copy(cute::Copy_Atom\u0026lt;CopyOp, T\u0026gt;{}, thread_layout_A, vector_layout_A); // ---- TiledMMA：SM80_16x8x16，mma_layout=(2M,2N,1K)，mma_tile=(32M,32N,16K) ---- // MMA atom SM80_16x8x16：单次 mma.sync.aligned.m16n8k16，A operand shape=(16M, 16K) // mma_layout (2,2,1)：warp 内沿 (M,N,K) 排列 2×2×1 个 atom，扩大计算覆盖范围 // mma_tile (32,32,16)：在 mma_layout 基础上再 tiling，最终每 warp 覆盖 (32M,32N,16K) // 总线程数 = 32（warp） × MMA_LAYOUT_M × MMA_LAYOUT_N × MMA_LAYOUT_K = 32×2×2×1 = 128 using MMATraits = cute::MMA_Traits\u0026lt;cute::SM80_16x8x16_F16F16F16F16_TN\u0026gt;; using MMAAtomShape = MMATraits::Shape_MNK; // (16M, 8N, 16K) constexpr auto mma_atom = cute::MMA_Atom\u0026lt;MMATraits\u0026gt;{}; constexpr auto mma_atom_shape = MMAAtomShape{}; constexpr int MMA_LAYOUT_M = 2, MMA_LAYOUT_N = 2, MMA_LAYOUT_K = 1; // CTA内 warp 排布 constexpr int NUM_MMA_TILE_M = 1, NUM_MMA_TILE_N = 2, NUM_MMA_TILE_K = 1; constexpr auto MMA_TILE_M = cute::get\u0026lt;0\u0026gt;(mma_atom_shape) * NUM_MMA_TILE_M * MMA_LAYOUT_M; // 16*1*2=32 constexpr auto MMA_TILE_N = cute::get\u0026lt;1\u0026gt;(mma_atom_shape) * NUM_MMA_TILE_N * MMA_LAYOUT_N; // 8*2*2=32 constexpr auto MMA_TILE_K = cute::get\u0026lt;2\u0026gt;(mma_atom_shape) * NUM_MMA_TILE_K * MMA_LAYOUT_K; // 16*1*1=16 constexpr auto mma_layout = cute::make_layout(cute::make_shape(cute::Int\u0026lt;MMA_LAYOUT_M\u0026gt;{}, cute::Int\u0026lt;MMA_LAYOUT_N\u0026gt;{}, cute::Int\u0026lt;MMA_LAYOUT_K\u0026gt;{})); constexpr auto mma_tile = cute::make_tile(cute::Int\u0026lt;MMA_TILE_M\u0026gt;{}, cute::Int\u0026lt;MMA_TILE_N\u0026gt;{}, cute::Int\u0026lt;MMA_TILE_K\u0026gt;{}); constexpr auto tiled_mma = cute::make_tiled_mma(mma_atom, mma_layout, mma_tile); // 128 threads total // ---- SMEM -\u0026gt; REG tiled copy via ldmatrix ---- // SM75_U16x8_LDSM_T：ldmatrix.sync.aligned.x4.trans（转置加载） // 一次加载 4 个 8×8 half 矩阵片段 = 64 字节，供 SM80_16x8x16 的 A operand 使用 // make_tiled_copy_A 根据 tiled_mma 自动推导与 MMA atom A operand 匹配的分片方案 using Copy_Atom_A = cute::Copy_Atom\u0026lt;cute::SM75_U16x8_LDSM_T, T\u0026gt;; constexpr auto smem_tiled_copy_A = cute::make_tiled_copy_A(Copy_Atom_A{}, tiled_mma); // block=128 线程（4 warp），grid=(M/bM, K/bK)=(32, 128) dim3 block(cute::size(tiled_mma)); // 128 threads dim3 grid(M / bM, K / bK); kernel_load\u0026lt;\u0026lt;\u0026lt;grid, block\u0026gt;\u0026gt;\u0026gt;(d_A, M, K, strideA, smem_layout_A, gmem_tiled_copy_A, smem_tiled_copy_A, tiled_mma); if (cudaError_t err = cudaDeviceSynchronize(); err != cudaSuccess) { std::cout \u0026lt;\u0026lt; \"CUDA error: \" \u0026lt;\u0026lt; cudaGetErrorString(err) \u0026lt;\u0026lt; std::endl; }}3.1. GEMM 分块考虑GMEM分块为(bM=128, bK=32)，M-major，每个CTA负责一个分块。线程布局为(16M, 8K)，M-major，一个线程负责加载一个128-bit=8 half的向量，线程沿M方向连续访问，M方向16个线程，16 * 8 = 128 half = 256B，填满两条L2 cache line（使用cp.async将数据从GMEM加载到SMEM）。3.2. Swizzle 设计分析Swizzle设计。按照uint128_t = 16B = 8 half简化为8-Bank，得到Swizzle\u0026lt;BBits=3, MBase=3, SShift=?\u0026gt;，即交换模式有$2^{B}=2^{3}=8$个（uint128_t），覆盖长度$2^{B+M}=2^{3+3}=64$个half。如果取SShift=3，则整个交换周期覆盖的长度为$2^{B+M+S}=2^{3+3+3}=512$个half。由于ldmatrix一个phase需要8列，又bM=128，可以取atom_shape=(128, 8)。此时取SShift=4，则整个交换周期覆盖的长度为$2^{B+M+S}=2^{3+3+4}=1024=128 \\times 8$个half，即覆盖整个atom_shape。在本例中，bM=128，所以取SShift=4，保证交换周期覆盖整个bM维度。该Swizzle可以保证ldmatrix的一个phase在K方向上8个uint128_t读取不冲突，也可以保证GMEM -\u0026gt; SMEM不冲突（16个线程从GMEM的M方向上取连续的128 half，其中每个线程取1 uint128_t = 8 half），并存到经过Swizzle之后的SMEM中（GMEM(m, k) =\u0026gt; SMEM(m, k')）。 ldmatrix一个phase拷贝，线程排布是K方向上连续，GMEM -\u0026gt; SMEM的线程排布是M方向上连续，Swizzle设计保证了两者都不冲突。3.3. GMEM -\u0026gt; SMEM 线程划分结果计算CTA的tile大小为(128M, 32K)，Thread-Value布局为((16M, 8K), (8M, 1K))，每个线程在M方向上拷贝uint128_t=8 half，即(8M, 1K)中的8M。Tiler大小为(128, 8)。TiledCopy Tiler_MN: (_128,_8) TiledLayout_TV: (_128,_8):(_8,_1)Copy_Atom ThrID: _1:_0 ValLayoutSrc: (_1,_8):(_0,_1) ValLayoutDst: (_1,_8):(_0,_1) ValLayoutRef: (_1,_8):(_0,_1) ValueType: 16b打印信息给出了两个信息，Tiler大小，线程布局。这个TiledCopy覆盖范围为(128M, 8K)。线程布局为一个线程负责拷贝连续的8 half，且线程在M方向上连续。\\[128 threads \\times 8 half/thread = 128M \\times 8K = 1024 half\\]Inverse TV-Layout展示如下:从上图看到S与D的线程访问模式一样（即layout一样），这是因为根据Copy_Traits定义：template \u0026lt;class S, class D\u0026gt;struct Copy_Traits\u0026lt;SM80_CP_ASYNC_CACHEALWAYS\u0026lt;S,D\u0026gt;\u0026gt;{ using SrcLayout = Layout\u0026lt;Shape\u0026lt;_1, Int\u0026lt;sizeof_bits\u0026lt;S\u0026gt;::value\u0026gt;\u0026gt;\u0026gt;; // S = uint128_t → 128 bits using DstLayout = Layout\u0026lt;Shape\u0026lt;_1, Int\u0026lt;sizeof_bits\u0026lt;D\u0026gt;::value\u0026gt;\u0026gt;\u0026gt;; // D = uint128_t → 128 bits};// using CopyOp = cute::SM80_CP_ASYNC_CACHEALWAYS\u0026lt;VectorType\u0026gt;;如注释行给出的实例参数，S与D的CopyOperation一样，所有布局也一样。测试代码链接：https://github.com/HPC02/cuda_perf/blob/master/src/study_codes/study_tests/test_tiled_copy.cu。另外，如果将vector_layout_A改为(16,1):(1,16)，得到：TiledCopy Tiler_MN: (_256,_8) TiledLayout_TV: (_128,_16):(_16,_1)Copy_Atom ThrID: _1:_0 ValLayoutSrc: (_1,_8):(_0,_1) ValLayoutDst: (_1,_8):(_0,_1) ValLayoutRef: (_1,_8):(_0,_1) ValueType: 16b3.4. tiled mma 及 smem tiled copyMMAOperation使用SM80_16x8x16_F16F16F16F16_TN，覆盖MNK=(16, 8, 16)，并且在MNK分别扩展为(2-layout, 2-layout * 2-tile, 1-layout)。即： 对矩阵A，在M方向上线程扩展为2倍，占用线程32*2=64个，覆盖(32M, 16K)，每个线程负责A中$32 \\times 16 \\div 64 = 8 half$。 对矩阵B，在N方向上线程扩展为2倍，占用线程32*2=64个。另外N方向上Atom重复2次，可以覆盖(2*16N, 16K)，每个线程负责B中$16 \\times 16 \\div 64 = 4 half$。其中，占用线程T0~T31、T64~T95。配置的TiledMMA，其Inverse TV-Layout如下图所示：由于MMA指令对数据在RF中的布局有特定格式要求，在创建SMEM TiledCopy，以及划分线程SMEM -\u0026gt; REG拷贝操作时，都需要使用到TiledMMA的布局。比如针对A tile： using Copy_Atom_A = cute::Copy_Atom\u0026lt;cute::SM75_U16x8_LDSM_T, T\u0026gt;; constexpr auto smem_tiled_copy_A = cute::make_tiled_copy_A(Copy_Atom_A{}, tiled_mma); auto thr_mma = tiled_mma.get_slice(threadIdx.x); auto mma_tCrA = thr_mma.partition_fragment_A(tensor_smem_A); // (MMA=(2,2), MMA_M, MMA_K) auto smem_thr_copy_A = smem_tiled_copy_A.get_slice(threadIdx.x); auto smem_tCsA = smem_thr_copy_A.partition_S(tensor_smem_A); // (CPY, CPY_M, CPY_K) SMEM 侧 auto smem_tCrA_view = smem_thr_copy_A.retile_D(mma_tCrA); // (CPY, CPY_M, CPY_K) REG 侧，与 mma_tCrA 共享存储 cute::copy(smem_tiled_copy_A, smem_tCsA, smem_tCrA_view);ThreadCopy负责线程执行MMA需要的REGs。另一方面由于源layout与目的layout不同，所以需要retile_D，将目的layout（RF layout）调整为与源layout（SMEM layout）一致。打印的layout信息如下：smem_tCsA (src):smem_ptr[16b](0xffffa6000000) o ((_8,_1),(_2,_2),_2):((_1,_0),(32,_64),_2048)mma_tCrA (dst):ptr[16b](0xffffa4fffc60) o ((_2,_2,_2),_4,_2):((_1,_2,_4),_8,_32)smem_tCrA_view (dst view):ptr[16b](0xffffa4fffc60) o ((_8,_1),_4,_2):((_1,_0),_8,_32)3.5. 有关Inverse TV-Layout图示的说明上面Inverse TV-Layout的图示中A只有64个线程，实际上CTA中128个线程都有从SMEM中加载数据，只不过可以分为两个64线程group，这两个group从SMEM上加载的位置相同，即SMEM被加载了两次，但是加载到不同的thread group的寄存器中，图示不能体现出来。tiled_mma infomation:TiledMMA ThrLayoutVMNK: (_32,_2,_2,_1):(_1,_32,_64,_0) PermutationMNK: (_32,_32,_16)MMA_Atom ThrID: _32:_1 Shape_MNK: (_16,_8,_16) LayoutA_TV: ((_4,_8),(_2,_2,_2)):((_32,_1),(_16,_8,_128)) LayoutB_TV: ((_4,_8),(_2,_2)):((_16,_1),(_8,_64)) LayoutC_TV: ((_4,_8),(_2,_2)):((_32,_1),(_16,_8))smem_tiled_copy_A infomation:TiledCopy Tiler_MN: (_32,_16) TiledLayout_TV: ((_4,_8,_2,_2),((_2,_2,_2),(_1,_1))):((_64,_1,_16,_0),((_32,_8,_256),(_0,_0)))Copy_Atom ThrID: _32:_1 ValLayoutSrc: (_32,_8):(_8,_1) ValLayoutDst: ((_4,_8),(_1,_2,_4)):((_16,_1),(_1,_8,_64)) ValLayoutRef: ((_4,_8),(_1,_2,_4)):((_16,_1),(_1,_8,_64)) ValueType: 16b以ThrLayoutVMNK为例，$t=V \\times 1 + M \\times 32 + N \\times 64 + K \\times 0$。这个128线程完整映射表： 线程 ID 范围 V M N K 0 .. 31 0-31 0 0 0 32 .. 63 0-31 1 0 0 64 .. 95 0-31 0 1 0 96 .. 127 0-31 1 1 0 由于mma指令时按照warp执行的，实际执行时，128个线程被拆分为4个warp执行，每个warp执行计算得到相应的C(M', N')。得到如下执行分配关系： 计算 发起线程 需要 A 数据 需要 B 数据 C[M, N0] = A × B[:,N0] threads 0-31, 32-63 (N-group 0) ✓ 在自己 REG 里 N0 片段 C[M, N1] = A × B[:,N1] threads 64-95, 96-127 (N-group 1) ✓ 在自己 REG 里 N1 片段 对矩阵A，N=0与N=1时，分别从同一个SMEM加载两次A tile，但是分别加载到不同的线程组（64线程）的寄存器中了。针对B tile的加载也是同理。 Operand 冗余维度 加载次数 原因 A N ×2 N=0 组和 N=1 组各 load 一次相同 SMEM A B M ×2 M=0 组和 M=1 组各 load 一次相同 SMEM B 4. 辅助用测试代码列表 GMEM -\u0026gt; SMEM，ldmatrix \u0026amp; mma，Swizzle完整流程测试代码：https://github.com/HPC02/cuda_perf/blob/master/src/study_codes/test_gmem_smem_swizzle/test_async_cp_ldmatrix.cu GMEM -\u0026gt; SMEM 的Swizzle配置测试代码：https://github.com/HPC02/cuda_perf/blob/master/src/study_codes/test_gmem_smem_swizzle/test_gemm_smem_swizzle.cu ldmatrix与ldmatrix.trans对比测试（以及打印latex）：https://github.com/HPC02/cuda_perf/blob/master/src/study_codes/study_tests/test_tiled_copy.cu mma打印信息以及latex：https://github.com/HPC02/cuda_perf/blob/master/src/study_codes/study_tests/test_mma.cu 另一个测试ldmatrix加载不同大小SMEM tile时bank conflict情况分析：https://github.com/HPC02/cuda_perf/blob/master/src/study_codes/test_ldmatrix/test_ldmatrix.cuA. 参考资料 Tensor Core MMA Swizzle Layout：杨轶凡博客 cute 之 Swizzle：知乎reed文章中，ldmatrix 相关内容 tensorcore中ldmatrix指令的优势是什么？ Graphene: An IR for Optimized Tensor Computations on GPUs Graphene: An IR for Optimized Tensor Computations on GPUs：来自 papaercache 的中文翻译 CuTe ldmatrix：毛磊博客 ldmatrix与swizzle（笔记） 实用 Swizzle 教程（一） cute代码仓库 – gemm_config for sm80A.1. GTC 资料 CUDA Techniques to MaximizeCompute and Instruction Throughput A Generalized Micro-kernel Abstraction for GPU Linear Algebra CuTe– CUDA Tensors Speaking Tensor Cores with CUTLASS 2024A.2. 相关工具 Excalidraw：白板工具，提供VSCode插件" },

  { "title": "CUTLASS-Cute 初步(4)：Swizzle", "url": "/blog/2025/Cute%E5%88%9D%E6%AD%A54-swizzle/", "categories": "CUDA",
  "tags": "CUDA", "date": "2025-02-26", "content":
  "Swizzle作用于SMEM的layout。给定layout范围内，Swizzle通过列异或操作（icol = irow ^ icol），周期性的coord重排，映射到新的物理地址offset。Swizzle定义了三个参数： $MBase$：以 $2^M$ 个一维坐标连续的元素为单位，将其当做一个元素； $SShift$：从Offset中提取的高位偏移，用于提取Offset的lead dimension； $BBits$：参与XOR的位数，用于提取一维index中的lead dimension、fast dimension中的部分bits。引用reed解释及图示，其输入为一个一维坐标的layout，通过swizzle将其拆分为二维坐标表示形式：给定义一个输入Offset:\u0026lt;LeadBits:FastBits\u0026gt;： 提取关系为：BBits+MBase -\u0026gt; FastBits，SShift+MBase -\u0026gt; LeadBits。 参与XOR操作的位宽为：BBits，即LeadBits中的低BBits位，FastBits中的高BBits位。BBits表示有$2^B$个交换模式，SShift表示交换模式的周期。通常$\\mid{S}\\mid \\ge B$，如果$\\mid{S}\\mid \\gt B$，则此交换模式重复$2^{\\mid{S}\\mid - B}$次，如果$\\mid{S}\\mid = B$，则只套用一次此交换模式。一般在设置Swizzle参数时，按输入的layout一行（准确的说是fast dimension）为周期进行swizzle，$2^{S+M}$ = 输入layout的列长度（此处仅指逻辑上的，实际完整的计算公式还需要考虑到元素存储字节数，具体见下面章Swizzle 参数设计规则）。比如 half 类型的layout (8, 32):(32, 1)，定义swizzle\u0026lt;3, 3, 3\u0026gt;，即 8 个元素形成新的最小单位（M），8 个最小单位为一行（B），所以swizzle从$8 \\times 8 = 64$个元素开始。见下面示例。B 为 8，则整个swizzle周期为 8 行。 设计Swizzle参数时，要求S \u0026gt;= B，否则不能提取到LeadBits。 设计Layout时，如果Layout的fast dimension长度小于 $2^{S+M}$，Swizzle不能完整的提取LeadBits，导致Swizzle失效或部分失效。 异或操作数学符号为 $\\oplus$。1. Cute Swizzle 示例定义 layout (8, 32):(32, 1)，定义swizzle\u0026lt;3, 2, 3\u0026gt;。定义的Swizzle含义如下： $2^{B+M}=2^{3+2}$ = 32：fast dimension 长度，32个elements。 $2^{M}=2^2=4$：4个element组成一个最小单位。 $2^{B}=2^3=8$：8个交换模式。 $2^{S}=2^3=8$：交换模式周期为8。代码如下：from cutlass import cutefrom cute_viz import render_layout_svg, render_swizzle_layout_svg@cute.jitdef test_swizzle_layout(): layout_2d = cute.make_layout((8, 32), stride=(32, 1)) sw = cute.make_swizzle(3, 2, 3) swizzled_layout = cute.make_composed_layout(sw, 0, layout_2d) render_layout_svg(layout_2d, \"out/original_layout.svg\") render_swizzle_layout_svg(swizzled_layout, \"out/swizzled_layout.svg\")test_swizzle_layout()结果如下：2. Swizzle 逻辑及规律class Swizzle { public: Swizzle(int num_bits, int num_base, int num_shft) : m_num_bits(num_bits), m_num_base(num_base), m_num_shft(num_shft) { CHECK2(m_num_bits \u0026gt;= 0, \"BBits must be positive.\"); CHECK2(m_num_base \u0026gt;= 0, \"MBase must be positive.\"); CHECK2(std::abs(m_num_shft) \u0026gt;= m_num_bits, \"abs(SShift) must be more than BBits.\"); } template \u0026lt;class Offset\u0026gt; auto apply(Offset offset) const noexcept { return offset ^ shiftr(offset \u0026amp; m_yyy_msk); // ZZZ ^= YYY } template \u0026lt;class Offset\u0026gt; auto operator()(Offset offset) const noexcept { return apply(offset); } private: template \u0026lt;class Offset\u0026gt; auto shiftr(Offset offset) const noexcept { return m_msk_sft \u0026gt;= 0 ? offset \u0026gt;\u0026gt; m_msk_sft : offset \u0026lt;\u0026lt; -m_msk_sft; } int m_num_bits; int m_num_base; int m_num_shft; int m_bit_msk = (1 \u0026lt;\u0026lt; m_num_bits) - 1; int m_yyy_msk = m_bit_msk \u0026lt;\u0026lt; (m_num_base + std::max(0, m_num_shft)); int m_zzz_msk = m_bit_msk \u0026lt;\u0026lt; (m_num_base - std::min(0, m_num_shft)); int m_msk_sft = m_num_shft;}; m_zzz_msk没有参与swizzle计算，在CuTe库中用作检查作用。Swizzle根据参数BBits、SShift，生成offset的掩码：高位部分提取掩码yyy_msk、低位部分提取zzz_msk，即分别对应行提取掩码、列提取掩码。此时，可以将MBase看作是BBits的一部分。直观展示如下： bits bits -- --0bxxxxxxxxxxxxxxxxYYxxxxxxxxxZZxxx \u0026lt;---------\u0026gt;--- shift base针对每个offset，经过swizzle映射之后，异或更新低位掩码对应的值：0bxxxxxxxxxxxxxxxxYYxxxxxxxxxAAxxx其中 AA = ZZ ^ YY。2.1. Swizzle 参数影响规律分析分别以行混淆周期，以及列混淆周期，这两个层次来分析。以一个 layout (32, 16):(16, 1)为例，分析swizzle\u0026lt;B=4, M=0, S4\u0026gt; 参数变动对结果的影响规律。经过offset \u0026amp; yyy_mk提取行号低四位（以及shiftr操作得到最终行号提取掩码），得到第 0 行、第 16 行（0x10）由于行号掩码提取过后的低 4 位为 0，导致swizzle无效。2.1.1. B 参数对周期的影响如果使用 swizzle\u0026lt;B=3, M=0, S=4\u0026gt;，导致高位掩码提取行号的低三位，行号为 0、8、16、24 时，得到的高位 YY 部分均为 0，swizzle 异或操作不生效。如果设 S = 3，则layout的左半部分（列 0 ~ 8）呈现 0、8、16、24 的行混淆周期。右半部分暂没有理清规律。2.1.2. S 参数以及列长度对周期的影响如果使用 S = 5，行号提取范围扩大到 32，由于 layout 的行号、列号范围对应掩码为 4 位，导致的结果为行号有效的掩码位为 0bxxxx0 \u0026gt;\u0026gt; 1，即最低一位被丢弃，且有效的掩码位为 4 位（注意 shiftr 的实现是提取高位之后右移S位）。最终的规律为：第 0、1 行维持不变，行周期变为 32 行，即第32、33行维持不变。中间的行，则每两行异或计算结果一致，即如果应用于解决 bank conflicts，此时只能消除一半的 bank conflicts。2.2. 测试代码from cutlass import cutefrom cute_viz import render_swizzle_layout_svg@cute.jitdef test_swizzle_layout(): layout_2d = cute.make_layout((32, 16), stride=(16, 1)) sw = cute.make_swizzle(4, 0, 4) swizzled_layout = cute.make_composed_layout(sw, 0, layout_2d) # render_layout_svg(layout_2d, \"out/original_layout.svg\") render_swizzle_layout_svg(swizzled_layout, \"out/swizzled_layout.svg\")test_swizzle_layout()2.3. Swizzle 与 Layout 的关系Swizzle 的实现与 Layout 没有关系，但是从其实现看，最终形成的 index，fast dimension （即内存连续的维度）位于低位，得到的结果就是对 fast dimension 形成XOR操作，即改变其映射顺序。演示代码片段如下： for (int i = 0; i \u0026lt; cute::size\u0026lt;0\u0026gt;(layout); i++) { for (int j = 0; j \u0026lt; cute::size\u0026lt;1\u0026gt;(layout); j++) { int idx = layout(i, j); int swizzled_idx = swizzle(idx); int bank_id = (swizzled_idx * element_size / 4) % 32; } }创建一个列主序的 Layout (32, 32):(1, 32)，XOR此时作用在行号上（即 fast dimension），但是从物理存储上看，其结果与行主序的 layout 是相同的。这个列主序 Layout 原始图示，以及 Swizzle 映射之后图示如下：另外，Swizzle就是一个复合映射，演示代码如下： constexpr auto smem_atom_layout_A = cute::make_layout(cute::make_shape(cute::Int\u0026lt;32\u0026gt;{}, cute::Int\u0026lt;8\u0026gt;{})); constexpr auto smem_atom_layout_A_swizzled = cute::composition(swizzle_A, smem_atom_layout_A); constexpr auto smem_shape_A = cute::make_shape(bM, bK, bP); // (bM, bK, bP) constexpr auto smem_layout_A = cute::tile_to_shape(smem_atom_layout_A_swizzled, smem_shape_A);3. Swizzle 参数设计规则假定矩阵每个元素大小为 S-byte，向量化访问的宽度为N个元素，shared memory 的 fast dimension 为X个元素。即： 符号 含义 S_elem 每个元素的大小（bytes） N 向量化访问的元素个数 X Fast dimension 的元素个数 $\\text{MBase} = log_{2}\\text{N}$，即向量化访问的元素个数。$\\text{SShift} = log_{2}\\text{X} - \\text{MBase}$，即 X = $2^{\\text{MBase} + \\text{SShift}}$，这样使得针对每一行的 swizzle 操作，掩码偏移对齐到行号的位置。即，将提取行号的掩码分为两部分：低 MBase 位不参与 swizzle，高 SShift 位参与 swizzle。$\\text{BBits} = log_{2}\\text{(32 * 4 / S)} - \\text{MBase}$。其原因为要确保 BBits 对应覆盖一次 shared memory 的访问字宽：128 字节，即 $\\text{S} \\times 2^\\text{MBase + BBits}$ = 32 * 4 = 128B。 即要求 32 个连续的 word 地址，经过 swizzle 之后，分别落入 shared memory 的32个bank中。3.1. Swizzle 参数设计示例１假定 half 类型（S_elem = 2 bytes）行主序矩阵，矩阵大小为 8 * 64。采用 128-bit 向量化访问指令，即每次访问 8 个 half 元素（4 个 word）。MBase = log2(8) = 3。SShift = log2(64) - MBase = 6 - 3 = 3。即其约束在于 shared memory 的 fast dimension 为 64 个元素，即掩码偏移量为 fast dimension 长度。BBits = log2(32 * 2) - MBase = 6 - 3 = 3。即要求 32 个连续的 word 地址，经过 swizzle 之后，分别落入 shared memory 的32个 bank 中。最终，得到 swizzle\u0026lt;3, 3, 3\u0026gt;。3.2. 一个错误的设计示例，以及修复输入 layout，其 Offset 组成为：\u0026lt;LeadBits\u0026gt;:\u0026lt;FastBits\u0026gt;。Swizzle 对输入 layout 的约束为： 输入 layout 的 fast dimension 长度，即连续内存访问的维度长度，应该等于 $2^{MBase + SShift}$。 如果输入 layout fast dimension 长度小于 $2^{MBase + SShift}$，则导致 Swizzle 只取到LeadBits的高位部分，低位部分取到FastBits里面去了。 输入 layout 的 lead dimension 长度小于 $2^{MBase + SShift}$，没有影响，即一个交换模式周期没有被完整应用。错误示例如下： constexpr auto bM = cute::Int\u0026lt;128\u0026gt;{}; constexpr auto bK = cute::Int\u0026lt;32\u0026gt;{}; constexpr auto MBase_A = constexpr_log2(sizeof(cute::uint128_t) / sizeof(T)); // log2(8) = 3 constexpr auto BBits_A = constexpr_log2(32 * 4 / sizeof(T)) - MBase_A; // log2(64) - 3 = 3 constexpr auto SShift_A = constexpr_log2(bM) - MBase_A; // log2(128) - 3 = 4 constexpr auto swizzle_A = cute::Swizzle\u0026lt;BBits_A, MBase_A, SShift_A\u0026gt;{}; constexpr auto smem_atom_layout_A = cute::make_layout(cute::make_shape(cute::Int\u0026lt;32\u0026gt;{}, cute::Int\u0026lt;8\u0026gt;{})); constexpr auto smem_atom_layout_A_swizzled = cute::composition(swizzle_A, smem_atom_layout_A); constexpr auto smem_layout_A = cute::tile_to_shape(smem_atom_layout_A_swizzled, cute::make_shape(bM, bK));上述示例中，输入 layout 的 fast dimension=32 =\u0026gt; 位宽为 5；LeadBits= 8 =\u0026gt; 位宽为 3。得到 Offset 位格式为\u0026lt;LeadBits=3:FastBits=5\u0026gt;。与 Swizzle 对齐的位宽组成要求为\u0026lt;X:4+3\u0026gt;，导致 Swizzle 只从lead dimension的高位部分获取2-Bit，低位2-Bit是从Offset的fast dimension中获取的；另外，经过BBits位与之后，LeadBits只有1-Bit起作用。其结果就是，Swizzle 之后的结果，shape 等于输入 layout 的 shape，但是没有达到预期的效果。测试代码：@cute.jitdef test_swizzle_layout2(): # 该swizzle不起作用 \"\"\"该swizzle不起作用, 因为fast dimension长度不满足要求\"\"\" layout_3d = cute.make_layout((32, 8)) # 列主序布局 sw = cute.make_swizzle(3, 3, 4) # 要求fast dimension 长度为 2^(4+3) = 128 swizzled_layout = cute.make_composed_layout(sw, 0, layout_3d) render_layout_svg(layout_3d, \"out/original_layout2.svg\") render_swizzle_layout_svg(swizzled_layout, \"out/swizzled_layout2.svg\")test_swizzle_layout2()修复代码如下： // constexpr auto bM = cute::Int\u0026lt;128\u0026gt;{}; constexpr auto smem_atom_layout_A = cute::make_layout(cute::make_shape(bM, cute::Int\u0026lt;8\u0026gt;{}));对应的测试代码：@cute.jitdef test_swizzle_layout3(): \"\"\"修正: fast dimension长度满足要求, swizzle生效\"\"\" \"\"\"MBase=3: 8个元素组成一组. \"\"\" layout_3d = cute.make_layout((128, 8)) # 列主序布局 sw = cute.make_swizzle(3, 3, 4) # 要求fast dimension 长度为 2^(4+3) = 128 swizzled_layout = cute.make_composed_layout(sw, 0, layout_3d) render_layout_svg(layout_3d, \"out/original_layout3.svg\") render_swizzle_layout_svg(swizzled_layout, \"out/swizzled_layout3.svg\")test_swizzle_layout3()Swizzle之后的Layout：相关测试代码：https://github.com/HPC02/cuda_perf/blob/master/src/study_codes/test_gmem_smem_swizzle/test_gemm_smem_swizzle.cu。4. Thread Block Swizzle由于 CUDA 调度实际上是以 block id 的顺序进行调度的，有时候，通过对 thread block 映射的tile进行重新排序，最大限度的利用L2 cache中个的数据，提升性能。GPU 硬件调度器通常优先调度 x，然后 y，最后 z 维度（单个SM调度thread block？还是所有SM -\u0026gt; thread block的顺序？）。 Thread block swizzle，以及 block 调度顺序，见 github issue：[QST]how to understand “block swizzling”，以及博客Nvidia Tensor Core-CUDA HGEMM Advanced Optimization。 关于使用 thread block swizzle 复用 L2 cache，见 NVIDIA 博客Optimizing Compute Shaders for L2 Locality using Thread-Group ID Swizzling。A. 参考资料 cute 之 Swizzle：来自知乎 Reed 文章 Tensor Core MMA Swizzle Layout：来自 Yang Yifan 博客。待学习 CuTe Swizzle：来自 Lei Mao 博客，其中段落Vectorized Memory Access讲述如何设计连续内存访问的 swizzle 参数。 CUTLASS CuTe GEMM细节分析（四）——谈谈Swizzle模板参数中关于B和S的一些误区：知乎 Anonymous 文章 淺談CUTLASS / CuTe的Swizzling Functor CUDA 013 - Swizzle 的工作原理：网页版 swizzle 可视化 SwizzleVis：Python 实现的 swizzle 可视化工具，网页查看结果A.1. 其他资料 github CUDA-Learn-Notes：很多学习用的 kernel 代码 NVIDIA Ampere GA102 GPU Architecture Whitepaper" },

  { "title": "CUTLASS-Cute 初步(3.1)：TiledCopy 以及 TiledMMA 配置示例", "url": "/blog/2025/Cute%E5%88%9D%E6%AD%A53.1-TileMMA%E9%85%8D%E7%BD%AE%E7%A4%BA%E4%BE%8B/", "categories": "CUDA",
  "tags": "CUDA", "date": "2025-02-26", "content":
  "cute_tiled_mma_preview.cu // Configure data type. using TA = cute::half_t; using TB = cute::half_t; using TC = cute::half_t; // Configure static \"shared memory\". // The \"shared memory\" is actually on host for preview purpose. // For tiled mma, the shared memory layout has to be static. constexpr int bM{128 * 2 / sizeof(TA)}; constexpr int bN{128 * 2 / sizeof(TB)}; constexpr int bK{32}; auto const blk_M = cute::Int\u0026lt;bM\u0026gt;{}; auto const blk_N = cute::Int\u0026lt;bN\u0026gt;{}; auto const blk_K = cute::Int\u0026lt;bK\u0026gt;{}; auto const smem_shape_A{cute::make_shape(blk_M, blk_K)}; auto const smem_shape_B{cute::make_shape(blk_N, blk_K)}; auto const smem_shape_C{cute::make_shape(blk_M, blk_N)}; auto const smem_stride_A{cute::make_stride(cute::Int\u0026lt;1\u0026gt;{}, blk_M)}; // Column-major auto const smem_stride_B{cute::make_stride(cute::Int\u0026lt;1\u0026gt;{}, blk_N)}; // Column-major auto const smem_stride_C{cute::make_stride(cute::Int\u0026lt;1\u0026gt;{}, blk_M)}; // Column-major auto const smem_layout_A{cute::make_layout(smem_shape_A, smem_stride_A)}; // (blk_M, blk_K) auto const smem_layout_B{cute::make_layout(smem_shape_B, smem_stride_B)}; // (blk_N, blk_K) auto const smem_layout_C{cute::make_layout(smem_shape_C, smem_stride_C)}; // (blk_M, blk_N) auto const size_a{blk_M * blk_K}; auto const size_b{blk_N * blk_K}; auto const size_c{blk_M * blk_N}; auto h_A = thrust::host_vector\u0026lt;TA\u0026gt;(size_a); auto h_B = thrust::host_vector\u0026lt;TB\u0026gt;(size_b); auto h_C = thrust::host_vector\u0026lt;TC\u0026gt;(size_c); // Make tensor for smem_A and smem_B. auto smem_tensor_A{cute::make_tensor(h_A.data(), smem_layout_A)}; auto smem_tensor_B{cute::make_tensor(h_B.data(), smem_layout_B)}; auto smem_tensor_C{cute::make_tensor(h_C.data(), smem_layout_C)};1. MMA_Atom 以及 TiledMMA 配置位于 SMEM 中的 tile 大小为 $M \\times N \\times K = 128 \\times 128 \\times 32$，其中： A 矩阵为 $M \\times K = 128 \\times 32$，row-major layout； B 矩阵为 $K \\times N = 32 \\times 128$，column-major layout； C 矩阵为 $M \\times N = 128 \\times 128$。1.1. MMA_Atom 配置MMA_Atom 使用的配置为 cute::SM80_16x8x16_F16F16F16F16_TN，使用一个 warp，即 32 个线程处理这个 MMA Atom。处理的 MNK 规模为：$M’ \\times N’ \\times K’ = 16 \\times 8 \\times 16$，其中： A sub-tile 为 $M’ \\times K’ = 16 \\times 16$； B sub-tile 为 $K’ \\times N’ = 16 \\times 8$； C sub-tile 为 $M’ \\times N’ = 16 \\times 8$。mma_atomMMA_Atom ThrID: _32:_1 Shape_MNK: (_16,_8,_16) LayoutA_TV: ((_4,_8),(_2,_2,_2)):((_32,_1),(_16,_8,_128)) LayoutB_TV: ((_4,_8),(_2,_2)):((_16,_1),(_8,_64)) LayoutC_TV: ((_4,_8),(_2,_2)):((_32,_1),(_16,_8))分配到线程，每个线程处理的元素数量为：A 矩阵为 2 x 2 x 2 = 8 个元素，B 矩阵为 2 x 2 = 4 个元素，得到 C 矩阵中 2 x 2 = 4 个元素。 📌 线程数计算：$ThrNum_{A} = 16 \\times 16 \\div 8 = 32$，$ThrNum_{B} = 16 \\times 8 \\div 4 = 32$，$ThrNum_{C} = 16 \\times 8 \\div 4 = 32$。1.2. Tiled MMA 配置 // Configure tiled MMA. using MmaTraits = cute::MMA_Traits\u0026lt;cute::SM80_16x8x16_F16F16F16F16_TN\u0026gt;; using MmaAtomShape = MmaTraits::Shape_MNK; auto const mma_atom = cute::MMA_Atom\u0026lt;MmaTraits\u0026gt;{}; auto const mma_atom_shape = MmaAtomShape{}; // Repeating the mma atom along the M, N, and K dimensions. // This increases the number of threads to process the tiled MMA. constexpr int MMA_LAYOUT_M{2}; constexpr int MMA_LAYOUT_N{2}; constexpr int MMA_LAYOUT_K{1}; auto mma_layout{cute::make_layout( cute::make_shape(cute::Int\u0026lt;MMA_LAYOUT_M\u0026gt;{}, cute::Int\u0026lt;MMA_LAYOUT_N\u0026gt;{}, cute::Int\u0026lt;MMA_LAYOUT_K\u0026gt;{}))}; // Repeating the mma processing along the M, N, and K dimensions. // This does not increase the number of threads to process the tiled MMA. // But the number of registers required for processing the tiled MMA increases. constexpr int NUM_MMA_TILE_M{1}; constexpr int NUM_MMA_TILE_N{2}; constexpr int NUM_MMA_TILE_K{1}; constexpr int MMA_TILE_M{cute::get\u0026lt;0\u0026gt;(mma_atom_shape) * MMA_LAYOUT_M * NUM_MMA_TILE_M}; constexpr int MMA_TILE_N{cute::get\u0026lt;1\u0026gt;(mma_atom_shape) * MMA_LAYOUT_N * NUM_MMA_TILE_N}; constexpr int MMA_TILE_K{cute::get\u0026lt;2\u0026gt;(mma_atom_shape) * MMA_LAYOUT_K * NUM_MMA_TILE_K}; auto mma_tile{cute::make_tile(cute::Int\u0026lt;MMA_TILE_M\u0026gt;{}, cute::Int\u0026lt;MMA_TILE_N\u0026gt;{}, cute::Int\u0026lt;MMA_TILE_K\u0026gt;{})}; auto tiled_mma{cute::make_tiled_mma(mma_atom, mma_layout, mma_tile)};在 M 维度上，MMA Atom 重复 2 次，在 N 维度上重复 2 次，在 K 维度上重复 1 次。一共需要 2 x 2 x 1 = 4 个 MMA Atom 来处理这个 tiled MMA。每个 Atom 由一个 warp（32 个线程）处理，整个 tiled MMA 由 4 个 warp（128 个线程）处理。即 ThrLayoutVMNK = (_32, _2, _2, _1):(_1, _32, _64, _0)。经过此配置后，能处理的 MNK 规模为 $(M’ \\times 2) \\times (N’ \\times 2) \\times (K’ \\times 1) = 32 \\times 16 \\times 16$。另外，通过配置 PermutationMNK（对应以前版本的 ValLayoutMNK），使得一个 tiled MMA 在 M/N/K 方向上处理更多的元素（即一个线程处理更多的元素）。这里配置 N 维度上乘以 2，得到该 tiled MMA 处理的 MNK 规模为 $32 \\times 32 \\times 16$，即 PermutationMNK: (_32, _32, _16)。tiled_mmaTiledMMA ThrLayoutVMNK: (_32,_2,_2,_1):(_1,_32,_64,_0) PermutationMNK: (_32,_32,_16)MMA_Atom ThrID: _32:_1 Shape_MNK: (_16,_8,_16) LayoutA_TV: ((_4,_8),(_2,_2,_2)):((_32,_1),(_16,_8,_128)) LayoutB_TV: ((_4,_8),(_2,_2)):((_16,_1),(_8,_64)) LayoutC_TV: ((_4,_8),(_2,_2)):((_32,_1),(_16,_8))1.3. Tiled MMA 划分总结以$M \\times N \\times K = 128 \\times 128 \\times 32$，以及 MMA Atom SM80_16x8x16_F16F16F16F16_TN 为例，MNK 三个维度分块划分（计算公式），可以按如下几种：1.3.1 划分方式一$\\frac{M}{M’} \\times \\frac{N}{N’} \\times \\frac{K}{K’}$ = $\\frac{128}{16} \\times \\frac{128}{8} \\times \\frac{32}{16}$ = $8 \\times 16 \\times 2$ = $256$，即 MMA Atom 在 M 维度重复 8 次，N 维度 重复16 次，K 维度重复 2 次。Atom 总共需要循环 256 次。1.3.2 划分方式二 $\\text{MMA_M} = \\frac{M}{M’ \\times \\text{MMA_TILE_M}} = \\frac{128}{16 \\times 2} = 4$ $\\text{MMA_N} = \\frac{N}{N’ \\times \\text{MMA_TILE_N}} \\times \\text{NUM_MMA_TILE_N} = \\frac{128}{8 \\times 2 \\times 2} \\times 2 = 8$ $\\text{MMA_K} = \\frac{K}{K’ \\times \\text{MMA_TILE_K}} = \\frac{32}{16 \\times 1} = 2$分别得到A_partition:(MMA, MMA_M, MMA_K) = (8, 4, 2)，B_partition:(MMA, MMA_N, MMA_K) = (4, 8, 2)，C_partition:(MMA, MMA_M, MMA_N) = (4, 4, 8)。即对A tile，Atom在M方向上重复4次，在K方向上重复2次；针对B tile，在N方向上重复8次，在K方向上重复2次；针对C tile，在M方向上重复4次，在N方向上重复8次。计算结果与代码中的打印信息一致（partition_fragment_A/B/C、partition_A/B/C，以及TiledCopy::partition_S）。或者其他配置方式，类似上面两种计算方式。1.3.3 总结一个 thread block 如何划分一个 tiled MMA，从性能上需要综合考虑以下几个因素： SM 上可用的寄存器数量，来确定一个 thread block 中可以有多少个线程来处理这个 tiled MMA。每个线程处理的元素数量越多，需要的寄存器数量就越多。 在寄存器够用的情况下，尽量让一个 thread block 中的线程数量能够充分利用 SM 上的计算资源（即 CUDA Core、Tensor Core）。每个 tiled MMA 需要多少个线程来处理，取决于 MMA Atom 的配置以及 tiled MMA 的配置。2. 内存分块以及 TiledCopy在将 thread block 对应的 tile 内存再次分解为线程 sub-tile 过程中，使用 CuTe 分块，有三种方式： 使用 partition 分块，得到 SMEM / GMEM slice。 使用 partition_fragment 分块，得到寄存器片段（register fragment）。 使用 TiledCopy / ldmatrix 进行分块以及传输。2.1. 使用 partition、partition_fragment 分块partition_A/B/C使用 partition 方法，获取原始 layout 的分块，即生成一个 slice。分块之后，数据还是在 SMEM / GMEM 中，且保留了原有的 tile 的 stride 信息。 auto thread_mma{tiled_mma.get_slice(THREAD_IDX)}; auto thread_layout_C_smem_tensor_A_no_tiled_copy{thread_mma.partition_A(smem_tensor_A)}; // (MMA, MMA_M, MMA_K) auto thread_layout_C_smem_tensor_B_no_tiled_copy{thread_mma.partition_B(smem_tensor_B)}; // (MMA, MMA_N, MMA_K) auto thread_layout_C_smem_tensor_C_no_tiled_copy{thread_mma.partition_C(smem_tensor_C)}; // (MMA, MMA_M, MMA_N)打印信息如下：thread_layout_C_smem_tensor_A_no_tiled_copyptr[16b](0x5c6c073e78c0) o ((_2,_2,_2),_4,_2):((_128,_8,_1024),_32,_2048)thread_layout_C_smem_tensor_B_no_tiled_copyptr[16b](0x5c6c073e98d0) o ((_2,_2),_8,_2):((_128,_1024),_16,_2048)thread_layout_C_smem_tensor_C_no_tiled_copyptr[16b](0x5c6c073eb8e0) o ((_2,_2),_4,_8):((_128,_8),_32,_2048)其中，含义如下： 维度 含义 MMA (第0维) 一次 tiled MMA 计算中该线程负责的元素 MMA_M (第1维) 沿 M 方向需要循环的次数 MMA_K (第2维) 沿 K 方向需要循环的次数 具体到 A/B/C 矩阵上，含义如下： 矩阵 shape 解释 A ((_2,_2,_2), _4, _2) MMA=8个元素, M循环4次, K循环2次 B ((_2,_2), _8, _2) MMA=4个元素, N循环8次, K循环2次 C ((_2,_2), _4, _8) MMA=4个元素, M循环4次, N循环8次 按线程切分之后，保留的第一个 mode，另外两个 mode 含义是（以 A 为例），MMA Atom 按 M 维度循环 4 次，K 维度循环 2 次。对于 B/C 矩阵类似。 📌 循环次数计算方式如下： SMEM tile：M * N * K = 128 * 128 * 32。 tiled_mma：M’ * N’ * K’ = 32 * 32 * 16。 得到计算公式： MMA_M = 128 / 32 = 4 MMA_N = 128 / 32 * 2 = 8 (由于配置了 permutation，使得 tiled MMA 在 N 维度上处理更多的元素，所以循环次数增加了) MMA_K = 32 / 16 = 2 💡 从打印信息看出，不论是直接使用 thread_layout_C_smem_tensor_A/B_no_tiled_copy，还是将其拷贝到寄存器，由于不连续，导致线程每次访问 2 * 2 * 2 = 8 个元素时，需要分开拷贝，即分 8 次访问 SMEM / GMEM 来加载数据到寄存器中。partition_fragment_A/B/Cpartition_fragment 则创建寄存器片段（register fragment），以复用寄存器数据。分块之后，数据在寄存器中，且不保留原有 tile 的 stride 信息，而是变为紧凑型布局。线程在做 gemm 之前，使用 copy 将数据从 SMEM 中加载到寄存器中。 auto thread_layout_C_register_tensor_A{thread_mma.partition_fragment_A(smem_tensor_A)}; // (MMA, MMA_M, MMA_K) auto thread_layout_C_register_tensor_B{thread_mma.partition_fragment_B(smem_tensor_B)}; // (MMA, MMA_N, MMA_K) auto thread_layout_C_register_tensor_C{thread_mma.partition_fragment_C(smem_tensor_C)}; // (MMA, MMA_M, MMA_N)打印信息如下：thread_layout_C_register_tensor_Aptr[16b](0x7ffc34e465f0) o ((_2,_2,_2),_4,_2):((_1,_2,_4),_8,_32)thread_layout_C_register_tensor_Bptr[16b](0x7ffc34e46670) o ((_2,_2),_8,_2):((_1,_2),_4,_32)thread_layout_C_register_tensor_Cptr[16b](0x7ffc34e466f0) o ((_2,_2),_4,_8):((_1,_2),_4,_16)总结使用 partition / partition_fragment 分块，分块数据可能不连续，导致需要多次访问 SMEM / GMEM。使用合适的 ldmatrix 可以一次将一次计算所需要的 sub-tile 数据拷贝到寄存器。 📌 partition_A/B/C 用于创建 slice view，partition_fragment_A/B/C 生成寄存器 tensor，即分配寄存器存储空间。2.2. 使用 TiledCopy / ldmatrix 进行分块以及传输前面使用的 partition / partition_fragment 方式，直接使用 TiledMMA/ThrMMA 分块得到的。使用 TiledCopy / ldmatrix 进行分块以及传输。设置 TiledCopy 需要的 Copy_Atom、CopyTraits（layout 信息），并使其传输的 sub-tile 大小与 tiled MMA 的计算需求一致。2.2.1. Copy Atom 配置针对 A、B，使用 cute::SM75_U16x8_LDSM_T 生成 Copy Atom 来复制 A、B 的 sub-tile，其对应的 PTX 为：ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {r0, r1, r2, r3}, [addr];即使用一个 warp（32 个线程）同时加载一个 sub-tile 的数据到寄存器中。 ldmatrix 只支持16位数据，不支持32位数据；另外 ldmatrix 是 PTX 指令，对应到 ncu 中看到的 LDSM 指令。x1 表示使用前 8 个线程，x2、x4 依次类推。m8n8 表示 load 操作的子块大小为 8x8，x4 则表示同时加载 4 个子块。cute::SM75_U16x8_LDSM_T 打印信息如下：copy_atom_ACopy_Atom ThrID: _32:_1 ValLayoutSrc: (_32,_8):(_8,_1) ValLayoutDst: ((_4,_8),(_1,_2,_4)):((_16,_1),(_1,_8,_64)) ValLayoutRef: ((_4,_8),(_1,_2,_4)):((_16,_1),(_1,_8,_64)) ValueType: 16bcopy_atom_BCopy_Atom ThrID: _32:_1 ValLayoutSrc: (_32,_8):(_8,_1) ValLayoutDst: ((_4,_8),(_1,_2,_4)):((_16,_1),(_1,_8,_64)) ValLayoutRef: ((_4,_8),(_1,_2,_4)):((_16,_1),(_1,_8,_64)) ValueType: 16b在前面 TiledMMA 的配置中，经过 AtomLayout 以及 Permutation，最终得到的 tiled MMA 处理的 sub-tile 大小为 $M’ \\times N’ \\times K’ = 32 \\times 32 \\times 16$。其中 A 矩阵的 sub-tile 大小为 $M’ \\times K’ = 32 \\times 16$，B 矩阵的 sub-tile 大小为 $K’ \\times N’ = 16 \\times 32$。使用SM75_U16x8_LDSM_T，其一次拷贝的 sub-tile 大小为 $32 \\times 8 = 256$，即对应上面打印信息中的 ValLayoutSrc: (_32, _8):(_8, _1)。因此，TiledCopy 的 Copy Atom 配置满足一次加载 执行一个 Tiled MMA 所需要的数据。另外，在 Tiled MMA 配置中，B 配置了 permutation，使得 B 矩阵的 sub-tile 大小与 A 矩阵大小一致，否则针对 B，就需要选择其他的 Copy Atom 来满足 tiled MMA 的计算需求。2.2.2. TiledCopy 以及 ThreadCopy 创建创建 TiledCopy，Copy_Atom 配置如上SM75_U16x8_LDSM_T，TV-Layout 则由上述的 Tiled MMA 给出。代码如下： auto copy_atom_A = cute::Copy_Atom\u0026lt;cute::SM75_U16x8_LDSM_T, TA\u0026gt;{}; auto copy_atom_B = cute::Copy_Atom\u0026lt;cute::SM75_U16x8_LDSM_T, TB\u0026gt;{}; auto smem_tiled_copy_A{cute::make_tiled_copy_A(copy_atom_A, tiled_mma)}; auto smem_tiled_copy_B{cute::make_tiled_copy_B(copy_atom_B, tiled_mma)};Tiled Copy 信息如下：smem_tiled_copy_ATiledCopy Tiler_MN: (_32,_16) TiledLayout_TV: ((_4,_8,_2,_2),((_2,_2,_2),(_1,_1))):((_64,_1,_16,_0),((_32,_8,_256),(_0,_0)))Copy_Atom ThrID: _32:_1 ValLayoutSrc: (_32,_8):(_8,_1) ValLayoutDst: ((_4,_8),(_1,_2,_4)):((_16,_1),(_1,_8,_64)) ValLayoutRef: ((_4,_8),(_1,_2,_4)):((_16,_1),(_1,_8,_64)) ValueType: 16bsmem_tiled_copy_BTiledCopy Tiler_MN: (_32,_16) TiledLayout_TV: ((_4,_8,_2,_2),((_2,_2),(_2,_1))):((_64,_1,_0,_8),((_32,_256),(_16,_0)))Copy_Atom ThrID: _32:_1 ValLayoutSrc: (_32,_8):(_8,_1) ValLayoutDst: ((_4,_8),(_1,_2,_4)):((_16,_1),(_1,_8,_64)) ValLayoutRef: ((_4,_8),(_1,_2,_4)):((_16,_1),(_1,_8,_64)) ValueType: 16b使用 Thread Copy 以及创建本线程 sub-tile分为两步，首先使用 ThrCopy\u0026lt;…\u0026gt;::partition_S 获取线程的 tensor，再使用 ThrCopy\u0026lt;…\u0026gt;::retile_D 获取调整 shape 之后的 tensor。代码如下： auto smem_thread_copy_A{smem_tiled_copy_A.get_slice(THREAD_IDX)}; auto smem_thread_copy_B{smem_tiled_copy_B.get_slice(THREAD_IDX)}; auto thread_layout_C_smem_tensor_A_tiled_copy{smem_thread_copy_A.partition_S(smem_tensor_A)}; auto thread_layout_C_smem_tensor_B_tiled_copy{smem_thread_copy_B.partition_S(smem_tensor_B)}; auto thread_layout_C_register_tensor_A_copy_view{smem_thread_copy_A.retile_D(thread_layout_C_register_tensor_A)}; auto thread_layout_C_register_tensor_B_copy_view{smem_thread_copy_B.retile_D(thread_layout_C_register_tensor_B)};thread_layout_C_smem_tensor_A/B_tiled_copy 作为源 tensor，他们的信息如下：thread_layout_C_smem_tensor_A_tiled_copyptr[16b](0x57b7b93248c0) o ((_8,_1),_4,_2):((_1,_0),_32,_2048)thread_layout_C_smem_tensor_B_tiled_copyptr[16b](0x57b7b93268d0) o ((_8,_1),_4,_2):((_1,_0),_32,_2048)作为目的 tensor 的 thread_layout_C_register_tensor_A/B（由 ThrMMA\u0026lt;…\u0026gt;::partition_fragment_A/B 获取），其信息如下：thread_layout_C_register_tensor_Aptr[16b](0x7ffc34e465f0) o ((_2,_2,_2),_4,_2):((_1,_2,_4),_8,_32)thread_layout_C_register_tensor_Bptr[16b](0x7ffc34e46670) o ((_2,_2),_8,_2):((_1,_2),_4,_32)所以需要对目的 tensor 进行调整，使得其 shape 与源 tensor 一致。调整之后的目的 tensor thread_layout_C_register_tensor_A/B_copy_view 的信息如下：thread_layout_C_register_tensor_A_copy_viewptr[16b](0x7ffd7a129350) o ((_8,_1),_4,_2):((_1,_0),_8,_32)thread_layout_C_register_tensor_B_copy_viewptr[16b](0x7ffd7a1293d0) o ((_8,_1),_4,_2):((_1,_0),_8,_32) 📌 通过 ThrCopy::partition_S/D、ThrCopy::partition_fragment_A/B等得到的分块，cute::copy只负责Shape(CPY,CPY_M, CPY_K)的维度；额外的维度需要手动切片，并在外部执行循环。ThrMMA得到的分片同理。2.3. 总结使用 partition / partition_fragment 分块，分块数据可能不连续，导致需要多次访问 SMEM / GMEM。使用合适的 ldmatrix 可以一次将一次计算所需要的 sub-tile 数据拷贝到寄存器。使用 TiledCopy / ldmatrix 的流程是： 配置 Copy Atom / Copy Traits，使得一次 tiled copy 的数据量满足一次 tiled MMA 计算的需求。 创建 TiledCopy，并获取线程对应的 ThreadCopy。 使用 ThreadCopy 的 partition_S 获取线程的 tensor。 使用 ThreadCopy 的 retile_D 获取调整 shape 之后的 tensor，使得其与 tiled copy 的源 tensor shape 一致。 使用 ThreadCopy 执行 copy 操作，将数据从 SMEM 中加载到寄存器中。 ldmatrix 导致的 smem bank 冲突，以及如何通过 swizzle 来规避，见博客系列中的：CUTLASS-Cute 初步(4.1)：MMA Swizzle -- MMA、ldmatrix、smem swizzle。参考及资料 CuTe Tiled MMA：Mao Lei 博客 CuTe General Matrix Multiplication：Mao Lei GitHub GEMM 代码 0t_mma_atom：CUTLASS-CuTe 0t_mma_atom 文档 待阅读 CuTe ldmatrix：Mao Lei 博客 ldmatrix 待阅读 ldmatrix指令例子和其带来的smem bank冲突 待阅读 ldmatrix与swizzle（笔记） 待阅读" },

  { "title": "CUTLASS-Cute 初步(3)：TiledCopy 以及 TiledMMA", "url": "/blog/2025/Cute%E5%88%9D%E6%AD%A53-TiledCopy-TiledMMA/", "categories": "CUDA",
  "tags": "CUDA", "date": "2025-02-26", "content":
  "1. Cute TiledCopy tiled_copy.cu：官方示例层次化的copy抽象，分为几个可组合的层次： CopyOperation：NVidia在不同的硬件架构、不同的存储层次之间数据搬运提供了不同的指令，如ldmatrix和LDS等，还有针对Ampere架构的cp.async等。 Copy_Traits：主要提供拷贝的metadata信息：源Thread-Value Layout、目标Thread-Value Layout等。 Copy_Atom：封装一个完整的最小数据搬运单元，包含CopyOperation和Copy_Traits。 TiledCopy：根据线程布局，重复使用Copy_Atom完成分块数据搬运计算。 make_tiled_copy{_A|B|C}：提供用户级别的API接口；1.1. CopyOperationCopyOperation封装了执行一次数据搬运指令，以及所需要的指令参数。其中指令参数（源参数、目的参数）描述了参数的类型以及个数，供API层级的copy函数使用。示例：struct SM75_U16x8_LDSM_T{ using SRegisters = uint128_t[1]; using DRegisters = uint32_t[4]; CUTE_HOST_DEVICE static void copy(uint128_t const\u0026amp; smem_src, uint32_t\u0026amp; dst0, uint32_t\u0026amp; dst1, uint32_t\u0026amp; dst2, uint32_t\u0026amp; dst3) {#if defined(CUTE_ARCH_LDSM_SM75_ACTIVATED) uint32_t smem_int_ptr = cast_smem_ptr_to_uint(\u0026amp;smem_src); asm volatile (\"ldmatrix.sync.aligned.x4.trans.m8n8.shared.b16 {%0, %1, %2, %3}, [%4];\\n\" : \"=r\"(dst0), \"=r\"(dst1), \"=r\"(dst2), \"=r\"(dst3) : \"r\"(smem_int_ptr));#else CUTE_INVALID_CONTROL_PATH(\"Trying to use ldmatrix without CUTE_ARCH_LDSM_SM75_ACTIVATED.\");#endif }};1.2. Copy_Traits执行数据拷贝操作，有如下需求：向量化指令，合并访存。另外一般是以一个warp为单位，即需要考虑SRC/DST是如何分配给一个warp中的线程的。典型的如ldmatrix指令：一个phase要求一个warp中一组8个线程（如T0~T7），分别提供源SMEM地址uint128_t，以及32个线程的目的REG（uint32_t）。Copy_Traits提供了源Thread-Value Layout、目的Thread-Value Layout等信息，即描述了线程如何访问数据的布局。示例：template \u0026lt;\u0026gt;struct Copy_Traits\u0026lt;SM75_U16x8_LDSM_T\u0026gt;{ // Logical thread id to thread idx (warp) using ThrID = Layout\u0026lt;_32\u0026gt;; // Map from (src-thr,src-val) to bit using SrcLayout = Layout\u0026lt;Shape \u0026lt; _32,_128\u0026gt;, Stride\u0026lt;_128, _1\u0026gt;\u0026gt;; // Map from (dst-thr,dst-val) to bit using DstLayout = Layout\u0026lt;Shape \u0026lt;Shape \u0026lt; _4, _8\u0026gt;,Shape \u0026lt;_16, _2, _4\u0026gt;\u0026gt;, Stride\u0026lt;Stride\u0026lt;_256,_16\u0026gt;,Stride\u0026lt; _1,_128,_1024\u0026gt;\u0026gt;\u0026gt;; // Reference map from (thr,val) to bit using RefLayout = DstLayout;};1.3. TiledCopy通过TiledCopy，将Copy_Atom扩展到更多线程。需要提供两个参数：Thread-Value Layout，以及Tiler_MN，其中**Tiler_MN表示最终处理的M、N维度的tile大小，TV_Layout则包含了需要的线程数量，以及线程如何访问数据的布局。TiledCopy实现如下：template \u0026lt;class Copy_Atom, class LayoutCopy_TV, // (tid,vid) -\u0026gt; coord [Need not be 2D...] class ShapeTiler_MN\u0026gt; // coord spacestruct TiledCopy : Copy_Atom{ // Layout information from the CopyAtom using AtomThrID = typename Copy_Atom::ThrID; // thrid -\u0026gt; thr_idx using AtomLayoutSrc = typename Copy_Atom::ValLayoutSrc; // (thr,val) -\u0026gt; offset using AtomLayoutDst = typename Copy_Atom::ValLayoutDst; // (thr,val) -\u0026gt; offset using AtomLayoutRef = typename Copy_Atom::ValLayoutRef; // (thr,val) -\u0026gt; offset using AtomNumThr = decltype(size\u0026lt;0\u0026gt;(AtomLayoutRef{})); using AtomNumVal = decltype(size\u0026lt;1\u0026gt;(AtomLayoutRef{})); // Layout information for the TiledCopy using Tiler_MN = ShapeTiler_MN; using TiledLayout_TV = LayoutCopy_TV; using TiledNumThr = decltype(size\u0026lt;0\u0026gt;(TiledLayout_TV{})); using TiledNumVal = decltype(size\u0026lt;1\u0026gt;(TiledLayout_TV{})); CUTE_STATIC_ASSERT_V(TiledNumThr{} % AtomNumThr{} == Int\u0026lt;0\u0026gt;{}, \"TiledCopy uses too few thrs for selected CopyAtom\"); CUTE_STATIC_ASSERT_V(TiledNumVal{} % AtomNumVal{} == Int\u0026lt;0\u0026gt;{}, \"TiledCopy uses too few vals for selected CopyAtom\"); // ....};从代码中，可以看到，TiledCopy中Thread数量比如为CopyAtom中线程数量的整数倍，Value数量也是CopyAtom中Value数量的整数倍。1.4. ThrCopy前面几个小节所描述的CopyOperation、Copy_Traits、Copy_Atom、TiledCopy等都是在编译期定义的抽象，描述了数据搬运的指令、线程访问数据的布局，以及如何将一个基本的搬运单元扩展到更多线程上。ThrCopy则是运行时根据线程 id 获取每个线程实际执行的数据搬运任务。通过调用TiledCopy::get_slice接口，生成一个ThrCopy对象。通过ThrCopy的接口partition_S/D，获取线程拷贝的SRC/DST数据的Layout，比如(CPY, CPY_M, CPY_K)。另外，待切分的Tensor可能其Shape不满足cute::copy的要求，使用ThrCopy::retile_S/D接口，重新切分成满足要求的Shape，retile前后Layout，其size(layout)与cosize(layout)保持一致，即指向同一块GMEM/SMEM且大小不变。比如从TiledMMA创建的TiledCopy，其DST的Layout与SRC数据的Layout不匹配，需要使用retile_D接口重新切分成与SRC Tensor一致的Shape。ThrCopy的定义如下：template \u0026lt;class TiledCopy, class ThrIdx\u0026gt;struct ThrCopy { auto partition_S(Tensor\u0026amp;\u0026amp; stensor); auto partition_D(Tensor\u0026amp;\u0026amp; dtensor); auto retile_S(Tensor\u0026amp;\u0026amp; stensor); auto retile_D(Tensor\u0026amp;\u0026amp; stensor);};1.5. make_tiled_copy{_A|B|C}使用make_tiled_copy可以直接使用Thread-Value Layout创建一个TiledCopy对象，其中Tiler_MN可以从TV-Layout推导。如果从一个TiledMMA创建TiledCopy，由于TiledMMA表示的是(M, N, K)三个维度（见下面章节TiledMMA的介绍），所以针对A/B/C，分别使用对应的make_tiled_copy_A/B/C接口创建TiledCopy对象。使用示例： using T = cute::half_t; using VectorType = cute::uint128_t; using CopyOp = cute::SM80_CP_ASYNC_CACHEALWAYS\u0026lt;VectorType\u0026gt;; constexpr auto thread_shape_A = cute::make_shape(cute::Int\u0026lt;16\u0026gt;{}, cute::Int\u0026lt;8\u0026gt;{}); // (16M, 8K) constexpr auto thread_stride_A = cute::make_stride(cute::Int\u0026lt;1\u0026gt;{}, cute::size\u0026lt;0\u0026gt;(thread_shape_A)); // (1, 16)，M-major constexpr auto thread_layout_A = cute::make_layout(thread_shape_A, thread_stride_A); constexpr auto NUM_ELEMENTS_A = sizeof(VectorType) / sizeof(T); constexpr auto vector_shape_A = cute::make_shape(cute::Int\u0026lt;NUM_ELEMENTS_A\u0026gt;{}, cute::Int\u0026lt;1\u0026gt;{}); // (8M, 1K) constexpr auto vector_stride_A = cute::make_stride(cute::Int\u0026lt;1\u0026gt;{}, cute::size\u0026lt;0\u0026gt;(vector_shape_A)); constexpr auto vector_layout_A = cute::make_layout(vector_shape_A, vector_stride_A); constexpr auto gmem_tiled_copy_A = cute::make_tiled_copy(cute::Copy_Atom\u0026lt;CopyOp, T\u0026gt;{}, thread_layout_A, vector_layout_A); // constexpr auto tiler_mn = decltype(gmem_tiled_copy_A)::Tiler_MN{}; cute::print(gmem_tiled_copy_A); std::cout \u0026lt;\u0026lt; \"-----------------------------\" \u0026lt;\u0026lt; std::endl; cute::print_latex(gmem_tiled_copy_A);打印信息如下：TiledCopy Tiler_MN: (_128,_8) TiledLayout_TV: (_128,_8):(_8,_1)Copy_Atom ThrID: _1:_0 ValLayoutSrc: (_1,_8):(_0,_1) ValLayoutDst: (_1,_8):(_0,_1) ValLayoutRef: (_1,_8):(_0,_1) ValueType: 16b含义如下： Tiler_MN: (_128,_8)：表示分块的M维度大小为128，K维度大小为8。 TiledLayout_TV: (_128,_8):(_8,_1)：表示线程访问数据的布局，其中(_128,_8)表示线程数量为128，且每个线程处理8个元素，Stride信息表明了每个线程处理的数据在内存中是连续的。 Tiler_MN和TiledLayout_TV，其size表示整个tile的大小，即M维度和K维度的总元素数量。1.6. cute::copy – 执行数据搬运指令copy 函数是拷贝的实际执行函数，完成线程指令的执行：void copy(TiledCopy const\u0026amp; copy, Tensor const\u0026amp; src, Tensor\u0026amp; dst);void copy_if(TiledCopy const\u0026amp; copy, PrdTensor const\u0026amp; pred, Tensor const\u0026amp; src, Tensor\u0026amp; dst);2. MMAAtom 以及 TiledMMA分块 MMA 抽象，将 MMA_Atom 分为几个可组合的层次： MMAOperation：是对 D=A*B + C 的 PTX 指令封装，以使用不同的数据类型以及 PTX 指令，包括使用 CUDA Core / Tensor Core。如 UniversalFMA\u0026lt;\u0026gt;、SM80_16x8x8_F32F16F16F32_TN。 MMA_Traits：和 Copy_Traits 类似，提供了 MMAOperation 类型没有提供，但是其使用者 MMA_Atom 却需要的起到桥梁作用的信息。如数据类型信息，TV layout 信息。 MMA_Atom：将 MMAOperation 和 MMA_Traits 结合，并提供 fragment 划分接口。 TiledMMA：根据 LayoutTile_TV 切分的线程布局，重复使用 MMA_Atom 完成分块矩阵乘加计算。 ThrMMA：完成实际的生成线程对应的 tensor。2.1. MMAOperation以SM80_16x8x8_F32F16F16F32_TN为例，封装了 SM80 架构下，16x8x8 大小的矩阵乘加指令 D=A * B + C，数据类型为 A:F16、B:F16、C:F32、D:F32。A 矩阵 row-major，B 矩阵 column-major。 BLAS 中约定 normal 矩阵为列优先。T(transpose) 表示使用转置矩阵，即 row-major 存储。 下图原图见 Thakkar_BLISRetreat2023.pdf 第 30 页。 SM80_16x8x8_F32F16F16F32_TN 对应的 inverse TV-Layout 如下： inverse TV-Layout 表示 element coordinate -\u0026gt; thread index的映射关系。SM80_16x8x8_F32F16F16F32_TN 对应的 MMA_Atom 信息如下：MMA_Atom ThrID: _32:_1 Shape_MNK: (_16,_8,_8) LayoutA_TV: ((_4,_8),(_2,_2)):((_32,_1),(_16,_8)) LayoutB_TV: ((_4,_8),_2):((_16,_1),_8) LayoutC_TV: ((_4,_8),(_2,_2)):((_32,_1),(_16,_8))对应的代码如下：print(MMA_Atom\u0026lt;SM80_16x8x8_F32F16F16F32_TN\u0026gt;{});// using MMA = MMA_Traits\u0026lt;SM80_16x8x8_F32F16F16F32_TN\u0026gt;;// print(\"ALayout: \"), print(typename MMA::ALayout{}), print(\"\\n\");// print(\"BLayout: \"), print(typename MMA::BLayout{}), print(\"\\n\");// print(\"CLayout: \"), print(typename MMA::CLayout{}), print(\"\\n\");MMA_Atom\u0026lt;SM80_16x8x8_F32F16F16F32_TN\u0026gt; mma;print_latex(mma);/* 或者如下写法TiledMMA tiled_mma = make_tiled_mma(SM80_8x8x4_F64F64F64F64_TN{});print_latex(tiled_mma);*/ 📌 SM80_16x8x8_F32F16F16F32_TN 使用一个 warp（32 个线程）处理 MNK 规模为 16 * 8 * 8 的一个 sub-tile。一个线程处理 A 中的 2 * 2 个数据，即LayoutATV 中的第二个 mode (_2, _2)，则线程数位 $ThrNum{A} = 16 \\times 8 \\div 4 = 32$。同理可以知道，每个线程处理 B、C 中多少个数据，以及需要的线程数。CUDA PTX 文档也给出了指令 m16n8k8 的布局信息：9.7.14.5.7. Matrix Fragments for mma.m16n8k8。2.2. MMA_TraitsMMA_Traits 提供数据类型信息，以及 TV layout 信息，比如需要根据 MMAOperation 中定义的指令，补充 A/B/C 的 layout 信息。需要提供的信息如下：using ElementDVal = // Logical A-value typeusing ElementAVal = // Logical B-value typeusing ElementBVal = // Logical C-value typeusing ElementCVal = // Logical D-value typeusing ElementAFrg = // A-type consumed by MMA (if ommitted, same as ElementAVal)using ElementBFrg = // B_type consumed by MMA (if ommitted, same as ElementBVal)using ElementCFrg = // C_type consumed by MMA (if ommitted, same as ElementCVal)using Shape_MNK = // Logical MxNxK shape of the MMAusing ThrID = // Logical thread id (tid) -\u0026gt; tidxusing ALayout = // (Logical thread id (tid), Logical value id (vid)) -\u0026gt; Flat MK-coordusing BLayout = // (Logical thread id (tid), Logical value id (vid)) -\u0026gt; Flat NK-coordusing CLayout = // (Logical thread id (tid), Logical value id (vid)) -\u0026gt; Flat MN-coord2.3. MMA_AtomMMA_Atom 封装了 MMAOperation 和 MMA_Traits。创建寄存器 fragment提供了创建寄存器 fragment 的接口 make_fragment_A/B/C：template \u0026lt;class CTensor\u0026gt;static constexpr auto make_fragment_C(CTensor\u0026amp;\u0026amp; ctensor);template \u0026lt;class ATensor\u0026gt;static constexpr auto make_fragment_A(ATensor\u0026amp;\u0026amp; atensor);template \u0026lt;class BTensor\u0026gt;static constexpr auto make_fragment_B(BTensor\u0026amp;\u0026amp; btensor);调用 FMA 指令提供 call 接口，调用 MMAOperation 指令：template \u0026lt;class TD, class DLayout, class TA, class ALayout, class TB, class BLayout, class TC, class CLayout\u0026gt;constexpr void call(Tensor\u0026lt;TD, DLayout\u0026gt;\u0026amp; D, Tensor\u0026lt;TA, ALayout\u0026gt; const\u0026amp; A, Tensor\u0026lt;TB, BLayout\u0026gt; const\u0026amp; B, Tensor\u0026lt;TC, CLayout\u0026gt; const\u0026amp; C) const; TODO: 调用之前进行一个 unpack 操作？2.4. TiledMMATiledMMA的模版参数表达了TiledMMA对MMA_Atom上的扩展逻辑：AtomLayoutMNK表示M、N、K方向上分别扩展的倍数（同时处理的MNK问题的规模也会相应倍数的增加）。TiledMMA::get_slice、TiledMMA::get_thread_slice函数功过给定线程 id 则获取线程对应到ThrMMA对象。template \u0026lt;class MMA_Atom, class AtomLayoutMNK, class PermutationMNK = Tile\u0026lt;Underscore,Underscore,Underscore\u0026gt;\u0026gt;struct TiledMMA : MMA_Atom { auto get_slice(ThrIdx const\u0026amp; thr_idx) const { auto thr_vmnk = thr_layout_vmnk_.get_flat_coord(thr_idx); return ThrMMA\u0026lt;TiledMMA, decltype(thr_vmnk)\u0026gt;{*this, thr_vmnk}; } auto get_thread_slice(ThrIdx const\u0026amp; thr_idx) const { return get_slice(thr_idx); } auto thrfrg_C(CTensor\u0026amp;\u0026amp; ctensor) const { .... }};三个模板参数的含义： 参数名 类型 说明 MMA_Atom 底层指令 定义单条 MMA 指令涉及的线程和值的布局 AtomLayoutMNK Layout\u0026lt;Shape\u0026lt;_2,_2,_1\u0026gt;\u0026gt; 在 M/N/K 方向上重复多少个 atom（分配更多线程） PermutationMNK Tile\u0026lt;_8, _16, _8\u0026gt;{} 每个线程在 M/N/K 方向上处理更多的值（不增加线程） AtomLayoutMNK决定如何将MMA_Atom复制到更多的线程上执行，且将MMA_Atom处理的线程扩展为size(AtomLayoutMNK)倍数。注意：AtomLayoutMNK是以倍数的形式给出的，且AtomLayoutMNK定义的是mma_layout。在cutlass / CuTe 3.4之前版本还有ValLayoutMNK参数，从3.4版本开始去掉该模板参数。PermutationMNK可以替代 ValLayoutMNK的功能。PermutationMNK 定义的是mma_tile，即定义这个TiledMMA最终处理的MNK规模，不影响线程的扩展（cAtomLayoutMNK已经定义了线程扩展的布局）。PermutationMNK 的展开讲述见下面章节。2.4.1. 四层 Layout 以及获取线程 fragment根据给定的 MMA_Atom、以及 AtomLayoutMNK 参数，生成一个四层 Layout 结构：using ThrLayoutVMNK = decltype(tiled_product(AtomThrID{}, AtomLayoutMNK{}));ThrLayoutVMNK thr_layout_vmnk_; Mode 0 (V): Threads within a single atom Mode 1 (M): Atom tiles in M dimension Mode 2 (N): Atom tiles in N dimension Mode 3 (K): Atom tiles in K dimension以 thrfrg_A 为例：// Tile a tensor or a layout from shape // (M,K,...) // to shape // ((ThrV,(ThrM,ThrK)),(FrgV,(RestM,RestK,...))) // where // ThrV: The threads local to an MMA. layout\u0026lt;0\u0026gt;(ThrLayoutVMNK): ThrV -\u0026gt; thread_idx // ThrM: The threads tiled in M. layout\u0026lt;1\u0026gt;(ThrLayoutVMNK): ThrM -\u0026gt; thread_idx // ThrK: The threads tiled in K. layout\u0026lt;3\u0026gt;(ThrLayoutVMNK): ThrK -\u0026gt; thread_idx // FrgV: The values local to an MMA. // RestM: The values tiled in M. // RestK: The values tiled in K.template \u0026lt;class ATensor\u0026gt;constexpr auto thrfrg_A(ATensor\u0026amp;\u0026amp; atensor) const;即得到的线程切分后的 subtile 布局为 ((ThrV,(ThrM,ThrK)),(FrgV,(RestM,RestK,…)))。2.4.2. TiledMMA 流程示意2.5. ThrMMATiledMMA 根据具体的线程 id 分解得到 ThrMMA 结构，提供 partition 函数接口，以及 partition_fragment 函数接口。如 Tensor C 为 BLK_M x BLK_N，则 partition_C 可以得到线程级别的任务，维度为 (MMA, MMA_M, MMA_N), MMA 表达了 TileMMA 一次能计算的单元，MMA_M, MMA_N 表达了 M 方向和 N 方向需要分块数量。partition_fragment 类函数是按照 partition 类函数返回的 Tensor 形状生成的对应的寄存器表示。template \u0026lt;class TiledMMA, class ThrVMNK\u0026gt;struct ThrMMA : TiledMMA { Tensor partition_C(Tensor C); Tensor partition_A(Tensor A); Tensor partition_B(Tensor B); Tensor partition_fragment_C(Tensor C); Tensor partition_fragment_A(Tensor A); Tensor partition_fragment_B(Tensor B);}2.6. Permutation：置换Permutation是一个Tiler，由三个独立的分量组成，分别作用于M、N、K维度。Permutation可以实现两种变换： 创建Tiler，即定义最终TiledMMA处理的MNK规模。 对M、N、K维度的逻辑坐标进行重新映射。以SM80_8x8x4_F64F64F64F64_TN为例，这个MMA_Atom的Inverse TV-Layout如下：TiledMMA ThrLayoutVMNK: (_32,_1,_1,_1):(_1,_0,_0,_0) PermutationMNK: (_,_,_)MMA_Atom ThrID: _32:_1 Shape_MNK: (_8,_8,_4) LayoutA_TV: ((_4,_8),_1):((_8,_1),_0) LayoutB_TV: ((_4,_8),_1):((_8,_1),_0) LayoutC_TV: ((_4,_8),_2):((_16,_1),_8)代码如下： using namespace cute; TiledMMA tiled_mma = make_tiled_mma(SM80_8x8x4_F64F64F64F64_TN{}); std::cout \u0026lt;\u0026lt; \"=== SM80_8x8x4_F64F64F64F64_TN ===\\n\"; print(tiled_mma), print(\"\\n\"); std::cout \u0026lt;\u0026lt; \"=== SM80_8x8x4_F64F64F64F64_TN latex format ===\\n\"; print_latex(tiled_mma); std::cout \u0026lt;\u0026lt; \"\\n\";2.6.1. PermutationMNK 的作用1：Tiler using namespace cute; TiledMMA tiled_mma = make_tiled_mma(SM80_8x8x4_F64F64F64F64_TN{}, Layout\u0026lt;Shape\u0026lt;_1, _1, _1\u0026gt;\u0026gt;{}, // AtomLayout Tile\u0026lt;_8, _16, _8\u0026gt;{}); // Tiler std::cout \u0026lt;\u0026lt; \"=== SM80_8x8x4_F64F64F64F64_TN ===\\n\"; print(tiled_mma), print(\"\\n\"); std::cout \u0026lt;\u0026lt; \"=== SM80_8x8x4_F64F64F64F64_TN latex format ===\\n\"; print_latex(tiled_mma); std::cout \u0026lt;\u0026lt; \"\\n\";TiledMMA ThrLayoutVMNK: (_32,_1,_1,_1):(_1,_0,_0,_0) PermutationMNK: (_8,_16,_8)MMA_Atom ThrID: _32:_1 Shape_MNK: (_8,_8,_4) LayoutA_TV: ((_4,_8),_1):((_8,_1),_0) LayoutB_TV: ((_4,_8),_1):((_8,_1),_0) LayoutC_TV: ((_4,_8),_2):((_16,_1),_8)2.6.2. PermutationMNK 的作用2：重新映射逻辑坐标通过对M、N、K这三个维度进行置换操作，即逻辑坐标重新映射，实现线程访问数据的合并，即合并访存。 using namespace cute; TiledMMA tiled_mma_perm = make_tiled_mma(SM80_8x8x4_F64F64F64F64_TN{}, Layout\u0026lt;Shape\u0026lt;_1, _1, _1\u0026gt;\u0026gt;{}, // AtomLayout Tile\u0026lt;_8, // Permutation on M, equivalent to 8:1, identity Layout\u0026lt;Shape\u0026lt;_2, _4, _2\u0026gt;, Stride\u0026lt;_1, _4, _2\u0026gt;\u0026gt;, // Permutation on N, size 16 _8\u0026gt;{}); std::cout \u0026lt;\u0026lt; \"=== SM80_8x8x4_F64F64F64F64_TN ===\\n\"; print(tiled_mma_perm), print(\"\\n\"); std::cout \u0026lt;\u0026lt; \"=== SM80_8x8x4_F64F64F64F64_TN latex format ===\\n\"; print_latex(tiled_mma_perm); std::cout \u0026lt;\u0026lt; \"\\n\";TiledMMA ThrLayoutVMNK: (_32,_1,_1,_1):(_1,_0,_0,_0) PermutationMNK: (_8,(_2,_4,_2):(_1,_4,_2),_8)MMA_Atom ThrID: _32:_1 Shape_MNK: (_8,_8,_4) LayoutA_TV: ((_4,_8),_1):((_8,_1),_0) LayoutB_TV: ((_4,_8),_1):((_8,_1),_0) LayoutC_TV: ((_4,_8),_2):((_16,_1),_8)2.6.3. 参考资料 0t_mma_atom：官方文档，讲了Permute操作，包括scatter permutation [QST] What is PermutationMNK in TiledMMA in CUTLASS 3.4 changes? 02_layout_algebra.md – Logical Divide 2-D Example 0t_mma_atom – TiledMMAs2.7. UniversalFMA 可以参考 Thakkar_BLISRetreat2023.pdf 第 26 页。UniversalFMA 是一个标量 FMA 操作的 MMAOperation 实现，定义如下：template \u0026lt;class D, class A = D, class B = A, class C = D\u0026gt;struct UniversalFMA { using DRegisters = D[1]; using ARegisters = A[1]; using BRegisters = B[1]; using CRegisters = C[1]; CUTE_HOST_DEVICE static constexpr void fma(D\u0026amp; d, A const\u0026amp; a, B const\u0026amp; b, C const\u0026amp; c) { // Forward to an ADL/cute free function for these types using cute::fma; fma(d, a, b, c); // 这里的实现就是d = a * b + c; }};template \u0026lt;class D, class A, class B, class C\u0026gt;struct MMA_Traits\u0026lt;UniversalFMA\u0026lt;D,A,B,C\u0026gt;\u0026gt; { using ValTypeD = D; using ValTypeA = A; using ValTypeB = B; using ValTypeC = C; // Logical shape of the MMA using Shape_MNK = Shape\u0026lt;_1,_1,_1\u0026gt;; // Logical thread id (tid) -\u0026gt; tidx using ThrID = Layout\u0026lt;_1\u0026gt;; // 只有一个thread参与 // (Logical thread id (tid), Logical value id (vid)) -\u0026gt; coord // (tid,vid) -\u0026gt; (m,k) using ALayout = Layout\u0026lt;Shape\u0026lt;_1,_1\u0026gt;\u0026gt;; // (tid,vid) -\u0026gt; (n,k) using BLayout = Layout\u0026lt;Shape\u0026lt;_1,_1\u0026gt;\u0026gt;; // (tid,vid) -\u0026gt; (m,n) using CLayout = Layout\u0026lt;Shape\u0026lt;_1,_1\u0026gt;\u0026gt;;};参考官方示例函数gemm_nt：https://github.com/NVIDIA/cutlass/blob/main/examples/cute/tutorial/sgemm_sm80.cu#L478，从中提取部分代码如下：using TA = float;using TB = float;using TC = float;TiledMMA mmaC = make_tiled_mma(UniversalFMA\u0026lt;TC, TA, TB\u0026gt;{}, Layout\u0026lt;Shape\u0026lt;_16, _16, _1\u0026gt;\u0026gt;{}); // 16x16x1 TiledMMAstd::cout \u0026lt;\u0026lt; \"\\nTiledMMA Layouts (UniversalFMA 16 16 1):\" \u0026lt;\u0026lt; std::endl;print(mmaC), print(\"\\n\");/*print(\"ALayout: \"), print(typename decltype(mmaC)::ALayout{}), print(\"\\n\");print(\"BLayout: \"), print(typename decltype(mmaC)::BLayout{}), print(\"\\n\");print(\"CLayout: \"), print(typename decltype(mmaC)::CLayout{}), print(\"\\n\");*/std::cout \u0026lt;\u0026lt; \"\\nMMA Atom Layout:\" \u0026lt;\u0026lt; std::endl;print_latex(mmaC);Inverse TV-Layout 如下：MMA_Atom 信息如下：TiledMMA ThrLayoutVMNK: (_1,_16,_16,_1):(_0,_1,_16,_0) PermutationMNK: (_,_,_)MMA_Atom ThrID: _1:_0 Shape_MNK: (_1,_1,_1) LayoutA_TV: (_1,_1):(_0,_0) LayoutB_TV: (_1,_1):(_0,_0) LayoutC_TV: (_1,_1):(_0,_0)A. 资料 A Generalized Micro-kernel Abstraction for GPU Linear Algebra：BLIS Retreat 2023 上的论文，介绍了 CUTLASS-Cute 中 TiledCopy 和 TiledMMA 的设计细节。待阅读 Introduction to CUDA Performance Optimization：CUDA Programming and Optimization 课程的 PPT，介绍了 CUDA 优化相关的知识，其中第 28-30 页介绍了 CUTLASS-Cute 中 TiledMMA 的设计细节。待阅读A.1. TiledCopy 资料 CuTe Tiled Copy：Mao Lei 博客 cute 之 Copy抽象：reed 知乎文章 CUTLASS 笔记 (4)：Tiled Copy：知乎杨远航文章 cute/tutorial/tiled_copy.cu：官方示例代码A.2. MMA Atom 资料 0t_mma_atom.md：官方文档，MMA Atom 文档 cute 之 MMA抽象：reed 知乎文章 CuTe Tiled MMA：Mao Lei 博客，如何配置 TiledMMA Thakkar_BLISRetreat2023.pdf MMA Atoms and TiledMMAA.3. 参考代码 sm80_mma_multistage.hpp：官方示例代码 sgemm_sm80.cu：官方示例代码A.4. Latex转换工具 TeXPage Aspose.TeX viewerA.5. Layout / TV-Layout 可视化 cutlass-viz cute_render cute-viz example_cute_tv_layout：使用cute-viz可视化的测试代码" },

  { "title": "CUTLASS-Cute 初步(2.1)：Tensor \u0026 Layout 实操笔记", "url": "/blog/2025/Cute%E5%88%9D%E6%AD%A52.1-Tensor-Layout-%E5%AE%9E%E6%93%8D%E7%AC%94%E8%AE%B0/", "categories": "CUDA",
  "tags": "CUDA", "date": "2025-02-26", "content":
  "1. local_tile给定一个 tiler，使用 local_tile 函数，将 Tensor 按 tiler 的 shape 切分成多个 tile。由于切分之后，每个 tile 保留输入 Tensor 的stride信息，以及rest mode的shape，故称为Inner-Partition。1.1. 示例一：常规使用（常规 CTA 切分）比如有一个 4x6 的 Tensor，将其切分并分配到 thread block，每个 thread block 获取到的 tile 大小为 (2,2)，当前 thread block 的 coord 为 (1,1)，就可以使用如下代码： constexpr int M = 4, N = 6; auto layout = cute::make_layout(cute::make_shape(M, N), cute::make_stride(N, cute::Int\u0026lt;1\u0026gt;{})); auto tensor = cute::make_tensor(h_data.data(), layout); constexpr auto tiler = cute::make_shape(cute::Int\u0026lt;2\u0026gt;{}, cute::Int\u0026lt;2\u0026gt;{}); constexpr auto coord = cute::make_coord(1, 1); // 第(1,1)块（0-indexed） auto tile = cute::local_tile(tensor, tiler, coord); cute::print(tile);输出：ptr[32b](0xaaaafc5faa38) o (_2,_2):(6,_1)图示如下：1.1.1. 输入 Tensor 轴数量多余于 tiler 的轴数量下面测试输入 Tensor 轴数量多余于 tiler 的轴数量。测试代码： constexpr int M = 4, N = 6, K = 8; auto layout = cute::make_layout(cute::make_shape(M, N, K), cute::make_stride(N * K, K, cute::Int\u0026lt;1\u0026gt;{})); auto tensor = cute::make_tensor(h_data.data(), layout); constexpr auto tiler = cute::make_shape(cute::Int\u0026lt;2\u0026gt;{}, cute::Int\u0026lt;2\u0026gt;{}); constexpr auto coord = cute::make_coord(1, 2); auto tile12 = cute::local_tile(tensor, tiler, coord); cute::print(tile12);输出：ptr[32b](0xaaab0eebc6f0) o (_2,_2,8):(48,8,_1)1.1.2. 总结local_tile的逻辑是：首先使用zipped_divide将输入 Tensor 切分，比如 输入 Tensor：(4,6,8):(48,8,1)，使用 tiler(2,2)切分，得到：zipped_divide得到的结果是((_2,_2),(2,3,8)):((48,8),(96,16,_1))。 其中(_2,_2)表示每个 tile 的 shape 是 (2,2)，(2,3,8)表示每个 tile 的 rest mode shape 是 (2,3,8)，(48,8)表示每个 tile 的 stride 是 (48,8)，(96,16,_1)表示每个 tile 的 rest mode stride 是 (96,16,_1)。最后根据 coord 获取对应的 tile。zipped_divide返回一个 rank-2 的 tensor：([tile mode], [rest mode]) 顶层 mode 内容 含义 tile mode (mode-0) (_2,_2):(48,8) 一个 tile 内部的坐标系：2×2 个元素，stride 保留自原 tensor rest mode (mode-1) (2,3,8):(96,16,_1) tile 外部的坐标系：有多少个 tile，以及未被切分的 trailing 维度 “rest”的字面意思是剩余，指”tile 坐标系之外剩下的所有东西”： 已切分维度的块编号（你选哪一块？→ 2 个 M 块 × 3 个 N 块） 未切分维度的完整范围（K=8 没有被 tiler 动过，直接挂在 rest 末尾） 如何切分，与 stride 没有关系，即选择前两个轴来切分，不是依照 stride 来选择轴（即与内存布局没有关系）。完整测试代码：https://github.com/HPC02/cuda_perf/blob/master/src/study_codes/study_tests/test_local_tile.cu中的 Case03。1.2. 示例二：在 coord 中使用 cute::_，得到该轴方向上的 tile 集合使用cute::_作用于 coord 中的某个轴，其含义是获取该轴方向上的 tile 集合，即在该轴方向上，生成新的 trailing mode。示例代码如下： constexpr int M = 4, N = 6; auto layout = cute::make_layout(cute::make_shape(M, N), cute::make_stride(N, cute::Int\u0026lt;1\u0026gt;{})); auto tensor = cute::make_tensor(h_data.data(), layout); constexpr auto tiler = cute::make_shape(cute::Int\u0026lt;2\u0026gt;{}, cute::Int\u0026lt;2\u0026gt;{}); constexpr auto coord = cute::make_coord(0, cute::_); auto tile00 = cute::local_tile(tensor, tiler, coord); // (tileM, tileN, k) cute::print(tile00);输出：ptr[32b](0xaaaafd59bed0) o (_2,_2,3):(6,_1,_2) 从结果看到，在第二个轴上使用 slice 操作符，shape 中的最后一个轴 3 表示在该轴上有 3 个 tile。完整测试代码：https://github.com/HPC02/cuda_perf/blob/master/src/study_codes/study_tests/test_local_tile.cu中的 Case04。1.3. 示例三：tiler 中使用 Step筛选 tiler 以及 coord，使用 cute::X 标记的轴不参与切分Step 本质上是一个轴选择器，筛选出需要的 tiler 以及 coord，使用 cute::X 标记的轴不参与切分。 constexpr int M = 4, N = 6, K = 8; constexpr auto problem_shape = cute::make_shape(M, N, K); auto stride_A = cute::make_stride(K, cute::Int\u0026lt;1\u0026gt;{}); auto layout_A = cute::make_layout(cute::select\u0026lt;0, 2\u0026gt;(problem_shape), stride_A); auto tensor_A = cute::make_tensor(h_data.data(), layout_A); constexpr auto tiler = cute::make_shape(cute::Int\u0026lt;2\u0026gt;{}, cute::Int\u0026lt;2\u0026gt;{}, cute::Int\u0026lt;4\u0026gt;{}); constexpr auto coord = cute::make_coord(0, 0, cute::_); // tile_A (tileM, tileK, k) auto tile_A = cute::local_tile(tensor_A, tiler, coord, cute::Step\u0026lt;cute::_1, cute::X, cute::_1\u0026gt;{}); cute::print(tile_A);输出：ptr[32b](0xaaaafe0c3860) o (_2,_4,2):(8,_1,_4) 即得到一个 MK 方向的 tile(2, 4) 的集合，K 方向上有两个 tile，这两个 tile 在 M 方向上坐标为 0（划分成 2*2 个分块）。另外，在 cute::Step\u0026lt;\u0026gt; 中使用 cute::X 标记的轴不参与切分，例如：// tiler rank=2，Step\u0026lt;_1,X\u0026gt; 表示只切第0维，第一维整体保留local_tile(tensor, make_shape(Int\u0026lt;2\u0026gt;{}), coord_m, Step\u0026lt;_1, X\u0026gt;{});2. local_partitionlocal_tile与local_partition都是基于zipped_divide（得到([tile_mode], [rest_mode])），区别在于：local_tile第二步是切入 rest_mode，而local_partition第二步是切入 tile_mode。即： local_tile：先切分得到 tile 和 rest，然后根据 coord 从 [rest_mode] 中选取一个 tile（保留[tile_mode]）。 local_partition：先切分得到 tile 和 rest，然后根据 tid 从 [tile_mode] 中选取一个 tile（保留[rest_mode]）。图示zipped_divide的结果：tensor_tiled 的结构: mode-0 (tile mode): 大小 64，标记内部元素编号 0..63 mode-1 (rest mode): 大小是 \"有几个这样的块\"访问方式：tensor_tiled(i, j) i = 块内第 i 个元素 j = 第 j 个块local_tile的行为：tensor_tiled( _, j=blockIdx ) ↑ ↑ 全取 固定\"哪个块\"local_partition的行为：tensor_tiled( i=tid, _ ) ↑ ↑ 固定\"我在 全取\"所有块\" 块内的位置\"另一个差异点是local_partition是交错分配数据到不同线程的（内存合并访问），而local_tile是连续分配数据到同一个线程的，如下图示：原始 tensor (BM=64, BK=16):┌────────────────────────────┐│ 0 1 2 ... │ ← BK=16 列│ 64 65 66 ... ││ ... ││ │ ← BM=64 行└────────────────────────────┘local_tile: 选出某个矩形块（连续区域） ┌──────┐ │ tile │ ← 这个 CTA 的全部数据 └──────┘local_partition: 从整个矩阵中取 tid 对应的行（交错分配） 行0 → tid=0 行1 → tid=1 ... 行63 → tid=63 每个 tid 得到 1 行 × 16 列 = 16 个元素示例代码： constexpr int M = 8, K = 6; auto layout = cute::make_layout(cute::make_shape(M, K), cute::make_stride(cute::Int\u0026lt;1\u0026gt;{}, M)); auto tensor = cute::make_tensor(h_data.data(), layout); constexpr auto thr_layout{cute::make_layout(cute::make_shape(cute::Int\u0026lt;4\u0026gt;{}, cute::Int\u0026lt;2\u0026gt;{}))}; auto partition = cute::local_partition(tensor, thr_layout, 1); // (2,3):(_4,16) cute::print(partition);由结果可知，local_partition将 tensor 沿着 mode-0 切分成四块，每块两行，交错分配给四个线程；沿着 mode-1 切分成两块，每块三列，交错分配给两个线程。是用 cute::Step\u0026lt;\u0026gt; 来选择切分轴，例如： auto partition = cute::local_partition(tensor, thr_layout, 1, cute::Step\u0026lt;cute::_1, cute::X\u0026gt;{}); // (2,6):(_4,8)此时，只沿着 M 轴切分成四块，每块两行，交错分配给四个线程；K 轴不切分，整体保留。2.1. local_tile 与 local_partition 对比总结 步骤 local_tile (inner) local_partition (outer) 步骤一 zipped_divide(tensor, tiler) → ([tile], [rest]) 同左 步骤二 用 coord 索引 [rest]，消掉块编号维度 用 tid 索引 [tile]，消掉块内位置维度 返回内容 第 coord 号块内部的所有元素 所有块中位置为 tid 的元素集合 语义 “我是第几个 CTA，取走我的 tile” “我是第几个线程，取走我的元素” “inner/outer”这两个别名也恰好反映了这个区别：inner 返回块的内部，outer 返回块的外部视角（跨块的分布）。参考代码见：https://github.com/HPC02/cuda_perf/blob/master/src/study_codes/study_tests/test_local_partition.cu。A. 资料 cute 之 Tensor 02_layout_algebra：官方文档" },

  { "title": "CUTLASS-Cute 初步(2)：Tensor \u0026 Layout Algebra", "url": "/blog/2025/Cute%E5%88%9D%E6%AD%A52-Tensor-Layout-Algebra/", "categories": "CUDA",
  "tags": "CUDA", "date": "2025-02-26", "content":
  "github – 测试代码 CuTe Tensors：官方文档0. 紧凑 stride 的生成：LayoutLeft 与 LayoutRight当调用 make_layout(shape) 而不显式指定 stride 时，CuTe 使用 LayoutLeft（紧凑列主序）自动生成 stride。也可以通过 LayoutRight（紧凑行主序）来构造。auto shape = make_shape(2, make_shape(2, 2));auto manual = make_layout(shape, make_stride(4, make_stride(2, 1)));auto left = make_layout(shape, LayoutLeft{}); // 等价auto right = make_layout(shape, LayoutRight{});print_layout(left);print_layout(right);LayoutLeft 生成 stride 的算法是”从左侧开始的 exclusive prefix product”——shape 从左到右做前缀乘，得到广义列主序 stride：(2,(2,2)):(_1,(2,4)) 0 1 2 3 +---+---+---+---+ 0 | 0 | 2 | 4 | 6 | +---+---+---+---+ 1 | 1 | 3 | 5 | 7 | +---+---+---+---+LayoutRight 从右侧开始做 exclusive prefix product，对于 depth=1 的 shape 就是行主序，但对层级 shape 的结果可能不符合直觉。 LayoutLeft、colexicographical order、column-major 本质上对应同一种遍历顺序（列优先）。LayoutRight 则对应 row-major。 理解 stride 是如何自动生成的，有助于理解后面的 coalesce 和 composition 操作为什么会有特定的行为。1. CuTe 中的 Tensor 划分 (Partitioning a Tensor)在 GEMM 计算是，需要对矩阵进行划分（分块），以便线程块（Thread Block）和线程（Thread）能够并行处理数据。常用的有 Inner-Partitioning、Outer-Partitioning，以及 TV-layout-Partition。1.1. Inner-PartitioningGEMM 计算，先需要按照 Thread Block 划分为若干 Tile，即给每个 Thread Block 分配一个 Tile。如下所示：Tensor A = make_tensor(ptr, make_shape(8,24)); // (8,24)auto tiler = Shape\u0026lt;_4,_8\u0026gt;{}; // (_4,_8)Tensor tiled_a = zipped_divide(A, tiler); // ((_4,_8),(2,3))在使用 tiler 对 A 进行切分之后，(_4, _8) 是第一个mode（first mode），(2, 3) 是第二个mode（second mode）。 第一个 mode (mode 0)：Tile 的 shape 第二个 mode (mode 1)：Tile 在全局中的排列 mode 是张量代数（tensor algebra）中的一个概念，表示张量的不同维度，或者叫逻辑轴。为了避免与 tensor 中的维度（dimension）混淆，故使用不同的术语。沿着第二个mode 切分之后，可以得到 2x3 个 Tile，每个 Tile 的 shape 是 4x8：Tensor cta_a = tiled_a(make_coord(_,_), make_coord(blockIdx.x, blockIdx.y)); // (_4,_8)由于保留了内部 tile 的 shape 信息，这种切分方式叫做 Inner-Partitioning。CuTe 使用 inner_partition(Tensor, Tiler, Coord) 函数来实现内部分块，在实际使用时，则使用另外一个代替的函数 local_tile(Tensor, Tiler, Coord)来进行 tile 划分（分配给 Thread Block）。1.2. Outer-Partitioning在 Thread Block 划分好 tile 之后，下一步就是将 sub-tile 分配给线程（Thread）：Tensor thr_a = tiled_a(threadIdx.x, make_coord(_,_)); // (2,3)这一步叫做 Outer-Partitioning，因为这种划分方式对 tile 进行划分，保留了分块的 shape 信息。在 CuTe 中，使用 outer_partition(Tensor, Tiler, Coord) 函数来实现外部分块，在实际使用时，则使用另外一个代替的函数 local_partition(Tensor, Layout, Idx) 来进行 sub-tile 划分（分配给 Thread）。2. Layout Algebra2.1. Coalesce：Layout 展平为一维一些二维的 layout，其二维索引 (m, n) 可以转换得到对应的一维索引 k。例如：layout (M, N) : (1, M)k = m + n * M例子：auto layout = Layout\u0026lt;Shape\u0026lt;_2,Shape \u0026lt;_1,_6\u0026gt;\u0026gt;, Stride\u0026lt;_1,Stride\u0026lt;_6,_2\u0026gt;\u0026gt;\u0026gt;{};auto result = coalesce(layout); // _12:_1有三种情况可以展平： (s0, _1) : (d0, d1) =\u0026gt; s0 * d0。忽略第二维的 stride。 (_1, s1) : (d0, d1) =\u0026gt; s1 * d1。忽略第一维的 stride。 (s0, s1) : (d0, d1) =\u0026gt; s0 + s1 * d0。当 d1 == d0 * s0 时，可以展平为一维。其他情况不能展平为一维 layout，比如 layout codomain 出现空洞，即不是连续映射的。不能展平的维度维持原有的 shape 和 stride 信息。 在 CuTe 中，以一个二维 layout 为例，可以使用layout(m, n) 索引的形式访问，也可以使用 layout(k) 的形式访问。使用一维索引 k 访问时，其应该等于使用 coalesce 展平后的 layout 进行访问。见https://github.com/NVIDIA/cutlass/blob/main/test/unit/cute/core/coalesce.cpp。2.1.1. By-mode Coalesce：展平部分轴例如 layout ((M, N), K, L) : ((1, M), M * N, M * N * K)，希望展平子 layout (M, N):(1, M)，保持其他轴的 shape 不变：coalesce(layout, Step\u0026lt;_1, _1, _1\u0026gt;{});Step\u0026lt;_1, _1, _1\u0026gt; 表示展平为三维的 layout，因为这个 Step\u0026lt;\u0026gt; 的 rank 是 3。CuTe 对每个轴重新计算，得到新的 layout：((M*N), K, L) : (1, M*N, M*N*K)一个完整的例子：// Coalesce within each top-level modeauto layout = make_layout(make_shape(Int\u0026lt;2\u0026gt;{}, make_shape(Int\u0026lt;1\u0026gt;{}, Int\u0026lt;6\u0026gt;{})), make_stride(Int\u0026lt;1\u0026gt;{}, make_stride(Int\u0026lt;6\u0026gt;{}, Int\u0026lt;2\u0026gt;{})));// layout is (2,(1,6)):(1,(6,2))auto result = coalesce(layout, Step\u0026lt;_1,_1\u0026gt;{}); // result is (2,6):(1,2)// Coalesced within each mode, preserving rank-2 structure2.2. Composition：Layout 复合首先，Layout 本质就是一个映射函数，Composition 即经过多次映射：R(c) := (A o B)(c) := A(B(c))计算过程如下：composition 有如下属性： 兼容性：compatible(B, R) - B的每个坐标都可以用作R的坐标，因为B定义了R的定义域 函数等价性：对于B定义域内的所有i，R(i) == A(B(i))。官方测试用例test composition2.2.1. By-mode Compositioncomposition 函数提供一个重载版本，第二个参数为 tiler，对部分维度（轴）进行复合操作。Tiler 可以是： 一个 Layout Tiler tuple Shape，会被解释为步长为 1 的 Tiler2.2.2. Composition：reshape \u0026amp; reorderingReshape// 20-element layout with stride 2auto layout_1d = make_layout(Int\u0026lt;20\u0026gt;{}, Int\u0026lt;2\u0026gt;{}); // 20:2// Reshape to 5x4 row-majorauto tiler = make_layout(make_shape(Int\u0026lt;5\u0026gt;{}, Int\u0026lt;4\u0026gt;{}), make_stride(Int\u0026lt;4\u0026gt;{}, Int\u0026lt;1\u0026gt;{})); // (5,4):(4,1)auto result = composition(layout_1d, tiler); // (5,4):(8,2)// Maps (i,j) coordinates to layout_1d using tiler pattern结果如下（下左为 tiler，下右为 result）：Extract Subtile// (12,(4,8)):(59,(13,1))auto a = make_layout(make_shape (12,make_shape ( 4,8)), make_stride(59,make_stride(13,1)));// \u0026lt;3:4, 8:2\u0026gt;auto tiler = make_tile(Layout\u0026lt;_3,_4\u0026gt;{}, // Apply 3:4 to mode-0 Layout\u0026lt;_8,_2\u0026gt;{}); // Apply 8:2 to mode-1// (_3,(2,4)):(236,(26,1))auto result = composition(a, tiler);结果如下：2.2.3. 1-D Index 以及 Composition 验证Layout A (6,2):(8,2)： 0 1 +----+----+ 0 | 0 | 2 | +----+----+ 1 | 8 | 10 | +----+----+ 2 | 16 | 18 | +----+----+ 3 | 24 | 26 | +----+----+ 4 | 32 | 34 | +----+----+ 5 | 40 | 42 | +----+----+Layout B (4,3):(3,1): 0 1 2 +----+----+----+ 0 | 0 | 1 | 2 | +----+----+----+ 1 | 3 | 4 | 5 | +----+----+----+ 2 | 6 | 7 | 8 | +----+----+----+ 3 | 9 | 10 | 11 | +----+----+----+Composed Layout C: 0 1 2 +----+----+----+ 0 | 0 | 1 | 2 | +----+----+----+ 1 | 3 | 4 | 5 | +----+----+----+ 2 | 6 | 7 | 8 | +----+----+----+ 3 | 9 | 10 | 11 | +----+----+----+验证代码：import cutlassfrom cutlass import cute@cute.jitdef compose_verify(): A = cute.make_layout((6, 2), stride=(8, 2)) B = cute.make_layout((4, 3), stride=(3, 1)) C = cute.composition(A, B) flat = cute.coalesce(B) for i in cutlass.range_constexpr(cute.size(flat)): print(f\"C({i}) = {C(i)}, \\tflat({i}) = {flat(i)}, \\tA(flat({i})) = {A(flat(i))}\")compose_verify()打印结果：C(0) = 0, flat(0) = 0, A(flat(0)) = 0C(1) = 24, flat(1) = 3, A(flat(1)) = 24C(2) = 2, flat(2) = 6, A(flat(2)) = 2C(3) = 26, flat(3) = 9, A(flat(3)) = 26C(4) = 8, flat(4) = 1, A(flat(4)) = 8C(5) = 32, flat(5) = 4, A(flat(5)) = 32C(6) = 10, flat(6) = 7, A(flat(6)) = 10C(7) = 34, flat(7) = 10, A(flat(7)) = 34C(8) = 16, flat(8) = 2, A(flat(8)) = 16C(9) = 40, flat(9) = 5, A(flat(9)) = 40C(10) = 18, flat(10) = 8, A(flat(10)) = 18C(11) = 42, flat(11) = 11, A(flat(11)) = 42A. 参考 CuTe Layout and Tensor Tutorial：deepwiki algegra 使用示例解析 CuTe Layout Algebra：deepwiki algebra 解析 cute_layout_algebra.ipynb：官方 CuteDSL notebookA.1. 学习参考资料 CuTe Layout Algebra。官方文档 deepwiki – cutlass – Layout Algebra cute Layout 的代数和几何解释。来自知乎 reed 文章 CuTe Layout and Tensor。来自Yifan Yang (杨轶凡) 博客 CuTe Layout Representation and Algebra：来自 arxiv 的论文 CuTe Layout 的范畴论基础 CuTe 02 - Layout 运算 Cute概念速通：待阅读三方学习测试代码： github – code for layout algebraA.2. 其他资料 A Generalized Micro-kernel Abstraction for GPU Linear Algebra：NVIDIA PPT Algebra – 2.12 Inverses and composition：数学理论：composition Categorical Foundations for CuTe Layouts Layout Algebra：三方实现的C++ Layout Algebra 库 On CuTe layoutsA.3. 工具 将 SVG 合并到 SVG：工具：合并svg图片" },

  { "title": "CUTLASS-Cute 初步(1)：Layout", "url": "/blog/2025/Cute%E5%88%9D%E6%AD%A51-Layout/", "categories": "CUDA",
  "tags": "CUDA", "date": "2025-02-26", "content":
  "github – Layout 测试代码1. Layout 核心概念CuTe（CUDA Tensor）是 CUTLASS 3.x 引入的底层张量抽象库，用于简化 BLAS 操作和内存布局管理。最核心的概念是 Layout。Layout = (Shape, Stride)，它是一个从逻辑坐标空间到一维内存索引空间的映射函数：offset = Σ (coord[i] * stride[i]) Shape：逻辑维度，定义坐标空间（”domain”，定义域）。 Stride：每个维度在内存中的步长，决定映射结果（”codomain”，值域）。以 1D layout 8:2 为例——8 个元素、stride 为 2： size(layout) = 8 —— 定义域大小（有多少个逻辑坐标） cosize(layout) = 16 —— 值域大小（映射到的最大 index + 1，即 layout(size-1) + 1）Stride 为 0 时（如 8:0），所有坐标映射到同一个地址，此时 size=8, cosize=1。 Layouts are functions from integers to integers. 这一总结来自官方文档。意味着：每个 Layout 都可以用 1-D 坐标去索引，无论它看起来是几维的。这是 CuTe 与 C++23 mdspan 的关键区别——mdspan 的 N-D 视图不接受 1-D 坐标，而 CuTe 的 Layout 原生支持。1.1. IntTuple、rank 与 depthShape 和 Stride 都是 IntTuple——要么是单个整数，要么是 IntTuple 的 tuple（可嵌套）。IntTuple 中的整数可以是编译期静态整数（用 Int\u0026lt;N\u0026gt;{} 或别名 _1、_2 等表示），也可以是运行期动态整数。Layout 的两个重要结构属性： 属性 含义 示例 rank 最外层有几个元素（mode） (4,2) rank=2；((2,2),2) rank=2 depth 括号嵌套的最大深度 (4,3) depth=1；(3,(6,2),8) depth=2 rank/depth 示例：Layout Shape rank depth 含义──────────── ──── ───── ────(8) 1 1 向量(4,2) 2 1 矩阵（行和列）(M,N,K) 3 1 3-D tensor((2,2),2) 2 2 矩阵，但第一 mode 是嵌套的(3,(6,2),8) 3 2 最深处 (6,2) 嵌套了 2 层2. 构造 LayoutLayout 通过 make_layout(shape, stride) 构造。如果不指定 stride，默认按 LayoutLeft（列优先）生成紧凑 stride。下面用 4 个递进式例子来展示 Layout 的表达能力。2.1. 例一：1D 向量auto v1 = make_layout(Int\u0026lt;8\u0026gt;{}); // 8:_1 — 静态 shape，stride 默认 1auto v2 = make_layout(8, 2); // 8:2 — 动态 shape，stride 为 28:1 就是连续排列的 8 个元素，index = coord。8:2 则每个坐标步进 2。2.2. 例二：2D 矩阵 — 列主序与行主序// 列主序（column-major）：沿列 stride=1，沿行 stride=Mauto col_major = make_layout(make_shape(2, 3), make_stride(1, 2));// 行主序（row-major）：沿行 stride=1，沿列 stride=Nauto row_major = make_layout(make_shape(2, 3), make_stride(3, 1));(2,3):(1,2) 列主序输出： 0 1 2 +---+---+---+ 0 | 0 | 2 | 4 | +---+---+---+ 1 | 1 | 3 | 5 | +---+---+---+(2,3):(3,1) 行主序输出： 0 1 2 +---+---+---+ 0 | 0 | 1 | 2 | +---+---+---+ 1 | 3 | 4 | 5 | +---+---+---+auto val = col_major(1, 2); // = 1*1 + 2*2 = 5所谓”列主序”就是最左侧 mode 的 stride=1，colexicographical 遍历等价于逐列访问。2.3. 例三：带嵌套 mode 的层级 Layout当 mode 本身也是多维时，使用嵌套 tuple 来表达：auto shape = make_shape(4, make_shape(2, 2)); // 4 行，(2,2) 的嵌套列auto stride = make_stride(4, make_stride(1, 2)); // stride_i=4, stride_j=1, stride_k=2auto layout = make_layout(shape, stride);print_layout(layout);输出：(4,(2,2)):(4,(1,2)) 0 1 2 3 +----+----+----+----+ 0 | 0 | 1 | 2 | 3 | +----+----+----+----+ 1 | 4 | 5 | 6 | 7 | +----+----+----+----+ 2 | 8 | 9 | 10 | 11 | +----+----+----+----+ 3 | 12 | 13 | 14 | 15 | +----+----+----+----+auto val = layout(2, make_coord(1, 0)); // = 2*4 + 1*1 + 0*2 = 9这里虽然 Shape 里嵌套了 (2,2)，但顶层 rank 仍然是 2（行 + 嵌套列）。print_layout 将其展示为 4×4 的表格，而嵌套的列 mode 被摊平为 4 列。2.4. 例四：改变 stride 对数据排布的影响保持 Shape 不变，将 stride 改为 (2,(1,8))：auto shape = make_shape(4, make_shape(2, 2));auto stride = make_stride(2, make_stride(1, 8));auto layout = make_layout(shape, stride);print_layout(layout);输出：(4,(2,2)):(2,(1,8)) 0 1 2 3 +----+----+----+----+ 0 | 0 | 1 | 8 | 9 | +----+----+----+----+ 1 | 2 | 3 | 10 | 11 | +----+----+----+----+ 2 | 4 | 5 | 12 | 13 | +----+----+----+----+ 3 | 6 | 7 | 14 | 15 | +----+----+----+----+auto val = layout(2, make_coord(1, 0)); // = 2*2 + 1*1 + 0*8 = 5比较例三和例四：同样的 Shape (4,(2,2))，只改变 stride，就得到了完全不同的数据排布。例三中每一行的 4 个元素在内存中是连续的 0,1,2,3；例四中每一行的前两个和后两个元素在内存中相隔 8，呈现”分块”效果。 Stride 分量 值 含义 stride_i 2 沿行移动一步，offset +2 stride_j 1 沿内层列移动一步，offset +1 stride_k 8 沿外层列移动一步，offset +8（跳到另一个 4 元素的 block） 关于层级 Layout 的几何直觉，可参考 https://note.gopoux.cc/hpc/cute/layout/3. 层级 Layout 与 coordinate3.1. 层级 Layout 的解读层级 Layout 的核心思想：用嵌套的 Shape/Stride 表达”layout of layouts”。以 (4,(2,4)):(2,(1,8)) 为例：上图中 a 和 b 分别是无嵌套的列主序和行主序。c 和 d 则带有层级结构。示例 c — 仅在列方向有嵌套： 内层 layout（红色框内）：(4,2):(2,1) —— 4 行 2 列，列主序 外层 layout：(1,4):(4,1) —— 1 行 4 列（每个”元素”是一个内层 layout） 合并 shape：(4, (2, 4))，stride：(2, (1, 8)) stride_i=2：内层行方向步长 stride_j=1：内层列方向步长 stride_k=8：外层列方向步长（等于一个完整内层 block 的 cosize） 示例 d — 行列方向均有嵌套： 内层 layout（红色框内）：(2,2):(2,2) 外层 layout：(2,4):(4,1) 合并 shape：((2,2), (2,4))，stride：((1,4), (2,8))3.2. 坐标访问与 slice访问层级 Layout 时，坐标也需对应嵌套结构：auto row_coord = make_coord(1, 3); // 内层行=1, 外层行=3auto col_coord = make_coord(2, 4); // 内层列=2, 外层列=4auto coord = make_coord(row_coord, col_coord);auto val = layout(coord);Slice 操作使用 cute::_（类似 Python 的 :）保留某个维度的所有元素：auto layout = make_layout( make_shape(make_shape(2, 4), make_shape(3, 5)), make_stride(make_stride(3, 6), make_stride(1, 24)));auto row_coord = make_coord(1, 1); // 固定行坐标auto col_coord = make_coord(cute::_, cute::_); // 保留所有列auto coord = make_coord(row_coord, col_coord); // 获取子块 Bauto sub_layout = cute::slice(coord, layout);4. Layout 的三种坐标空间每个 Layout 天然拥有三套坐标系统： 1-D 坐标：单个整数 R-D 坐标：rank 维坐标（如 2D layout 的 (m,n)） h-D 坐标：层级坐标（自然坐标，与 Shape 的嵌套结构一致）这三种坐标可以互相转换，前提是 Shape 兼容（size 相等且定义域相互包含）。4.1. colexicographical order 坐标转换CuTe 使用 colexicographical order（余字典序，从右往左变化更快）进行坐标映射。以 shape (3,(2,3)) 为例： 1-D 2-D Natural   1-D 2-D Natural 0 (0,0) (0,(0,0))   9 (0,3) (0,(1,1)) 1 (1,0) (1,(0,0))   10 (1,3) (1,(1,1)) 2 (2,0) (2,(0,0))   11 (2,3) (2,(1,1)) 3 (0,1) (0,(1,0))   12 (0,4) (0,(0,2)) 4 (1,1) (1,(1,0))   13 (1,4) (1,(0,2)) 5 (2,1) (2,(1,0))   14 (2,4) (2,(0,2)) 6 (0,2) (0,(0,1))   15 (0,5) (0,(1,2)) 7 (1,2) (1,(0,1))   16 (1,5) (1,(1,2)) 8 (2,2) (2,(0,1))   17 (2,5) (2,(1,2)) 转换算法：从左到右逐级取模和除法。以 shape (M, N) 从 1-D 转到 2-D 为例：coord = (index % M, (index / M) % N)以 shape (M, (N, K)) 从 1-D 转到 h-D 为例：index = 7，M=3： mode_0 = 7 % 3 = 1， 7 / 3 = 2 mode_1 = (2 % 2 = 0， 2 / 2 = 1 % 3 = 1) → (1, (0, 1))坐标转换流程： idx2crd crd2idx1-D ──────────→ natural (h-D) coord ──────────→ index ↑ idx2crd │R-D ──────────→──────┘C++ 示例：auto shape = Shape\u0026lt;_3,Shape\u0026lt;_2,_3\u0026gt;\u0026gt;{};print(idx2crd( 16, shape)); // (1,(1,2))print(idx2crd(_16{}, shape)); // (_1,(_1,_2))print(idx2crd(make_coord( 1,5), shape)); // (1,(1,2))print(idx2crd(make_coord(_1{},5), shape)); // (_1,(1,2))print(idx2crd(make_coord( 1,make_coord(1, 2)), shape)); // (1,(1,2))print(idx2crd(make_coord(_1{},make_coord(1,_2{})), shape)); // (_1,(1,_2))可以看到无论用哪种坐标输入 idx2crd，只要指的是同一个逻辑位置，得到的 natural 坐标都等价。 坐标转换只依赖 Shape，不依赖 Stride。而调用 layout(coord) 计算 index 时才需要 Stride： 操作 依赖 Shape 依赖 Stride 1-D ↔ R-D ↔ h-D 坐标转换（idx2crd） ✅ ❌ 坐标兼容性判断（compatibility） ✅ ❌ 坐标 → 内存 index（crd2idx / layout(coord)） ✅ ✅ A. 参考及资料 reed – cute 之 Layout Yifan Yang (杨轶凡) – CuTe Layout and Tensor CUTLASS CUTE 1 Layout Algebra 01_layout.md：CUTLASS/CuTe 官方文档 CuTe-Copy for GPUMode PPT：NVIDIA 官方 PPT CUTLASS: A CUDA C++ Template Library for Accelerating Deep Learning：YouTube 视频A.1. 更多学习资料 AI Kernel Learning — Quick Start Learn CUTLASS the hard way 系列" },

  { "title": "使用 Nsight Compute 进行 kernel 性能分析", "url": "/blog/2025/ncu-%E6%80%A7%E8%83%BD%E5%88%86%E6%9E%90/", "categories": "CUDA",
  "tags": "CUDA", "date": "2025-02-25", "content":
  "测试用例：github – cuda_perf 编译时，加上 -lineinfo 参数，Nsight Compute 分析时，可以看到具体的 C++/cu 代码。 另外一篇性能分析文章：CUDA 架构(1.1)：Hopper架构及性能分析(ncu) + 性能优化。1. 启动配置使用Nsight Compute运行被测试CUDA程序，启动时，指定metrics为full：1.1. 命令行方式# 完整分析，输出到文件ncu --set full -o gemm_sm80_profile ./ncu_gemm_sm80 4096 4096 4096# 查看 memory 相关指标ncu --set memory ./ncu_gemm_sm80 4096 4096 4096# 查看 compute 相关指标ncu --set compute ./ncu_gemm_sm80 4096 4096 4096# 查看 roofline# 1. Memory Throughput(GB/s)、L2 Cache Throughput(GB/s)、Compute (SM) Throughput(TFLOPS/s)# 结果以百分比形式给出，即与理论峰值的比值# 2. SM Active Cycles / Elapsed Cycles，他们的比值表示 SM 的利用率，即可以看出等待内存的时间占比ncu --set roofline ./ncu_gemm_sm80 4096 4096 4096# 查看 occupancy 相关指标ncu --section Occupancy ./ncu_gemm_sm80.exe 4096 4096 4096# 查看 shared memory bank conflictncu --metrics l1tex__data_bank_conflicts_pipe_lsu_mem_shared,l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum ./ncu_gemm_sm80 4096 4096 40962. 分析Bank Conflicts被测试CUDA程序运行结束后，打开Nsight Compute的结果页面（Details），进入Memory Workload Analysis章节，在章节标题右侧，选择Memory Tables，查看Shared Memory部分：Details页面里显示的Shared Memory的Bank Conflicts信息，由于其测量数据来源于硬件计数器，还可能包括仲裁冲突(arbitration conflicts)： LDGSTS的Fill Return（全局/共享内存加载的返回数据） TMA的Fill Return（Tensor Memory Accelerator的返回数据） Tensor Core对共享内存的读取使用Source页面的L1 Wavefonts Shared Excessive结果，应该更准确：两个页面分析结果为什么会有差异？ 指标 数据来源 测量内容 用途 Details Page 硬件计数器 Bank Conflict + 仲裁冲突 实际性能影响 Source Page 代码分析 纯粹的 Bank Conflict 代码优化 如何识别可优化的Bank Conflict： 使用源页面（Source View）的Excessive计数器 这些是由地址发散和真实的Bank Conflict引起的 这些可以通过代码优化来消除哪些冲突无法修复： 详情页面中显示的仲裁冲突无法通过代码优化解决 这些是系统级别的问题（Tensor Core、TMA、全局内存的更高优先级访问） 这些是硬件行为的自然结果A. 参考资料 Profiling Guide：官方文档。 How to Understand and Optimize Shared Memory Accesses using Nsight Compute：NVIDIA 课程 Some question about GTC2020 cutlass’s talk about conflict free load shared memory Questions on L1 Bank Conflict statistic discrepancies between Details and Source pages Shared memory bank conflicts and nsight metric Requests, Wavefronts, Sectors Metrics: Understanding and Optimizing Memory-Bound Kernels with Nsight Compute。视频资料。 NVIDIA Kernel级性能分析工具Nsight Compute入门详解：详细讲解 Nsight Compute 中各个页面中的参数。" },

  { "title": "CUDA入门：Bank Conflict", "url": "/blog/2025/CUDA%E5%85%A5%E9%97%A8-Bank-Conflict/", "categories": "CUDA",
  "tags": "CUDA", "date": "2025-02-25", "content":
  "使用到的测试代码：bank_conflict.cu1. Bank Conflicts (Shared Memory)1.1. Bank 划分针对Shared Memory的访问，CUDA使用bank机制，将shared memory的访问（读/写）映射到不同的bank，以实现并行访问。bank以4字节为单位，共32个bank。这样，一个时钟周期内，可以并行访问32个不同的bank，即访问128字节的数据。映射公式bank index = (address /4) % 32。 ⭐ 每次发起共享内存事务（transation）时，可以从这 32 个 bank 中分别读取一个 32 位数据。以 32 位的字为单位索引，则 bank 以地址的低 5 位进行划分，与高位没有关系。图示transaction：Thread（在CUDA Core中） ↓ 访问Shared Memory ↓[Bank系统处理] ← Transaction 在这里发生 ↓返回数据到Thread举例：warp中定义的shared memory如何映射到banks：__shared__ float s[64];如上变量，其映射如下：1.2. Bank Conflicts在一次transaction的时候，如果，当warp中的不同线程访问到同一个bank中的不同地址时，就会产生Bank Conflicts，导致访问串行化：需要分成多次transaction。有N个线程访问同一个bank，称为N-way Bank Conflicts。 所谓 Bank Conflicts，只与transaction有关，即其由 Shared Memory 访问控制器相关。引用https://forums.developer.nvidia.com/t/how-to-understand-the-bank-conflict-of-shared-mem/260900/2：When you store (or load) more than 4 bytes per thread, which is like saying more than 128 bytes per warp, the GPU does not issue a single transaction. The largest transaction size is 128 bytes. If you request 16 bytes per thread, then warp wide that will be a total of 512 bytes per request (warp-wide). The GPU will break that up into 4 transactions (in that case: T0-T7 make up a transaction, T8-T15 are a transaction, and so on), each of which is 128 bytes wide. The determination of bank conflicts is made per transaction, not per request or per warp or per instruction.如下情况，会产生Bank Conflicts： 一次transaction中，warp中的多个线程，访问同一个bank中的不同地址； 一次transaction中，warp中的多个线程，访问shared memory的下一个128字节，被映射到同一个bank。此时也是属于上一种情况：同一个bank中的不同地址。如下情况，Bank Conflicts不会产生： warp中的线程，访问地址唯一对应到bank簇的每个bank，不论是顺序，还是错位； warp中的多个线程，访问同一个bank中的相同地址–使用boardcast分发相同地址数据到多个线程； warp中的线程，单个线程一次访问多个bank，但其他线程不访问这些bank。此时，生成多次transaction。 总结来说：一次内存事务（SMEM 读/写），传输大小为128字节，并对应到32个bank。如果SMEM中有两个不同地址落入到同一个bank，则产生bank conflict，导致访问串行化。另外一个概念是：bank conflict，是在warp级别上发生的（可以理解为：warp是一个最小的调度单元，其对SMEM的一次访问对应到一次内存事务）。2. Bank Conflicts 示例 以下示例来自博客Notes About Nvidia GPU Shared Memory Banks。如下示例，产生32路Bank Conflicts：const int num_iters = 10000; // 全局常量__global__ void all_conflicts() { __shared__ float s[32][32]; [[maybe_unused]] int warp_id = threadIdx.y; int lane_id = threadIdx.x; // thread 在 warp 中的 id float *ptr = \u0026amp;s[lane_id][0]; int addr = (int)(uintptr_t)ptr \u0026amp; 0xFFFF; [[maybe_unused]] float r1; // 声明输出变量 for (int j = 0; j \u0026lt; num_iters; j++) { // num_iters 定义为 100'000 asm volatile(\"ld.volatile.shared.f32 %0, [%1];\" : \"=f\"(r1) : \"r\"(addr)); }}// launched withall_conflicts\u0026lt;\u0026lt;\u0026lt;1, dim3(32, 8)\u0026gt;\u0026gt;\u0026gt;();// Gride size: 1（即只有一个Block）// Block size: dim3(32, 8)（即有8个warp，每个warp 32个线程）由于是 32-way Bank Conflicts，则每个 warp 产生的 bank conflicts 次数是 10000 * 31，所有 warp 总共是 8 * 10000 * 31 = 2,480,000 次。与 Nsight Compute 测量结果吻合。2.1. conflict free 代码参考__global__ void conflict_free_kernel() { __shared__ float s[8][32]; int warp_id = threadIdx.y; int lane_id = threadIdx.x; float *ptr = \u0026amp;s[warp_id][lane_id]; int addr = (int)(uintptr_t)ptr \u0026amp; 0xFFFF; [[maybe_unused]] float r1; // 声明输出变量 for (int j = 0; j \u0026lt; num_iters; j++) { asm volatile(\"ld.volatile.shared.f32 %0, [%1];\" : \"=f\"(r1) : \"r\"(addr)); }}// conflict_free_kernel\u0026lt;\u0026lt;\u0026lt;1, dim3(32, 8)\u0026gt;\u0026gt;\u0026gt;();// Gride size: 1（即只有一个Block）// Block size: dim3(32, 8)（即有8个warp，每个warp 32个线程）2.2. 矢量读写指令使用矢量指令ld.shared.v4可以读取4个连续的32位数据。如下代码，一个线程读取s[4*i], s[4*i+1], s[4*i+2], s[4*i+3]（分为四个transaction）。会产生四路Bank Conflicts：__global__ void vectorized_loads() { __shared__ float sh[8][128]; int warp_id = threadIdx.y; int lane_id = threadIdx.x; float4* ptr = reinterpret_cast\u0026lt;float4*\u0026gt;(\u0026amp;sh[warp_id][lane_id * 4]); int addr = (int)(std::uintptr_t)ptr \u0026amp; 0xFFFF; float4 r; for (int j = 0; j \u0026lt; num_iters; j++) { asm volatile (\"ld.volatile.shared.v4.f32 {%0,%1,%2,%3}, [%4];\" : \"=f\"(r.x), \"=f\"(r.y), \"=f\"(r.z), \"=f\"(r.w) : \"r\"(addr)); }}// vectorized_loads\u0026lt;\u0026lt;\u0026lt;1, dim3(32, 8)\u0026gt;\u0026gt;\u0026gt;(); 上图只给出了一半的线程访问情况。编号(Lane)为0的线程与编号为8的线程，访问的shared memory中的数据映射到了同一个bank 0；同时，编号为16，以及24，同样映射到了bank 0。不过，由于其每个线程一次访问 4 个 32 位数据，其平均访问时间折算下来，与 32 位加载相当。要想避免Bank Conflicts，可以错开（interleave）冲突的线程访问的顺序，比如： 线程0：s[0] -\u0026gt; s[1] -\u0026gt; s[2] -\u0026gt; s[3] 线程8：s[33] -\u0026gt; s[34] -\u0026gt; s[35] -\u0026gt; s[32]3. 避免 Bank Conflicts 的方法3.1. Paddingwarp内多个线程访问同一bank会引发冲突，导致串行化访问。 通过在二维共享内存数组的列数上 +1 padding，可打破映射冲突：从第二行开始，Shared Memory中的数据到bank的映射偏移一个bank，且每行累积。示意图：示例代码：__shared__ float sData[BLOCKSIZE][BLOCKSIZE + 1]; // +1 避免bank冲突int x = threadIdx.x;int y = threadIdx.y;sData[x][y] = matrix[y * col + x];__syncthreads();matrixTest[y * col + x] = sData[x][y];当warp中的线程步长间距为128字节（32个32位数据）时，适用于padding，例如： 场景 冲突原因 Padding方案 列访问 行步长=32*4 列+1 步长访问 步长是32*k 改变数组维度 结构体数组 字段偏移相同 结构体+padding 斜向访问 特定步长产生周期 适当增加维度 3.2. SwizzleSwizzle是通过重新排列线程访问顺序，来避免Bank Conflicts。假设有32×32的二维数组，原本按列访问产生冲突：// 原始访问（产生冲突）int x = threadIdx.x;int y = threadIdx.y;float val = s[x][y]; // 同列线程映射到同一bank使用Swizzle变换：// Swizzle：对线程索引进行XOR操作int x = threadIdx.x;int y = threadIdx.y;int swizzled_x = x ^ (y % 32); // 用XOR改变x坐标float val = s[swizzled_x][y]; // 现在不同线程映射到不同bank内存布局对比：原始内存：[0,0] [1,0] [2,0] ... [31,0] \u0026lt;- 映射到bank 0,1,2...31[0,1] [1,1] [2,1] ... [31,1] \u0026lt;- 映射到bank 0,1,2...31（重复）...Swizzle后的访问顺序：线程[0,0]访问 s[0^0][0] = s[0,0]线程[1,0]访问 s[1^0][0] = s[1,0]线程[2,0]访问 s[2^0][0] = s[2,0]...线程[0,1]访问 s[0^1][1] = s[1,1] \u0026lt;- 同列不同线程，映射不同bank线程[1,1]访问 s[1^1][1] = s[0,1]...内存物理位置不变，但访问顺序改变了常见Swizzle操作：// 方法1：XOR swizzle（最常用）int swizzled = x ^ (y \u0026amp; (WARP_SIZE - 1));// 方法2：位移swizzleint swizzled = (x + y) % 32;// 方法3：混合操作int swizzled = (x + (y \u0026gt;\u0026gt; 4)) % 32;Swizzle更多资料： CUTLASS CuTe GEMM细节分析（三）——Swizzle\u0026lt;B, M, S\u0026gt;模板参数的取值 issue – how to understand “block swizzling”：Swizzle可以提升L2 cache命中率 issue – Swizzling the shared memory4. 概念总结Lane ID：每个线程在warp中的编号，范围为0-31，物理固定。sector：L1TEX / L2 缓存的最小传输单位是 sector，大小为 32 字节。一条缓存 cacheline = 128 字节 = 4 个 sector。一个 warp（32 线程）发出内存请求时，硬件会把所有线程的访问地址合并（coalesce），看需要覆盖几个 sector。wavefront：在 L1TEX语境下，表征 L1TEX 处理一次内存请求，需花费的次数。举例如下：情形一：4 sectors，1 wavefront（理想）32 个线程访问连续的 128 字节（完全合并） → 覆盖 4 个 sector，但地址都在同一个 cacheline → 1 个 wavefront 处理完毕，消耗 1 个周期情形二：4 sectors，2+ wavefronts（有冲突/分散）32 个线程访问的地址跨越多个 cacheline 或访问了 shared memory 中有 bank conflict 的位置→ 同样是 4 个 sector 的数据量→ 但 L1TEX 无法一次处理，需要拆成 2+ 个 wavefront→ 消耗 2+ 个周期内存事务（Memory Transaction）：每次访问共享内存时，硬件会将访问请求打包成内存事务。一个 warp 内的多个线程可能会访问同一个 bank，从而产生 bank conflicts，导致内存事务被串行化。transaction 与 wavefront 应该是同一个意思。Bank Conflicts相关名词：当发生 bank conflict 时，warp 需要额外的一个 cycle 来重新提交 shared memory 的访问指令到 LSU 单元，该指令需要在 MIO 中排队，这种排队会导致访问延迟增加，此时 warp 可能处于等待数据返回的状态，warp state 标识为 Stall Short Scoreboard。如果 MIO 队列满，此时 warp 先需要等待 MIO 队列处于非空的状态，此时 warp state 标识为 Stall MIO Throttle。A. 学习资料 CCUDA 编程手册系列第五章: 性能指南 NVIDIA – Performance Optimization: Paulius Micikevicius Programming Guidelines and GPU Architecture Reasons Behind Them cuda程序优化-2.访存优化 Optimize GEMM step by step。代码：https://github.com/gty111/GEMM_MMA，优化步骤见不同分支。" },

  { "title": "CUDA性能概述：影响因素及优化方法", "url": "/blog/2025/CUDA%E6%80%A7%E8%83%BD%E6%A6%82%E8%BF%B0/", "categories": "CUDA",
  "tags": "CUDA", "date": "2025-02-24", "content":
  "性能分析依据Roofline 模型，该模型根据算术强度（AI）划分两种性能区间（以A100为例）： 内存受限（Memory-bound）：当AI低于13 FLOPs/Byte时，性能由内存带宽决定。 计算受限（Compute-bound）：当AI高于13 FLOPs/Byte时，性能由计算能力决定。1. 性能影响及优化分类1.1. 内存优化 global memory – 数据重用：例如将大块内存数据加载到shared memory，以减少对global memory的访问次数，从而提高AI。 shared memory – Bank Conflicts：比如矩阵计算C = A * B中的共享内存访问冲突，将B矩阵加载到shared memory之后，进行转置，避免Bank Conflicts。 提升global memory带宽利用率。由于warp访问global memory时，是以128字节（32x4字节）为单位进行对齐访问，若warp内线程访问的地址不连续或未对齐，会导致多次transaction，从而降低带宽利用率。建议确保warp内线程访问的地址连续且对齐，以实现单次transaction访问完整的128字节数据。参考：https://www.olcf.ornl.gov/wp-content/uploads/2020/04/04-CUDA-Fundamental-Optimization-Part-2.pdf。 与带宽利用率有关的另一个概念是全局内存合并访问（Global memory coalescing）。即warp内线程访问的地址应连续且对齐，以实现单次transaction访问完整的128字节数据，从而提升带宽利用率。线程内避免跨步长访问。细节参考：https://cseweb.ucsd.edu/classes/wi12/cse260-a/Lectures/Lec09.pdf。1.2. 指令延迟优化 线程块并发（Occupancy）与延迟隐藏：合理配置共享内存和寄存器使用量，使SM可同时调度多个线程块，提高Warp的调度选择范围，从而隐藏内存访问延迟。 线程分叉（Thread Divergence）：Warp中线程执行路径不一致会导致序列化执行，降低吞吐率，建议使用分支无关的代码（如min/max替代if-else）以避免分歧。另外，如果数据分块不能完全分配到32个线程，可使用C+=A*0替代多余的条件分支。1.3. CPU-GPU 交互优化 使用Stream实现计算与数据传输重叠（Overlap）：使用Stream+异步传输+异步启动kernel。 并发启动多个kernel。 传输大块内存。 pinned memory：固定分配出一块内存给CPU/GPU交互使用，禁用内存页管理不会被换出。此时GPU驱动可以直接使用DMA，传输速度接近理论值。2. 参考资料 Introduction to CUDA Performance Optimization：认真看，全面介绍CUDA性能优化，以及以及性能参数计算。 NVIDIA – Fundamental Optimizations in CUDA：pdf文档。3. CUDA、CuTe 及其他资料收集 CUTLASS: Fast Linear Algebra in CUDA C++。强烈推荐，内容对应 GTC2018；详细讲解了在 cutlass 背景下针对 GEMM 优化的分块、内外积转换、缓存，以及相关基础概念和分层设计。 GEMM优化博客。强烈推荐，结合性能指标分析从一个最原始的 GEMM 开始优化，注重性能分析 nv blog 2025 cutlass。强烈推荐，内容对应 GTC 2023，介绍 cute，结合官方文档的 CuTe系列 理解 Layout 设计抽象。 nv 2025 blog cutlass 3 介绍。 3.x 系列引入的特性和抽象分层，对应 GTC 2023，是上文的补充第二部分。 GTC 2018：CUTLASS: Software Primitives for Dense Linear Algebra at All Levels and Scales within CUDA。初次介绍CUTLASS的GTC视频。 GTC 2019：PROGRAMMING TENSOR CORES: NATIVE VOLTA TENSOR CORES WITH CUTLASS。Volta、Turing架构下的cutlass优化。 GTC 2021：Accelerating Convolution with Tensor Cores in CUTLASS。卷积支持。 GTC 2023：Developing Optimal CUDA Kernels on Hopper Tensor Cores。cutlass 3.x 引入 cute 抽象 GTC 2024：CUTLASS: A Performant, Flexible, and Portable Way to Target Hopper Tensor Cores。卷积支持、epilog tree GTC 2025：Programming Blackwell Tensor Cores with CUTLASS。TBC，blackwell 特性。 The Present and Future of CUTLASS Tensor Core Programming3.1. 其他博客 reed CUTE系列 杨远航 CUTE 系列笔记 写给大家看的 CuTe 教程 – tiled mma，系列 tri dao 对 flash-attention 3 的介绍" },

  { "title": "CUDA 架构及对应的计算能力CC", "url": "/blog/2025/CUDA%E6%9E%B6%E6%9E%84%E5%8F%8A%E5%85%B6%E8%AE%A1%E7%AE%97%E8%83%BD%E5%8A%9BCC/", "categories": "CUDA",
  "tags": "CUDA", "date": "2025-02-23", "content":
  "1 计算能力 5.x1.1 体系架构计算能力 5.x 的 GPU 设备使用 Maxwell 架构（2014），下图展示了该架构下的流多处理器（SM）的结构示意图（以 GeForce GTX 980 为例）。https://developer.nvidia.com/blog/maxwell-most-advanced-cuda-gpu-ever-made/在 Maxwell 架构下，一个 SM 中包括： 128 个用于算术运算的 CUDA Core（相比于 Kepler 架构的 192 个有所缩减，但是峰值利用效率从 66.7\\% 提升到了几乎 100\\%）； 32 个用于单精度浮点运算的特殊函数单元； 4 个 warp 调度器； 由所有功能单元共享的常量内存的缓存，可加快从驻留在设备内存中的常量内存空间的读取速度； 24 KB 的纹理/L1 缓存，用于缓存来自全局内存的读取； 64 KB（计算能力 5.0）或 96 KB（计算能力 5.2）的共享内存。还有一个由所有 SM 共享的 L2 缓存（设备级别的，不属于单个 SM），用于缓存对本地或全局内存的访问，包括临时寄存器溢出。应用程序可以通过检查设备属性 l2CacheSize 来查询 L2 缓存大小。值得注意的是，在 Maxwell 架构中，Texture Cache 和普通全局内存数据的缓存（L1）合并成了 Unified Cache，大小为 24 KB，用户通过 const __restrict__ 修饰指针，或者使用 __ldg() 函数读取全局内存数据，就可以直接利用 Unified Cache，大部分情况下不需要手工写纹理提取的代码，就可以达到纹理内存的加速效果。1.2 全局内存Maxwell 架构下，对于全局内存访问的缓存通常在 L2 缓存中，但这也不是绝对的。如果编译器检测到某些数据在 Kernel 的整个执行过程中内是只读的，那么它会自行优化使用 __ldg() 函数来读取这些数据，将全局内存的数据缓存在 Unified Cache 中，当然用户也可以显式使用 __ldg() 函数完成这一行为。除此之外，编译器对数据是否满足只读条件的检测并不是绝对的，用户可以使用 const __restrict__ 修饰指针从而增加编译器检测到只读条件的概率。对于计算能力 5.0 的设备，如果数据在 Kernel 的整个生命周期内不是只读的，那么该数据不能缓存 Unified Cache 中；而对于计算能力为 5.2 的设备，默认情况下全局内存的数据不会缓存在 Unified Cache 中，但可以使用以下机制启用缓存： 使用带有适当修饰符的内联汇编执行读取，参考 PTX 参考手册； 使用 -Xptxas -dlcm=ca 编译标志进行编译，在这种情况下，所有读取都被缓存，除了使用带有禁用缓存的修饰符的内联汇编执行的读取场景； 使用 -Xptxas -fscm=ca 编译标志进行编译，在这种情况下，所有读取都被缓存，包括使用内联汇编执行的读取，无论使用何种修饰符。在计算能力 5.2 的设备上使用上面列出的三种机制之一启用缓存时，除非 Kernel 启动的时候 block 内消耗了太多的 SM 寄存器资源，否则将会把所有 Kernel 中全局内存的读取缓存到 Unified Cache 中。1.3 共享内存共享内存在物理上被划分为连续的 32 个 bank，这些 bank 宽度相同且能被同时访问，Maxwell 架构中 bank 宽度为 32bit，bank 带宽为每个时钟周期 32bit。如果来自一个 warp 的多个线程访问到共享内存中相同 32bit 内的某个 byte 时，不会产生 bank conflict，读取操作会在一次广播中完成，如果是写入操作，则每个地址仅会被其中一个线程写入，具体是哪个线程未定义；而如果访问的是同一个 bank 的不同 32bit 的地址，则会产生 bank conflict。我们可以把共享内存的 32 个 bank 想象为由很多层组成，每个 bank 每层有 32bit，假设同一 warp 内的不同线程访问到某个 bank 的同一层的数据，此时不会发生 bank conflict，但如果同一 warp 内的不同线程访问到某个 bank 的不同层的数据，此时将产生 bank conflict。2 计算能力 6.x2.1 体系架构计算能力 6.x 的 GPU 设备使用 Pascal 架构（2016），Pascal 架构和 Maxwell 架构实际上并无的本质区别，可以认为是对 Maxwell 架构的修补版本，而且这代 GPU 的发布正好迎来了 AI 和深度学习的浪潮。下图展示了该架构下的流多处理器（SM）的结构示意图（以 GeForce GTX 1080 为例）。 https://developer.nvidia.com/blog/inside-pascal/在 Pascal 架构下，一个 SM 中包括： 64（计算能力 6.0）或 128（计算能力 6.1 和 6.2）个用于算术运算的 CUDA Core； 16（计算能力 6.0）或 32（计算能力 6.1 和 6.2）个用于单精度浮点运算的特殊函数单元； 2（计算能力 6.0）或 4（计算能力 6.1 和 6.2）个 warp 调度器； 由所有功能单元共享的常量内存的缓存，可加快从驻留在设备内存中的常量内存空间的读取速度； 24 KB（计算能力 6.0 和 6.2）或 48 KB（计算能力 6.1）的纹理/L1 缓存，用于缓存来自全局内存的读取； 64 KB（计算能力 6.0 和 6.2）或 96 KB（计算能力 6.1）的共享内存。还有一个由所有 SM 共享的 L2 缓存（设备级别的，不属于单个 SM），用于缓存对本地或全局内存的访问，包括临时寄存器溢出。应用程序可以通过检查设备属性 l2CacheSize 来查询 L2 缓存大小。在 Pascal 架构中，纹理缓存和 L1 缓存依然保留了 Maxwell 架构下的组织形式，两者合并成了 Unified Cache，从硬件层面上支持了纹理提取和表面提取的各种寻址模式和数据过滤策略。2.2 全局内存Pascal 架构下全局内存的行为方式与 Maxwell 架构相同。2.3 共享内存Pascal 架构下共享内存的行为方式与 Maxwell 架构相同。3 计算能力 7.x3.1 体系架构计算能力 7.x 的 GPU 设备可以分成两种不同架构，即计算能力 7.0 的 Volta 架构、计算能力 7.5 的 Turing 架构，这两个架构某种程度上的可以看成同一种架构，相当于 Maxwell 架构中的计算能力 5.0 和 5.2 的关系，而不是看成 Turing（7.5）是 Volta（7.0）的下一代全新架构。只是在 Turing 架构中新引入了支持整数和浮点的并发执行能力以及 Tensor Core 的增强，除此之外并无本质区别。其中，基于 Volta 架构的 GPU 产品主要是 Tesla 和 Jetson 系列，下图展示了该架构下的流多处理器（SM）的结构示意图（以 Tesla GV100 为例）。而基于 Turing 架构的 GPU 产品主要是 Tesla、Quadro、和 GeForce 系列，其流多处理器（SM）的结构与 Volta 架构基本一致，不再展示 SM 示意图，有兴趣的读者可自行查阅产品白皮书。 https://www.nas.nasa.gov/hecc/support/kb/basics-on-nvidia-gpu-hardware-architecture_704.html https://www.olcf.ornl.gov/wp-content/uploads/2018/12/summit_workshop_Volta-Architecture.pdf在 Volta 和 Turing 架构下，一个 SM 中包括： 64 个用于单精度浮点算术运算的 FP32 Core； 32 个用于双精度浮点算术运算的 FP64 Core； 64 个用于整数算术运算的 INT32 Core； 8 个用于深度学习矩阵运算的混合精度 Tensor Core； 16 个用于单精度浮点运算的特殊函数单元； 4 个 warp 调度器； 由所有功能单元共享的常量内存的缓存，可加快从驻留在设备内存中的常量内存空间的读取速度； 一个统一的 128 KB（Volta 架构）或 96 KB（Turing 架构）的 纹理/L1 缓存/共享内存。在 Volta 和 Turing 架构中，L1 缓存、纹理缓存和共享内存三者在物理上被统一成了一个有 128 KB 或 96 KB 的数据缓存，用户可以使用 CUDA Runtime API 来指定其中共享内存所占的容量大小。3.2 独立线程调度Volta 架构在 warp 内的线程之间引入了独立线程调度（Independent Thread Scheduling），此功能引入了之前不可用的 warp 内同步模式，并简化了移植 CPU 代码时的代码更改。然而，如果用户对之前架构的 warp 同步模式进行了假设（即隐含地假设同一 warp 中的线程在每个指令上都是同步的代码），那么独立线程调度机制也可能导致 warp 内实际参与执行代码的线程集合与预期的线程集合截然不同。对于使用到 warp 内置函数（如 __shfl*、__any、__all、__ballot）的应用程序，开发人员有必要将这些函数替换为具有 *_sync 后缀的新函数。新的 warp 内置函数新增了线程掩码参数，明确定义哪些通道（warp 内的线程）参与 warp 内置函数运算。关于 warp 内置函数的详细信息，请参阅【CUDA编程】束内表决函数（Warp Vote Function）和【CUDA编程】束内洗牌函数（Warp Shuffle Functions）。由于这些 warp 内置函数是从 CUDA 9.0 开始引入的，所以为了提高代码兼容性，建议通过如下预处理器宏有条件地执行代码：#if defined(CUDART_VERSION) \u0026amp;\u0026amp; CUDART_VERSION \u0026gt;= 9000// *_sync intrinsic#endif要注意的是，这些内置函数可用于所有架构，而不仅仅是 Volta 或 Turing 架构，并且在大多数情况下，单个代码库就足以满足所有架构的需求。但是对于 Pascal 和更早的架构，mask 中的所有线程在聚合时必须执行相同的 warp 内置函数，并且 mask 中所有值的并集必须等于 warp 内置函数的掩码。以下代码模式在 Volta 架构上合法，但在 Pascal 或更早的架构上不合法。if (tid % warpSize \u0026lt; 16) { ... float swapped = __shfl_xor_sync(0xffffffff, val, 16); ...} else { ... float swapped = __shfl_xor_sync(0xffffffff, val, 16); ...}对于 Pascal 或更早的架构的正确写法应该是在 __shfl_xor_sync 中分别传入 0x0000ffff、0xffff0000 作为 mask，这样才能保证实际活动线程与 mask 一致。在 Volta 架构下，旧版本 __ballot(1) 函数的替代品是 __activemask() 函数，用来返回活动状态的线程掩码。要注意的是，即使在单个代码路径中，warp 中的线程也可能发散，也就是说 warp 发散并不仅限于代码路径这种显式发散的场景，也可能隐式地出现发散。因此，__activemask() 和 __ballot(1) 可能只返回当前代码路径上的活动线程子集。比如以下代码示例：// Sets bit in output[] to 1 if the correspond element in data[i]// is greater than ‘threshold’, using 32 threads in a warp.for(int i=warpLane; i\u0026lt;dataLen; i+=warpSize) { unsigned active = __activemask(); unsigned bitPack = __ballot_sync(active, data[i] \u0026gt; threshold); if (warpLane == 0) output[i/32] = bitPack;}上面代码本来在 Pascal 或更早的架构是 warp 聚合的代码，但从 Volta 架构后就不保证聚合了。比如：当 dataLen 不是 warpSize 的整数倍时会有隐式的 warp 发散。但是 __activemask() 只针对当前活动的线程，并不一定包括所有要经过这个地方的 lane，就是说 Volta 架构后 CUDA 不保证 warp 只在循环条件下发散，如果有其他原因导致发散，那可能计算出的结果就会和预期不同。比如说有 mask 为 0xffff0000 的 16 个线程要执行 __activemask()，Pascal 或更早的架构会对这 16 个线程统一返回 active = 0xffff0000。但从 Volta 架构后，可能这 16 个线程是分两组过去的，可能前 8 个线程得到 active = 0xff000000，后 8 个线程得到 active = 0x00ff0000。官方文档中给出的一个正确写法如下：for(int i=warpLane; i-warpLane\u0026lt;dataLen; i+=warpSize) { unsigned active = __ballot_sync(0xFFFFFFFF, i \u0026lt; dataLen); if (i \u0026lt; dataLen) { unsigned bitPack = __ballot_sync(active, data[i] \u0026gt; threshold); if (warpLane == 0) output[i/32] = bitPack; }}新的代码中使用了非发散的循环条件和束内表决函数 __ballot_sync() 安全地表决出当前代码路径下的满足 i \u0026lt; dataLen 线程的掩码。Volta 架构引入了 warp 内同步函数 __syncwarp()，应用程序中针对全局内存和共享内存的读写操作，如果假设了某个操作对同一 warp 中的其他线程可见，那么必须在相应的读写操作之间插入同步指令 __syncwarp()，显式地让 warp 内的线程都到达该屏障点。如果没有显式使用 __syncwarp() 进行同步，那么任何关于 warp 内线程步调一致执行的假设都是不可靠并错误的。下面给出一个 block 内规约场景使用到 warp 内同步函数 __syncwarp() 的代码示例：__shared__ float s_buff[BLOCK_SIZE];s_buff[tid] = val;__syncthreads();// Inter-warp reductionfor (int i = BLOCK_SIZE / 2; i \u0026gt;= 32; i /= 2) { if (tid \u0026lt; i) { s_buff[tid] += s_buff[tid+i]; } __syncthreads();}// Intra-warp reduction// Butterfly reduction simplifies syncwarp maskif (tid \u0026lt; 32) { float temp; temp = s_buff[tid ^ 16]; __syncwarp(); s_buff[tid] += temp; __syncwarp(); temp = s_buff[tid ^ 8]; __syncwarp(); s_buff[tid] += temp; __syncwarp(); temp = s_buff[tid ^ 4]; __syncwarp(); s_buff[tid] += temp; __syncwarp(); temp = s_buff[tid ^ 2]; __syncwarp(); s_buff[tid] += temp; __syncwarp();}if (tid == 0) { *output = s_buff[0] + s_buff[1];}__syncthreads();尽管 _syncthreads() 一直被认为是同步 block 中的所有线程，但在 Pascal 或更早的架构中只支持在 warp 级别强制同步。因此在某些情况下，只要每个 warp 中至少有一些线程到达屏障，这就会认为 block 所有线程均到达屏障点，这在独立线程调度引入之前是可行的。但从 Volta 架构开始，CUDA 内置的 __syncthreads() 和 PTX 指令 bar.sync（及其派生类）在 block 内每个线程中强制执行，因此在 block 中所有未退出的线程到达之前，屏障不会成功。3.3 全局内存Volta 和 Turing 架构下全局内存的行为方式与 Maxwell 架构相同。3.4 共享内存前面介绍过，在 Volta 和 Turing 架构中，L1 缓存、纹理缓存和共享内存三者在物理上被统一成了一个数据缓存，其中共享内存容量可以使用 CUDA Runtime API 进行设置。对于 Volta 架构，统一数据缓存大小为 128 KB，共享内存容量可以设置为 0、8、16、32、64 或 96 KB；而对于 Turing 架构，统一数据缓存大小为 96 KB，共享内存容量可以设置为 32 或 64 KB。与 Kepler 架构不同，驱动程序会自动为每个 Kernel 配置共享内存容量以避免共享内存占用瓶颈，从而允许该 Kernel 与已启动的 Kernel 并发执行。在大多数情况下，驱动程序的默认行为会提供最佳性能。由于驱动程序并不保证能知道 Kernel 的全部共享内存负载，所以有些情况下应用程序显式提供共享内存容量配置会取得更好的性能，比如，对于某些使用很少或不使用共享内存的 Kernel 来说，给其设置更大的共享内存容量，一定程度上可能增加它与其他 Kernel（比如需要使用较多共享内存）的并发程度。在 Volta 和 Turing 架构中提供了新的 cudaFuncSetAttribute() API，使得用户可以自定义一个首选的共享内存容量（即在统一数据缓存中的划分比例），注意这个 API 是 per-Kernel 级的。与 Kepler 架构引入的传统 cudaFuncSetCacheConfig() API 相比，cudaFuncSetAttribute() 放宽了启动时对首选共享容量的要求，也就是说，旧版 API 将共享内存容量视为内核启动的硬性要求，导致不同共享内存容量配置的 Kernel 会先执行一个 Kernel，完成后重设共享内存容量，然后再执行下一个 Kernel。而使用新的 API 时，用户指定的划分比例只是一个提示，驱动程序可能根据实际情况选择与此不同的配置以获取最佳性能。下面给出使用 cudaFuncSetAttribute() API 的代码示例：// Device code__global__ void MyKernel(...){ __shared__ float buffer[BLOCK_DIM]; ...}// Host codeint carveout = 50; // prefer shared memory capacity 50% of maximum// Named Carveout Values:// carveout = cudaSharedmemCarveoutDefault; // (-1)// carveout = cudaSharedmemCarveoutMaxL1; // (0)// carveout = cudaSharedmemCarveoutMaxShared; // (100)cudaFuncSetAttribute(MyKernel, cudaFuncAttributePreferredSharedMemoryCarveout, carveout);MyKernel \u0026lt;\u0026lt;\u0026lt;gridDim, BLOCK_DIM\u0026gt;\u0026gt;\u0026gt;(...);上面的代码中把 MyKernel 的首选共享内存容量占比设定为 50\\%，注释中还提供了几个其他的参考枚举值。如果选择的整数百分比不能完全映射到支持的容量（比如 Volta 架构的设备支持 0、8、16、32、64 或 96 KB 的共享容量），则自动使用下一个更大的容量。例如，在上面的示例中，最大容量 96 KB 的 50\\% 是 48 KB，这不是受支持的共享内存容量，因此，实际会向上舍入为 64 KB。在 Volta 和 Turing 架构中，允许单个线程块能够寻址全部 96 KB 的共享内存。为了保证不同架构下代码兼容性，静态共享内存分配仍被限制为 48 KB，如需寻址超过 48 KB 的共享内存，需要使用动态共享内存的方式，并且使用 cudaFuncSetAttribute() API 显式指定，示例代码如下。// Device code__global__ void MyKernel(...){ extern __shared__ float buffer[]; ...}// Host codeint maxbytes = 98304; // 96 KBcudaFuncSetAttribute(MyKernel, cudaFuncAttributeMaxDynamicSharedMemorySize, maxbytes);MyKernel \u0026lt;\u0026lt;\u0026lt;gridDim, blockDim, maxbytes\u0026gt;\u0026gt;\u0026gt;(...);3.5 Tensor Core与上一代 Pascal 架构相比，Volta 架构引入了 Tensor Core，并在 Turing 架构中进一步增强，以满足神经网络中数以百万计的神经元所需的训练、推理性能。每个 Tensor Core 执行以下运算：$D = A \\times B + C$其中 $A$、$B$、$C$ 和 $D$ 是 $4\\times4$ 矩阵，矩阵 $A$、$B$ 元素类型为 FP16，而累加矩阵 $C$ 和 $D$ 的元素类型可以是 FP16 或 FP32。在实际应用中，Tensor Core 被用来执行更大的 2D 或更高维的矩阵运算，这些大矩阵的乘法运算都是由这些 $4\\times4\\times4$ 的矩阵运算组成的。在 CUDA 9 C++ API 中提供了 Warp-Level 矩阵运算 API 用来调动 Tensor Core 进行运算。该 API 支持专门的矩阵加载、矩阵乘法、累加以及矩阵存储操作，以有效地使用 Tensor Core 的性能。除了直接使用 CUDA C++ Warp-Level 矩阵运算 API 以外，cuBLAS 和 cuDNN 等官方库也已经完成底层代码更新以利用 Tensor Core 进行深度学习应用，用户可以直接使用这些官方库，而不必基于 Warp-Level 矩阵运算 API 重新编写代码。4 计算能力 8.x4.1 体系架构计算能力 8.x 的 GPU 设备采用 Ampere 架构，下图展示了该架构下的流多处理器（SM）的结构示意图（以 Tesla GA100 为例）。 https://developer.nvidia.com/blog/nvidia-ampere-architecture-in-depth/ https://en.wikipedia.org/wiki/Ampere_(microarchitecture)在 Ampere 架构下，一个 SM 中包括： 64（计算能力 8.0）或 128（计算能力 8.6、8.7、8.9）个用于单精度浮点算术运算的 FP32 Core； 32（计算能力 8.0）或 2（计算能力 8.6、8.7、8.9）个用于双精度浮点算术运算的 FP64 Core； 64 个用于整数算术运算的 INT32 Core； 4 个混合精度第三代 Tensor Core，支持 fp16、__nv_bfloat16、tf32、 sub-byte、fp64 等数据类型的矩阵运算（计算能力 8.0、8.6、8.7）； 4 个混合精度第四代 Tensor Core，支持 fp8、fp16、__nv_bfloat16、tf32、 sub-byte、fp64 等数据类型的矩阵运算（计算能力 8.9）； 16 个用于单精度浮点运算的特殊函数单元； 4 个 warp 调度器； 由所有功能单元共享的常量内存的缓存，可加快从驻留在设备内存中的常量内存空间的读取速度； 一个统一的 192 KB（计算能力 8.0、8.7）或 128 KB（计算能力 8.6、8.9）的 纹理/L1 缓存/共享内存。同 Volta 和 Turing 架构一样，在 Ampere 架构中 L1 缓存、纹理缓存和共享内存三者在物理上被统一成了一个有 192 KB 或 128 KB 的数据缓存，用户可以使用 CUDA Runtime API 来指定其中共享内存所占的容量大小。4.2 全局内存Ampere 架构下全局内存的行为方式与 Maxwell 架构相同。4.3 共享内存同 Volta 和 Turing 架构一样，共享内存容量可以使用 CUDA Runtime API 进行设置，只是可设置的共享内存容量值不同。对于计算能力为 8.0、8.7 的设备，统一数据缓存大小为 192 KB，共享内存容量可以设置为 0、8、16、32、64、100、132 或 164 KB；而计算能力为 8.6、8.9 的设备，统一数据缓存大小为 128 KB，共享内存容量可以设置为 0、8、16、32、64 或 100 KB。计算能力 8.0、8.7 的设备允许单个线程块寻址多达 163 KB 的共享内存，而计算能力 8.6、8.9 的设备允许多达 99 KB 的共享内存。Ampere 架构中每个线程块的最大共享内存量小于每个 SM 上可用的最大共享内存分区，未提供给线程块使用的 1 KB 共享内存将保留给系统使用。超过 48 KB 的共享内存使用，需要使用动态共享内存的方式，并且使用 cudaFuncSetAttribute() API 显式指定。5 计算能力 9.x5.1 体系架构当前最新的 GPU 设备计算能力最高只到 9.0，采用 Hopper 架构，下图展示了该架构下的流多处理器（SM）的结构示意图（以 Tesla GH100 为例）。 https://developer.nvidia.com/blog/nvidia-hopper-architecture-in-depth/ https://medium.com/@nimritakoul01/nvidia-hopper-h100-tensor-core-gpu-architecture-8c356614dbea在 Hopper 架构下，一个 SM 中包括： 128 个用于单精度浮点算术运算的 FP32 Core； 64 个用于双精度浮点算术运算的 FP64 Core； 64 个用于整数算术运算的 INT32 Core； 4 个混合精度第四代 Tensor Core，支持 fp8、fp16、__nv_bfloat16、tf32、 INT8、fp64 等数据类型的矩阵运算； 16 个用于单精度浮点运算的特殊函数单元； 4 个 warp 调度器； 由所有功能单元共享的常量内存的缓存，可加快从驻留在设备内存中的常量内存空间的读取速度； 一个统一的 256 KB 的纹理/L1 缓存/共享内存。在 Hopper 架构中 L1 缓存、纹理缓存和共享内存三者在物理上被统一成了一个有 256 K 的数据缓存，用户可以使用 CUDA Runtime API 来指定其中共享内存所占的容量大小，从 Volta 架构开始，连续 4 个大版本一直保持该形式，只是缓存容量有所变化。5.2 全局内存Hopper 架构下全局内存的行为方式与 Maxwell 架构相同。5.3 共享内存与 Ampere 架构类似，共享内存容量可以使用 CUDA Runtime API 进行设置，只是可设置的共享内存容量值不同。对于计算能力为 9.0 的设备（如 NVIDIA H100 Tensor Core GPU），统一数据缓存的大小为 256 KB，共享内存容量可以设置为 0、8、16、32、64、100、132、164、196 或 228 KB。计算能力 9.0 的设备允许单个线程块寻址多达 227 KB 的共享内存，超过 48 KB 的共享内存使用，需要使用动态共享内存的方式，并且使用 cudaFuncSetAttribute() API 显式指定。5.4 加速专用计算Hopper 架构还引入了以下特性，用来加速矩阵乘加（MMA）计算： MMA 指令的异步执行； 应用于跨 warp 大矩阵的 MMA 指令； 在 warp 组之间动态重新分配寄存器容量，以支持更大的矩阵； 直接从共享内存访问的操作数矩阵这些功能仅在 CUDA 编译工具链中通过内联 PTX 提供，并已经在 CUDA 官方库的底层代码中更新。建议用户通过 CUDA-X 库（如 cuBLAS、cuDNN 或 cuFFT）利用上述特性。同时也建议用户在编写 Kernel 代码时通过 CUTLASS 库利用这些特性，CUTLASS 是一个基于 CUDA C++ 模板抽象的集合，使得用户可以在 CUDA 内的所有级别和尺度上进行高性能矩阵乘法（GEMM）和相关计算。最近几代架构资料 NVIDIA GPU Architecture: from Pascal to Turing to Ampere CUDA GPU Compute Capability：不同计算能力CC对应的 GPU 设备列表。 Matching CUDA arch and CUDA gencode for various NVIDIA architectures：不同GPU及CC对应的 gencode 列表。 NVIDIA Hopper Architecture In-Depth" },

  { "title": "CUDA 架构(1.1)：Hopper架构及性能分析(ncu) + 性能优化", "url": "/blog/2025/CUDA%E6%9E%B6%E6%9E%841.1-Hopper%E5%8F%8A%E6%80%A7%E8%83%BD%E4%BC%98%E5%8C%96/", "categories": "CUDA",
  "tags": "CUDA", "date": "2025-02-23", "content":
  "1. GPU thread hierarchy, SIMT, and Warp divergence1.1. Thread Hierarchy: Grid \u0026amp; Blocksgrid是thread blocks的集合，代表这次启动kernel的全部工作。thread block之间完全独立，不能共享数据，不能通过__syncthreads同步，不能假设他们的执行顺序，不能假设两个thread block同时启动及运行。有关thread block与SM： 若干个thread block可以同时在一个SM上运行（可以达到Hide Latency），前提是它们的资源需求（寄存器、共享内存等）满足SM的资源限制（Occupancy）。 同一个thread block的所有线程必须在同一个SM上运行，因为它们需要共享资源（如共享内存）和进行同步。\\[占用率Occupancy = \\frac{\\text{SM上实际驻留的Warp数}}{\\text{SM支持的最大Warp数}}\\] 资源 SM上限 如何影响 最大Block数 32 硬性上限，不管Block多小 最大Warp数 64 即最多2048个线程同时驻留 最大线程数 2048 同上 寄存器文件 65536个32-bit寄存器 每个线程用的寄存器越多，能驻留的线程越少 共享内存 最大228 KB 每个Block用的shared memory越多，能放的Block越少 计算Occupancy举例：场景1：每个Block有256个线程，每个线程用32个寄存器，每个Block用0 shared memory线程数限制： 2048 / 256 = 8 个BlockWarp数限制： 64 / (256/32) = 64/8 = 8 个Block寄存器限制： 65536 / (256 × 32) = 65536/8192 = 8 个BlockBlock数限制： 32→ 结果：min(8, 8, 8, 32) = 8 个Block可以驻留即哪个资源限制最紧，该资源就形成瓶颈。1.2. Thread Hierarchy: ClustersHopper架构引入了Thread Block Clusters，以及硬件上的分布式共享内存（Distributed Shared Memory, DSMEM）。 Cluster可配置为1/2/4/8/16个thread block。 Cluster保证Cluster内的Thread Block同时被调度到不同的SM上执行，以保证他们之间的通信，以及避免等待死锁。 Clusters：将多个thread block组织成一个cluster，允许它们共享数据和进行更高效的通信。每个cluster内的thread block可以通过DSM进行数据交换。应用场景： GEMM：每个block负责一个tile，多个block组成一个cluster，cluster内的block通过DSM共享tile数据，不用重复从全局内存访问。 Reduction：每个block负责一部分数据，多个block组成一个cluster，cluster内的block通过DSM共享中间结果。1.3. SIMT Architecture有条件分支的时候，warp内每个线程需要执行所有的分支路径，即线程分叉（warp diverge），导致部分clock浪费。图中，标记红叉的部分是浪费的时钟周期，这部分计算生成输出。在使用ncu分析时，从Warp State Statistics中查看线程分叉指标，以及从Source页面查看一些指标。2. Warp scheduling and Kernel profiling at a glance2.1. Warp Scheduler StatisticsWarp Slot状态（在Warp Scheduler Statistics页面中查看）： 状态 含义 性能影响 Unused Slot空闲，没有Warp Occupancy低，浪费硬件能力 Active Warp驻留在SM上 越高越好（更多Warp可供调度） Stalled Active但在等待某些东西 正常现象，但占比过高说明有瓶颈 Eligible 万事俱备，等待被选中 越多越好（调度器有的选） Selected 本周期被选中发射指令 受限于Scheduler数量（每SM 4个） 含义：这个Warp Slot上有一个Warp驻留在SM上，它的寄存器、状态等都保存在硬件中。Active = Stalled + Eligible（包括Selected）。Stall Reasons原因分类：Stall来源列表： Warp State / Stall Reason 含义 常见触发场景 Stall Math Pipe Throttle 需要的数学/计算流水线（FP/INT/SFU 等）繁忙，当前 warp 发不进去（结构性拥塞） 大量 FMA/FP64/INT/特殊函数 sin/exp/rsqrt 等，或者指令高度集中在同一类计算管线 Stall LG Throttle Local/Global 内存管线（LD/ST）相关资源/队列饱和，导致新的 LD/ST 不能及时发射 频繁 global load/store、local memory（寄存器溢出导致的本地内存访问）、大量 scatter/gather 访问 Stall MIO Throttle 内存 I/O 相关的“杂项/特殊”管线拥塞（通常是某类内存/原子/特殊访存通路资源受限），warp 因结构性资源不足而等待 原子操作密集、某些特殊内存指令组合（不同于纯 LG/纯 TEX 的那类）、高频内存系统操作 Stall Long Scoreboard 等待长延迟指令的结果（典型：global memory、L2/DRAM miss），数据没回来导致后续指令依赖无法就绪 读取全局内存后立刻使用（依赖链短）、cache miss 多、访问不合并 Stall Short Scoreboard 等待相对短延迟指令的结果（例如 shared/常量/某些管线较短的依赖），仍未就绪 shared memory 访问后紧接使用、shared bank conflict、常量/纹理命中但仍有固定延迟 Stall Barrier 等待同步点（barrier）满足（如 __syncthreads()），本 warp 已到但其他 warp 未到 线程块内工作不均衡、分支导致部分 warp 慢、sync 过频繁 Stall Branch Resolving 控制流相关原因导致等待（分支目标/谓词/重汇合等处理阶段），本周期无法发射 分支多、分支发散严重、间接跳转/复杂控制流 Stall Wait 等待某类“等待/依赖管理”相关的内部条件完成（NCU 将其归为 Wait 类等待） 某些需要等待条件满足的指令序列（与依赖/同步/内存子系统管理相关） Stall Not Selected warp 其实是 Eligible，但当周期调度器选了别的 warp（发射槽位有限） Eligible warps 很多、调度竞争激烈；通常是“系统健康”的表现之一 Selected （积极状态） Warp 被调度器选中，在本周期发射了一条指令。 无阻塞。这是性能良好的标志，表明调度器能持续找到可执行的 Warp。该值应尽可能高。若 Selected 低而 Stall 高，说明存在瓶颈。 2.2. ncu 瓶颈分析其中，Latency Bound是最糟糕的情况。既没有充分利用计算资源，也没有跑满内存带宽。程序的执行速度受限于各种操作的延迟（如内存读取延迟、分支延迟、同步延迟）。由于没有足够的并行工作（Warp）来隐藏这些延迟，导致GPU的执行单元大部分时间处于空闲等待状态。优化手段包括： 增加 occupancy（减少寄存器/共享内存占用、调整 block 大小）、增加并行度； 减少同步与串行化：减少 __syncthreads()、减少热点原子； 降低依赖链：插入独立计算、流水化/双缓冲，或增加流水深度； 减少发散：数据重排、warp 内路径一致；Comp \u0026amp; Mem Bound，最理想的状态，性能已接近硬件的Roofline，意味着已经成功地将程序的计算和内存利用都推向了硬件的极限。系统同时处于计算和带宽的饱和状态。2.3. ncu 中相关 Page 页面首先，使用GPU Speed of Light页面总体区分是否存在性能优化空间，以及性能瓶颈的类型。Speed Of Light (SOL) Throughput给出是compute bound还是memory bound，或者latency bound。如上图中，红圈中的占比数值对应上面表格中的蓝色条。百分比小的性能指标，表示数据量小，或者复用率高。Scheduler statistics页面给出warp调度状态统计信息：Stall占比、Eligible(就绪)占比、Select(被选中)占比等。Warp State Staticstics则给出了具体Stall占比，即性能瓶颈的来源。2.4. Source Page 中的性能指标 📝在Source Page页面中，查看瓶颈的对应的代码/汇编，比如： Compute Bound：查看Instructions Executed（指令执行栏），并通过菜单栏定位到占比最多的指令； Memory Bound：查看L1 Tag Requests Global/L1 Wavefronts Shared/L2 Theotetical Sectors Global等栏目； Latency Bound：查看Warp Stall Samplining(Not issued Samplies)栏目； 📝对一些Stall比较长的计算指令，使用Scoreboard Dependencies查看其依赖的指令和资源，从而分析瓶颈原因。（Scoreboard即记分牌，有专门的硬件记录指令的状态，比如在等待从内存加载数据）3. Roofline 模型Roofline模型是一个性能分析工具，用于评估程序的性能瓶颈。它将程序的性能限制分为两类：计算限制（Compute Bound）和内存限制（Memory Bound）。通过分析程序的算术强度（Arithmetic Intensity）和硬件的计算能力与内存带宽，可以确定程序是受计算资源限制还是内存带宽限制。其中： $\\pi$：理论计算性能峰值，单位为FLOPS或FLOP/s（每秒浮点运算次数）。 $\\beta$：理论内存带宽峰值，单位为Bytes/s（每秒字节数）。 $I$：算术强度，单位为FLOP/Byte（每字节的浮点运算次数）。 $P$：当前$I$能达到的性能，单位为FLOPS或FLOP/s。Roofline限制线描述的是： 当数据搬运的速度比较小的时候，此时计算在等待数据，性能受内存带宽限制（Memory Bound），性能随算术强度$I$线性增长。 当数据搬运速度达到一定程度时，计算速度达到理论峰值，此时性能受计算资源限制（Compute Bound），性能达到平台的计算峰值$\\pi$，不再随算术强度$I$增加。用公式描述Roofline的这两段瓶颈区如下：\\[P = \\min(\\beta \\cdot I, \\pi)\\]或者：\\[P =\\begin{cases}\\beta \\cdot I \u0026amp; \\text{if } I \u0026lt; I_{max} \u0026amp; \\text{Memory Bound} \\\\\\pi \u0026amp; \\text{if } I \\geq I_{max} \u0026amp; \\text{Compute Bound}\\end{cases}\\]\\[I_{max} = \\frac{\\pi}{\\beta}\\] 补充理解：当实际点落在斜坡线上时，说明程序的数据搬运比较慢：没有合并访存，L1/L2 Cache Miss比较高等原因。3.1. Roofline演示1定位性能瓶颈类型： 落在带宽斜坡 → 内存受限（Memory Bound） 落在算力平顶 → 计算受限（Compute Bound） 两者都远未达到 → 延迟受限（Latency Bound）图中，红色点表示Latency Bound，即FLOPS上不去，瓶颈可能来自： 内存访问延迟，如L1/L2 Cache Miss较高，没有合并访存； 指令延迟，如指令依赖，或者分支发散严重（此时issue Slot Utilization较低）； SM occupancy较低，无法隐藏指令/内存访问的延迟； 同步开销，如__syncthreads()过多，原子操作等； 存在非FMA计算；3.2. Roofline 瓶颈优化3.2.1. Memory Bound 时重点查看 NCU 指标 问题信号 对应优化方向 Global Load/Store Efficiency \u0026lt; 100% 未合并访问，检查 SoA/AoS、对齐、stride L1 Hit Rate 过低 数据局部性差，考虑 tiling / __ldg() L2 Hit Rate 过低 工作集超出 L2；考虑 L2 persistence / 减小 tile Shared Memory Efficiency \u0026lt; 100% 存在 bank conflict DRAM Throughput 接近峰值但 kernel 仍慢 已达带宽极限，需减少访存量（fusion / 算法改进） Sectors/Request \u0026gt; 1（理想值为 1） 未对齐或未合并 Stall Long Scoreboard 高 全局内存访问延迟未隐藏 3.2.2. Compute Bound 时重点查看 NCU 指标 问题信号 对应优化方向 Branch Efficiency \u0026lt; 100% 分支发散 Eligible Warps Per Cycle 过低 ILP 不足或 occupancy 过低 Issue Slot Utilization \u0026lt; 50% 指令调度不饱和 FP32/FP16/Tensor Pipe Utilization 不均衡 未使用合适的计算管线（如该用 Tensor Core 没用） Register Spill (Local Memory) \u0026gt; 0 寄存器溢出 3.2.3. Latency Bound 时重点查看 NCU 指标 问题信号 对应优化方向 Occupancy (Achieved vs Theoretical) 差距大 寄存器 / Shared Memory 用量过高 Stall Reasons 面板 高 Stall Not Selected / Stall Barrier 同步开销过大 Stall Long Scoreboard 高 全局内存延迟未隐藏 Stall Short Scoreboard 高 Shared Memory / L1 延迟未隐藏，special math instructions (MUFU)、dynamic branching (BRX/JMX) Stall LG Throttle（LG 指令队列满） 高   A. 资料 CUDA Techniques to Maximize Compute and Instruction Throughput CUDA Techniques to Maximize Compute and Instruction Throughput(pdf) HOW TO UNDERSTAND AND OPTIMIZE SHARED MEMORY ACCESSES USING NSIGHT COMPUTE(S41723)(pdf)A.1. Roofline 模型资料 Roofline Model与深度学习模型的性能分析 手把手建立Roofline模型（CPU） Roofline模型 CUDA 算子迭代优化方案（NCU 驱动） Profiling Guide：官方文档 GTC Silicon Valley-2019: Performance Analysis of GPU-Accelerated Applications using the Roofline Model GTC Silicon Valley-2019: Performance Analysis of GPU-Accelerated Applications using the Roofline Model(pdf)" },

  { "title": "CUDA 架构", "url": "/blog/2025/CUDA%E6%9E%B6%E6%9E%84/", "categories": "CUDA",
  "tags": "CUDA", "date": "2025-02-23", "content":
  "1. 硬件结构英伟达CUDA/GPU架构演变，以及不同架构的硬件能力：硬件层次结构如下（以Fermi架构为例）： 一个GPU中包含若干个SM（Streaming Multiprocessor，流多处理器），对应上图中左边； 一个SM中包含32个CUDA Core（也叫SP），对应上图中右边； 一个CUDA Core中包含一个ALU，一个FPU。 有些SM中还包含Tensor Core，与CUDA Core协同参与计算。1.1. 缓存层级一个CUDA Core内部，包含： Register File：16K 32-bit寄存器文件； L0 I-Cache一个SM内部，包含： Shared Memory：一个SM内部的threads可访问（一个Thread Block内的所有线程可访问）； L1 Cache：一个SM内部的threads可访问； Constant Cache； Register File：编译时分配给每个thread使用的寄存器文件。 注意：高级的 NVIDIA GPU 中，包含若干个 Subcore，比如 Ampere 架构的 GPU 中包含4个 Subcore，每个 Subcore 包含 32 个 CUDA Core 和 1 个 Tensor Core。如下图所示：https://developer.nvidia.com/blog/nvidia-ampere-architecture-in-depth/整个GPU内部，包含： L2 Cache：所有SM共享访问； Global Memory：所有SM共享访问。内存访问速度示意图： 注1：Shared Memory与L1 Cache共享片上内存，通过cudaFuncSetAttribute(kernel_name, cudaFuncAttributePreferredSharedMemoryCarveout, carveout);提示驱动分配多少给Shared Memory，但是尽量不要使用这个函数。资料： CUDA: GPU内存架构示意。文章结尾有一些内存优化相关链接。 huggingface – The Ultra-Scale Playbook:Training LLMs on GPU Clusters – A primer on GPUs 英伟达GPU架构总结2. 编程模型软件/硬件层次结构对应关系： 层级 说明 硬件对应 Grid 所有要执行的Block的集合 整个GPU Block 一组线程，一起在同一个SM上执行 一个SM（流多处理器） Thread 执行kernel代码的最小单位 一个CUDA Core 代码执行过程：myKernel\u0026lt;\u0026lt;\u0026lt;100, 256\u0026gt;\u0026gt;\u0026gt;(data, n) // 启动kernel↓ 产生 Grid，包含100个BlockBlock 0: ├── Thread 0 执行myKernel ├── Thread 1 执行myKernel ├── ... └── Thread 255 执行myKernelBlock 1: ├── Thread 0 执行myKernel ├── ... └── Thread 255 执行myKernel...（总共100个Block）在调度时，一个Block中的Theads只会分配到一个SM中执行，如果资源不允许，则需要执行多次调度循环，才能执行完这个Block中的所有Threads。另一方面，多个Block可以分配到同一个SM中执行。即，如果资源允许，或者当前Block中的Theads由于访问延迟而阻塞时，SM可以调度其他Block中的Threads来执行。（比如一个Warp的寄存器写后读会产生24个时钟延迟，则需要分配24个Warp来掩盖延迟）同一SM上可运行的Block数量 = min( ⌊最大thread数 / 每个Block的thread数⌋, ⌊共享内存大小 / 每个Block占用的shared memory⌋, ⌊寄存器总数 / (每个thread占用寄存器数 × 每个Block的thread数)⌋, 硬件限制（通常8-16个Block）)查看SM资源利用率：nvprof --metrics achieved_occupancy ./program# occupancy = (实际运行的warp数) / (理论最大warp数) Fermi同时执行最多16个kernel。2.1. kernel 索引 变量 类型 说明 threadIdx uint3 当前thread在其Block内的索引（0-based） blockIdx uint3 当前Block在Grid中的索引（0-based） blockDim dim3 当前Block的维度/大小（thread数量） gridDim dim3 当前Grid的维度/大小（Block数量） \u0026lt;\u0026lt;\u0026lt;grid, block\u0026gt;\u0026gt;\u0026gt;中，grid表示Grid的维度/大小（Block数量），block表示Block的维度/大小（Thread数量）。一维示例及执行分析：__global__ void kernel1D(float *data) { // blockDim.x = 256（启动时指定） // blockIdx.x = 0, 1, 2, ... （当前Block在Grid中的位置） // threadIdx.x = 0, 1, 2, ..., 255（当前thread在Block中的位置） // 计算全局线程索引 int globalIdx = blockIdx.x * blockDim.x + threadIdx.x; data[globalIdx] = data[globalIdx] * 2;}int main() { kernel1D\u0026lt;\u0026lt;\u0026lt;100, 256\u0026gt;\u0026gt;\u0026gt;(data); // 100个Block，每个Block 256个thread}kernel\u0026lt;\u0026lt;\u0026lt;100, 256\u0026gt;\u0026gt;\u0026gt;(data)Grid: Block 0: blockIdx.x=0 ├── Thread 0: threadIdx.x=0, globalIdx=0 ├── Thread 1: threadIdx.x=1, globalIdx=1 ├── ... └── Thread 255: threadIdx.x=255, globalIdx=255 Block 1: blockIdx.x=1 ├── Thread 0: threadIdx.x=0, globalIdx=256 ├── Thread 1: threadIdx.x=1, globalIdx=257 ├── ... └── Thread 255: threadIdx.x=255, globalIdx=511 ... Block 99: blockIdx.x=99 ├── Thread 0: threadIdx.x=0, globalIdx=25344 └── ...二维示例：__global__ void kernel2D(float *matrix, int width) { // blockDim.x = 16, blockDim.y = 16 // blockIdx.x = 0,1,2,... blockIdx.y = 0,1,2,... // threadIdx.x = 0,...,15 threadIdx.y = 0,...,15 // 计算全局行列索引 int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; int idx = y * width + x; matrix[idx] = matrix[idx] * 2;}int main() { dim3 blockDim(16, 16); // 16x16=256个thread/block dim3 gridDim(10, 10); // 10x10=100个block kernel2D\u0026lt;\u0026lt;\u0026lt;gridDim, blockDim\u0026gt;\u0026gt;\u0026gt;(matrix, width);}3. CUDA 中一些重要概念3.1. 线程束分叉 Warp Divergence当存在条件分支时，不同的线程执行不同的代码分支，需要串行执行两个分支。此时，当一些线程直线分支1时，其他线程的执行被浪费（执行结果不写入），导致一些clock浪费。3.2. Wavefront一次 Wavefront 是指一次打包的 load / store 操作。可能由于资源原因，或者地址不对齐等原因，导致 warp 的一次 request 可能需要分为多个 wavefront 来完成。比如理想情况下，warp 从 SMEM 一次请求 128 字节数据只需要一个 wavefront；如果发生 bank conflict，则需要多个 wavefront 来完成。官方文档描述见Nsight Compute – Metrics Guide。并参考https://forums.developer.nvidia.com/t/reuse-of-l1-shared-memory-during-execution-of-consecutive-wavefronts/288610。3.3. N-way Bank Conflict示例：访问一个列主序矩阵(float)，warp 内每个线程使用指令 LDS.128 或 STS.128 一次访问一个 float4。则一次访问（32 * 16B = 512B）需要分成 4 个内存事物，分别是：T0 ~ T7、T8 ~ T15、T16 ~ T23、T24 ~ T31。由于是列主序，T0 ~ T7 访问同一个bank 0，这个叫 8-way bank conflict，需要拆分成 8 个wavefront，产生 7 个 bank conflicts。warp 的一个 request，总共产生的 bank conflict 数量是：bank_conflicts = 7 * 4 = 283.3. 计算强度 Arithmetic Intensity在性能瓶颈分析过程中，需要确定瓶颈是计算能力还是内存带宽。计算强度定义为每次内存访问所执行的计算量，通常以 FLOPS/Byte 表示。定义公式：计算强度 = 算术运算次数 / 访问的内存字节数以 N _ N 矩阵乘法为例，执行 N^3 次乘加运算，以及 N^2 _ (N-1)次加法，访问 3 * N^2 个元素（A、B、C 矩阵），每个元素假设为 4 字节（float 类型），则计算强度为：计算强度 = N^3 / (3 * N^2 * 4) = N / 12 (FLOPS/Byte)Nsight Compute 相关资料 详细介绍：【CUDA调优指南】缓存\u0026amp;访存流程 深入理解 roofline 模型4. 参考资料 NVidia GPU指令集架构-寄存器 CUDA Refresher: The CUDA Programming Model 深入解析 NVIDIA Hopper 架構 CUDA C++ Best Practices Guide CUDA编程：基础与实践 pdf CUDA blogs4.1. 手册等 CUDA GPU Compute Capability：不同架构的计算能力对照表" },

  { "title": "HPC 零散笔记集合", "url": "/blog/2025/HPC%E7%AC%94%E8%AE%B0-%E9%9B%B6%E6%95%A3%E6%A6%82%E5%BF%B5/", "categories": "HPC",
  "tags": "HPC, CUDA, AI", "date": "2025-02-23", "content":
  "1. 概念：带宽(Bandwidth) vs 延迟(Latency)定义如下： 带宽(Bandwidth)：单位时间内能传输的数据量，通常以 GB/s 或 TB/s 表示。 延迟(Latency)：发出一次内存访问请求到数据返回所需的时间，通常以 ns（纳秒）或 时钟周期 表示。在不同的应用场景下，强调的性能指标不同。比如在AI训练中，需要持续搬运大量数据，比如大量密集的GEMM计算，即此时是计算密集型。另外，CPU/GPU 的算力极强。这种情况下，带宽经常成为性能瓶颈。而针对推理场景，由于其访问数据是随机访问的 KV-Cache，细粒度、低复用，带宽利用率低。这种情形下，等待数据的时间，即延迟成为性能瓶颈。推理与训练的对比如下： 对比项 训练 推理（Decode 阶段） Batch size 大（数百~数千） 小（1~几十） 数据访问模式 连续大块读写 逐 token 随机小块访问 计算强度 高（Compute-bound） 低（Memory-bound） KV Cache 访问 无 每步都要读取历史 KV 带宽 vs 延迟区别总结： 指标 带宽 延迟 衡量什么 数据吞吐量（多宽） 访问响应时间（多快） 单位 GB/s, TB/s ns, 时钟周期 决定因素 总线宽度、内存并行度 物理距离、内存层级 训练瓶颈 ✅ 主要瓶颈 次要（大 batch 可掩盖） 推理瓶颈 次要 ✅ 主要瓶颈（小 batch / decode） 代表硬件 HBM（大容量高带宽） 片上 SRAM（小容量低延迟） 2. nvidia-smi 工具使用查询以及设置功耗、频率：# 查询 GPU 功耗nvidia-smi -q -d POWER# 查询 GPU 支持的频率nvidia-smi -q -d SUPPORTED_CLOCKS# 设置 GPU 以最高速度运行nvidia-smi -q -d PERFORMANCE# 一秒间隔查询 GPU 使用情况nvidia-smi dmon# 一秒间隔查询 GPU 进程nvidia-smi pmon参考：nvidia-smi：控制您的GPU" },

  { "title": "CUDA 笔记集合", "url": "/blog/2025/CUDA%E7%AC%94%E8%AE%B0%E9%9B%86%E5%90%88/", "categories": "HPC",
  "tags": "CUDA, CuTe", "date": "2025-02-23", "content":
  "1. cutlass/CuTe GEMM 中矩阵的存储方式 NT / TN / NN / TT1.1. 背景BLAS 的约定是：所有矩阵一律按 column-major 存储，然后用 transA/transB 标志告诉 BLAS 要不要对它做转置：\\[C = \\alpha \\cdot op(A) \\cdot op(B) + \\beta \\cdot C\\]其中： 当transX为N时：$op(X) = X$，当transX为T时：$op(X) = X^T$。 乘法要求 $op(A)$ 是 $M \\times K$，$op(B)$ 是 $K \\times N$。1.2. GEMM 命名含义CuTe给矩阵做了一个约定：A(M, K)，B(N, K)，C(M, N)，即： A 矩阵：(M，K) – M 行 K 列 B 矩阵：(N，K) – N 行 K 列（不同于 BLAS 及其他典型约定） C 矩阵：(M，N) – M 行 N 列 即 CuTe 对 B 的约定，默认即为转置形式，即$B^T$，正好与 BLAS 约定形成转置关系。由于 CuTe 中对 A/B/C 的约束，导致在调用 BLAS 的时候，通过设置主序来表达转置关系： 针对 A，如果是 N，则使用 CuTe 表示的时候，A 是 (M, K)，column-major；如果是 T，则使用 CuTe 表示的时候，A 是 (M, K)，row-major。 针对 B，如果是 N，则使用 CuTe 表示的时候，B 是 (N, K)，row-major；如果是 T，则使用 CuTe 表示的时候，B 是 (N, K)，column-major。 即 B 在 CuTe 中的表示，与 A 在 CuTe 中的表示，转换规律正好相反。BLAS的约定与cutlass/CuTe的GEMM实现中矩阵A/B的主序关系如下： BLAS A Majorness A Layout B Majorness B Layout 解释 NT M-jajor (M,K):(1,ldA) N-major (N,K):(1,ldB) A(M,K) 列主序；B(N,K) 列主序 TN K-jajor (M,K):(ldA,1) K-major (N,K):(ldB,1) A(M,K) 行主序；B(N,K) 行主序 NN M-jajor (M,K):(1,ldA) K-major (N,K):(ldB,1) A(M,K) 列主序；B(N,K) 行主序 TT K-jajor (M,K):(ldA,1) N-major (N,K):(1,ldB) A(M,K) 行主序；B(N,K) 列主序 官方文档中相关描述见：https://github.com/NVIDIA/cutlass/blob/main/media/docs/cpp/cute/0x_gemm_tutorial.md#the-full-tensors-shapes-strides-and-data。gemm_nttemplate \u0026lt;class TA, class TB, class TC, class Alpha, class Beta\u0026gt;static cudaError_t gemm_nt(int m, int n, int k, Alpha alpha, TA const* A, int ldA, TB const* B, int ldB, Beta beta, TC* C, int ldC, cudaStream_t stream)nt含义是： 此时，BLAS 约定 $A$ 是 (M, K)，column-major。此时 CuTe 约定格式与 BLAS 约定一致。 此时，BLAS 的约定 $B^T$ 是 (K, N)。正好与 CuTe 约定的 B 矩阵形状一致，且是 column-major。最终得到：strider_A = cute::make_stride(cute::_1, ldA); // column-majorshape_A = cute::make_shape(M, K);stride_B = cute::make_stride(cute::_1, ldB); // column-majorshape_B = cute::make_shape(N, K);stride_C = cute::make_stride(cute::_1, ldC); // column-majorshape_C = cute::make_shape(M, N);gemm_tntemplate \u0026lt;class TA, class TB, class TC, class Alpha, class Beta\u0026gt;static cudaError_t gemm_tn(int m, int n, int k, Alpha alpha, TA const* A, int ldA, TB const* B, int ldB, Beta beta, TC* C, int ldC, cudaStream_t stream) $A^T$是(M, K)。按照 CuTe 约定，CuTe 要表达 $A^T$，只能是 (M, K) + row-major。 由于 CuTe 对 B 的约定导致其存储格式与 BLAS 约定形成转置关系，B 的的存储格式是 (N, K) + row-major。得到：stride_A = cute::make_stride(ldA, cute::_1); // row-majorshape_A = cute::make_shape(M, K);stride_B = cute::make_stride(ldB, cute::_1); // row-majorshape_B = cute::make_shape(N, K);stride_C = cute::make_stride(cute::_1, ldC); // column-majorshape_C = cute::make_shape(M, N);gemm_nntemplate \u0026lt;class TA, class TB, class TC, class Alpha, class Beta\u0026gt;static cudaError_t gemm_nn(int m, int n, int k, Alpha alpha, TA const* A, int ldA, TB const* B, int ldB, Beta beta, TC* C, int ldC, cudaStream_t stream) $A$ 是 (M, K)，column-major。CuTe 约定格式与 BLAS 约定一致。 $B$ 是 (N, K)，row-major。由于 CuTe 对 B 的约定导致其存储格式与 BLAS 约定形成转置关系，B 的的存储格式是 (N, K) + row-major。得到：stride_A = cute::make_stride(cute::_1, ldA); // column-majorshape_A = cute::make_shape(M, K);stride_B = cute::make_stride(ldB, cute::_1); // row-majorshape_B = cute::make_shape(N, K);stride_C = cute::make_stride(cute::_1, ldC); // column-majorshape_C = cute::make_shape(M, N);gemm_tttemplate \u0026lt;class TA, class TB, class TC, class Alpha, class Beta\u0026gt;static cudaError_t gemm_tt(int m, int n, int k, Alpha alpha, TA const* A, int ldA, TB const* B, int ldB, Beta beta, TC* C, int ldC, cudaStream_t stream) $A^T$ 是 (M, K)，row-major。CuTe 约定格式与 BLAS 约定形成转置关系，因此 A 的存储格式是 (M, K) + row-major。 $B^T$ 是 (K, N)，column-major。CuTe 约定格式与 BLAS 约定一致，即 CuTe 默认表达 BLAS 的转置形式。得到：stride_A = cute::make_stride(ldA, cute::_1); // row-majorshape_A = cute::make_shape(M, K);stride_B = cute::make_stride(cute::_1, ldB); // column-majorshape_B = cute::make_shape(N, K);stride_C = cute::make_stride(cute::_1, ldC); // column-majorshape_C = cute::make_shape(M, N);1.3. 内存访问效率分析–访存合并在划分 A/B 的的过程中，一般按照 M 方向划分 tile（针对 A），或者按照 N 方向划分 tile（针对 B）。比如如下 Thread-Value Layout 划分：// Define thread layouts.auto const thread_shape_A{cute::make_shape(cute::Int\u0026lt;16\u0026gt;{}, cute::Int\u0026lt;8\u0026gt;{})}; // (THR_M, THR_K)auto const thread_shape_B{cute::make_shape(cute::Int\u0026lt;16\u0026gt;{}, cute::Int\u0026lt;8\u0026gt;{})}; // (THR_N, THR_K)auto const thread_shape_C{cute::make_shape(cute::Int\u0026lt;32\u0026gt;{}, cute::Int\u0026lt;4\u0026gt;{})}; // (THR_M, THR_N)auto const thread_stride_A{cute::make_stride(cute::Int\u0026lt;1\u0026gt;{}, cute::size\u0026lt;0\u0026gt;(thread_shape_A))}; // column-majorauto const thread_stride_B{cute::make_stride(cute::Int\u0026lt;1\u0026gt;{}, cute::size\u0026lt;0\u0026gt;(thread_shape_B))}; // column-majorauto const thread_stride_C{cute::make_stride(cute::Int\u0026lt;1\u0026gt;{}, cute::size\u0026lt;0\u0026gt;(thread_shape_C))}; // column-majorauto const thread_layout_A{cute::make_layout(thread_shape_A, thread_stride_A)}; // (THR_M, THR_K)auto const thread_layout_B{cute::make_layout(thread_shape_B, thread_stride_B)}; // (THR_N, THR_K)auto const thread_layout_C{cute::make_layout(thread_shape_C, thread_stride_C)}; // (THR_M, THR_N)此时，得到：$\\text{thr_id} = m \\times 1 + k \\times 16$。即，同一个 warp 内的 32 个连续线程，它们的 thread ID 沿第一个维度（M 或 N）连续变化。所以：加载 A tile 时：warp 内线程沿 M 维度连续 → 如果 A 在 M 维度内存连续（column-major, stride=1），就是合并访存 ✅加载 B tile 时：warp 内线程沿 N 维度连续 → 如果 B 在 N 维度内存连续（column-major, stride=1），就是合并访存 ✅对这几种形式的 GEMM，可以得到其访存能否合并： 变体 A 的 M 维度 B 的 N 维度 加载 A 加载 B gemm_nt stride=1 (连续) stride=1 (连续) uint128_t合并 ✅ uint128_t合并 ✅ gemm_nn stride=1 (连续) stride=ldB (不连续) uint128_t合并 ✅ 逐元素拷贝 ❌ gemm_tn stride=ldA (不连续) stride=ldB (不连续) 逐元素拷贝 ❌ 逐元素拷贝 ❌ gemm_tt stride=ldA (不连续) stride=1 (连续) 逐元素拷贝 ❌ uint128_t合并 ✅ 参考代码来源：https://github.com/leimao/CUTLASS-Examples/blob/main/examples/cute_general_matrix_multiplication/cute_general_matrix_multiplication_tensor_core_gmem_tiled_copy_smem_tiled_copy_tiled_mma_sm80_pipeline.cu。更多资料： Row-Major VS Column-Major: https://leimao.github.io/blog/Row-Major-VS-Column-Major/2. nvcc 打印的 ptxas 信息含义以如下一条打印信息为例（以CTA中的线程数 128 为例）：ptxas info : Used 168 registers, used 1 barriers, 49152 bytes smem, 419 bytes cmem[0] 字段 含义 REG per thread 每个线程占用 168 个寄存器。SM 中共 65536 个寄存器。 寄存器是否够用，以及可以驻留的CTA数量计算公式：65536 / 128 / 168 SMEM per CTA 每个 CTA 使用了 49152 字节的共享内存。SMEM 容量 164KB。SMEM 可以驻留的CTA数量：164KB / 49152B ≈ 3 CTA 3. CUDA 循环展开 #pragma unroll循环展开指令，分为部分展开，以及完全展开。循环展开，可以节省一部分循环语句的开销，部分循环展开可能会产生一些额外的开始/结束的开销。循环展开，就是将循环体（block），展开变为更大的执行体（block）。循环展开之后，编译器可以更好的重组原本在不同iteration中的指令，将加载指令提前，从而提高指令的吞吐量（throughput）。不好的方面是，提前加载导致寄存器占用时间拉长，导致寄存器压力，可能导致SM占用率（occupancy）。另外，循环体中如果有局部变量（寄存器），展开之后，会占用更多的寄存器，同样导致occupancy下降。完全展开，会将索引替换为常量嵌入到指令中，将一些下标访问的数组（这些数据只能放在local memory中），优化为使用寄存器。另一方面，完全展开之后，可能会消除一些条件判断语句。比如如下代码将原本只能存放在local memory中的数组优化为寄存器：// 展开前：编译器不知道 i 的值，arr 只能放内存（local memory）for (int i = 0; i \u0026lt; 4; i++) { arr[i] = arr[i] * 2;}// 完全展开后：索引全是常量，编译器可以把 arr[0]~arr[3] 各分配一个寄存器arr[0] = arr[0] * 2; // → reg0 = reg0 * 2arr[1] = arr[1] * 2; // → reg1 = reg1 * 2arr[2] = arr[2] * 2; // → reg2 = reg2 * 2arr[3] = arr[3] * 2; // → reg3 = reg3 * 2如下代码，优化掉分支逻辑：// 展开前for (int i = 0; i \u0026lt; 4; i++) { if (i == 0) { special_init(); } process(i);}// 完全展开后，编译器可以做常量传播// i=0: if (0 == 0) → true → 保留 special_init()special_init();process(0);// i=1: if (1 == 0) → false → 整个 if 被消除（死代码消除）process(1);// i=2, i=3 同理process(2);process(3);循环展开，导致指令膨胀，指令膨胀过大导致ICache存不下的时候，会导致i-cache miss增加，降低指令吞吐量。 参考自NVIDIAl论坛：Understanding unrolling and concurrent memory operations" },

  { "title": "C++ 右值引用，万能引用，完美转发", "url": "/blog/2025/cpp-rvalue/", "categories": "Cpp",
  "tags": "Cpp", "date": "2025-02-18", "content":
  "主要概念： 引用的本质在C++内部实现是一个常指针。 左值引用，右值引用。右值引用限制了其只能接收右值，可以利用这个特性从而提供重载。 template 万能引用，引用折叠。 完美转发：std::forward 。完美转发 std::forward模板的万能引用只是提供了能够接收同时接收左值引用和右值引用的能力，但是引用类型的唯一作用就是限制了接收的类型，后续使用中都退化成了左值，我们希望能够在传递过程中保持它的左值或者右值的属性, 如果不使用forward，直接按照下面的方式写就会导致问题。void RFn(int\u0026amp;\u0026amp; arg){}template\u0026lt;typename T\u0026gt;void ProxyFn(T\u0026amp;\u0026amp; arg){ RFn(arg);}void main(){ ProxyFn(1);}会发现右值版本不能传过去, [int]无法到[int\u0026amp;\u0026amp;]，就导致参数不匹配。 C++: 左值引用(\u0026amp;), 右值引用(\u0026amp;\u0026amp;),万能引用(template \u0026amp;\u0026amp;)详解 与 完美转发(forward) 实现剖析 说一下C++左值引用和右值引用" },

  { "title": "IO多路复用与linux epoll", "url": "/blog/2025/IO%E5%A4%9A%E8%B7%AF%E5%A4%8D%E7%94%A8%E4%B8%8Eepoll/", "categories": "Systems",
  "tags": "Linux, Networking", "date": "2025-02-07", "content":
  "定义：I/O多路复用是指在单线程中同时监视多个文件描述符的状态变化（如可读、可写、异常等），当其中一个或多个文件描述符发生状态变化时，内核会通知应用程序进行相应的处理。1. epoll 与 select/poll 区别epoll的实现代码fs/eventpoll.c，其分为三个接口函数： epoll_create: 创建一个epoll实例，返回一个文件描述符。 epoll_ctl: 向epoll实例中添加、修改或删除需要监视的文件描述符。 epoll_wait: 等待epoll实例中监视的文件描述符发生状态变化，并返回就绪的文件描述符列表。具体流程分为两个阶段：注册阶段（epoll_ctl EPOLL_CTL_ADD）： 为目标 fd 创建 epitem，插入 ep-\u0026gt;rbr 红黑树。 调用 ep_item_poll(epi, \u0026amp;epq.pt, 1)，内部通过 ep_ptable_queue_proc 分配一个 eppoll_entry，将 ep_poll_callback 注册为回调函数，并通过 add_wait_queue 将其挂入目标 fd（如 socket）自身的等待队列，完成事件监听挂钩。等待/触发阶段（epoll_wait）： epoll_wait 调用 ep_poll，先检查 ep-\u0026gt;rdllist 是否有就绪事件：若有则直接调用 ep_try_send_events 传递事件并返回；若无则将当前进程加入 ep-\u0026gt;wq（epoll 实例自身的等待队列），挂起进程等待唤醒。 当目标 fd（如 socket）收到数据后，内核驱动层（网络栈/设备驱动）调用 wake_up() 唤醒该 fd 自身的等待队列，从而触发挂在其上的回调 ep_poll_callback。 ep_poll_callback 将对应 epitem 加入 ep-\u0026gt;rdllist（若此时正在向用户空间传输事件，则暂存于 ep-\u0026gt;ovflist 溢出链表），然后调用 wake_up(\u0026amp;ep-\u0026gt;wq) 唤醒 epoll_wait 中挂起的进程。 进程被唤醒，ep_poll 调用 ep_send_events，后者先通过 ep_start_scan 调用 list_splice_init(\u0026amp;ep-\u0026gt;rdllist, txlist) 将 rdllist 中的 epitem 原子地转移到临时链表 txlist，同时将 ep-\u0026gt;ovflist 从 EP_UNACTIVE_PTR 置为 NULL，以接收此传输期间新到的事件。 ep_send_events 遍历 txlist 中每个 epitem，调用 ep_item_poll（内部调用 vfs_poll）获取 fd 上最新的 events（防止使用过时状态），再将 events 与 epoll_event.data 通过 epoll_put_uevent 拷贝到用户空间。LT 模式下若事件仍未处理完，epitem 会被重新加回 ep-\u0026gt;rdllist，保证下次 epoll_wait 仍能返回；ET 模式则不重新加入。 ep_done_scan 将 ovflist 中传输期间到达的新事件合并回 ep-\u0026gt;rdllist，并在 rdllist 非空时再次调用 wake_up(\u0026amp;ep-\u0026gt;wq) 通知等待者。epoll_wait流程如下：图中涉及两个不同的等待队列：目标 fd（如 socket）自身的等待队列（ep_poll_callback 注册于此，收到数据时被触发）；以及 ep-\u0026gt;wq（epoll 实例的等待队列，epoll_wait 调用者挂于此处休眠）。两者概念不同，需加以区分。1.1. epoll 与 poll/select 特点总结select/poll给定一个fd数组，每次都是从用户态传递到内核态，内核态遍历该数组，填充有事件发生的fd。然后内核态将该数组传回用户态，用户态再次遍历该数组，处理有事件发生的fd。因此，select/poll的时间复杂度为O(n)。epoll只在有事件发生时才通知用户程序，只将就绪的fd返回给用户程序。epoll有如下特点： 红黑树: 内核中epoll_ctl将需要监听的文件描述符(针对网络通信就是套接字)时，保存在红黑树中。添加/删除/索引的时间复杂度为O(log n)。 rdllist: ep-\u0026gt;rdllist是内核存储的就绪事件列表，当有事件发生时比如套接字数据可读，驱动将fd对应的epitem加入到rdllist中（通过fd等待队列上的回调函数ep_poll_callback）。epoll_wait通过ep_start_scan将rdllist转移到临时txlist后扫描处理，更新对应的epoll_event数据，并返回给用户程序。2. ET模式与LT模式区别LT模式：当fd就绪时，epoll_wait 会一直返回该fd，直到事件被处理。例如，如果一个socket连接有数据可读，epoll_wait 会每次都返回该socket fd，直到数据被完全读取。ET模式: 当fd从未就绪变为就绪时，epoll_wait 只会返回一次该fd。例如，如果一个socket有数据可读（kernel中缓冲区由空变为非空），epoll_wait 只会在数据第一次到达时返回该套接字，之后即使有更多数据到达，也不会再次返回。ET的这个特性使得其要求应用程序必须一次性读取所有数据，否则可能会错过后续的数据到达事件，不能实时处理数据。因此，使用ET模式时，必须将套接字设置为非阻塞模式，并在事件处理函数中循环读取数据，直到返回EWOULDBLOCK错误。以下是ET模式下的事件处理示例：while (1) { ssize_t n = read(fd, buffer, sizeof(buffer)); if (n == -1) { if (errno == EAGAIN || errno == EWOULDBLOCK) { break; // 数据读完 } // 处理其他异常情形 } // 处理读取的数据} 明显的，ET模式相较于LT模式，减少了系统调用次数，提高了性能。3. 设置socket为非阻塞 Blocking read Non-blocking read ET模式下，使用阻塞模式socket，如果数据量较大，需要多次read，最后一次可能没有数据可读，此时read将一直阻塞。使用非阻塞模式socket，read返回EWOULDBLOCK即代表数据读完。4. epoll 调用流程图graph TD A[启动服务器] --\u0026gt; B[创建 epoll 实例] B --\u0026gt; C[创建监听套接字] C --\u0026gt; D[设置监听套接字为非阻塞] D --\u0026gt; E[绑定监听套接字到指定端口] E --\u0026gt; F[监听连接请求] F --\u0026gt; G[将监听套接字添加到 epoll 实例中] G --\u0026gt; H[进入事件循环] H --\u0026gt; I[调用 epoll_wait 等待事件发生] I --\u0026gt; J{有事件发生吗?} J --\u0026gt;|是| K[处理就绪事件] J --\u0026gt;|否| I K --\u0026gt; L{事件类型} L --\u0026gt;|新连接| M[接受新连接] M --\u0026gt; N[设置新连接套接字为非阻塞] N --\u0026gt; O[将新连接套接字添加到 epoll 实例中] L --\u0026gt;|可读事件| P[读取数据] P --\u0026gt; Q{读取成功吗?} Q --\u0026gt;|是| R[处理读取的数据] Q --\u0026gt;|否| S[关闭套接字] R --\u0026gt; T[将套接字修改为可写事件] L --\u0026gt;|可写事件| U[写入数据] U --\u0026gt; V[将套接字修改为可读事件] S --\u0026gt; H O --\u0026gt; H T --\u0026gt; H V --\u0026gt; H5. libEventlibEvent是一个跨平台的事件通知库，提供了统一的接口来处理不同操作系统上的事件驱动机制（如select、poll、epoll等）。它封装了底层的事件处理细节，使得开发者可以更方便地编写高性能的网络应用程序。libEvent可以处理的事件类型包括（还有很多高级功能没有列出来，比如处理复杂的事件组合、优先级事件等）： I/O事件：如套接字可读、可写等。 定时器事件：在指定时间后触发的事件。 信号事件：当特定信号发生时触发的事件，使用evsigsel。libEvent基本使用流程为，首先创建一个事件基础设施（event_base），然后为每个需要监视的事件创建一个事件对象（event），并将其添加到基础设施中。最后，调用event_base_dispatch进入事件循环，等待事件发生并处理。资源： PcapNG Playback：基于libEvent实现的时间戳回放 libevent：博客文章 23.libevent：博客文章参考 Linux下的I/O复用与epoll详解 gitee – epoll sample Linux I/O 模型详解 Blocking I/O, Nonblocking I/O, And Epoll Linux I/O multiplexing examples Network Programming: How is epoll implemented?：见Nelson Elhage的回答 Linux 内核的 epoll 实现" },

  { "title": "Jekyll Chirpy 主题配置优化记录", "url": "/blog/2025/Jekyll-Chirpy%E4%B8%BB%E9%A2%98%E9%85%8D%E7%BD%AE%E4%BC%98%E5%8C%96%E8%AE%B0%E5%BD%95/", "categories": "Computer",
  "tags": "jekyll, Website", "date": "2025-02-03", "content":
  "概述本文记录了将 Jekyll 博客从 chirpy-starter 迁移到完整 jekyll-theme-chirpy 主题后遇到的问题及解决方案，以及后续的一系列优化改进。主要涵盖： GitHub Pages 部署错误修复 数学公式渲染引擎切换（MathJax → KaTeX） 目录（TOC）展开配置优化一、部署错误修复问题背景从 chirpy-starter 切换到 jekyll-theme-chirpy 完整主题后，GitHub Actions 部署失败，错误信息为：Error: Can't find stylesheet to import. ╷3 │ @import 'colors/typography-light';根本原因chirpy-starter 使用预编译的静态资源，而 jekyll-theme-chirpy 完整主题需要在构建时编译 SCSS 和 JavaScript 文件。GitHub Actions 工作流缺少 Node.js 环境和前端资源构建步骤。解决方案1. 添加 Node.js 环境和构建步骤修改 .github/workflows/pages-deploy.yml，在 Jekyll 构建之前添加：- name: Setup Node.js uses: actions/setup-node@v4 with: node-version: \"20\"- name: Install dependencies run: npm install- name: Build frontend assets run: npm run build提交记录：847b43c - fix: add Node.js build step for jekyll-theme-chirpy2. 解决 package-lock.json 缺失问题初次构建时遇到 npm ci 要求 package-lock.json 的错误。由于 .gitignore 忽略了该文件，改用 npm install：- name: Install dependencies run: npm install # 改为 npm install提交记录：251cdcb - fix: use npm install instead of npm ci3. 解决配置文件冲突git pull --rebase 时 _config.yml 发生合并冲突，保留用户自定义配置（如 avatar 路径），删除冲突标记。提交记录：c56b96e - fix: resolve _config.yml merge conflict二、数学公式渲染优化从 MathJax 迁移到 KaTeX迁移动机 性能提升：KaTeX 渲染速度比 MathJax 快 5-10 倍 一致性：VS Code Markdown 预览使用 KaTeX，切换后本地预览与网站渲染一致 转义字符处理：KaTeX 对 \\text{row\\_offset} 等转义字符的处理更符合预期实施步骤1. 添加 KaTeX CDN 资源修改 _data/origin/basic.yml 和 _data/origin/cors.yml：# basic.ymlkatex: css: https://cdn.jsdelivr.net/npm/katex@0.16.9/dist/katex.min.css js: https://cdn.jsdelivr.net/npm/katex@0.16.9/dist/katex.min.js auto-render: https://cdn.jsdelivr.net/npm/katex@0.16.9/dist/contrib/auto-render.min.js# cors.ymlkatex: https://cdn.jsdelivr.net2. 在页面头部加载 KaTeX CSS修改 _includes/head.html，在 \u0026lt;/head\u0026gt; 之前添加：{% if page.math %}\u0026lt;link rel=\"stylesheet\" href=\"{{ site.data.origin[origin].katex.css | relative_url }}\" integrity=\"sha384-n8MVd4RsNIU0tAv4ct0nTaAbDJwPJzDEaqSD1odI+WdtXRGWt2kTvGFasHpSy3SV\" crossorigin=\"anonymous\"/\u0026gt;{% endif %}3. 在页面底部加载 KaTeX JavaScript修改 _includes/js-selector.html，替换 MathJax 部分：{% if page.math %}\u0026lt;!-- KaTeX --\u0026gt;\u0026lt;script defer src=\"{{ site.data.origin[origin].katex.js | relative_url }}\" integrity=\"sha384-VQ8d8WVFw0yHhCk5E8I86oOhv48xLpnDZx5T9GogA/Y84DcCKWXDmSDfn13bzFZY\" crossorigin=\"anonymous\"\u0026gt;\u0026lt;/script\u0026gt;\u0026lt;script defer src=\"{{ site.data.origin[origin].katex['auto-render'] | relative_url }}\" integrity=\"sha384-+XBljXPPiv+OzfbB3cVmLHf4hdUFHlWNZN5spNQ7rmHTXpd7WvJum6fIACpNNfIR\" crossorigin=\"anonymous\"\u0026gt;\u0026lt;/script\u0026gt;\u0026lt;script defer src=\"{{ '/assets/js/data/katex.js' | relative_url }}\"\u0026gt;\u0026lt;/script\u0026gt;{% endif %}4. 创建 KaTeX 配置文件新建 assets/js/data/katex.js 处理 kramdown 生成的 MathJax 格式标签：document.addEventListener(\"DOMContentLoaded\", function () { // Convert kramdown's MathJax format to KaTeX format document.querySelectorAll(\"script[type='math/tex']\").forEach(function (el) { const texText = el.textContent; const span = document.createElement(\"span\"); katex.render(texText, span, { throwOnError: false, displayMode: false, }); el.parentNode.replaceChild(span, el); }); document.querySelectorAll(\"script[type='math/tex; mode=display']\").forEach(function (el) { const texText = el.textContent; const div = document.createElement(\"div\"); katex.render(texText, div, { throwOnError: false, displayMode: true, }); el.parentNode.replaceChild(div, el); }); // Auto-render for regular KaTeX delimiters renderMathInElement(document.body, { delimiters: [ { left: \"$$\", right: \"$$\", display: true }, { left: \"$\", right: \"$\", display: false }, ], throwOnError: false, });});5. 启用数学公式支持修改 _config.yml 的 defaults 配置，为所有文章启用数学公式：defaults: - scope: path: \"\" type: posts values: math: true提交记录： 90b1eb3 - render math formular using MathJax c69851b - update: add Katex and use Katex for math render三、目录（TOC）展开优化问题描述默认情况下，TOC（Table of Contents）只在滚动到某个二级标题时才展开其下的三级、四级标题，影响导航体验。优化目标让所有标题级别（h2、h3、h4）默认完全展开，方便用户快速定位内容。实施方案1. 修改桌面端 TOC 配置编辑 _javascript/modules/components/toc/toc-desktop.js：static options = { tocSelector: '#toc', contentSelector: '.content', ignoreSelector: '[data-toc-skip]', headingSelector: 'h2, h3, h4', orderedList: false, scrollSmooth: false, collapseDepth: 6, // 新增：展开所有标题级别 headingsOffset: 32};2. 更新编译后的 JavaScript由于本地环境未安装 Node.js，直接修改编译后的文件 assets/js/dist/post.min.js：// 找到 tocbot 配置部分，添加 collapseDepth:6P(Mn, \"options\", { tocSelector: \"#toc\", contentSelector: \".content\", ignoreSelector: \"[data-toc-skip]\", headingSelector: \"h2, h3, h4\", orderedList: !1, scrollSmooth: !1, collapseDepth: 6, // 添加此行 headingsOffset: 32,});提交记录：812ebd7 - enable unroll all sub-directories in dir list配置说明 collapseDepth: 6：展开前 6 级标题 由于博客使用 h2-h4（3 级标题），设置为 6 确保所有标题完全展开 用户刷新页面（Ctrl+F5）即可看到效果四、经验总结1. 主题选择建议 chirpy-starter：适合快速搭建，使用预编译资源，无需本地构建 jekyll-theme-chirpy：适合深度定制，需要配置 Node.js 构建环境2. CI/CD 最佳实践 完整主题需在 GitHub Actions 中添加 Node.js 环境 优先使用 npm install 而非 npm ci（如果不追踪 package-lock.json） 构建顺序：Node.js 环境 → npm install → npm build → Jekyll build3. 数学公式渲染选择 特性 MathJax KaTeX 渲染速度 较慢 快（5-10倍） 功能完整度 非常完整 覆盖常用功能 包大小 较大（~200KB） 较小（~100KB） VS Code 预览 ✗ ✓ 转义字符处理 需额外配置 原生支持 推荐：对于技术博客，KaTeX 是更好的选择。六、支持 KaTeX 服务器端渲染（PR #2603）背景参考上游仓库 PR #2603，该 PR 添加了对 KaTeX 服务器端渲染的支持，通过 jektex 插件在构建时渲染数学公式，而不是在浏览器端通过 JavaScript 渲染。实施方案1. 添加数学引擎配置修改 _config.yml，添加数学引擎选择配置：# Math equation rendering engine.math: # Choose engine for rendering math equations. # mathjax — client-side rendering (loads JavaScript library) # katex — server-side rendering via jektex plugin (faster, no JS required) engine: # [mathjax | katex]2. 添加 jektex 配置在 _config.yml 的 kramdown 配置后添加 jektex 配置：# Jektex configuration for server-side KaTeX renderingjektex: cache_dir: \".jektex-cache\" # Cache directory for rendered equations ignore: [\"**/*\"] # Ignore all by default (enable when math.engine is katex) silent: false # Show rendering progress macros: [] # Global LaTeX macros (e.g., [[\"\\\\\\\\Q\", \"\\\\\\\\mathbb{Q}\"]])3. 更新 .gitignore添加 jektex 缓存目录到 .gitignore：# Misc.jektex-cache4. 条件加载 CSS修改 _includes/head.html，根据引擎类型加载 CSS：{% if page.math %} {% assign math_engine = site.math.engine | default: 'mathjax' %} {% if math_engine == 'katex' %}\u0026lt;!-- KaTeX CSS for server-side rendering --\u0026gt;\u0026lt;link rel=\"stylesheet\" href=\"{{ site.data.origin[type].katex.css | relative_url }}\" /\u0026gt;{% endif %} {% endif %}5. 条件加载 JavaScript修改 _includes/js-selector.html，只在使用 MathJax 时加载 JS：{% if page.math %} {% assign math_engine = site.math.engine | default: 'mathjax' %} {% if math_engine == 'mathjax' %}\u0026lt;!-- MathJax --\u0026gt;\u0026lt;script src=\"{{ '/assets/js/data/mathjax.js' | relative_url }}\"\u0026gt;\u0026lt;/script\u0026gt;\u0026lt;script async src=\"https://cdnjs.cloudflare.com/polyfill/v3/polyfill.min.js?features=es6\"\u0026gt;\u0026lt;/script\u0026gt;\u0026lt;script id=\"MathJax-script\" async src=\"{{ site.data.origin[type].mathjax.js | relative_url }}\"\u0026gt;\u0026lt;/script\u0026gt;{% endif %} {%- comment -%} KaTeX is rendered server-side via jektex plugin, only CSS needed {%- endcomment -%} {% endif %}6. 添加 KaTeX 样式修改 _sass/base/_base.scss，添加 KaTeX 溢出处理：/* KaTeX */.katex-display { overflow: auto hidden;}7. 添加 jektex 依赖修改 jekyll-theme-chirpy.gemspec，添加 jektex 依赖：spec.add_runtime_dependency \"jektex\", \"~\u0026gt; 0.1.1\"使用方法选择 MathJax（默认）不需要配置，或显式设置：math: engine: mathjax选择 KaTeX在 _config.yml 中设置：math: engine: katexjektex: ignore: [\"*.xml\"] # 处理 markdown 文件，忽略 feed.xml然后运行：bundle add jektexbundle install特性对比 特性 MathJax KaTeX (jektex) 渲染方式 客户端 JavaScript 服务器端构建时 页面加载速度 较慢 快 LaTeX 支持 更完整 常用功能完整 \\label 和 \\eqref ✓ ✗ 需要 JavaScript ✓ ✗ 缓存支持 ✗ ✓ 注意事项 兼容性：KaTeX 不支持某些 LaTeX 特性，如 \\label 和 \\eqref 构建时间：首次构建时，KaTeX 会渲染所有公式并缓存，后续构建会快很多 默认行为：如果不配置 math.engine，默认使用 MathJax（保持向后兼容）提交记录：参考 PR #26034. 配置文件管理 关键配置文件建议纳入版本控制 合并冲突时优先保留自定义配置 及时提交配置变更，便于回溯五、参考资源 Jekyll Chirpy Theme KaTeX Documentation tocbot Configuration GitHub Actions - setup-node附录：完整文件清单本次优化涉及的主要文件：.github/workflows/pages-deploy.yml # CI/CD 配置_config.yml # Jekyll 全局配置_data/origin/basic.yml # CDN 资源配置_data/origin/cors.yml # CORS 配置_includes/head.html # 页面头部（KaTeX CSS）_includes/js-selector.html # 脚本加载器（KaTeX JS）assets/js/data/katex.js # KaTeX 配置脚本（新建）_javascript/modules/components/toc/toc-desktop.js # TOC 桌面端配置assets/js/dist/post.min.js # 编译后的 JS（TOC 配置）通过这一系列优化，博客成功完成了主题迁移，并在性能和用户体验上都有明显提升。后续可以继续探索更多主题定制选项。链接 Chirpy Jekyll Theme – Tutorial Chirpy Jekyll Theme Markdown 渲染器配置选项其他主题 al-folio：另一个流行的 Jekyll 主题，适合学术博客，支持 LaTeX 数学公式和丰富的功能。 Academic Pages：适合学术博客，支持 LaTeX 数学公式和丰富的功能。" },

  { "title": "Chirpy主题的安装与使用指南", "url": "/blog/2025/Chirpy%E6%A8%A1%E6%9D%BF%E5%AE%89%E8%A3%85%E6%8C%87%E5%8D%97/", "categories": "Computer",
  "tags": "Website", "date": "2025-02-02", "content":
  "Chirpy主题的前置与安装指南设置Github Pages和Jekyll Github pages的官方文档: GitHub Pages 文档 使用 Jekyll 设置 GitHub Pages 站点 配置 GitHub Pages 站点的自定义域部署Chirpy主题 Chirpy主题的官方文档: Jekyll Theme Chirpy Wiki 官方指南 Chirpy 及其汉化版 pansong291 官方部署指南 以及 示例 【避坑篇】使用Github Pages搭建个人主页or博客网站【上】 【快速部署+客制化】Github Pages+Jekyll Chirpy 速搭个人主页 Jekyll博客搭建教程（上篇）（下篇） 搭建个人博客：Jekyll + Github Pages + VSCode 使用Jekyll + Github Pages搭建静态网站 如何从chirpy-starter升级为完整的jekyll-theme-chirpy: Upgrade GuideChirpy主题的使用与进阶 撰写文章 自定义图标 Chirpy主题的个性化自定义 Chirpy主题的安装、编写、图标、头像 Chirpy主题的进阶使用 对Chirpy进行简单美化 对于Chirpy各种奇怪bug的研究Chirpy主题的开发 Chirpy官方语法说明书 Development \u0026amp; Test EnvironmentsChirpy主题中文章中可使用的命令文章的头信息(Front Matter)文章作为.md文件存在/_posts中, 每个.md文件需要有一段头信息(frontmatter), 可参照撰写新帖子, 其中更改作者author可见Author Information. 关于更深入的内容可参照Front Matter 和 Posts.事实上, 在如下Front Matter中仅有title与date是必选项.---layout: # 在Chirpy模板中默认为postpermalink: # 最终导向的URLtitle: TITLEdescription: DESCRIPTION # 可选author: # 默认是social.namedate: YYYY-MM-DD HH:MM:SS +0800 # 时间和时区可选categories: [TOP_CATEGORIE, SUB_CATEGORIE]tags: [TAG]excerpt_separator: # 摘录段落published: # 是否公开toc: true # 目录comment: true # 评论区pin: true # 置顶math: true # 数学模式, 由MathJax支持mermaid: true # Markdown的一种扩展, 用于画图cdn: url # 只能作为img的根目录,所以对我作用不大---类型提示markdown中引文以\u0026gt; 在行首出现, Chirpy提供了一些提示(但我不一定记得用):\u0026gt; 显示 `tip` 类型提示的例子。\u0026gt; {: .prompt-tip }\u0026gt; 显示 `info` 类型提示的例子。\u0026gt; {: .prompt-info }\u0026gt; 显示 `warning` 类型提示的例子。\u0026gt; {: .prompt-warning }\u0026gt; 显示 `danger` 类型提示的例子。\u0026gt; {: .prompt-danger } 显示 tip 类型提示的例子。 显示 info 类型提示的例子。 显示 warning 类型提示的例子。 显示 danger 类型提示的例子。内嵌pdf和视频图片我用\u0026lt;iframe\u0026gt;标签在网站中内嵌pdf文件,一般是如下形式：\u0026lt;iframe src=\"https://cdn.jsdelivr.net/gh/user/repo@version/file/\" frameborder=\"0\" width=\"100%\" height=\"600px\"\u0026gt; \u0026lt;/iframe\u0026gt;其中用了js Delivr的CDN，也可以放其他链接。对于内嵌视频,详见Write a New Post. 对于平台视频可使用如下语法:{% include embed/{Platform}.html id='{ID}' %}其中{Platform}是平台名称(支持Youtube, Twitch和Bilibili), {ID}是视频ID.对于一般的视频和音频可以直接用:{% include embed/video.html src='{URL}' %}{% include embed/audio.html src='{URL}' %}内嵌图片可使用![My helpful screenshot](/assets/img/postimg/screenshot.jpg)脚注在正文插入脚注使用命令正文[^footnote],正文且可以写反向脚注(一般在正文最后)[^footnote]: 这是一个脚注添加评论功能 在 Chirpy 上整合 Giscus Hugo 博客引入 Giscus 评论系统定制字体，默认展开所有目录修改assets\\css\\jekyll-theme-chirpy.scss文件,:/* ============================================================ *//* 自定义字体样式：改善中英文混排效果 *//* ============================================================ *//* 优化主文本字体：中文字体和英文字体混排 */body,html { /* 字体栈优先级： 1. 系统中文字体（如有） 2. 西方字体用于英文显示 3. 降级字体 */ font-family: -apple-system, BlinkMacSystemFont, \"Segoe UI\", \"Roboto\", \"Oxygen\", \"Ubuntu\", \"Cantarell\", \"Fira Sans\", \"Droid Sans\", \"Helvetica Neue\", \"Source Han Sans SC\", \"Noto Sans CJK SC\", \"Microsoft YaHei\", \"WenQuanYi Zen Hei\", sans-serif; font-size: 16px; line-height: 1.65;}/* 段落文本优化 */p { word-spacing: 0.05em; text-align: justify;}/* 列表文本优化 */li { word-spacing: 0.05em;}/* 标题字体优化 */h1,h2,h3,h4,h5,h6 { font-family: -apple-system, BlinkMacSystemFont, \"Segoe UI\", \"Roboto\", \"Oxygen\", \"Ubuntu\", \"Cantarell\", \"Fira Sans\", \"Droid Sans\", \"Helvetica Neue\", \"Source Han Sans SC\", \"Noto Sans CJK SC\", \"Microsoft YaHei\", sans-serif; line-height: 1.4; letter-spacing: 0.01em; font-weight: 600;}/* 代码和代码块字体保持等宽 */code,pre,.highlight { font-family: \"Cascadia Code\", \"Fira Code\", \"Source Code Pro\", \"Courier New\", \"Courier\", monospace; font-size: 0.95em; line-height: 1.5;}/* 内联代码 */:not(pre) \u0026gt; code { font-size: 0.9em; padding: 0.2em 0.4em;}/* 表格文本优化 */table td,table th { word-spacing: 0.02em;}/* 块引用文本 */blockquote { font-style: italic; border-left: 4px solid var(--primary-color, #6c63ff);}/* ============================================================ *//* 自定义样式：强制展开 TOC 所有子目录 *//* 覆盖 tocbot 的折叠样式，使所有子目录始终展开 *//* ============================================================ */.toc-list .is-collapsed { max-height: none !important;}.toc-list .is-collapsible { max-height: none !important;}" },

  { "title": "QEMU 学习资料", "url": "/blog/2025/qemu-study/", "categories": "Linux",
  "tags": "Linux, QEMU", "date": "2025-01-18", "content":
  "qemu\u0026amp;kvm学习笔记 qemu笔记 Qemu KVM(Kernel Virtual Machine)学习笔记" },

  { "title": "Memory Segmentation Cheet Sheets", "url": "/blog/2025/memory-segmentation-cheet-sheets/", "categories": "Linux",
  "tags": "Linux", "date": "2025-01-17", "content":
  "reference Memory Segmentation Cheet Sheets" },

  { "title": "std::pmr -- 内存池", "url": "/blog/2025/std-pmr-memory-pool/", "categories": "Cpp",
  "tags": "Cpp", "date": "2025-01-15", "content":
  "1. 介绍使用C++ 17的多态内存管理器(PMR)，可以实现一个简单的内存池。根据选择(std::pmr::memory_resource)，可以在内存不够的时候，向upstream申请内存。标准内存资源列表： memory_resource派生类 效率 线程安全 内存 std::pmr::synchronized_pool_resource() 效率低（内部需要上锁） 线程安全 更少碎片化 std::pmr::unsynchronized_pool_resource() 效率较高(内部不需要上锁) 非线程安全 更少碎片化 std::pmr::monotonic_buffer_resource() 效率最高 非线程安全 “只进不出”（从不释放、可传递进可选的缓冲区） 两个返回指向单例全局内存资源指针的函数： 函数名 特点 std::pmr::new_delete_resource() 默认的内存资源（转发给传统 new/delete） std::pmr::null_memory_resource() “永远拒绝” 2. 示例 基本示例：std::pmr::monotonic_buffer_resource的基本用法：src/hello_prm。 使用std::pmr::polymorphic_allocator初始化对象（alloc + construct）：src/polymorphic_allocator。 benchmark：不同memory_resource实现的性能对比：src/benchmark。用户自定义类使用PMR分配之后的析构：using user_class_alloc_traits = std::allocator_traits\u0026lt;std::pmr::polymorphic_allocator\u0026lt;UserClass\u0026gt;\u0026gt;;user_class_alloc_traits::destroy(userclass_allocator, userclass);user_class_alloc_traits::deallocate(userclass_allocator, userclass, 1);3. 参考 cppreference – pmr benchmark test source code C++17 the complete guide – Chap 29. 多态内存资源(PMR)" },

  { "title": "C++面向对象三个概念——重载、覆盖和隐藏", "url": "/blog/2025/cpp-override-overload-hide/", "categories": "Cpp",
  "tags": "Cpp", "date": "2025-01-03", "content":
  "1. overload 重载 同名函数，参数个数或类型不同； 相同作用域，即同一个类。2. override 覆盖 不在一个作用域，即父类与子类； 子类函数与基类函数同名，参数个数和类型相同； 基类使用virtual关键字，子类使用override关键字。例外的一个点是协变：基类返回基类指针，子类返回子类指针。此时也是override。#include \u0026lt;iostream\u0026gt;class Base {public: virtual Base* clone() const { std::cout \u0026lt;\u0026lt; \"Base::clone()\" \u0026lt;\u0026lt; std::endl; return new Base(*this); }};class Derived : public Base {public: // 使用协变返回类型，返回类型是 Base 的派生类型 Derived* Derived* clone() const override { std::cout \u0026lt;\u0026lt; \"Derived::clone()\" \u0026lt;\u0026lt; std::endl; return new Derived(*this); }};int main() { Derived d; Base* ptr = \u0026amp;d; Base* newPtr = ptr-\u0026gt;clone(); // 调用 Derived::clone() delete newPtr; return 0;}3. hide 隐藏作用域不同，同名函数不能生成override，而是覆盖。子类中的成员函数、数据成员，将覆盖基类中的同名成员函数、数据成员。如果要调用基类中的同名函数，需要使用基类名::函数名，或者使用using引入。#include \u0026lt;iostream\u0026gt;using namespace std;class Base {public: void func() { cout \u0026lt;\u0026lt; \"Base::func()\" \u0026lt;\u0026lt; endl; } void func(int) { cout \u0026lt;\u0026lt; \"Base::func(int)\" \u0026lt;\u0026lt; endl; }};class Derived : public Base {public: void func() { cout \u0026lt;\u0026lt; \"Derived::func()\" \u0026lt;\u0026lt; endl; } // 隐藏基类所有 func 版本 void callBaseFunc() { Base::func(); } // 显式调用基类的 func()};int main() { Derived d; d.func(); // 调用 Derived::func() // d.func(1); // 编译错误，基类 func(int) 被隐藏 d.Base::func(1); // 显式调用基类的 func(int) return 0;}4. 引用 c++三大概念要分清–重载，隐藏（重定义），覆盖（重写）" },

  { "title": "编译FFMPEG使能 NVIDIA 硬件解码", "url": "/blog/2024/%E7%BC%96%E8%AF%91FFMPEG%E4%BD%BF%E8%83%BDNvidia%E7%A1%AC%E4%BB%B6%E8%A7%A3%E7%A0%81/", "categories": "Linux",
  "tags": "Linux, FFmpeg, NVIDIA", "date": "2024-12-31", "content":
  "1. 依赖于 NVIDIA NVIDIA 显卡驱动(.run文件安装) CUDA Toolkit(.run文件安装)安装 NVIDIA 驱动(run文件)# 禁用 nouveau 开源驱动cat \u0026gt; /etc/modprobe.d/blacklist-nouveau.conf \u0026lt;\u0026lt;EOFblacklist nouveauoptions nouveau modeset=0EOF# Update initramfsupdate-initramfs -u# 需要重启生效reboot# Install compilation tools and Xorg dependenciesapt install -y make gcc linux-headers-$(uname -r) pkg-config xserver-xorg xorg-dev# Install Vulkan and GLVND development librariesapt install -y libvulkan1 libglvnd-devsystemctl stop gdm.service./NVIDIA-Linux-x86_64-xxx.run安装CUDA：# https://developer.nvidia.com/cuda-toolkit-archivewget https://developer.download.nvidia.com/compute/cuda/12.6.3/local_installers/cuda-repo-debian12-12-6-local_12.6.3-560.35.05-1_amd64.debsudo dpkg -i cuda-repo-debian12-12-6-local_12.6.3-560.35.05-1_amd64.debsudo cp /var/cuda-repo-debian12-12-6-local/cuda-*-keyring.gpg /usr/share/keyrings/sudo add-apt-repository contribsudo apt-get updatesudo apt-get -y install cuda-toolkit-12-6查看硬件编解码实时状态:watch -n 1 nvidia-smi -i 0 -q -d UTILIZATION资料: Debian 12安装Nvidia显卡驱动 download – nvidia 525.183.01 driver for linux CUDA Toolkit 12.6 Update 3 Downloads2. FFMPEG 编译依赖库安装sudo apt-get -y install autoconf automake build-essential libass-dev libfreetype6-dev libsdl2-dev libtheora-dev libtool libva-dev libvdpau-dev libvorbis-dev libxcb1-dev libxcb-shm0-dev libxcb-xfixes0-dev pkg-config texinfo zlib1g-devsudo apt-get -y install yasm nasmsudo apt-get -y install libx264-dev libx265-dev libvpx-dev libdav1d-dev # videosudo apt-get install libfdk-aac-dev libmp3lame-dev libopus-dev libopus-dev # audio3. 编译 FFMPEGgit clone https://github.com/FFmpeg/nv-codec-headers.gitcd nv-codec-headersmakesudo make install PREFIX=/usr # 如果安装到其他目录，需要将路径天骄到PKG_CONFIG_PATH./configure --extra-cflags=\"-I/usr/local/cuda/include\" --extra-ldflags=\"-L/usr/local/cuda/lib64\" --extra-libs=-lpthread --extra-libs=-lm --enable-shared --enable-gpl --enable-libfreetype --enable-libmp3lame --enable-libopus --enable-libvorbis --enable-libx264 --enable-libx265 --enable-nonfree --enable-cuda --enable-cuvid --enable-nvenc --enable-ffnvcodec --enable-cuvid # --enable-libfdk_aac --enable-libnppmake -j$(nproc)sudo make install4. 测试# https://github.com/omen23/ffmpeg-ffnvcodec-explanationffmpeg -hide_banner -encoders | grep nvencffmpeg -hide_banner -decoders | grep cuvidffmpeg -hide_banner -hwaccels# https://d2axc7bbtmotmv.cloudfront.net/chrome_test/Temple.mp4ffmpeg -c:v h264_cuvid -i input output.mkv # 使用 nvidia cuvid 解码# 拉流，使用cuda硬件解码ffmpeg -hwaccel cuda -i https://d2axc7bbtmotmv.cloudfront.net/chrome_test/Temple.mp4 out.mp4 out.mp4# 拉流，使用nvdec硬件解码ffmpeg -c:v h264_cuvid -i https://d2axc7bbtmotmv.cloudfront.net/chrome_test/Temple.mp4 out.mp4 out.mp4# 使用nvenc硬件编码ffmpeg -i input.mp4 -c:v h264_nvenc output.mp45. 参考 FFMPEG – HWAccel Intro – CUDA (NVENC/NVDEC) FFMPEG – Compile FFmpeg for Ubuntu, Debian, or Mint ffmpeg NVIDIA编解码一：ffmpeg编译安装 ffmpeg NVIDIA编解码三：英伟达硬编码 LINUX下，ffmpeg增加NVIDIA硬件编解码的步骤及解决办法6. 附加VAAPI相关的一些可能的设置选项:sudo apt install xserver-xorg-video-nvidia# export LD_PRELOAD=/usr/lib/x86_64-linux-gnu/libEGL.soexport VDPAU_DRIVER=nvidiaexport LIBVA_DRIVER_NAME=nvidia # vdpauexport LIBVA_DRIVERS_PATH=/usr/lib/x86_64-linux-gnu/dri # nvidia_drv_video.soexport LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/usr/local/libexport DRI_PRIME=1export VK_DRIVER_FILES=/usr/lib/x86_64-linux-gnu/libvulkan.soexport VK_ICD_FILENAMES=/usr/share/vulkan/icd.d/nvidia_icd.jsonVK_ADD_DRIVER_FILES=/run/opengl-driver/share/vulkan/icd.d/nvidia_icd.x86_64.jsonLIBVA_DRI3_DISABLE=1Nvidia相关的一些可能的环境变量:export __NV_PRIME_RENDER_OFFLOAD=1export __GLX_VENDOR_LIBRARY_NAME=nvidiaexport __NV_PRIME_RENDER_OFFLOAD=1export __VK_LAYER_NV_optimus=NVIDIA_onlyexport __GLX_VENDOR_LIBRARY_NAME=nvidiaexport __EGL_VENDOR_LIBRARY_NAME=nvidia一些工具软件:sudo apt install libvulkan1 libvulkan-dev vulkan-tools# https://gitlab.freedesktop.org/vdpau/libvdpau# https://github.com/KhronosGroup/Vulkan-Toolsvdpauinfovainfovulkaninfo --summary7. 附加2 Chromium 一些资料 Chromium sources7.1 安装Chromiumapt install software-properties-common apt-transport-https ca-certificates curl -ycurl -fSsL https://dl.google.com/linux/linux_signing_key.pub | sudo gpg --dearmor | sudo tee /usr/share/keyrings/google-chrome.gpg \u0026gt;\u0026gt; /dev/nullecho deb [arch=amd64 signed-by=/usr/share/keyrings/google-chrome.gpg] http://dl.google.com/linux/chrome/deb/ stable main | sudo tee /etc/apt/sources.list.d/google-chrome.listapt install google-chrome-stable7.2 一些启动命令尝试# export GOOGLE_API_KEY=\"no\"# export GOOGLE_DEFAULT_CLIENT_ID=\"no\"# export GOOGLE_DEFAULT_CLIENT_SECRET=\"no\"# export VK_ICD_FILENAMES=/usr/share/vulkan/icd.d/nvidia_icd.json# FFmpegVideoDecoder InitializeMediaLibrary GpuVideoDecoder# headless./chrome --ignore-gpu-blocklist --disable-gpu-driver-bug-workarounds --enable-features=AcceleratedVideoDecodeLinuxZeroCopyGL,AcceleratedVideoDecodeLinuxGL,VaapiIgnoreDriverChecks,VaapiOnNvidiaGPUs,AcceleratedVideoEncoder,UseOzonePlatform,UseMultiPlaneFormatForHardwareVideo,PlatformHEVCDecoderSupport --ozone-platform-hint=auto --enable-gpu --enable-logging=stderr --vmodule=*/media/* --v=0 --use-gl=angle --use-angle=vulkan --headless=new --remote-debugging-address=0.0.0.0 --remote-debugging-port=9221 --allow-chrome-scheme-url# headful./chrome --ignore-gpu-blocklist --disable-gpu-driver-bug-workarounds --enable-features=AcceleratedVideoDecodeLinuxZeroCopyGL,AcceleratedVideoDecodeLinuxGL,VaapiIgnoreDriverChecks,VaapiOnNvidiaGPUs,AcceleratedVideoEncoder,UseOzonePlatform,UseMultiPlaneFormatForHardwareVideo,PlatformHEVCDecoderSupport --ozone-platform-hint=auto --enable-gpu --enable-logging=stderr --vmodule=*/media/* --v=0 --use-gl=angle --use-angle=vulkan# --disable-features=UseSkiaRenderer,UseChromeOSDirectVideoDecoder# --disable-gpu-compositing --in-process-gpu Chromium docs – VA-API Disable VA-API on NVIDIA GPUs for ChromeOS and Linux Hardware-accelerated video decode on chromium w/ NVIDIA+VDPAU Hardware accelarated video decoding in chromium How To Enable Hardware Acceleration on Chrome, Chromium \u0026amp; Puppeteer on AWS in Headless mode blog – vdpau7.3 定制化 Chromium/FFmpeg# media/ffmpeg/scripts/build_ffmpeg.py# media/ffmpeg/scripts/robo_lib/config.py /etc/lsb-release -\u0026gt; /etc/os-release# third_party/ffmpeg/configure# media/ffmpeg/scripts/generate_gn.py# third_party/ffmpeg/chromium/config/Chrome/linux/x64/config_components.h# third_party/ffmpeg/ffmpeg_generated.gniexport PATH=`pwd`/third_party/llvm-build/Release+Asserts/bin:$PATH./media/ffmpeg/scripts/build_ffmpeg.py linux x64 --branding Chrome./third_party/ffmpeg/chromium/scripts/copy_config.shpython3 ./media/ffmpeg/scripts/generate_gn.py# build_ffmpeg.pyconfigure_flags['Chrome'].extend([ '--enable-decoder=aac,h264', '--enable-demuxer=aac', '--enable-parser=aac,h264', '--enable-ffnvcodec', '--enable-cuvid', '--enable-nvdec', '--enable-nvenc', '--enable-vdpau', '--enable-hwaccel=h264_vdpau', '--enable-hwaccel=h264_nvdec', '--extra-cflags=-isystem/usr/include', '--extra-cxxflags=-isystem/usr/include' ])# gn args out/releaseis_debug = falseis_component_build = falsetarget_cpu = \"x64\"enable_nacl = falsesymbol_level = 0is_official_build = truechrome_pgo_phase = 0ffmpeg_branding = \"Chrome\"proprietary_codecs = true# enable_linux_installer = trueautoninja -j4 -C out/release -v chrome# gn args out/debuggoogle_api_key=\"AIzaSyDxKL42zsPjbke5O8_rPVpVrLrJ8aeE9rQ\"google_default_client_id=\"595013732528-llk8trb03f0ldpqq6nprjp1s79596646.apps.googleusercontent.com\"google_default_client_secret=\"5ntt6GbbkjnTVXx-MSxbmx5e\"is_debug = trueis_component_build = truetarget_cpu = \"x64\"enable_nacl = falsesymbol_level = 2is_official_build = falsechrome_pgo_phase = 0ffmpeg_branding = \"Chrome\"proprietary_codecs = truedisable_fieldtrial_testing_config=truemedia_use_ffmpeg=trueexclude_unwind_tables=falseenable_ffmpeg_video_decoders=trueenable_widevine=falsertc_use_h264=trueblink_symbol_level=0v8_symbol_level=0enable_stack_trace_line_numbers = trueenable_backup_ref_ptr_support=falseenable_dangling_raw_ptr_checks=falseenable_dangling_raw_ptr_feature_flag=false#use_partition_alloc_as_malloc=false#use_allocator_shim=false#use_partition_alloc=falseuse_qt=false# run following if# use_sysroot=false# sudo ./build/install-build-deps.sh# sudo apt install qtbase5-dev 重新配置chrome中ffmpeg插件 chromium – debugging Tips for debugging on Linux How to get a stack trace at runtime Some tricks about debugging Chromium Chromium 编译与调试笔记" },

  { "title": "gcc 向量化相关选项", "url": "/blog/2024/gcc-vectorize-options/", "categories": "Cpp",
  "tags": "Cpp, GCC", "date": "2024-11-19", "content":
  "1. alias选项strict aliasing是编译器优化中依赖的一个假设，即不同类型的指针，指向不同的内存区域。基于该假设，gcc编译器可以进行一些优化。gcc优化-O2默认开启该选项(-fstrict-aliasing)。使用该选项，需要保证不同类型指针的内存区域不重叠，否则会导致未定义行为。例如:int a = 10;int* p1 = \u0026amp;a;float* p2 = (float*)\u0026amp;a;如果不能保证，则使用-fno-strict-aliasing选项。该选项导致性能下降，例如每次可能会从内存中读取数据，而不是寄存器。要保证代码类型转换安全，使用编译选项:add_compile_options(-Wstrict-aliasing) # -Werror -Wall参考: Casting does not work as expected when optimization is turned on2. vectorize选项 -ftree-vectorize: 整个代码中，可能的向量化优化。 -ftree-loop-vectorize: 循环中的向量化优化。 -fopt-info-vec-missed: 显示没有向量化的循环。 -fopt-info-vec-optimized: 显示已向量化的循环。另外，在本地，想要充分优化，设置:set(CMAKE_C_FLAGS_RELEASE \"-O3 -march=native\")set(CMAKE_CXX_FLAGS_RELEASE \"-O3 -march=native\")编写代码过程中，影响向量化的因素有: exception: 异常处理会影响向量化。尽可能使用noexcept或const。3. 调试–看IRadd_compile_options(-fdump-tree-dse) #查看 dead store elimination 之后的 IR4. vectorize 更多资料入门资料: Intel – Vectorization codebook step by step – Crunching Numbers with AVX and AVX2更多资料: Vectorization part1. Intro. Auto-vectorization in GCC Automatic Vectorization" },

  { "title": "perf性能分析(7) -- Top-down 分析方法及VTune工具", "url": "/blog/2024/perf-tools-07-top-down-method/", "categories": "Performance",
  "tags": "Performance, VTune", "date": "2024-11-10", "content":
  "现代性能分析，使用针对pipeline的分析办法(取代CPU cycles分析)。这源于现代CPU架构的复杂性。现代CPU处理指令架构，分为前端 Front-end，后端 Back-end两部分。阻碍指令执行的因素，从硬件看，源于前端或后端的Stall。1. CPU 流水线Intel CPU流水线一般分为5级。其中解码(ID)，意思是将指令操作分解为多个uOp(即拆分为多个更低级的硬件操作)，如ADD eax, [mem1]，可以拆分成两个微指令：从内存读取数据，再执行ADD操作。2. CPU 架构及流水线的执行过程执行过程：前端执行完IF -\u0026gt; ID之后，然后在一个名为allocation的过程中（下图中星标处），uOps被输送到后端。后端监控uOp的操作数(data operand)何时可用，并在可用的执行单元中执行uOps。当uOp执行完成之后，称之为执行完成(retirement)，且将uOp的结果被写会寄存器或者内存（经过Store Buffer写入内存）。大多数uOps都会完全通过流水线并退出，但有些投机指令uOps可能会在退出前被取消–如预测错误的分支。在Intel处理器中，一个core一般有四个执行端口，即每个cycle最多可以执行四个uOps。在处理器架构中，有一个抽象概念：pipeline slot（流水线槽），用来表示用于执行一个uOp所需要的硬件资源。在每个时钟周期，有四个流水线槽可用，流水线槽可以是空的，也可以是被uOp填充。流水线槽在Allocation阶段（上图中的星号标记处），将uOp从前端分配到后端执行单元。PMU监控流水线槽的状态，在每个时钟周期，衡量流水线槽的利用率（是否填充有uOp），并对流水线槽进行分类，确定是前端瓶颈还是后端瓶颈。3. Top-down 分析方法从性能分析的角度看，一条微指令在流水线中的性能指标可以分为： 退出(Retiring) – Micro Sequencer(微指令调度器)可能会成为瓶颈，例如调度浮点指令。 分支预测错误(Bad Speculation) – 分支预测错误，或者memory ordering violation(多核多线程共享数据情形)，导致Machine Clears(清除流水线)。 前端瓶颈(Front-End Bottleneck) 后端瓶颈(Back-End Bottleneck)Top-down分析思想根据上述的流水执行阶段及过程，首先从Top-level分类分析步骤：然后，继续细分（Breakdown），确定是哪个阶段中的哪个资源导致的stall： 相比以前基于事件的度量的方式，基于Top-down分析方法，可以更准确地定位性能瓶颈的根源，指导开发者进行针对性的优化。3.1. Frontend Bound – 前端瓶颈前端主要职责为读取指令，解码之后，发送给后端。遇到分支指令，需要经过预测器预测下一个指令的地址，这意味着会出现由于分支预测错误并清除流水线导致的ICache Miss而引起前端阻塞。而在取指/解码过程中，如果代码的局部性不好，由于ICache Miss，或者iTLB Miss，引起前端Stall。过高的miss率导致前端瓶颈。 L1/L2 Cache分为DCache和ICache。TLB也分为iTLB、dTLB和Second-Level TLB，地址从iTLB/dTLB中找不到则再从Second-Level TLB中查找。3.2. Back-End Bound – 后端瓶颈后端主要职责为执行指令，包括ALU、FPU、Memory等。后端瓶颈分为： Core Bound 除法指令过多，因为除法指令执行时间长； Execution Port Utilzation过高，导致执行单元饱和； 指令的数据依赖，即依赖上一个指令的结果，导致Stall； Memory Bound L1/L2/L3 Cache Miss引起的Stall，需要改进数据访问的局部性，或者减小数据的访问规模（如分块处理数据），或者False-Sharing引起的Cache Miss； Memory Bound：分为Memory BandWidth和Memory Latency，前者是指内存带宽不足，后者是指内存访问延迟过高； 3.3. 优化指令的处理能力（Throughput）例如，改用向量指令（如AVX）。或者使用更高效的算法。4. 一些优化手段4.1. Frontend 减少代码的footprint，如-fomit-frame-pointer 调整代码布局：如是用-fprofile-generate -fprofile-use 调整代码布局：如用__attribute__((hot)) 分支预测优化：如loop unrolling，特别是小的循环，如小于64次循环 分支预测优化：是用if代替三目运算符；避免if-elses结构，switch-case排序4.2. Back-End 减少function call，如inline 多线程避免false-sharing，使用内存对齐 gcc优化：如__builtin_expect4.3. 示例#define likely(x) __builtin_expect(!!(x), 1) //gcc内置函数, 帮助编译器分支优化 #define unlikely(x) __builtin_expect(!!(x), 0) if(likely(condition)) { // 这里的代码执行的概率比较高 } if(unlikely(condition)) { // 这里的代码执行的概率比较高 }#define CACHE_LINE __attribute__((aligned(64)))struct S1 { int r1; int r2; int r3; S1(): r1(1),r2(2),r3(3){} } CACHE_LINE;5. Hazard 介绍5.1. StructuralHazards – 结构性冲突结构性冲突本质是CPU中硬件资源的竞争，比如流水线中，前后指令之间都需要经过译码器，访问内存，形成对译码器的争用。5.2. DataHazards – 数据依赖五级流水线：取指IF -\u0026gt; 解码ID -\u0026gt; 执行EX -\u0026gt; 内存访问MEM -\u0026gt; 写回WB。Data Hazard是指后一条指令的操作数，依赖于前一条指令的结果。操作数依赖分为三种关系： 先写后读(Write-after-Read) – Data Denpendency 先读后写(Read-after-Write) – Anti-Dependency 写后写(Write-after-Write) – Output DependencyCPU处理Data Hazard办法有两种： 插入NOP指令，流水线停顿(Pipeline Stall)，或者叫流水线冒泡(Pipeline Bubbling)。 Operand Forwarding – 操作数转发。Operand Forwarding：在第一条指令的执行阶段完成之后，直接将结果数据传输给到下一条指令的 ALU。然后，下一条指令不需要再插入两个 NOP 阶段，就可以继续正常走到执行阶段。这样的解决方案，我们就叫作操作数前推(Operand Forwarding)，或者操作数旁路(Operand Bypassing)。其实更合适的名字应该叫操作数转发。这里的 Forward，其实就是我们写 Email 时的“转发”（Forward）的意思。5.3. ControlHazards – 控制依赖主要使用分支预测。5.4. 流水线 – 乱序执行更详细资料：cnblogs – 计算机组成原理——原理篇 处理器（中）6. 使用VTune工具进行Top-down分析对一个应用程序进行profile，一般首先使用Hotspots，获取测量得到的分类性能指标。以官方matrix_mul示例为例（以Windows为例，VTune自带样例的路径为C:\\Users\\Administrator\\Documents\\VTune\\Samples）：然后，使用Microarchitecture Exploration，再次进行测试。测试完成之后，在Bottom-up页面查看Top-down Tree，获取更细粒度的性能指标： 下图为局部展开的Front-End Bound、Back-End Bound。 从图中看，Front-End Bound为44.0%，即流水线槽（Pipeline Slot）由于前端没能及时提供指令而产生44.0%的空闲。 Back-End Bound为100.0%，意味着后端由于内存带宽限制、指令的数据依赖（比如指令依赖上一条指令的计算结果）、L1/L2/L3缓存访问延迟严重，或者执行单元饱和，导致流水线停顿。 综合前后端的瓶颈分析，前端应该是由于后端的瓶颈导致的。 6.1. Back-End Bound在Top-down的顶层分类中，使用流水线槽（Pipeline-Slot）来测量瓶颈：Front-End Bound、Back-End Bound、Bad Speculation、Retiring。在深入到Back-End之后，使用PMU来测量后端的性能指标值，比如测量一个指令的时钟周期或者事件，进而度量其Stall来源并算比例值。其测量值，与使用流水线槽度量及计算得到的值会有差异。另外，Back-End Bound又分为两个子类：Memory Bound与Core Bound，分别度量来自内存/Cache的Stall，以及来自Execution Port的Stall。6.2. 微架构优化方法在进行性能优化的时候，一般关注应用程序的顶层热点函数，即占用CPU最多的函数。VTune工具展示的profile结果也是热点 + 层次定位，比如Bottom-up页面，可以按照占用CPU时间排序查看，即查看热点函数以及其耗时占比。按照最优优化顺序，首先优化程序算法，以及并行化，即优化大头。之后，再进行微架构性能profile，并按照Top-down方法，针对性的优化瓶颈。在Top-down视图中，VTune会高亮显示Top-down分类中的瓶颈类型（如上面Bottom-up页面所示）。下表给出了不同分类应程序，其Top-down分类中各指标的合理范围，超出该范围5%则在VTune的profile结果页面中高亮显示： 这个阈值表是基于对Intel实验室中的一些工作负载的分析总结出来的，作为一般指导原则。 在Top-down分析过程中，没有高亮的指标，不需要花费时间优化，优化这一部分一般不会有太多的改善。 大多数未调优的应用程序都是后端瓶颈的。 解决后端问题通常与解决延迟源有关，延迟源会导致执行完成所需的时间超过必要时间。 6.3. 第三方优化示例 https://github.com/yao-matrix/mblog/blob/main/mips.md https://weedge.github.io/perf-book-cn/zh/chapters/7-Overview-Of-Performance-Analysis-Tools/7-1_Intel_Vtune_cn.htmlA. 资料 自顶向下的微架构分析方法：官方文档，介绍了Top-down分析方法的原理和步骤。 pdf – A Top-Down Method for Performance Analysis and Counters Architecture cnblogs –C/C++ 性能优化背后的方法论：TMAM 调优指南: Xeon E5 v3 pdf – intel lectures: Intel_VTune_Amplifier-jackson：PPT文档 github – pmu-tools User Interface Reference：VTune官方参考文档，该章节的子章节详细介绍了VTune的每个页面。A.1. 更多阅读 Yasin’s Publications：Intel Yasin 教授的论文列表 视频：From Top-down Microarchitecture Analysis to Structured Performance Optimizations Software Optimizations Become Simple with Top-Down Analysis on Intel Skylake - Ahmad Yasin @ IDF’15 译：内存分析 pdf – perf books cn – 第二部分：源代码优化" },

  { "title": "C++ 中 auto 和 decltype 的用法 (update 20241106)", "url": "/blog/2024/cpp-auto-and-decltype/", "categories": "Cpp",
  "tags": "Cpp", "date": "2024-11-06", "content":
  "1. auto 占位符1.1 规则auto推导的原则为：保持原有变量的类型(如cv限定)，大致分两种情况: auto: auto含义是创建了一个新的变量: 表达式为T或者T\u0026amp;或者const T\u0026amp; – auto推导为T – 即新变量的类型去除cv限定 (如果原有表达式有cv限定); 表达式为T* const或者T* – auto推导为T* – 新变量去除cv限定; 表达式为const T*或者const T* const – auto推导为const T*，即保持指针指向的内存区域的const属性。 auto\u0026amp;: auto\u0026amp;含义是alias，故auto\u0026amp;推导的结果是原有类型的的引用，不能少任何一个限定符，如: const T* const推导为const T* const \u0026amp;； const T推导为const T\u0026amp;； 1.2 一些应用场景 范围循环std::vector\u0026lt;int\u0026gt; numbers = {1, 2, 3, 4, 5};for (const auto\u0026amp; num : numbers) { std::cout \u0026lt;\u0026lt; num \u0026lt;\u0026lt; std::endl; // num 的类型被推导为 const int\u0026amp;} lambda 表达式auto add_func = [](int a, int b) -\u0026gt; int { return a + b;}; 结构化绑定（c++17新特性）std::pair\u0026lt;int, double\u0026gt; myPair = std::make_pair(42, 3.14);auto [first, second] = myPair; // 结构化绑定，first 和 second 的类型被推导为 int 和 double 返回类型后置template\u0026lt;typename T, typename U\u0026gt;auto add(T t, U u) -\u0026gt; decltype(t + u) { return t + u;}2. decltypedecltype 类型推导规则：2.1 decltype(entity)如果 entity 是一个不被括号包围的标识符、类访问表达式，那么decltype(entity)与 entity 类型一致。2.2 decltype((expression))如果expression是一个表达式，计算结果为类型T，那么： 如果 expression 为 xvalue（将亡值），那么 decltype 的结果是 T\u0026amp;\u0026amp;。 如果 expression 为 lvalue，那么 decltype 的结果是 T\u0026amp;。 如果 expression 为 prvalue（纯右值），那么 decltype 的结果是 T。2.3 说明及示例注意第一点中强调了 entity 是一个不被括号包围的标识符。因为当一个标识符被括号包围时，它就是一个左值表达式了，对应上面第二大点的第二小点。比如说 int x = 0;，x 是一个标识符，所以 decltype(x) 的结果为 int。但是 (x) 就是一个左值表达式，decltype((x)) 的结果就是 int\u0026amp;。3. 参考 C++11特性：decltype关键字 c++11-17 模板核心知识（九）—— 理解decltype与decltype(auto) decltype specifier" },

  { "title": "perf性能分析(6) -- perf实战(1) -- 分支预测", "url": "/blog/2024/perf-tools-06-perf-practice-01/", "categories": "Performance",
  "tags": "Performance, VTune", "date": "2024-11-05", "content":
  "使用Debug模式编译 sort 和 unsort 代码:#include \u0026lt;algorithm\u0026gt;#include \u0026lt;ctime\u0026gt;#include \u0026lt;iostream\u0026gt;int main(int argc, char* argv[]) { constexpr int kArrLen = 1024 * 1024; int* data = new int[kArrLen]; for (int c = 0; c \u0026lt; kArrLen; c++) data[c] = std::rand() % 256; // std::sort(data, data + kArrLen); // 是否排序 long long sum = 0; for (int i = 0; i \u0026lt; 30000; i++) { for (int c = 0; c \u0026lt; kArrLen; c++) { if (data[c] \u0026gt;= 128) { // 故意選 256 一半 sum += data[c]; } } } delete[] data; std::cout \u0026lt;\u0026lt; \"sum = \" \u0026lt;\u0026lt; sum \u0026lt;\u0026lt; std::endl;}1. sort/unsort 性能测试及比较: perf statThe CPU is Intel(R) Core(TM) i5-8260U CPU @ 1.60GHz.perf stat ./sort 测试结果:$ perf stat ./sortsum = 3016043160000 Performance counter stats for './sort': 63,995.28 msec task-clock # 1.000 CPUs utilized 1,039 context-switches # 16.236 /sec 7 cpu-migrations # 0.109 /sec 1,151 page-faults # 17.986 /sec 240,342,649,070 cycles # 3.756 GHz 473,524,864,500 instructions # 1.97 insn per cycle 63,163,134,842 branches # 986.997 M/sec 7,402,185 branch-misses # 0.01% of all branches 64.008866998 seconds time elapsed 63.983827000 seconds user 0.010999000 seconds sysperf stat ./unsort 测试结果:$ perf stat ./unsortsum = 3016043160000 Performance counter stats for './unsort': 206,441.99 msec task-clock # 1.000 CPUs utilized 484 context-switches # 2.344 /sec 44 cpu-migrations # 0.213 /sec 1,150 page-faults # 5.571 /sec 771,936,412,315 cycles # 3.739 GHz 473,410,337,438 instructions # 0.61 insn per cycle 63,144,272,701 branches # 305.869 M/sec 15,732,825,641 branch-misses # 24.92% of all branches 206.439026771 seconds time elapsed 206.440760000 seconds user 0.002000000 seconds sys1.1. Result comparisonunsort的分支预测失败率为24.92%, 且CPU IPC为0.61, 小于1.0.sort分支预测失败率为0.01%, CPU IPC为1.97.2. analysis using perf record and perf ananotate2.1. program unsort运行 perf record 和 perf annotate 命令:perf record ./unsort # 结束后生产 perf.data 文件perf annotate # 自动定位到最耗时的地方if语句比较失败跳转至循环体, 循环体占用24.44%的CPU时间. 原因是分支预测失败 – if语句在for循环中失败次数较多.2.2. program sort分支预测失败率为.051% + 4.31%, 且CPU IPC为1.97. 内存load操作: 41.68%. add指令: 10.54%.3. When build with Release如果使用-O2或者-O3编译，编译器使用SIMD指令，且消除了if语句。其结果就是，sort与unsort的性能差距不大，甚至sort的性能更好: pcmpgtd – SIMD比较指令，得到位掩码结果，存在结果寄存器中； pand – SIMD按bit进行and操作；基本等效于：sum += data[c] \u0026amp; -(data[c] \u0026gt;= 128);references 在 Linux 上使用 Performance 做效能分析(入門篇) Category:Architecture/x86-assembly StackOverflow: Why is processing a sorted array faster than processing an unsorted array? StackOverflow: Why is processing an unsorted array the same speed as processing a sorted array with modern x86-64 clang? 知乎 – 分支对代码性能的影响和优化 Modern Microprocessors – A 90-Minute Guide!" },

  { "title": "perf性能分析(5) -- linux perf 工具介绍", "url": "/blog/2024/perf-tools-05-perf-introduce/", "categories": "Performance",
  "tags": "Performance, VTune, TBB", "date": "2024-10-30", "content":
  "1. perf 介绍perf及子命令可以测量/记录系统性能，可以记录的性能数据项繁多。包括CPU/PMU等硬件数据，以及software counter/tracepoint等系统内核采集的数据。可以关注的几类： CPU / PMU (Performance Monitoring Unit)数据。包括: dTLB, iTLB, cache 计数以及miss计数；branch及branch miss计数。 memory 延时、阻塞； bus延时、阻塞； front end/back end阻塞； virtual memory相关: TLB相关。 pipeline相关。查看perf命令及子命令帮助信息：man perfman perf-topman perf-statman perf-recordman perf-report查看perf所有子命令:$ perf usage: perf [--version] [--help] [OPTIONS] COMMAND [ARGS] The most commonly used perf commands are: annotate Read perf.data (created by perf record) and display annotated code archive Create archive with object files with build-ids found in perf.data file bench General framework for benchmark suites buildid-cache Manage build-id cache. buildid-list List the buildids in a perf.data file c2c Shared Data C2C/HITM Analyzer. config Get and set variables in a configuration file. daemon Run record sessions on background data Data file related processing diff Read perf.data files and display the differential profile evlist List the event names in a perf.data file ftrace simple wrapper for kernel's ftrace functionality inject Filter to augment the events stream with additional information iostat Show I/O performance metrics kallsyms Searches running kernel for symbols kvm Tool to trace/measure kvm guest os list List all symbolic event types mem Profile memory accesses record Run a command and record its profile into perf.data report Read perf.data (created by perf record) and display the profile script Read perf.data (created by perf record) and display trace output stat Run a command and gather performance counter statistics test Runs sanity tests. top System profiling tool. version display the version of perf binary probe Define new dynamic tracepoints See 'perf help COMMAND' for more information on a specific command.查看子命令的帮助信息，如perf stat：$ perf stat -h Usage: perf stat [\u0026lt;options\u0026gt;] [\u0026lt;command\u0026gt;] -a, --all-cpus system-wide collection from all CPUs -A, --no-aggr disable aggregation across CPUs or PMUs -B, --big-num print large numbers with thousands' separators -C, --cpu \u0026lt;cpu\u0026gt; list of cpus to monitor in system-wide -D, --delay \u0026lt;n\u0026gt; ms to wait before starting measurement after program start (-1: start with events disabled) # ......2. Eventsperf 记录的性能数据项，称为events。主要分为软件 Events 和硬件 Events。软件 Events 比如有：context-switchs, minor-fault等等。硬件 Events 主要记录micro-architecture相关性能数据，由CPU/PMU提供。如果硬件没有提供，该对应该event不可用。使用 perf list 查看 perf 支持的 events：$ perf listList of pre-defined events (to be used in -e or -M): branch-instructions OR branches [Hardware event] branch-misses [Hardware event] bus-cycles [Hardware event] cache-misses [Hardware event] cache-references [Hardware event] cpu-cycles OR cycles [Hardware event] instructions [Hardware event] ref-cycles [Hardware event] alignment-faults [Software event] bpf-output [Software event] cgroup-switches [Software event] context-switches OR cs [Software event] cpu-clock [Software event] cpu-migrations OR migrations [Software event] dummy [Software event] emulation-faults [Software event] major-faults [Software event] minor-faults [Software event] page-faults OR faults [Software event] task-clock [Software event] duration_time [Tool event] user_time [Tool event] system_time [Tool event]cpu: L1-dcache-loads OR cpu/L1-dcache-loads/ L1-dcache-load-misses OR cpu/L1-dcache-load-misses/ L1-dcache-stores OR cpu/L1-dcache-stores/ L1-icache-load-misses OR cpu/L1-icache-load-misses/ LLC-loads OR cpu/LLC-loads/ LLC-load-misses OR cpu/LLC-load-misses/ LLC-stores OR cpu/LLC-stores/ LLC-store-misses OR cpu/LLC-store-misses/ dTLB-loads OR cpu/dTLB-loads/ dTLB-load-misses OR cpu/dTLB-load-misses/ dTLB-stores OR cpu/dTLB-stores/ dTLB-store-misses OR cpu/dTLB-store-misses/ iTLB-loads OR cpu/iTLB-loads/ iTLB-load-misses OR cpu/iTLB-load-misses/ branch-loads OR cpu/branch-loads/ branch-load-misses OR cpu/branch-load-misses/ node-loads OR cpu/node-loads/ node-load-misses OR cpu/node-load-misses/ node-stores OR cpu/node-stores/ node-store-misses OR cpu/node-store-misses/ branch-instructions OR cpu/branch-instructions/ [Kernel PMU event] branch-misses OR cpu/branch-misses/ [Kernel PMU event] bus-cycles OR cpu/bus-cycles/ [Kernel PMU event] cache-misses OR cpu/cache-misses/ [Kernel PMU event] cache-references OR cpu/cache-references/ [Kernel PMU event] cpu-cycles OR cpu/cpu-cycles/ [Kernel PMU event] instructions OR cpu/instructions/ [Kernel PMU event] mem-loads OR cpu/mem-loads/ [Kernel PMU event] mem-stores OR cpu/mem-stores/ [Kernel PMU event] # ......3. 使用 perf stat 记录性能数据，并输出到终端使用 perf stat 命令，可以运行被测试程序，并在程序结束之后，统计各不同Events的计数，并打印出来。也可以使用-p参数，指定运行程序的进程号。$ perf stat ls -l /usr/bin/ls Performance counter stats for 'ls -lh /home/hxf0223/': 1.65 msec task-clock # 0.757 CPUs utilized 0 context-switches # 0.000 /sec 0 cpu-migrations # 0.000 /sec 112 page-faults # 67.878 K/sec 4,892,473 cycles # 2.965 GHz 3,714,068 instructions # 0.76 insn per cycle 681,289 branches # 412.897 M/sec 20,522 branch-misses # 3.01% of all branches 0.002179855 seconds time elapsed 0.000000000 seconds user 0.002198000 seconds sys可以指定需要记录的events (-e参数)，以及repeat次数(-r参数)：perf stat -r 6 -e cache-misses ls -lh ~/ Performance counter stats for 'ls -lh /home/hxf0223' (6 runs): 20,763 cache-misses ( +- 14.14% ) 0.0015583 +- 0.0000369 seconds time elapsed ( +- 2.37% )4. 记录性能数据到文件，以及分析：perf record, perf report, perf annotate4.1 perf record 记录性能数据到文件使用 perf record 命令，运行被测试程序，并记录测量数据到perf.data数据文件。$ perf record ./test_grain_size./test_grain_size. Process ID: 6226Duration: 232.74 secondsSum: 1.25e+07. loop num: 10000[ perf record: Woken up 934 times to write data ][ perf record: Captured and wrote 283.898 MB perf.data (7441535 samples) ]$ ls -lh perf.data-rw------- 1 \u0026lt;groupname\u0026gt; \u0026lt;username\u0026gt; 284M Oct 31 21:25 perf.data4.2 perf report 查看性能数据使用 perf report 命令，查看perf.data数据文件。使用如下命令，直接打开perf.data文件：perf report4.3 perf annotate 显示源码级别的性能数据perf annotate -i perf.data如何编译使用-ggdb，则可以显示源码级别的性能数据。参考资料 PerfTool pdf Blog of Aleksandar Milenkovic (PerfTool pdf作者) edu: The Laboratory for Advanced Computer Architectures and Systems" },

  { "title": "perf性能分析(4) -- linux perf 工具基本使用(1)", "url": "/blog/2024/perf-tools-04-perf-usage/", "categories": "Performance",
  "tags": "Performance, VTune, TBB", "date": "2024-10-28", "content":
  "1. perf stat 基本使用perf stat 基本功能 – 统计： cycles 数, IPC (instructions per cycle)。IPC \u0026gt;= 1 表示指令执行效率高 分支切换次数(branchs), 分支预测失败次数(branch-misses)，以及失败比例 上下文切换次数(context switches)，以及每秒切换次数 CPU迁移次数(migrations)，以及每秒迁移次数 缺页次数(page faults)，以及每秒缺页次数$ sudo perf stat -p 8460 Performance counter stats for process id '8460': 1,763,985.38 msec task-clock # 7.974 CPUs utilized 5,976 context-switches # 3.388 /sec 19 cpu-migrations # 0.011 /sec 7 page-faults # 0.004 /sec 5,402,549,520,366 cycles # 3.063 GHz 1,249,986,676,725 instructions # 0.23 insn per cycle 247,349,019,761 branches # 140.222 M/sec 217,175,635 branch-misses # 0.09% of all branches 221.220644606 seconds time elapsed参考资料 redhat – Red Hat Enterprise Linux 8 监控和管理系统状态和性能 redhat – 监控和管理系统状态和性能(网页版) 调试技术之perf实战笔记" },

  { "title": "Linux 性能及统计工具 (1)", "url": "/blog/2024/Linux-perf-stat-commands/", "categories": "Performance",
  "tags": "Linux, Performance, Bash", "date": "2024-10-26", "content":
  "1. Virtual Memory Statistics – vmstatvmstat 检测cpu、系统内存(包括 slab)、进程、块设备IO等使用情况： CPU相关：用户时间 / 系统时间 / 空闲时间占比。每秒中断数量 / 上下文切换数量。活动进程数量 / 阻塞进程数量，fork进程数量。 内存相关：active/inactive内存，buff/cache内存，swap使用及每秒交换量。 IO：每秒读写的块数量。1.1. 基本用法$ vmstat -S M # 内存以兆为单位显示procs -----------memory---------- ---swap-- -----io---- -system-- -------cpu------- r b swpd free buff cache si so bi bo in cs us sy id wa st gu 1 0 0 14431 60 749 0 0 357 33 165 0 0 0 100 0 0 0[procs]r – 运行队列中进程数量，如果该值大于核数，说明系统可能存在CPU瓶颈。b – 阻塞的进程数量，如果该值较大，说明系统可能存在IO瓶颈。[memory]swpd – swap space，swpd显示当前被换到物理磁盘上的物理内存的大小。该值大于0表示物理内存不足。free – 空闲物理内存的大小。buff – 块设备的缓冲区占用的大小，buff、cache大代表系统有效利用了内存，能够提高IO性能(在有大量IO时，或者文件读写情况下)。cache – 文件缓存占用的大小。[swap]si – 每秒从物理磁盘读入到虚拟内存的大小。含义见swpd。so – 每秒从虚拟内存写到物理磁盘的大小。含义见swpd。[io]bi – 块设备每秒接收的块数量。bo – 块设备每秒发送的块数量。[system]in – interrupt，每秒中断数，包括时钟中断。cs – count per second，每秒上下文切换数量。包括系统调用、进程切换、线程切换等。[cpu]us – 用户进程执行时间百分比。sy – 内核进程执行时间百分比。id – 空闲时间百分比。一般 us + sy + id =100。wa – 等待IO时间百分比。st – 被信号中断的进程时间百分比。gu – 被引起的页错误的进程时间百分比。1.2 查看统计信息 – 内存及CPU相关$ vmstat -s -S M 15881 M total memory 946 M used memory 804 M active memory 252 M inactive memory 14421 M free memory 61 M buffer memory 749 M swap cache 4095 M total swap 0 M used swap 4095 M free swap 1707 non-nice user cpu ticks 82 nice user cpu ticks 1977 system cpu ticks 4068162 idle cpu ticks 176 IO-wait cpu ticks 0 IRQ cpu ticks 15 softirq cpu ticks 0 stolen cpu ticks 0 non-nice guest cpu ticks 0 nice guest cpu ticks 672044 K paged in 66205 K paged out 0 pages swapped in 0 pages swapped out 438059 interrupts 452000 CPU context switches 1729905791 boot time 3331 forks1.3 其他选项用法vmstat -a # 显示 active/inactive 内存vmstat -d # 显示磁盘IO统计信息vmstat -f # fork 进程统计信息vmstat -m # 内核 SLAB 相关统计信息vmstat 2 # 每两秒输出一次统计信息vmstat 1 5 -t # 每隔一秒输出一次统计信息，持续5秒，并显示时间戳2. top 命令$ toptop - 11:16:35 up 1:53, 5 users, load average: 3.15, 0.77, 0.26Tasks: 252 total, 2 running, 250 sleeping, 0 stopped, 0 zombie%Cpu(s):100.0 us, 0.0 sy, 0.0 ni, 0.0 id, 0.0 wa, 0.0 hi, 0.0 si, 0.0 stMiB Mem : 15881.5 total, 14023.4 free, 1340.6 used, 816.5 buff/cacheMiB Swap: 4096.0 total, 4096.0 free, 0.0 used. 14541.0 avail Mem PID USER PR NI VIRT RES SHR S %CPU %MEM TIME+ COMMAND # ... 第一行 启动时间，持续时间： 11:16:35 up 1:53 系统1分钟、5分钟、15分钟的CPU负载值：load average: 3.15, 0.77, 0.26 第三行 CPU占用：us 用户进程占比；sy 内核CPU占比；wa(IO wait) 等待IO的CPU时间占比；hi 硬件中断CPU时间占比；si 软件中断CPU时间占比；st 虚拟机偷取(steal)CPU时间占比。 2.1 查看进程下所有线程信息$ top -Hp \u0026lt;PID\u0026gt;# ......3. uptime 命令$ uptime 11:16:35 up 1:53, 5 users, load average: 3.15, 0.77, 0.26更多资料 Velocity2015_LinuxPerfTools pdf" },

  { "title": "有用的 GCC 编译选项收集", "url": "/blog/2024/%E6%9C%89%E7%94%A8%E7%9A%84GCC%E7%BC%96%E8%AF%91%E9%80%89%E9%A1%B9%E6%94%B6%E9%9B%86/", "categories": "GCC",
  "tags": "GCC, Linux, Performance", "date": "2024-10-24", "content":
  "1. 编译选项 -fverbose-asm-fverbose-asm 将编译信息(编译选项等等)、C源码中的变量名，以注释的形式嵌入到汇编代码中，便于分析。" },

  { "title": "内存分析需要理解的几个概念", "url": "/blog/2024/Linux-%E5%86%85%E5%AD%98%E7%AE%A1%E7%90%86%E8%A6%81%E7%90%86%E8%A7%A3%E7%9A%84%E6%A6%82%E5%BF%B5/", "categories": "Linux",
  "tags": "Linux, Performance", "date": "2024-10-21", "content":
  "“真实的” 内存空闲率 = (free + shared + buffers + cached)/ Total = 5860 M / 7983M X 100 % = 73.4 %1. Linux内存分类 匿名内存：存储用户计算过程中间的数据，与物理磁盘上的文件无关； File-Backed内存：用作磁盘高速缓存，其物理内存与物理磁盘上的文件对应；1.1 Shmem包括： shared memory. tmpfs：所使用的内存（基于内存的文件系统）。Linux将空闲内存用于缓存，并且在需要的时候，回收这些内存。 devtmpfs：指的是/dev文件系统，/dev目录下所有的文件占用的空间也属于共享内存。1.2 buffers \u0026amp; cache buffers： 用于块设备I/O缓存； cached：page cache，用于文件系统（读取文件时，kernel创建的缓存）；使用top命令查看时，buffers与cached通常合并在一起显示，他们都是kernel使用的缓存。在需要的时候，可以被回收。2. 查看内存使用情况2.1 内存利用率详细信息cat /proc/meminfoMemTotal: 8174352 kBMemFree: 376952 kBBuffers: 527412 kBCached: 5178924 kBSwapCached: 60 kBActive: 3061760 kBInactive: 4066588 kBActive(anon): 1112780 kBInactive(anon): 314156 kBActive(file): 1948980 kBInactive(file): 3752432 kBUnevictable: 6724 kBMlocked: 6724 kBSwapTotal: 16779884 kBSwapFree: 16777400 kBDirty: 376 kBWriteback: 0 kBAnonPages: 1428844 kBMapped: 64632 kBShmem: 644 kBSlab: 557384 kBSReclaimable: 338272 kBSUnreclaim: 219112 kBKernelStack: 4024 kBPageTables: 12440 kBNFS_Unstable: 0 kBBounce: 0 kBWritebackTmp: 0 kBCommitLimit: 20867060 kBCommitted_AS: 2406484 kBVmallocTotal: 34359738367 kBVmallocUsed: 111536 kBVmallocChunk: 34359455060 kBHugePages_Total: 0HugePages_Free: 0HugePages_Rsvd: 0HugePages_Surp: 0Hugepagesize: 2048 kBDirectMap4k: 6384 kBDirectMap2M: 2080768 kBDirectMap1G: 6291456 kB2.2 查看进程内存占用有用的信息： VmData：data段大小； VmExe：text段大小； Vmlib：共享库占用内存空间大小； VmRSS：物理内存使用量； VMSwap：交换分区使用量；cat /proc/30029/statusName: javaState: S (sleeping)Tgid: 30029Pid: 30029PPid: 29983TracerPid: 0Uid: 54322 54322 54322 54322Gid: 54323 54323 54323 54323FDSize: 8192Groups: 10 54323VmPeak: 2754032 kBVmSize: 2678836 kBVmLck: 0 kBVmHWM: 1337912 kBVmRSS: 1337512 kBVmData: 2575692 kBVmStk: 1012 kBVmExe: 60 kBVmLib: 101564 kBVmPTE: 3048 kBThreads: 98SigQ: 0/63825SigPnd: 0000000000000000ShdPnd: 0000000000000000SigBlk: 0000000000000004SigIgn: 0000000000000001SigCgt: 1000000180005cceCapInh: 0000000000000000CapPrm: 0000000000000000CapEff: 0000000000000000CapBnd: ffffffffffffffffCpus_allowed: ffffffffCpus_allowed_list: 0-31Mems_allowed: 00000000,Mems_allowed_list: 0voluntary_ctxt_switches: 12468nonvoluntary_ctxt_switches: 192.3 pmap 命令执行 ps aux 后输出的各进程的 RSS (resident set size), 表示进程占用内存的大小，单位是KB。 需要注意的是, RSS 值实际上是基于 pmap 命令，表示该进程正在使用的物理内存的总和。pmap -x 3002930029: /slview/jdk150/jdk1.5.0_06/bin/java -com.apache.TestAddress Kbytes RSS Dirty Mode Mapping0000000008048000 60 48 0 r-x-- java0000000008057000 8 8 8 rwx-- java0000000009f1d000 23184 23140 23140 rwx-- [ anon ]000000004d1f1000 108 96 0 r-x-- ld-2.5.so000000004d20c000 4 4 4 r-x-- ld-2.5.so000000004d20d000 4 4 4 rwx-- ld-2.5.so000000004d214000 1356 548 0 r-x-- libc-2.5.so000000004d367000 8 8 8 r-x-- libc-2.5.so00007f581e51d000 16 16 0 r--s- huanan-product-2.6.1-snapshots.jar00007f581e521000 24 24 0 r--s- dt.jar00007f581e527000 36 36 0 r--s- gnome-java-bridge.jar00007f581e530000 32 32 8 rw-s- 1322800007f581e538000 4 4 4 rw--- [ anon ]00007f581e539000 4 4 0 r---- [ anon ]00007f581e53a000 8 8 8 rw--- [ anon ]00007fffe9eb7000 84 32 32 rw--- [ stack ]00007fffe9fff000 4 4 0 r-x-- [ anon ]ffffffffff600000 4 0 0 r-x-- [ anon ](部分省略)---------------- ------ ------ ------total kB 2484196 36180 268803. 参考资料 How Linux Works(三): 内存管理 Linux 0.11 源码阅读笔记-内存的基础概念kernel内存分配学习文档: 深度剖析 Linux 伙伴系统的设计与实现 Linux 堆内存管理深入分析 [Linux内存管理 五、物理内存空间布局及管理](https://uniondong.github.io/docs/linux/linux_memory_manage/%E4%BA%94%E7%89%A9%E7%90%86%E5%86%85%E5%AD%98%E7%A9%BA%E9%97%B4%E5%B8%83%E5%B1%80%E5%8F%8A%E7%AE%A1%E7%90%86/)" },

  { "title": "perf性能分析(3) -- Intel VTune 配合 linux perf 使用", "url": "/blog/2024/perf-tools-03-Intel-VTune-and-perf-usage-03/", "categories": "Performance",
  "tags": "Performance, VTune, TBB", "date": "2024-10-20", "content":
  "1. 配置1.1 安装 perfsudo apt-get install linux-tools-common linux-tools-generic linux-tools-`uname -r`1.2 设置系统相关设置项以允许 perf 采集# 允许非特权用户进行内核分析和访问 CPU 事件echo 0 | sudo tee /proc/sys/kernel/perf_event_paranoidsudo sh -c 'echo kernel.perf_event_paranoid=0 \u0026gt;\u0026gt; /etc/sysctl.d/local.conf'# 启用内核模块符号解析以供非特权用户使用echo 0 | sudo tee /proc/sys/kernel/kptr_restrictsudo sh -c 'echo kernel.kptr_restrict=0 \u0026gt;\u0026gt; /etc/sysctl.d/local.conf'# 生效系统设置sudo sysctl -p其中： 第一项设置，允许非特权用户进行内核分析和访问 CPU 事件。 第二项设置(kptr_restrict)，可以使得perf工具可以访问到内核指针，即允许内核符号(kallsyms)被映射到用户层。2. 开始 perf 测量及收集数据在被测量的程序中，添加如下代码获取自身PID:#include \u0026lt;unistd.h\u0026gt; // getpidconst pid_t pid = getpid();std::cout \u0026lt;\u0026lt; argv[0] \u0026lt;\u0026lt; \". Process ID: \" \u0026lt;\u0026lt; pid \u0026lt;\u0026lt; std::endl;使用ggdb，并使用fno-omit-frame-pointer等编译选项，遍已完成之后，启动程序。随后启动perf命令：perf record -F 599 -e cycles,cache-misses -ag -p 12738 -- sleep 200其中： -e cycles,cache-misses 表示采集的类型，使用 perf list 可以列出所有可用的事件，如CPU相关的事件，cache相关的事件，以及是硬件采集 (PMC) 还是软件采集； -p 12738 表示进程 PID； sleep 200 表示采集持续时间；采集完成之后，将生成的perf.data 重命名为 data.perf， .perf 文件是 VTune 可以识别的文件格式。为防止采样频率与代码中的某些周期性代码因同频而导致每次采集到相同的地方，故设置采集频率不能为10的倍数，也不能是2的幂次方，可尽量避免采集误导性数据。参考Using perf On Arm platforms3. VTune 分析数据使用VTune Profiler 导入经过重命名的 data.perf ： Menu -\u0026gt; Import Result...; Import raw trace data， 选择 data.perf 文件； Import 按钮；4. 编译配置Windows下配置如下： 确保Debug Information Format设置为Program Database (/Zi); 确保Optimization选项设置为Maximum Optimizations (Favor Size) (/O1); 将Optimization Diagnostic Level设置为Level 2 (/Qopt-report:2);参考： Intel® VTune™ Profiler 分析 C++ 程序的常见性能瓶颈( Windows 平台)5. 参考资料 现代CPU性能分析与优化 – Linux Performance perf Examples6. 更多学习资料 Performance Analysis and Tuning on Modern CPU 中文翻译 Using Linux perf at Netflix" },

  { "title": "perf性能分析(2) -- Intel VTune 配置与使用(2)", "url": "/blog/2024/perf-tools-02-Intel-VTune-usage-02/", "categories": "Performance",
  "tags": "Performance, VTune, TBB, Cpp", "date": "2024-10-16", "content":
  "测试代码：test_tbb_perf_vtune_profiler注意：编译选项需要添加”-g”，以便于VTune Profiler可以显示源码信息。1. 测试原始来源VTune Profiler 进行性能分析：使用VTune Profiler测试TBB overhead。2. 资料 Intel TBB API 使用教程：Intel® oneAPI Threading Building Blocks Too long TBB’s shedule time when using parallel_deterministic_reduce3. 测试过程及优化3.1 reduce 针对一些比如遍历求和操作，他们之间没有顺序要求，可以改用并行的 reduce。前提是数据的构造代价小，如稀疏矩阵拷贝代价就比较大。 计算的先后顺序有关的，比如针对浮点的乘加操作，先后顺序变化影响计算精度，此时使用parallel_deterministic_reduce。官方解释是：合并顺序是预先定义好的，确保每次调用deterministic_reduce的结果相同。3.2 几个不同的优化方式Intel TBB 动态划分任务，以及把任务提交给线程执行，都需要消耗时间。优化包括： 调整 grain size 减少调度的开销。 使用静态划分static_partitioner减少调度开销。3.2.1 调整 grain size通过设置grain size，可以大致设定 TBB 每个任务要处理的数据量，即划分粒度：tbb::blocked_range\u0026lt;double*\u0026gt;(v, v + n, 1000)设置grain size之前，显示的热点 call stack 如下图（100次循环）：设置grain size等于10000，Intel TBB 内部调度时间明显减少（10000次循环）：3.2.2 使用静态划分static_partitioner通过设置使用static_partitioner，即预先划分好任务，减少调度开销，具有与blocked_range类似的效果，但其控制的方式不同。static_partitioner适用于任务均衡的计算场景。tbb::task_arena ta(8); double sum = ta.execute([\u0026amp;]() { return tbb::parallel_deterministic_reduce( tbb::blocked_range\u0026lt;double*\u0026gt;(v, v + n), 0.0, [](const tbb::blocked_range\u0026lt;double*\u0026gt;\u0026amp; r, double value) -\u0026gt; double { return std::accumulate(r.begin(), r.end(), value); }, std::plus\u0026lt;double\u0026gt;(), tbb::static_partitioner{}); }); return sum;" },

  { "title": "给 shared_ptr 添加自定义 deleter 的几种方式", "url": "/blog/2024/shared_ptr-deleters-usage/", "categories": "Cpp",
  "tags": "Cpp", "date": "2024-10-15", "content":
  "1. 使用函数#include \u0026lt;iostream\u0026gt;#include \u0026lt;memory\u0026gt;void deleter(Sample* ptr) { std::cout \u0026lt;\u0026lt; \"delete function called\" \u0026lt;\u0026lt; std::endl; delete ptr;}std::shared_ptr\u0026lt;Sample\u0026gt; sp(new Sample, deleter);2. 使用仿函数#include \u0026lt;iostream\u0026gt;#include \u0026lt;memory\u0026gt;struct Deleter { void operator()(Sample* ptr) { std::cout \u0026lt;\u0026lt; \"deleter function called\" \u0026lt;\u0026lt; std::endl; delete ptr; }};std::shared_ptr\u0026lt;Sample\u0026gt; sp(new Sample, Deleter{});3. 使用 lambda 表达式#include \u0026lt;iostream\u0026gt;#include \u0026lt;memory\u0026gt;std::shared_ptr\u0026lt;Sample\u0026gt; sp(new Sample, [](Sample* ptr) { std::cout \u0026lt;\u0026lt; \"lambda function called\" \u0026lt;\u0026lt; std::endl; delete ptr;});4. 使用 std::default_delete#include \u0026lt;iostream\u0026gt;#include \u0026lt;memory\u0026gt;std::shared_ptr\u0026lt;Sample\u0026gt; sp(new Sample, std::default_delete\u0026lt;Sample\u0026gt;{});" },

  { "title": "perf性能分析(1) -- Intel VTune 配置与使用(1)", "url": "/blog/2024/perf-tools-01-Intel-VTune/", "categories": "Performance",
  "tags": "Performance, VTune, TBB, Cpp", "date": "2024-10-15", "content":
  "1. Intel VTune 配置1.1 使能 ptrace使能进程跟踪 (ptrace) 功能 (attach ID)，以便 VTune 可以监控到进程的运行情况。sudo vim /etc/sysctl.d/10-ptrace.conf# set kernel.yama.ptrace_scope = 0# 使配置生效sudo sysctl --system -a -p | grep yama # 应用配置，或者也可以选择重启电脑1.2 安装 Sampling Driverscd ~/intel/oneapi/vtune/latest/sepdk/src./build-driver# 添加 vtune 组并将你的用户添加到该组# 创建一个新的 shell，或者重新启动系统sudo groupadd vtunesudo usermod -a -G vtune `whoami`# 安装 sep 驱动程序sudo ./insmod-sep -r -g vtune1.3 检查软硬件配置: 查看VTune Profiler 可以做哪些profiling~/intel/oneapi/vtune/latest/bin64python3 self_check.py最终会出来一个结果， 显示 VTune Profiler 可以做哪些 profiling, 哪些 profiling 不能做。2. 参考及资料 Intel® VTune™ Profiler 分析 C++ 程序的常见性能瓶颈( Windows 平台) perf Examples 现代CPU性能分析与优化 – Intel Vtune 自顶向下的微架构分析方法 Intel® VTune™ Profiler Performance Analysis Cookbook VTune Documentation \u0026amp; Code Samples HOW: Analysis Types intel-TBB使用笔记 Intel TBB API 官方教程 oneTBB Developer Guide3. TBB 调度实现 TBB并发库代码学习 C++ TBB 并行编程教程 – 控制用于执行的线程数量 oneTBB Developer Guide" },

  { "title": "Plugin 的创建及使用", "url": "/blog/2024/Plugin-Boost-DLL/", "categories": "Cpp",
  "tags": "Cpp, Plugin", "date": "2024-10-15", "content":
  "使用 BOOST_DLL_ALIAS 定义插件接口。 使用 import_alias 导入插件接口。 使用 creator / Factory 模式，提供创建插件实例的接口。Demo Code: test_plugin_dll1. 实现插件接口1.1 DSO/DLL原型定义class DIInterface { public: DIInterface() = default; virtual ~DIInterface() = default; //virtual std::shared_ptr\u0026lt;DIInterface\u0026gt; clone() = 0; virtual void loadInfo(const std::string\u0026amp; dataFile) = 0; virtual QWidget* getWidget() = 0; virtual void unload() = 0; protected: QWidget* widget_;};1.2 接口定义及实现std::shared_ptr\u0026lt;test::plugin::DIInterface\u0026gt; diLoader(const std::string\u0026amp; infoFile) { std::shared_ptr\u0026lt;test::plugin::DIInfoChaoke\u0026gt; ptr = std::make_shared\u0026lt;test::plugin::DIInfoChaoke\u0026gt;(); ptr-\u0026gt;loadInfo(infoFile); return ptr;}BOOST_DLL_ALIAS(diLoader, // 被封装成插件接口的函数名 diLoaderAlias); // 别名，可用于创建插件实例2. 插件加载// static const char* kDILoaderFnName = \"diLoaderAlias\";const std::string dir_path = boost::dll::program_location().parent_path().string();const auto chaoke_dll_path = fs::path(dir_path) / \"di_plugin_chaoke.dll\";lib_ = boost::dll::shared_library(chaoke_dll_path.string());try { auto creator = boost::dll::import_alias\u0026lt;di_loader_proto_t\u0026gt;(chaoke_dll_path.string(), kDILoaderFnName, boost::dll::load_mode::append_decorations); auto loader_proto = lib_.get_alias\u0026lt;di_loader_proto_t\u0026gt;(kDILoaderFnName); instance_ = loader_proto(dataFilePath);} catch (const std::exception\u0026amp; e) { SPDLOG_WARN(\"Failed to load dll: {}\", e.what()); return false;}3. 插件卸载在卸载DSO/DLL之后，由于类的成员函数代码内存被释放，故其vtable所指向的内存（即成员方法）变成非法地址，在调用其成员方法函数，以及析构函数时，会导致程序崩溃。所以在卸载DLL/DSO之前，需要先释放外部持有的资源： 调用unload接口，释放资源，比如释放QWidget; 释放外部持有的shared_ptr;4. 技巧：将boost::dll::shared_library生命周期与Plugin实例对象生命周期绑定定义如下deletor：struct library_holding_deleter { library_holding_deleter(std::shared_ptr\u0026lt;boost::dll::shared_library\u0026gt; libDLL) : lib_(libDLL) {} void operator()(DIInterface* p) const { delete p; } std::shared_ptr\u0026lt;boost::dll::shared_library\u0026gt; lib_;};在创建DIInterface实例时，传入library_holding_deleter（make_shared不支持传入自定义deleter）：// auto lib_dll = std::make_shared\u0026lt;boost::dll::shared_library\u0026gt;(chaoke_dll_path.string());// library_holding_deleter deletor(lib_dll);std::shared_ptr\u0026lt;test::plugin::DIInterface\u0026gt; diLoader(const std::string\u0026amp; infoFile, library_holding_deleter deletor) { std::shared_ptr\u0026lt;test::plugin::DIInfoChaoke\u0026gt; ptr(new test::plugin::DIInfoChaoke, deletor); ptr-\u0026gt;loadInfo(infoFile); return ptr;}5. 资料 Boost DLL – Plugin 跨平台的 plugin 開發函式庫：Boost DLL - 進階 Boost Plugin Loader" },

  { "title": "性能优化学习资料", "url": "/blog/2024/%E6%80%A7%E8%83%BD%E4%BC%98%E5%8C%96%E5%AD%A6%E4%B9%A0%E8%B5%84%E6%96%99/", "categories": "Performance",
  "tags": "Performance, Linux, CPU", "date": "2024-10-10", "content":
  "原作者 Linux Performance 原作者 github perf-tools pdf – Linux Performance Tools 中文翻译 pdf – Linux Performance Tools github – Performance Ninja Class ARMVirtualization: Performance and Architectural Implications blog – Sherlock’s blog" },

  { "title": "STL 图解", "url": "/blog/2024/STL%E5%9B%BE%E8%A7%A3/", "categories": "Cpp",
  "tags": "Cpp", "date": "2024-10-10", "content":
  "STL 包含五种主要组件： 算法（algorithm）：定义计算过程。 容器（container）：管理一组内存位置。 迭代器（iterator）：提供算法遍历容器的方法。 函数对象（function object）：将函数封装在对象中，供其他组件使用。 适配器（adaptor）：调整组件以提供不同的接口。从实现来看还需要包含： 分配器（allocator）：用于处理容器对内存的分配与释放请求。以下分析适用于 GCC9 。1. 源码阅读1.1. ::template在 __rebind 函数体中，在 :: 后面有个 template 关键字，这是用于告诉编译器 template 后面的 \u0026lt; 不是比较符号，而是模板参数符号。就是类似于 _Tp 前面的 typename 是告诉编译器 :: 后面的是类成员函数，而不是 static 函数。using type = typename _Tp::template rebind\u0026lt;_Up\u0026gt;::other;2. allocator1分配器是负责封装堆内存管理的对象。2.1. 分配器上图左侧。C++的默认的内存分配器 std::allocator，继承至 __gnu_cxx::new_allocator。2.1.1. __gnu_cxx::new_allocator（1）对传入类型进行了类型萃取。（2）rebind 重新绑定，定义 other 类型，用于萃取器萃取类型。（3）封装实现分配对象内存、初始化对象、析构对象、释放对象内存，底层使用 new 和 delete。 address() : 用于获取分配地址 allocate() : 用于分配内存 deallocate() : 用于释放内存 max_size() : 获取最大可分配数量 construct() : 调用已分配内存对象的构造函数 destroy() : 调用析构函数2.1.2. std::allocator（1）偏特化处理 void（2）偏特化处理 const 和 volatile，重新形成有效的分配器类型（3）类模板 std::allocator rebind 重新绑定，定义 other 类型，用于萃取器获取类型。2.2. 萃取器上图右侧。类 __gnu_cxx::__alloc_traits 继承于类std::allocator_traits，再继承于 std::__allocator_traits_base，用于获取内存分配器 allocator 的各个属性。2.2.1. std::__allocator_traits_base 用于获取内存分配器 allocator 的各个属性。2.2.2. std::allocator_traits 私有函数模板 _S 开头是封装对应名称分配器 _Alloc 原生的函数模板 共有函数模板是封装自身对应名称 _S 开头的私有函数模板 有个特化版本是当 _Alloc 是 std::allocator，起别名后重复调用通用模板的共有函数模板2.2.3. __gnu_cxx::__alloc_traits 全部继承 std::allocator_traits，重载 construct 和 destroy 非标准类型指针 _S 开头的静态函数是封装父类萃取类型3. iterator2迭代器是指向容器内元素的对象（如指针）。3.1. 迭代器类型上图中间。3.1.1. std::iterator_tag/// Marking input iterators.struct input_iterator_tag {};/// Marking output iterators.struct output_iterator_tag {};/// Forward iterators support a superset of input iterator operations.struct forward_iterator_tag : public input_iterator_tag {};/// Bidirectional iterators support a superset of forward iterator/// operations.struct bidirectional_iterator_tag : public forward_iterator_tag {};/// Random-access iterators support a superset of bidirectional/// iterator operations.struct random_access_iterator_tag : public bidirectional_iterator_tag {};STL库支持的迭代器类型是所有编程语言中最全面的，共有五种： InputIterator : 输入迭代器。支持对容器元素的逐个遍历，以及对元素的读取 (input)； OutputIterator : 输出迭代器。支持对容器元素的逐个遍历，以及对元素的写入 (output)。 ForwardIterator : 前向迭代器。向前逐个遍历元素。可以对元素读取； BidirectionalIterator : 双向迭代器。支持向前向后逐个遍历元素，可以对元素读取。 RandomAccessIterator : 随机访问迭代器。支持O(1)时间复杂度对元素的随机位置访问，支持对元素的读取。这些是空类型用于区分不同的迭代器，区别不在于它们所包含的内容，而在于它们的是什么类型，然后可以基于不同迭代器类型支持的不同操作。3.1.2. std::iterator此类只定义嵌套的 typedef，子类迭代器类可以继承这个类以节省一些工作，然后用于特化和重载。3.1.3. std::iterator_traits此类只定义嵌套的 typedef，简单地从 Iterator 转发嵌套的 typedef 参数。 提供指针和指向常量的指针的特化版本。3.2. 迭代器函数上图左侧。3.2.1. std::distance计算迭代器之间的距离，通过一层类型判断确定 iterator_tag，然后转发给具体函数 __distance()。包括单向、随机。3.2.2. std::advance向前或向后移动迭代器，通过一层类型判断确定 iterator_tag，然后转发给具体函数 __advance()。包括单向、双向、随机。 next() 向后移动 prev() 向前移动3.3. 衍生迭代器上图右侧。3.3.1. std::reverse_iterator反向迭代器。 记录当前迭代器 重载操作运算符，++ 内部实现为 --，-- 内部实现为 ++ 等反向操作 该迭代器全局的各种重载操作运算符函数模板3.3.2. std::back_insert_iterator尾部插入迭代器。 记录一个容器，调用容器自己实现的 push_back() 重载操作运算符，返回自身解引用 该迭代器全局的插入函数模板3.3.3. std::front_insert_iterator头部插入迭代器。 记录一个容器，调用容器自己实现的 push_front() 重载操作运算符，返回自身解引用 该迭代器全局的插入函数模板3.3.4. std::insert_iterator插入迭代器。 记录一个容器，调用容器自己实现的 insert() 重载操作运算符，返回自身解引用 该迭代器全局的插入函数模板3.3.5. __gnu_cxx::__normal_iterator这个迭代器适配器是一个普通的适配器，因为它不会改变迭代器参数的任何运算符的语义。它的主要目的是将不是类的迭代器（例如指针）转换为类迭代器。_Container 参数单独存在（不同容器），因此使用此模板的不同容器可以实例化不同的类型，即使 _Iterator 参数相同。 记录当前迭代器 重载操作运算符 该迭代器全局的各种重载操作运算符函数模板3.3.6. std::move_iterator类模板 move_iterator 是一个迭代器适配器，其行为与基础迭代器相同，只是其解引用运算符隐式将基础迭代的解引用运算符返回的值转换为右值引用。可以使用移动迭代器调用一些通用算法，以移动代替复制。 记录当前迭代器 重载操作运算符 该迭代器全局的各种重载操作运算符函数模板4. vectorstd::vector 是封装动态数组的顺序容器。连续存储元素，这意味着不仅可通过迭代器，还能用指向元素的常规指针访问元素。4.1. 基类std::_Vector_base 是基类，通过萃取获取分配类型和分配类型指针。typedef typename __gnu_cxx::__alloc_traits\u0026lt;_Alloc\u0026gt;::template rebind\u0026lt;_Tp\u0026gt;::other _Tp_alloc_type;typedef typename __gnu_cxx::__alloc_traits\u0026lt;_Tp_alloc_type\u0026gt;::pointer pointer;具体保存数据结构。struct _Vector_impl_data { pointer _M_start; // 起始指针 pointer _M_finish; // 实际空间结束指针 pointer _M_end_of_storage; // 分配空间结束指针}真正数据结构是 _Vector_impl，该结构多重继承于 _Vector_impl_data 和 _Tp_alloc_type（具体类型）实现封装管理内存调整（增加、收缩、删除）。4.2. 具体类std::vector 实现具体成员函数，实际都是对以下指针进行操作，部分进行迭代器封装。_M_impl._M_start_M_impl._M_finish_M_impl._M_end_of_storage例如 begin() end() empty() 函数。/** * Returns a read/write iterator that points to the first * element in the %vector. Iteration is done in ordinary * element order. */iterator begin() _GLIBCXX_NOEXCEPT { return iterator(this-\u0026gt;_M_impl._M_start);}/** * Returns a read/write iterator that points one past the last * element in the %vector. Iteration is done in ordinary * element order. */iterator end() _GLIBCXX_NOEXCEPT { return iterator(this-\u0026gt;_M_impl._M_finish);}/** * Returns true if the %vector is empty. (Thus begin() would * equal end().) */_GLIBCXX_NODISCARD bool empty() const _GLIBCXX_NOEXCEPT { return begin() == end();}5. liststd::list 是支持常数时间从容器任何位置插入和移除元素的容器。不支持快速随机访问。它通常实现为双向链表。5.1. 节点类型节点类型分为以下三个： _List_node_base : 基类节点，包含前向、后向指针。 _List_node_header : 头节点，继承于 _List_node_base，额外包含节点个数。 _List_node : 数据节点，继承于 _List_node_base，额外包含数据。5.2. 基类std::_List_node 是基类，该类的真正结构是 _List_impl，继承于 _List_node\u0026lt;_Tp\u0026gt;，内部包含 _List_node_header 头节点。std::_List_node 中声明 _List_impl 成员变量，操作都是对 _List_impl 成员变量中数据操作。5.3. 具体类std::list 实现具体成员函数，实际都是对以下指针进行操作，部分进行迭代器封装。_M_impl._M_node例如 begin() end() empty() 函数。 /** * Returns a read/write iterator that points to the first element in the * %list. Iteration is done in ordinary element order. */ iterator begin() _GLIBCXX_NOEXCEPT { return iterator(this-\u0026gt;_M_impl._M_node._M_next); } /** * Returns a read/write iterator that points one past the last * element in the %list. Iteration is done in ordinary element * order. */ iterator end() _GLIBCXX_NOEXCEPT { return iterator(\u0026amp;this-\u0026gt;_M_impl._M_node); } /** * Returns true if the %list is empty. (Thus begin() would equal * end().) */ _GLIBCXX_NODISCARD bool empty() const _GLIBCXX_NOEXCEPT { return this-\u0026gt;_M_impl._M_node._M_next == \u0026amp;this-\u0026gt;_M_impl._M_node; }6. dequestd::deque（double-ended queue，双端队列）是有下标顺序容器，它允许在它的首尾两端快速插入及删除。另外，在 deque 任一端插入或删除不会使指向其余元素的指针或引用失效。deque 的元素不是相接存储的：典型实现用单独分配的固定尺寸数组的序列，外加额外的序列，这表示下标访问必须进行二次指针解引用，与之相比 vector 的下标访问只进行一次。6.1. 基类std::_Deque_base 是基类，通过萃取获取分配类型和分配类型指针。真正数据结构是 _Deque_impl，该结构继承于 _Tp_alloc_type（具体类型）实现封装管理内存调整（初始化、增加、删除）。struct _Deque_impl{ _Map_pointer _M_map; // 固定尺寸数组关联的序列 size_t _M_map_size; // 序列大小 iterator _M_start; // 起始迭代器 iterator _M_finish; // 终止迭代器}6.2. 具体类std::deque 实现具体成员函数，实际都是对以下指针进行操作，部分进行迭代器封装。_M_impl._M_start_M_impl._M_finish例如 push_front() push_back()函数。 /** * @brief Add data to the front of the %deque. * @param __x Data to be added. * * This is a typical stack operation. The function creates an * element at the front of the %deque and assigns the given * data to it. Due to the nature of a %deque this operation * can be done in constant time. */ void push_front(const value_type\u0026amp; __x) { if (this-\u0026gt;_M_impl._M_start._M_cur != this-\u0026gt;_M_impl._M_start._M_first) { _Alloc_traits::construct(this-\u0026gt;_M_impl, this-\u0026gt;_M_impl._M_start._M_cur - 1, __x); --this-\u0026gt;_M_impl._M_start._M_cur; } else _M_push_front_aux(__x); } /** * @brief Add data to the end of the %deque. * @param __x Data to be added. * * This is a typical stack operation. The function creates an * element at the end of the %deque and assigns the given data * to it. Due to the nature of a %deque this operation can be * done in constant time. */ void push_back(const value_type\u0026amp; __x) { if (this-\u0026gt;_M_impl._M_finish._M_cur != this-\u0026gt;_M_impl._M_finish._M_last - 1) { _Alloc_traits::construct(this-\u0026gt;_M_impl, this-\u0026gt;_M_impl._M_finish._M_cur, __x); ++this-\u0026gt;_M_impl._M_finish._M_cur; } else _M_push_back_aux(__x); }7. 适配器7.1. stackstd::stack 类是容器适配器，它给予程序员栈的功能——特别是 FILO（先进后出）数据结构。适配器可以为标准容器 std::vector（包括 std::vector\u0026lt;bool\u0026gt;）、std::deque 和 std::list。如果没有为特定的 stack 类特化指定容器类，那么就会使用标准容器 std::deque。实现方式一般是声明适配器变量，然后封装调用适配器容器函数来实现自己函数功能。例如 top() push()函数。template\u0026lt;typename _Tp, typename _Sequence = deque\u0026lt;_Tp\u0026gt; \u0026gt;class stack{ // 适配器容器 deque\u0026lt;_Tp\u0026gt; _Sequence c; /** * Returns a read/write reference to the data at the first * element of the %stack. */ reference top() { __glibcxx_requires_nonempty(); return c.back(); } /** * @brief Add data to the top of the %stack. * @param __x Data to be added. * * This is a typical %stack operation. The function creates an * element at the top of the %stack and assigns the given data * to it. The time complexity of the operation depends on the * underlying sequence. */ void push(value_type\u0026amp;\u0026amp; __x) { c.push_back(std::move(__x)); }}7.2. queuestd::queue 类是容器适配器，它给予程序员队列的功能——尤其是 FIFO （先进先出）数据结构。适配器可以为标准容器 std::deque 和 std::list。实现方式一般是声明适配器变量，然后封装调用适配器容器函数来实现自己函数功能。例如 push() pop()函数。template\u0026lt;typename _Tp, typename _Sequence = deque\u0026lt;_Tp\u0026gt; \u0026gt;class queue{ // 适配器容器 deque\u0026lt;_Tp\u0026gt; _Sequence c; /** * @brief Add data to the end of the %queue. * @param __x Data to be added. * * This is a typical %queue operation. The function creates an * element at the end of the %queue and assigns the given data * to it. The time complexity of the operation depends on the * underlying sequence. */ void push(value_type\u0026amp;\u0026amp; __x) { c.push_back(std::move(__x)); } /** * @brief Removes first element. * * This is a typical %queue operation. It shrinks the %queue by one. * The time complexity of the operation depends on the underlying * sequence. * * Note that no data is returned, and if the first element's * data is needed, it should be retrieved before pop() is * called. */ void pop() { __glibcxx_requires_nonempty(); c.pop_front(); }}7.3. priority_queuepriority_queue 是容器适配器，它提供常数时间的（默认）最大元素查找，对数代价的插入与提取。适配器可以为标准容器 std::vector（包括 std::vector\u0026lt;bool\u0026gt;）和 std::deque。实现方式一般是声明适配器变量，然后封装调用适配器容器函数来实现自己函数功能。例如 top() pop()函数。template\u0026lt;typename _Tp, typename _Sequence = vector\u0026lt;_Tp\u0026gt;, typename _Compare = less\u0026lt;typename _Sequence::value_type\u0026gt; \u0026gt;class priority_queue{ // 适配器容器 vector\u0026lt;_Tp\u0026gt; _Sequence c; _Compare comp; /** * Returns a read-only (constant) reference to the data at the first * element of the %queue. */ const_reference top() const { __glibcxx_requires_nonempty(); return c.front(); } /** * @brief Removes first element. * * This is a typical %queue operation. It shrinks the %queue * by one. The time complexity of the operation depends on the * underlying sequence. * * Note that no data is returned, and if the first element's * data is needed, it should be retrieved before pop() is * called. */ void pop() { __glibcxx_requires_nonempty(); std::pop_heap(c.begin(), c.end(), comp); c.pop_back(); }}8. rb_tree红黑树（英语：Red–black tree）是一种自平衡二叉查找树，红黑树是每个节点都带有颜色属性的二叉查找树，颜色为红色或黑色。在二叉查找树强制一般要求以外，对于任何有效的红黑树我们增加了如下的额外要求： 节点是红色或黑色。 根是黑色。 所有叶子都是黑色（叶子是NIL节点）。 每个红色节点必须有两个黑色的子节点。（或者说从每个叶子到根的所有路径上不能有两个连续的红色节点。）（或者说不存在两个相邻的红色节点，相邻指两个节点是父子关系。）（或者说红色节点的父节点和子节点均是黑色的。） 从任一节点到其每个叶子的所有简单路径都包含相同数目的黑色节点。8.1. set, multiset3std::set 是关联容器，含有 Key 类型对象的已排序集。用比较函数 比较 (Compare) 进行排序。搜索、移除和插入拥有对数复杂度。set 通常以红黑树实现。std::multiset 是含有 Key 类型对象有序集的容器。与 set 不同，它允许多个 Key 拥有等价的值。用关键比较函数 Compare 进行排序。搜索、插入和移除操作拥有对数复杂度。8.2. map, multimapstd::map 是有序键值对容器，它的元素的键是唯一的。用比较函数 Compare 排序键。搜索、移除和插入操作拥有对数复杂度。map 通常实现为红黑树。multimap 是关联容器，含有键值对的已排序列表，同时容许多个元素拥有同一键。按照应用到键的比较函数 Compare 排序。搜索、插入和移除操作拥有对数复杂度。拥有等价键的键值对的顺序就是插入顺序，且不会更改。9. hashtable散列表（Hash table，也叫哈希表），是根据键（Key）而直接访问在内存储存位置的数据结构。也就是说，它通过计算出一个键值的函数，将所需查询的数据映射到表中一个位置来让人访问，这加快了查找速度。这个映射函数称做散列函数，存放记录的数组称做散列表。9.1. unordered_set, unordered_multisetunordered_set 是含有 Key 类型唯一对象集合的关联容器。搜索、插入和移除拥有平均常数时间复杂度。在内部，元素并不以任何特别顺序排序，而是组织进桶中。元素被放进哪个桶完全依赖其值的哈希。这允许对单独元素的快速访问，因为哈希一旦确定，就准确指代元素被放入的桶。不可修改容器元素（即使通过非 const 迭代器），因为修改可能更改元素的哈希，并破坏容器。unordered_multiset 是关联容器，含有可能非唯一 Key 类型对象的集合。搜索、插入和移除拥有平均常数时间复杂度。不要求此容器的迭代顺序稳定。9.2. unordered_map, unordered_multimapunordered_map 是关联容器，含有带唯一键的键-值 pair 。搜索、插入和元素移除拥有平均常数时间复杂度。元素在内部不以任何特定顺序排序，而是组织进桶中。元素放进哪个桶完全依赖于其键的哈希。这允许对单独元素的快速访问，因为一旦计算哈希，则它准确指代元素所放进的桶。unordered_multimap 是无序关联容器，支持等价的键（一个 unordered_multimap 可含有每个键值的多个副本）和将键与另一类型的值关联。 unordered_multimap 类支持向前迭代器。搜索、插入和移除拥有平均常数时间复杂度。不要求此容器的迭代顺序稳定。参考转载自 STL 图解 深入剖析STL内存分配器 allocator 及其萃取器 \u0026#8617; c++ STL迭代器（iterator) \u0026#8617; C++ 参考手册 \u0026#8617;" },

  { "title": "C++11 新特性", "url": "/blog/2024/C++11%E6%96%B0%E7%89%B9%E6%80%A7/", "categories": "Cpp",
  "tags": "Cpp", "date": "2024-10-10", "content":
  "C++11 新特性" },

  { "title": "OpenCL Buffer Objects", "url": "/blog/2024/OpenCL-Buffer-Objects/", "categories": "OpenCL",
  "tags": "OpenCL, Cpp", "date": "2024-09-28", "content":
  "1. clCreateBuffer 分配内存创建 OpenCL 内存对象函数原型为：clCreateBuffer(cl_context, // 上下文 cl_mem_flags, // 内存对象的性质，见下表 size_t, // 内存对象数据块大小 void *, // host_ptr 主机数据内存地址（可以为空） cl_int *);针对不同场景需求，OpenCL提供了不同的内存对象创建标志位。1.1. CL_MEM_USE_HOST_PTRauto src_matrix_ptr = aligned_malloc\u0026lt;int, 4096\u0026gt;(kMatrixSize * kMatrixSize);// fill the matrix with data...cl_mem clsrc = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, cl_buff_size, src_matrix_ptr, NULL);由于是 host malloc 分配的内存，runtime 会分配一个相应的 buffer，kernel开始执行时，将从host_ptr拷贝到OpenCL的buffer中。应该避免使用该标志位。所谓开始执行时，是指NDRangeKernel创建的命令，在device上执行时，状态变为Ready的时候。针对与host共享物理内存的device，如果host_ptr已经是地址对齐的，那么runtime应该不用分配内存了。如果host_ptr没有内存对齐，则runtime将进行拷贝操作。（要考虑物理内存页要连续？？）。Use this when an aligned buffer already exists on the host side. It must be aligned to a 4096 byte boundary and be a multiple of 64 bytes or you don't actually get zero copy. Below clCreateBuffer takes a host address and returns a corresponding device address.1.2. CL_MEM_ALLOC_HOST_PTRcl_mem clsrc = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, cl_buff_size, NULL, NULL);int* map_cl_src = (int*)clEnqueueMapBuffer(queue, clsrc, CL_TRUE, CL_MAP_WRITE, 0, cl_buff_size, 0, NULL, NULL, \u0026amp;status);memcpy(map_cl_src, src_matrix_ptr, cl_buff_size);在`device`内存空间分配`buffer`。通过使用clEnqueueMapBuffer提供对host的访问入口。这是最佳的内存分配方式，避免runtime内存拷贝操作。1.3. CL_MEM_ALLOC_HOST_PTR | CL_MEM_COPY_HOST_PTR添加 CL_MEM_COPY_HOST_PTR 标志位，runtime 会将 host_ptr 的数据拷贝到 OpenCL 的 buffer 中。适用于`host_ptr`没有内存地址对齐。1.4. 场景选择 如果可能，尽量使用 CL_MEM_ALLOC_HOST_PTR 标志位，避免runtime内存拷贝操作。 如果host_ptr已经是地址对齐的，可以使用 CL_MEM_USE_HOST_PTR 标志位。 如果host_ptr没有内存地址对齐，可以使用 CL_MEM_ALLOC_HOST_PTR | CL_MEM_COPY_HOST_PTR 标志位。1.5. 参考资料 编程与调试 C++ – OpenCL \u0026amp; CUDA 初探 – 十三、内存问题探讨 [Intel Community – Why is it needed to use CL_MEM_ALLOC_HOST_PTR CL_MEM_COPY_HOST_PTR instead of just CL_MEM_COPY_HOST_PTR for Intel HD Graphics?](https://community.intel.com/t5/OpenCL-for-CPU/Why-is-it-needed-to-use-CL-MEM-ALLOC-HOST-PTR-CL-MEM-COPY-HOST/m-p/1070056) Arm Guide to OpenCL Programming – 8.3 Memory Allocation opencl-sdk_developer-guide-core-xeon_2018 Getting the Most from OpenCL™ 1.2: How to Increase Performance by Minimizing Buffer Copies on Intel® Processor Graphics" },

  { "title": "OpenCL 学习资源", "url": "/blog/2024/OpenCL-Learn-Resources/", "categories": "OpenCL",
  "tags": "OpenCL, Cpp", "date": "2024-09-26", "content":
  "pdf 文档 AMD_OpenCL_Programming_Optimization_Guide Arm Guide to OpenCL Programming NVIDIA_OpenCL_Best-Practices-Guide opencl-sdk_developer-guide-processor-graphics_2019.4 NVIDIA OpenCL SDK Code Samples更多待整理 知乎 – OpenCL学习资料整理SDK AMD-APP-SDKInstaller" },

  { "title": "OpenCL 同步操作", "url": "/blog/2024/OpenCL-Synchronization/", "categories": "OpenCL",
  "tags": "OpenCL, Cpp", "date": "2024-09-25", "content":
  "1. Barrier1.1 clEnqueueBarrierWithWaitList// Provided by CL_VERSION_1_2cl_int clEnqueueBarrierWithWaitList( cl_command_queue command_queue, cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event);cl_int clWaitForEvent( cl_uint num_events, const cl_event *event_list);用于在OpenCL命令队列中插入一个同步点。其作用对象限于一个 command queue。 如果 event_wait_list 为空，则需要该同步点命令(clEnqueueBarrierWithWaitList)之前的命令全部执行完成，才能执行其之后的命令。 如果 event_wait_list 不为空，则需要等到所有事件都完成（CL_COMPLETE），才能执行其之后的命令。如果 event 不为空，可以用这个 event 阻塞host直到该命令执行完成(CL_COMPLETE)。1.2 work-group barrier作用于同一个 work-group 内的所有 work-item。void work_group_barrier(cl_mem_fence_flags flags);void work_group_barrier(cl_mem_fence_flags flags, memory_scope scope);flags 含义： CLK_LOCAL_MEM_FENCE：The barrier function will either flush any variables stored in local memory or queue a memory fence to ensure correct ordering of memory operations to local memory. CLK_GLOBAL_MEM_FENCE：The barrier function will queue a memory fence to ensure correct ordering of memory operations to global memory. This can be useful when work-items, for example, write to buffer or image objects and then want to read the updated data.参考： Barriers in OpenCL OpenCL 1.2 man – Barrier2. 同步点 – synchronization points2.1 clFinishcl_int clFinish(cl_command_queue command_queue);clFinish 也会在命令队列里面添加一个同步点，与 clEnqueueBarrierWithWaitList 不同的是，clFinish 对 host 是阻塞的。2.2 eventevent 也是一种同步点： 用于同步设备与 host，或者 同一个 context 创建的不同 command queue 之间的同步。所有 clEnqueueXXXX 函数，都可以设置 event 列表。 event 可以适用于同一个 context(或者 shared context) 创建的不同 command queue 之间的同步； event 效率应该比较慢，因为要经过 host （个人理解）；2.2.1 场景：不同 command queue 之间的同步针对 out-of-order 命令队列里面的同步，使用 event 对象特别合适，执行如下操作并同步：cl_event k_events[2]{};err = clEnqueueNDRangeKernel(commands, kernel1, 1, NULL, \u0026amp;global, \u0026amp;local, 0, NULL, \u0026amp;k_events[0]);err = clEnqueueNDRangeKernel(commands, kernel2, 1, NULL, \u0026amp;global, \u0026amp;local, 0, NULL, \u0026amp;k_events[1]);err = clEnqueueNDRangeKernel(commands, kernel3, 1, NULL, \u0026amp;global, \u0026amp;local, 2, k_events, NULL);2.2.2 场景：device 代码等待 host 发送event host 创建 event 对象使用 clCreateUserEvent 创建 event 对象：cl_event clCreateUserEvent(cl_context context, cl_int*errcode_ret); host在合适的时候设置 event 状态cl_int clSetUserEventStatus(cl_event event, cl_intexecution_status);2.3 其他阻塞操作其他几个 clEnqueueXXXX 命令，如果将同步参数设置为 CL_TRUE，则为阻塞执行 clEnqueueReadBuffer，clEnqueueReadBufferRect，clEnqueueReadImage clEnqueueWriteBuffer，clEnqueueWriteBufferRect，clEnqueueWriteImage clEnqueueSVMMemcpy，clEnqueueSVMMap clEnqueueMapBuffer其他阻塞命令： clReleaseCommandQueue clWaitForEvents 所有 阻塞命令 都会调用 clFlush，将命令队列中的命令提交(submit)到设备。 可以使用 event 对象实现不同命令队列之间的同步，不过需要显式/隐式调用 clFlush 提交命令到对应的命令队列/设备。 一个命令队列只与一个设备关联。3. 参考 OpenCL API – clEnqueueBarrierWithWaitList OpenCL API – Flush and Finish OpenCL: A Hands-on Introduction4. 更多资料 OpenCL 3.0 man pages OpenCL reference Guide Intel Tools for OpenCL™ Applications CUDA Toolkit 4.0论文资源： Google Scholar IEEE Xplore ACM Digital Library" },

  { "title": "OpenCL 平台模型、执行模型", "url": "/blog/2024/OpenCL-model/", "categories": "OpenCL",
  "tags": "OpenCL, Cpp", "date": "2024-09-20", "content":
  "1. 平台模型关键词： OpenCL Device CU – Compute Unit PE – Processing Element 2. 内存模型 github – OpenCL Guide –Memory Model3. 执行模型3.1 ContextContext 是针对Host端编程而产生的概念，表示设备的执行环境，包含： Devices：一个或多个OpenCL物理设备； Memory Objects：Host 端和/或 Device 端可见的内存对象； Program Objects: 包含源码及编译后的二进制代码； kernel Objects：Device 端执行的函数对象；3.1.1 命令队列 Command Queue一个Command Queue 对应一个Device。一个 Command Queue 中的命令包含如下三种类型： Kernel 相关命令：执行 Kernel 函数； Memory 相关命令： host \u0026lt;–\u0026gt; device 数据传输； host \u0026lt;–\u0026gt; device memory map / unmap； Memory Objects 之间数据传输； 同步相关命令；除了主机端往命令队列中添加命令外，在设备端，Kernel执行的时候，也可以往设备端的命令队列中添加命令，比如启动 Child kernel。如下图中的Ended表示所有该命令中的所有Work Group执行完毕，但可能Child kernel还没有执行完毕，以及更新global memory中的数据。关于每个状态的解释，参考 OpenCL 3.0 Spec – 3.2. Execution Model Queued：初始状态； Submitted：提交到Device，还没还有放入设备端的 work pool – 比如需要的资源没有准备好，或者work pool 满； Ready：命令提交到work pool，等待被调度； Running：已经被调度器调度到 CU 开始执行； Ended：所有 work group 执行完毕； Complete：Child kernel 执行完毕，global memory 中的数据更新完毕； 个人理解：一个Command Queue 同时存在于主机端，以及设备端 ？？3.2 NDRange – 索引空间表示一维 / 二维 / 三维索引空间：global index, group index, local index。OpenCL 软件调度将全局 work items 按 group 为单位，分配给 CU（一个 CU 包含多个 PE）。CU执行完当前 group 后，再调度下一个 group 到 CU 上执行。3.3 Work-Item index 关系划分好 work group size 之后，可以相互换算global index 和 local index，以及 group index。例如matrix大小为 $G_{x}$ x $G_{y}$，将其划分为 $W_{x}$ x $W_{y}$ 个工作组， 每个工作组的大小为 $L_{x}$ x $L_{y}$，则：\\[\\begin{cases}L_{x} = G_{x} / W_{x} \\\\L_{y} = G_{y} / W_{y}\\end{cases}\\]根据工作项ID（$l_x$, $l_y$）可以计算出全局ID（$g_x$, $g_y$）：\\[\\begin{cases}g_x = w_x * L_{x} + l_x \\\\g_y = w_y * L_{y} + l_y\\end{cases}\\]相反的，根据全局ID，计算出工作项ID，以及工作组ID，只需要分别进行取余、除法运算即可。4. 编程模型4.1 编程模型4.2 编程流程5. 参考资料 一文说清OpenCL框架 OpenCL 平台模型 - 执行模型 - 内存模型 - 编程模型 OpenCL 3.0 Spec6. 附加资料 Emulating Command Buffer Extensions with OpenCL Layers github – Intercept Layer for OpenCLTM Applications" },

  { "title": "Git加速资源", "url": "/blog/2024/git%E5%8A%A0%E9%80%9F%E8%B5%84%E6%BA%90/", "categories": "Git",
  "tags": "Git", "date": "2024-09-14", "content":
  "GitHub Proxy 加速 GitHub Proxy GitHub 加速链接转换使用清华镜像加速：git config --global url.\"https://mirrors.tuna.tsinghua.edu.cn/git/\".insteadOf https://github.com/使用GitClone加速：# 方法一(替换URL)git clone https://gitclone.com/github.com/tendermint/tendermint.git# 方法二(设置git参数)git config --global url.\"https://gitclone.com/\".insteadOf https://git clone https://github.com/tendermint/tendermint.git" },

  { "title": "总结：git 不常用命令", "url": "/blog/2024/git-commands/", "categories": "Git",
  "tags": "Git", "date": "2024-09-14", "content":
  "# 远程分支与本地分支有不相关的提交，合并远程分支git pull origin main --allow-unrelated-histories# 删除远程分之git push origin --delete \u0026lt;branch_name\u0026gt;# 更新远程分支列表git remote update origin --prune# 删除submodule# Remove the submodule entry from .git/configgit submodule deinit -f .\\3rd\\xz-v5.8.1\\# Remove the submodule directory from the superproject's .git/modules directoryrm -rf .git/modules/3rd/xz-v5.8.1# Remove the entry in .gitmodules and remove the submodule directory located at path/to/submodulegit rm -f 3rd/xz-v5.8.1删除git记录中的大文件及blob# 执行命令之前，保证仓库没有待提交的更改# 清理git gc# 根据文件名查找大文件的完整路径git rev-list --objects --all | grep \u0026lt;filename\u0026gt;# 使用完整大文件路径删除大文件git filter-branch --force --index-filter 'git rm -rf --cached --ignore-unmatch \u0026lt;full_path_filename\u0026gt;' --prune-empty --tag-name-filter cat -- --all# 列出前5个大文件# git rev-list --objects --all | grep \"$(git verify-pack -v .git/objects/pack/*.idx | sort -k 3 -n | tail -5 | awk '{print$1}')\"# 推送到远程仓库git push origin --force --all# 清除缓存rm -rf .git/refs/original/git reflog expire --expire=now --allgit gc --prune=now# 查看当前Git对象统计信息git count-objects -vgitlab无法强行推送问题解决需要owner设置如下：仓库的Settings -\u0026gt; Repository -\u0026gt; Protected branches 改成unprotected。" },

  { "title": "OpenCL 端编程流程及主要概念实践", "url": "/blog/2024/OpenCL-basic-coding-flow/", "categories": "OpenCL",
  "tags": "OpenCL, Cpp", "date": "2024-09-09", "content":
  "0. OpenCL 概念 平台 platform：OpenCL 实现的顶层容器，通常对应于一个 OpenCL 的实现厂商； 设备 device：执行 OpenCL 程序的硬件设备，可以是 CPU、GPU、FPGA，或其他计算加速设备； 上下文 context：管理设备和资源的的环境，一个上下文可以包括多个 device； 命令队列 command queue：向设备发送命令的队列，一个命令队列与一个给定的 device 相关联； 程序 program：CL 代码及其编译后的二进制，包含一个或多个 kernel； 内核 kernel：在设备上执行的函数，这是 OpenCL 程序的核心； 工作项 work item：kernel 执行的一个实例，类似于线程； 工作组 work group：工作项的集合，集合内的 work item 共享一个 Local Memory，以及进行同步；1. 编程流程编程步骤如下：一个示例源码：opencl_002_array_add2. OpenCL 内存模型 kernel 函数中，使用关键字 __global 标示的变量，存储在上图中的 Global Memory 中；__local 标示的变量，存储在 Local Memory 中。 OpenCL 也分 WorkGroup，使用__local修饰的变量，存储在Local Memory中，仅限于同一个 WorkGroup 中的 Work Item 可以共享访问该变量。 kernel 函数中定义的变量，存储在 Private Memory 中，仅限于 Work Item 内可访问。 kernel 函数也可以使用值传参，以及 指针传参 ，一般不推荐值传参。3. 概念解释：work group、work item 与 设置 index类似于 CUDA 中的 warp 概念以及 thread 概念，OpenCL 中也有 get_global_id() 和 get_local_id() 这两个函数，用来获取当前 work item 的全局和局部索引，用于表示当前任务的index。使用 clEnqueueNDRangeKernel 时，需要设置维度参数，函数原型如下：cl_int clEnqueueNDRangeKernel( cl_command_queue command_queue, // 命令队列 cl_kernel kernel, // 要执行的内核 cl_uint work_dim, // 工作维度，范围是1到3 const size_t *global_work_offset, // 全局工作项的偏移 const size_t *global_work_size, // 全局工作项的大小 const size_t *local_work_size, // 局部工作项的大小 cl_uint num_events_in_wait_list, // 依赖的事件数量 const cl_event *event_wait_list, // 依赖事件的列表 cl_event *event // 返回的事件); work_dim: 工作维度，表示 kernel 函数的执行次数，可以是 1, 2,3； global_work_offset: 全局工作项的偏移量，可以设为 NULL，表示从 (0,0,0) 开始； global_work_size: 全局工作项的大小，例如对于一个 1024x1024 的矩阵/图像，设置为 (1024, 1024)； local_work_size：指定每个 work group 分配的 work item 数量；例如如下代码设置 global_work_size，local_work_size：size_t global_work_size[2] = {1024, 1024}; // 1024x1024 的全局工作区size_t local_work_size[2] = {16, 16}; // 16x16 的局部工作区// 启动 kernelcl_int err = clEnqueueNDRangeKernel( queue, kernel, 2, // 2 维 NULL, // 全局偏移量设为 NULL global_work_size,// 全局工作区大小 local_work_size, // 局部工作区大小 0, // 没有依赖的事件 NULL, // 没有依赖的事件列表 NULL // 不需要返回的事件句柄);3.1 例子：矩阵转置work item index 演示代码，测试代码 003_opengl_matrix_transpose。kernel 部分：__kernel void matrixTransposeSimple(__global float* input, __global float* output, const uint width, const uint height) { uint gdx = get_global_id(0); uint gdy = get_global_id(1); output[gdy * width + gdx] = input[gdx * height + gdy];}work item index 演示代码 – C++部分代码：// 5. 准备数据，并创建 cl buffers Eigen::MatrixXf dst_matrix = Eigen::MatrixXf::Zero(kMatrixSize, kMatrixSize); Eigen::MatrixXf src_matrix = Eigen::MatrixXf::Random(kMatrixSize, kMatrixSize); auto src_matrix_ptr = src_matrix.data(), dst_matrix_ptr = dst_matrix.data(); const size_t cl_buff_size = kMatrixSize * kMatrixSize * sizeof(cl_float); cl_mem clsrc = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, cl_buff_size, src_matrix_ptr, NULL); cl_mem cldst = clCreateBuffer(context, CL_MEM_READ_WRITE, cl_buff_size, NULL, NULL); // 6. 设置 kernel 参数，并执行 kernel cl_int dimx = kMatrixSize, dimy = kMatrixSize; const auto err1 = clSetKernelArg(kernel, 0, sizeof(cl_mem), \u0026amp;clsrc); // param 0: source matrix const auto err2 = clSetKernelArg(kernel, 1, sizeof(cl_mem), \u0026amp;cldst); // param 1: destination matrix const auto err3 = clSetKernelArg(kernel, 2, sizeof(cl_int), \u0026amp;dimx); // param 2: width const auto err4 = clSetKernelArg(kernel, 3, sizeof(cl_int), \u0026amp;dimy); // param 3: height size_t global_work_size[] = {kMatrixSize, kMatrixSize}, local_work_size[] = {16, 16}; const auto err6 = clEnqueueNDRangeKernel(queue, kernel, 2, // 0, global_work_size, local_work_size, // 0, 0, 0); const auto err7 = clFinish(queue);3.2 如何设置 local_work_size在kernel中，有如下函数，分别获取指定维度上的local index，group index，global index：size_t get_global_id(uint D); // 获取全局索引，D=0,1,2size_t get_local_id(uint D); // 获取局部索引, D=0,1,2size_t get_group_id(uint D); // 获取组索引，D=0,1,2使用CL_KERNEL_WORK_GROUP_SIZE获取work group的最大尺寸：size_t max_work_group_size{}, max_work_group_size2{};clGetKernelWorkGroupInfo(kernel, device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), \u0026amp;max_work_group_size, NULL);clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(size_t), \u0026amp;max_work_group_size2, NULL);SPDLOG_INFO(\"Max work group size: {} / {}\", max_work_group_size, max_work_group_size2); // 输出 256 / 256在调用clEnqueueNDRangeKernel时，参数local_work_size设置为null，kernel将自动选择合适的local work size。FIXME：使用clinf 查看 Intel UHD Graphics 620信息，显示推荐的local work size为 8, 16, 32。 Affect of local_work_size on performance and why it is4. 参考及资料 OpenCL矩阵转置 OpenCL Matrix Transpose opencl(十八)—-矩阵转置、矩阵乘法 Some Basic Usage of Eigen with C++ 编程与调试 C++ – OpenCL \u0026amp; CUDA 初探 Intel – opencl-sdk_developer-guide-processor-graphics_2019.4 Eigenvalues and Eigenvectors" },

  { "title": "使用 FFmpeg 从视频中提取音频", "url": "/blog/2024/get-audio-from-vidio-using-ffmpeg/", "categories": "Tools",
  "tags": "Tools, FFmpeg", "date": "2024-09-08", "content":
  "1. 提取完整音频ffmpeg -i sample.mp4 -q:a 0 -map a sample.mp32. 提取特定时段的音频ffmpeg -i sample.mp4 -ss 00:03:05 -t 00:00:45.0 -q:a 0 -map a sample.mp3 ss 选项指定开始时间戳，使用 t 选项指定编码持续时间，例如从3分钟到5秒钟，持续45秒。 时间戳必须采用 HH：MM：SS.xxx 格式或以秒为单位。 如果你不指定 t 选项，它将会结束。 FFmpeg 提取视频的音频3. 直接提取音频流ffmpeg -i input.flv -vn -codec copy out.m4a其中，-i的意思是input，后接输入源。-codec的意思是直接复制流。 利用FFmpeg无损提取视频中源音频流" },

  { "title": "Windows 环境编译 VTK", "url": "/blog/2024/building-vtk-under-windows/", "categories": "Cpp",
  "tags": "Cpp, CAD, VTK", "date": "2024-09-05", "content":
  "1. 依赖项从OpenCascade官网下载编译好的包，提取里面的依赖包。下载地址：OpenCascade批量编译 VTK – 同时编译Debug/Release版本使用CMake生成完成之后，打开Visual Studio进行编译。选择生成 -\u0026gt; 批生成，选取如下 Debug Install、Release Install即可，并开始编译。编译时间较长。" },

  { "title": "spdlog 使用", "url": "/blog/2024/spdlog-usage/", "categories": "Cpp",
  "tags": "Cpp", "date": "2024-09-04", "content":
  "1. 使用时的编译选项if(CMAKE_BUILD_TYPE STREQUAL \"Release\") add_compile_definitions(-DSPDLOG_ACTIVE_LEVEL=SPDLOG_LEVEL_WARN) add_compile_definitions(-DNDEBUG)else() add_compile_definitions(-DSPDLOG_ACTIVE_LEVEL=SPDLOG_LEVEL_TRACE)endif()2. 日志级别，日志格式#include \u0026lt;spdlog/spdlog.h\u0026gt;#ifdef NDEBUG spdlog::set_level(spdlog::level::warn); // disable spdlog for performance test#else spdlog::set_level(spdlog::level::info); spdlog::set_pattern(\"%H:%M:%S.%e %t %s %! %v\");#endif// back to default format// spdlog::set_pattern(\"%+\");// alignment: 左对齐, 右对齐spdlog::info(\"{:\u0026gt;8} aligned, {:\u0026lt;8} aligned\", \"right\", \"left\");3. 使用技巧3.1 打印 std::vector#include \"spdlog/spdlog.h\"#include \"spdlog/fmt/bundled/ranges.h\"logger-\u0026gt;info (\"vector: {}\", fmt::join(vec, \", \"));SPDLOG_INFO(\"vector: {}\", vec);" },

  { "title": "Mesh 及相关开源仓库收集", "url": "/blog/2024/mesh-tools/", "categories": "CAD",
  "tags": "CAD", "date": "2024-09-04", "content":
  "meshlab Gmsh github – netgen github –carve : A fast, robust constructive solid geometry library. cgal : The Computational Geometry Algorithms Library. github – Clipper2 : An open source freeware library for clipping and offsetting lines and polygons. github – cock : A 3D boolean/CSG library. draco : An open-source library for compressing and decompressing 3D geometric meshes and point clouds geogram : A programming library of geometric algorithms github – libigl : A simple C++ geometry processing library MMG : Robust, open source \u0026amp; multidisciplinary software for remeshing. Qhull : Engine for convex hulls, Delaunay triangulations, Voronoi diagrams computations. Quartet : A tetrahedral mesh generator that does isosurface stuffing with an acute tetrahedral tile. TetGen : Tetrahedral mesh generation engine. Triangle : A two-Dimensional quality mesh generator and Delaunay triangulator. Geometric Tools CFD中文网" },

  { "title": "总结：using 几种使用场景", "url": "/blog/2024/keywork-using-of-cpp/", "categories": "Cpp",
  "tags": "Cpp", "date": "2024-09-04", "content":
  "1. using 声明 (using declaration)将命名空间中的某一个名字 (变量或函数) 引入到当前作用域中，使得当前作用域访问该名字，不需要使用命名空间，以及全局限定符::。{ using std::map; map\u0026lt;int, std::string\u0026gt; the_map; //ok}map\u0026lt;int, std::string\u0026gt; the_map2; //error1.1 使用 using 声明，子类可以使用父类中的私有(private)成员class Base{protected: int bn1; int bn2;};class Derived: private Base{public: using Base::bn1;};class DerivedAgain: public Derived{};int main() { Derived d; d.bn1 = 1; // ok d.bn2 = 2; // error, 'bn2' is a private member of 'Base' DerivedAgain da; da.bn1 = 3; // ok return 0;}2. 其他使用场景 引入命名空间(using directive)； 别名 alias，如: using alias_class_t = myClass；2.1 template + using 使用方法template\u0026lt;typename Val\u0026gt;using int_map_t = std::map\u0026lt;int, Val\u0026gt;;int main() { int_map_t\u0026lt;int\u0026gt; imap; return 0;}" },

  { "title": "复习：RVO NRVO and std::move", "url": "/blog/2024/RVO-NRVO-and-std_move/", "categories": "Cpp",
  "tags": "Cpp", "date": "2024-09-04", "content":
  "RVO 以及 NRVO RVO：返回值优化，其功能为：消除子函数返回的临时对象导致的拷贝构造。 NRVO：其功能为：消除子函数中返回的局部对象导致的拷贝构造。class Object { //...};Object getObjRVO() { return Object(); // RVO}Object getObjNRVO() { Object localObj; return localObj; // NRVO}std::move使用场景： 局部对象赋值给长生命周期对象时，使用std::move。前提是该类里面有非trival成员，如std::string，或支持移动构造的自定义类成员； std::vector等容器使用emplace_back代替push_back，此时针对局部对象使用std::move + emplace_back可以避免拷贝构造 – 代之以移动构造； std::thread线程不可复制，只能所有权转移。如将线程对象添加到std::vector中，则需要使用std::move转移所有权。不要使用std::move： 对于返回值优化的函数，不要使用std::move。" },

  { "title": "CRTP：使用笔记", "url": "/blog/2024/CRTP-introduce/", "categories": "Cpp",
  "tags": "Cpp, Template", "date": "2024-09-02", "content":
  "1. 动态继承运行时时间损耗 每个 virtual 方法，都需要通过指针查找到虚函数入口（间接寻址），且可能引起I-Cache cache miss； virtual 方法，不能被优化为inline，针对一些短小的函数，时间损耗较大；2. CRTP 使用举例：子类也是 template 模板类template \u0026lt;typename DerivedT\u0026gt;class HoleDetectorBase { public: using underlying_type = HoleDetectorBase\u0026lt;DerivedT\u0026gt;; //...};template \u0026lt;typename condPairFuncT, typename condGHIPairFuncT\u0026gt;class HoleDetector2 : public HoleDetectorBase\u0026lt;HoleDetector2\u0026lt;condPairFuncT, condGHIPairFuncT\u0026gt;\u0026gt; { public: using base_t = HoleDetectorBase\u0026lt;HoleDetector2\u0026lt;condPairFuncT, condGHIPairFuncT\u0026gt;\u0026gt;; using typename base_t::underlying_type; //...}; 继承写法为：HoleDetector2\u0026lt;condPairFuncT, condGHIPairFuncT\u0026gt;，即带上template参数。2.1 添加在基类中封装子类的函数template \u0026lt;typename DerivedT\u0026gt;class HoleDetectorBase {//...DerivedT\u0026amp; underlying() { return static_cast\u0026lt;DerivedT\u0026amp;\u0026gt;(*this); }DerivedT const\u0026amp; underlying() const { return static_cast\u0026lt;DerivedT const\u0026amp;\u0026gt;(*this); }//...}; Note：基类中不能引用子类中的子类型，因为对父类而言，子类是incomplete type。3. 子类中使用 CRTP 父类的成员方法子类中使用父类的成员函数，需要加上base_t::前缀，或者this指针。auto groups = base_t::group_holes(valid_holes);4. 参考资料 unique_ptr 与抽象类的多态 An Implementation Helper For The Curiously Recurring Template Pattern The Curiously Recurring Template Pattern in C++ The cost of dynamic (virtual calls) vs. static (CRTP) dispatch in C++ C++CRTP概念与应用和concept" },

  { "title": "Intel TBB 并行计算", "url": "/blog/2024/intel-tbb-parallel/", "categories": "TBB, Cpp",
  "tags": "TBB, Cpp, Performance, HPC", "date": "2024-08-28", "content":
  "1. TBB 简介Intel TBB主要功能模块： 并行算法 任务调度 并行容器 同步原语 内存分配器1.1. 并行算法 parallel_for parallel_reduce parallel_scan parallel_do parallel_sort parallel_invoke pipeline, parallel_pipeline1.2. 并行容器 concurrent_vector concurrent_hash_map concurrent_queue1.3. 编译及链接参考之前文档Intel TBB malloc 使用 (windows)(2024-08-13)。2. 并行计算头文件包含：#include \u0026lt;oneapi/tbb/parallel_sort.h\u0026gt;#include \u0026lt;tbb/parallel_for.h\u0026gt;#include \u0026lt;tbb/tbb.h\u0026gt;2.1 sortstd::vector\u0026lt;dataxy_info_t\u0026gt; data_xy_info;//...tbb::parallel_sort(data_xy_info.begin(), data_xy_info.end(), [](const dataxy_info_t\u0026amp; a, const dataxy_info_t\u0026amp; b) { return (a.x_ \u0026lt; b.x_); });对150万数据排序，时间对比：[2024-08-28 12:44:37.924] [info] 02. Sorted 1503894 samples in 80.45 ms (TBB sort)[2024-08-28 12:44:42.610] [info] 11. Sorted 1503894 samples in 206.20 ms (std sort)2.2 parallel_for_each// void parse_dataxy_info(dataxy_info_t\u0026amp; info);tbb::parallel_for_each(data_xy_info.begin(), data_xy_info.end(), parse_dataxy_info);2.3. 任务分块：blocked_range 以及 partitioner分块是通过分区（partitioner）和粒度（grainsize）控制的。Intel TBB定义了几种partitioner，即不同的分区策略，不同分区策略对grainsize的控制有所不同。默认使用auto_partitioner。blocked_range的构造函数形式为：blocked_range\u0026lt;T\u0026gt;(begin,end,grainsize)。结合使用blocked_range以及partitioner，如下是不同分区策略的一些区别：使用auto_partitioner以及affinity_partitioner，会根据执行情况调整每个线程的负载，可能会有一些额外开销，static_partitioner则不会在执行过程中调整线程的负载，可能会有一些负载不均的情况。2.3.1. affinity_partitioneraffinity_partitioner除了自动选择grainsize（粒度）之外，还会记录每个线程处理的数据块，以便在后续的parallel_for等并行算法中优先分配给同一线程，减少数据迁移，提高缓存命中率。对于内存数据需要重复访问的情况，affinity_partitioner尤其有用。如下图示中，上半部分显示一个线程加载并处理一块数据，以及对应的Cache区域。下半部分示意两个线程处理的内存区域有重叠，即Cache区域有重叠。如果内存重叠区域的线程都有回写的需求，可能会出现False Sharing或者叫Cache Coherence Traffic，导致线程1的Cache被线程2逐出导致Cache Miss。相关官方文档： Bandwidth and Cache Affinity Partitioner Summary2.3.2. 示例：基于分块的 parallel_for 如果需要使用到任务共享写变量，需要添加锁，或者使用原子变量，如下面示例中使用到的原子变量。TBB不保证线程安全。{ std::atomic_int overflow_acc{}, underflow_acc{}; auto statistic_abnormal_func = [\u0026amp;](const sample_data_t\u0026amp; sample) { // func body: 统计异常数据 if (sample.mapped_pos_.y_ \u0026lt; 0) { underflow_acc++; } else if (sample.mapped_pos_.y_ \u0026gt; 255) { overflow_acc++; } }; tbb::parallel_for( tbb::blocked_range\u0026lt;size_t\u0026gt;(0, ds.ds_.size()), [\u0026amp;](const tbb::blocked_range\u0026lt;size_t\u0026gt;\u0026amp; r) { for (size_t i = r.begin(); i != r.end(); i++) { statistic_abnormal_func(ds.ds_[i]); } }, tbb::auto_partitioner()); if (overflow_acc \u0026gt; 0 || underflow_acc \u0026gt; 0) { //SPDLOG_WARN(\"samples overflow {}. underflow {} in dataset\", overflow_acc.load(), underflow_acc.load()); samples_overflow_acc2 += overflow_acc.load(); samples_underflow_acc2 += underflow_acc.load(); }}2.3.3. 示例：基于分块的 parallel_reduce归并主要用于查找最大值，合并计算（如求和）等场景。其原理如下：参考资料：intel-TBB使用笔记3. 任务调度及线程池针对具体并行任务，设定并行度，使用task_arena。需要设置全局并行度的，使用global_control。3.1 task_arena#include \u0026lt;oneapi/tbb/info.h\u0026gt;#include \u0026lt;oneapi/tbb/parallel_for.h\u0026gt;#include \u0026lt;oneapi/tbb/task_arena.h\u0026gt;#include \u0026lt;cassert\u0026gt;oneapi::tbb::task_arena arena(4); // 4 threadsarena.execute([] { oneapi::tbb::parallel_for( /* ... */ [] { // This arena is limited with for threads assert(oneapi::tbb::this_task_arena::max_concurrency() == 4); });});3.2 global_control#include \u0026lt;oneapi/tbb/info.h\u0026gt;#include \u0026lt;oneapi/tbb/task_arena.h\u0026gt;#include \u0026lt;oneapi/tbb/global_control.h\u0026gt;#include \u0026lt;cassert\u0026gt;oneapi::tbb::global_control global_limit(oneapi::tbb::global_control::max_allowed_parallelism, 2);// the default parallelismoneapi::tbb::parallel_for( /* ... */ [] { // No more than two threads is expected; however, tbb::this_task_arena::max_concurrency() can return a bigger value int thread_limit = oneapi::tbb::global_control::active_value(oneapi::tbb::global_control::max_allowed_parallelism); assert(thread_limit == 2);});3.3 设置并行线程 Stack Size#include \u0026lt;oneapi/tbb/global_control.h\u0026gt;oneapi::tbb::global_control global_limit(tbb::global_control::thread_stack_size, 16 * 1024 * 1024); Migrating from tbb::task_scheduler_init4. TBB原理：任务划分TBB采用递归分块的方式划分Task，直到任务足够小（小于grainsize）时才执行。如下图所示：在执行时，按照深度优先的顺序执行Task，以获取最大的局部性和缓存效率。另一方面在，在负载均衡调整的过程中，Work Steal则采样广度优先的方式，以避免与被Steal的线程Cache冲突。如下图所示：资料：PPT–Parallel Programming with Intel® Threading Building BlocksA. 资料 TBB并发库代码学习：博客，介绍了一些内部实现 Intel Thread Building Blocks (TBB)：中文博客，其中任务调度讲解比较好 Intel CommunityA.1. 官方文档 oneTBB Developer Guide Pro TBB Pro TBB：Pro TBB 代码示例 Today’s TBB pdf Source Code" },

  { "title": "Intel TBB malloc 内存分配器", "url": "/blog/2024/intel-tbb-malloc/", "categories": "TBB, Cpp",
  "tags": "TBB, Cpp, Performance, HPC", "date": "2024-08-28", "content":
  "1. TBB Malloc 使用入门有两种方式使用TBB Malloc：Run-Time替换，Link-Time替换。替换的函数(routines)包括： routines Linux MacOS Windows global C++ new / delete √ √ √ C库：malloc / calloc / realloc / free √ √ √ C库(C11)：aligned_alloc √ - - POSIX：posix_memalign √ √ - Run-Time替换方法： 平台 替换方法 Linux export LD_PRELOAD=$TBBROOT/lib/intel64/gcc4.8/libtbbmalloc_proxy.so.2 MacOS export DYLD_INSERT_LIBRARIES=$TBBROOT/lib/intel64/gcc4.8/libtbbmalloc_proxy.dylib Link-Time替换方法： 平台 替换方法 Linux \u0026amp; MacOS -L$TBBROOT/lib/intel64/gcc4.8 -ltbbmalloc_proxy Windows tbbmalloc_proxy.lib /INCLUDE:\"__TBB_malloc_proxy\" Link-Time替换，需要添加编译flags。不添加这些编译flags，可能导致malloc等这些函数被内联为汇编，即失去了符号，也就没有办法替换了。需要添加的flags如下：Linux/MacOS平台下，添加如下编译flags（适用于Linux/MacOS）：-fno-builtin-malloc-fno-builtin-calloc-fno-builtin-realloc-fno-builtin-freeWindows平台下，icc编译器添加如下编译flags：- /Qfno-builtin-malloc- /Qfno-builtin-calloc- /Qfno-builtin-realloc- /Qfno-builtin-free 替换Proxy入口代码文件为src/tbbmalloc_proxy/proxy.cpp。1.1. 直接使用TBB malloc（Compile-Time） Allocation Routine Deallocation Routine 对应系统运行时库 scalable_malloc, scalable_calloc, scalable_realloc scalable_free C Standard library scalable_aligned_malloc, scalable_aligned_realloc scalable_aligned_free Microsoft C runtime 与C++ STL对应的allocator：一个简单使用示例：#include \u0026lt;vector\u0026gt;#include \u0026lt;algorithm\u0026gt;#include \u0026lt;execution\u0026gt;#include \u0026lt;tbb/scalable_allocator.h\u0026gt;std::vector\u0026lt;int, tbb::scalable_allocator\u0026lt;int\u0026gt;\u0026gt; v{…};std::sort(std::execution::par, v.begin(), v.end());1.2. Huge PagesTBB malloc支持Huge Pages，可以通过设置环境变量TBB_MALLOC_USE_HUGE_PAGES来启用，或者在代码中设置：scalable_allocation_mode( TBBMALLOC_USE_HUGE_PAGES,1);Huage Pages可以减少malloc调用，减少TLB Miss，提高性能，尤其是在分配大块内存时。1.3. memory_pool_allocatorTBB malloc提供了一个memory_pool_allocator，它通过预先分配一大块内存，避免TBB Malloc模块的管理开销。管理器需要查找空闲块、分割块、记录元数据（这块内存多大、是否在用等）。其约束为，每次分配的大小为固定的P，因为它是通过简单的基于 P 取模的偏移量获取内存块的。适用于作为STL容器的分配器。一个示例如下：#define TBB_PREVIEW_MEMORY_POOL 1#include \"oneapi/tbb/memory_pool.h\"#include \u0026lt;list\u0026gt;int main() { oneapi::tbb::memory_pool\u0026lt;std::allocator\u0026lt;int\u0026gt;\u0026gt; my_pool; typedef oneapi::tbb::memory_pool_allocator\u0026lt;int\u0026gt; pool_allocator_t; std::list\u0026lt;int, pool_allocator_t\u0026gt; my_list(pool_allocator_t{my_pool}); my_list.emplace_back(1);}当memory_pool中内存用尽之后，将向底层Alloc申请内存进行扩容。新增的内存切分成若干FreeBlock，并添加到pool内部的空闲链表。扩容大小由extMemPool-\u0026gt;granularity决定，扩容的内存可能来自其他线程丢弃的的orphaned blocks，也可能来自底层Alloc分配的内存。其类声明如下：//! Thread-safe growable pool allocator for variable-size requeststemplate \u0026lt;typename Alloc\u0026gt;class memory_pool : public pool_base另外，定义了一个fixed_pool，内存池耗尽之后不会扩展。1.4. 局部替代 new \u0026amp; delete当某些模块需要自定义allocator时，可以通过局部替代new和delete来实现：#include \u0026lt;tbb/parallel_for.h\u0026gt;#include \u0026lt;tbb/tbb_allocator.h\u0026gt;// No retry loop because we assume that// scalable_malloc does all it takes to allocate the memory,// so calling it repeatedly will not improve the situation at all// No use of std::new handler because it cannot bedone in portable and// thread-safe way We throw std::bad alloc() when scalable mallocreturns NULL// (we return NULL if it is a no-throw implementation)void *operator new(size_t size) throw(std::bad_alloc) { if (size == 0) size = 1; if (void *ptr = scalable_malloc(size)) return ptr; throw std::bad_alloc();}void *operator new[](size_t size) throw(std::bad_alloc) { return operator new(size);}void *operator new(size_t size, const std::nothrow_t \u0026amp;) throw { if (size == 0) size = 1; if (void *ptr = scalable_malloc(size)) return ptr; return NULL;}void *operator new[](size_t size, const std::nothrow_t \u0026amp;) throw { return operator new(size, std::nothrow);}void operator delete(void *ptr) throw() { if (ptr != 0) scalable_free(ptr);}void operator delete[](void *ptr) throw() { operator delete(ptr); }void operator delete(void *ptr, const std::nothrow_t \u0026amp;) throw() { if (ptr != 0) scalable_free(ptr);}void operator delete[](void *ptr, const std::nothrow_t \u0026amp;) throw() { operator delete(ptr, std::nothrow);}int main(int argc, char **argv) { const size_t size = 1000; const size_t chunk = 100; // scalable malloc will be called to allocate // the memory for this array of integers int *p = new int[size]; tbb::parallel_for(size_t{0}, size, [=](size_t chunk) { // scalable_malloc will be called to allocate the memory for this // array of integers int *p = new int[chunk]; // scalable_free will be called to deallocate the memory for this // array of integers delete[] p; }); return 0;} 代码示例来自Pro TBB第七章：Scalable Memory Allocation。2. TBB scalable allocator 架构分析TBB scalable allocator基于前后端两层分层架构：分为Frontend和Backend两层。Frontend负责处理基于线程的内存分配请求，Backend负责物理内存分配和全局内存管理（线程分配及回收）。每个线程都有一个自己独立的本地缓存（local cache）。当线程请求内存时，首先检查本地缓存中分配，如果有空闲块，则直接返回本地的空闲块，这个查询及分配操作是无锁的。另外，本地维护的空闲链表，被分类为不同大小的内存块（Size Class），每个Size Class维护一个空闲链表，比如8字节、16字节、2KB等。Backend负责物理内存分配/释放、处理碎片。当某个线程的缓存空了，会向Backend请求内存块，当线程本地的缓存太多时，会将多余的内存块返回给Backend。由于Backend是共享的，所有访问Backend需要加锁。另外，Huge Pages的支持也是在Backend实现的。这套前端/后端的分层架构，特点为： Thread-Local Storage (TLS): 利用 TLS 存储每个线程的分配器状态，避免全局锁。 Cross-Thread Recycling: 如果线程 A 释放了内存，而线程 B 需要内存，分配器需要能安全地将 A 释放的块转交给 B。TBB malloc 使用一种延迟回收或集中式回收的机制来处理这种跨线程转移，同时尽量减少锁竞争。 Object Caching: 释放的内存不会立即还给 OS，而是保留在缓存中，以便下次快速重用。 Frontend代码路径：src/tbbmalloc/frontend.cpp。Backend代码路径：src/tbbmalloc/backend.cpp。tbbmalloc相关的部分高层代码路径：include/oneapi/tbb/memory_pool.hsrc/tbb/allocator.cppsrc/tbbmalloc_proxy/proxy.cpp。2.1. Linux系统符号替换ELF 格式中，符号有三种绑定类型（st_bind）： 类型 说明 STB_GLOBAL 强符号，全局可见，重复定义时链接报错 STB_WEAK 弱符号，可被强符号覆盖，未覆盖时使用弱版本 STB_LOCAL 局部符号，仅文件内可见 dlopen以及动态链接器只加载找到的第一个强符号，在proxy.cpp中，定义这些符号（部分定义截图）：并调用dlsym(RTLD_NEXT, ...)保留系统原始的symbol，用作初始化阶段bootstrap调用：inline void InitOrigPointers(){ // race is OK here, as different threads found same functions if (!origFuncSearched.load(std::memory_order_acquire)) { orig_free = dlsym(RTLD_NEXT, \"free\"); orig_realloc = dlsym(RTLD_NEXT, \"realloc\"); orig_msize = dlsym(RTLD_NEXT, \"malloc_usable_size\"); orig_libc_free = dlsym(RTLD_NEXT, \"__libc_free\"); orig_libc_realloc = dlsym(RTLD_NEXT, \"__libc_realloc\"); origFuncSearched.store(true, std::memory_order_release); }}另外，使用__attribute__((alias))将libc库中的__libc_*等函数替换为对应的routines： alias(“sym”) 让当前符号成为 sym 的别名，两个名字指向同一段代码。即将__libc_free等符号截获并替换为free等符号。2.2. TBB scalable allocator 的实现初始化的时候，Backend以1MB为单位，从系统申请内存，并将这些内存切分成Block（每个Block大小为16KB，内存对齐为16KB），并放在global heap of free blocks中。这个时候，申请的global heap of free blocks常驻内存，不会被释放，以保证这些内存的复用。ExtMemoryPool::backend实现了Backend。当线程TLS内存不够时，再向Backend从global heap of free blocks中申请内存块（Block），此时Backend需要从系统申请新的内存块（以1MB为单位）。另外，还有一个Global Heap of Abandoned Blocks。线程退出时，TLSData 中尚未释放的 Slab 会被孤立化（orphaned），放入全局废弃队列（shareOrphaned/privatizeOrphaned），供其他线程认领复用。2.2.1. Size ClassFrontend分配的最小单位是Object，并分成不同Size Class的Object： 小对象(Small)：≤64字节，共8个bin 分隔对象(Segregated)：64~1024字节，共16个bin 拟合对象(Fitting)：1792字节、2688字节、…8128字节，共5个bin 大对象(Large)：\u0026gt;8128字节（即大于fittingSize5） 巨型对象(Huge)：\u0026gt;4MB，直接从操作系统分配–mmap/munmapSize Class划分定义如下：2.2.2. bin（桶） 以及 Memory Block / SlabObject并不是直接从整块内存中分配的，而是从Block中分配的。每个Block包含多个Object，并且每个Size Class维护一个空闲链表（free list）来管理这些Block。由于Object有不同的Size Class，所以按Size Class分类，创建对应的链表，即有若干个Block链表，每个链表管理一个Size Class的Block。当线程请求内存时，从Size Class对应的的Block空闲链表查找，如果有空闲的Block，则从中分配一个Object；如果没有，则向Backend请求新的Block。不同Size Class的Block链表，存在对应的bin(桶)中，一共定义了32个bin。每个线程有一个TLSData，里面存了一个Bin bin[numBlockBins]。每个bin中定义了一个activeBlk，以及一个mailbox。整个Frontend核心数据结构可以示意如下：TLSData // 每个线程独立一份（Thread Local Storage）└── bin[numBlockBins] // 32个bin，按Size Class分类 ├── bin[0] // Size Class: 小对象 8字节 │ ├── activeBlk ──────────────► Block // 当前活跃Block的指针 │ ├── mailbox // 跨线程回收的Object队列 │ └── BlockList (LIFO) │ ├── Block │ │ ├── BlockHeader // 元数据：size class、已分配数等 │ │ └── Object[N] // N个等大的Object槽位 │ └── ... ├── bin[1] // Size Class: 16字节 │ └── ... └── ... // 共32个binObject的分配：在一个bin中Block的索引顺序为Full Block(后) =\u0026gt; Active Block =\u0026gt; Empty enough block(前)，bin中有activeBlk信息，当需要分配Object时，查询activeBlk中是否有空槽位（Slot），有则直接返回这个Slot，如果没有则移动到下一个索引继续查找，这样顺序索引速度快。当一个Block中Object槽位被释放足够多时（满足Empty enough block条件），这个Block移动到Active Block的前面。Object的释放：每个Block包含一个BlockHeader（LocalBlockFields/GlobalBlockFields），里面有Size Class等信息，所以释放Object需要的信息直接从Block中获取，而不需要在Object自身存储这些信息。当一个Block中的所有被释放之后，这个Block归还到global heap of free blocks中。（Empty enough block是指有足够多的槽位空位–低于1/4，并不是完全槽位被释放。） 使用Block / Object结构以及等分设计，是的内存更紧凑，针对小Object具有良好的内存局部性以及缓存局部性，即相同Size Class的内存对象都分配在一起。实际实践上，附近的变量更可能被访问到，这样就能更好地利用CPU缓存，提升性能。2.2.3. 生产-消费模型：跨线程回收由于Frontend是线程本地的，如果允许其他线程访问则需要加锁。为了避免锁竞争，在Block内，定义了两个链表：freeList以及publicFreeList，并定义了两个接口函数，用于释放public free list、将public free list里面的Object回收到free list：void freePublicObject(FreeObject *objectToFree);void privatizePublicFreeList(bool reset = true);public free list的访问性能略有下降，其使用atomic操作保证线程安全：std::atomic\u0026lt;FreeObject*\u0026gt; publicFreeList;2.2.4. 碎片整理：Coalescing(合并)Coalescing流程主要发生在Backend，目的是将相邻的空闲内存块合并，以减少碎片，并在整个内存区域（MemRegion）空闲时将其归还给操作系统。2.2.4.1. 核心数据结构 GuardedSize其中定义枚举：enum State { LOCKED, // 块正在被使用 COAL_BLOCK, // 块正在参与合并，block is coalescing now MAX_LOCKED_VAL = COAL_BLOCK, LAST_REGION_BLOCK, // 区域末尾的标记块，used to mark last block in region // values after this are \"normal\" block sizes MAX_SPEC_VAL = LAST_REGION_BLOCK // 正常的大小值（\u0026gt; MAX_SPEC_VAL）：块空闲 };2.2.4.2. 核心数据结构：FreeBlock myL：保护本块大小的锁 leftL：保护左侧邻居大小的锁（存在于右侧块头部，供合并时快速访问） nextToFree：用于组成合并延迟队列的链表指针2.2.4.3. 核心类：CoalRequestQ — 延迟合并请求队列当合并因邻居块被锁定而无法立即进行时，当前块会被放入这个无锁队列，等待后续处理。coalescQ 的存在是为了避免死锁与过度竞争：当两个线程同时释放相邻块时，有一方会检测到邻居正在合并（COAL_BLOCK 状态），并将自己的块放入 coalescQ，由后续任意线程的分配/清理操作来处理。2.2.4.4. 核心合并函数：doCoalescdoCoalesc()尝试将一个块与其左右邻居进行合并，返回合并后的大块；若无法立即合并，则将块放入coalescQ并返回nullptr。2.2.4.5. 其他核心处理函数函数coalescAndPutList()对合并后的块列表逐一处理， 场景 A：整个区域已空（memRegion != nullptr 且块大小等于区域的初始块大小） 场景 B：合并后的块放回 bin 最终解锁函数scanCoalescQ()：延迟合并队列。scanCoalescQ() 负责将 coalescQ 中积压的块取出并重新尝试合并，在以下时机被调用： 分配时（genericGetBlock()）：每次尝试获取块前，先扫描延迟队列 等待块释放时（BackendSync::waitTillBlockReleased()）：监测 inFlyBlocks，若有进展则扫描队列 软限制清理时（releaseCachesToLimit()）：超过内存软限制时扫描 全局清理时（Backend::clean()）：清理 advance regions 前先扫描2.3. Frontend: handling large objectsTODOA. 资料 Memory Allocation：官方文档 Scalable Memory Allocation for Parallel Algorithms：PPT资料，介绍TBB Malloc的使用 scalable_allocators：Intel TBB scalable_malloc benchmark The_Foundations_for_Scalable_Multi-Core_Software_in_Intel_Threading_Building_Blocks：早期讲述TBB scalable allocator设计的文章（Scalable Memory Allocator Architecture），介绍了当时的设计思路和实现细节 COMP522-2019-TBB-OpenMP：设计思路对应的PPT资料 How to Use oneTBB for Efficient Memory Allocation in C++ Applications：Intel 官方文章，介绍了tbbmalloc架构（不够完整清晰，作为补充材料） On the Impact of Memory Allocation on High-Performance Query Processing：性能测试对比论文" },

  { "title": "信号处理资料：Butterworth 和 Chebyshev 滤波器", "url": "/blog/2024/dsp-filter-butterworth-and-chebyshev/", "categories": "Algorithm",
  "tags": "Cpp, Algorithm", "date": "2024-08-28", "content":
  "滤波器资料 Digital Signal Processing github – related sources功率密度谱 (PSD) 资料 功率密度谱(Power Spectral Density)笔记 white noise filtering" },

  { "title": "使用 mmap 读取文件", "url": "/blog/2024/read-file-using-mmap/", "categories": "Cpp",
  "tags": "Cpp, Performance", "date": "2024-08-27", "content":
  "关于mmap介绍，见之前文章 总结：内存访问优化(2024-08-13)。使用跨平台支持的三方库github – mio。1. 使用 mio 映射内存读取文件相关头文件：#include \u0026lt;chrono\u0026gt;#include \u0026lt;filesystem\u0026gt;#include \u0026lt;fstream\u0026gt;#include \u0026lt;string\u0026gt;#include \"3rd_utils.h\"#include \"spdlog/spdlog.h\"#include \"typedef.h\"#include \u0026lt;mio/mmap.hpp\u0026gt;#include \u0026lt;system_error\u0026gt; // for std::error_codenamespace {sample_dataset_t loadDataXYFromFileMM(const std::string\u0026amp; filename) { constexpr size_t sample_size = sizeof(double) * 2 + sizeof(int32_t); sample_dataset_t ds; if (!std::filesystem::is_regular_file(filename)) { SPDLOG_ERROR(\"File not found: {}\", filename); return ds; } std::error_code error; mio::mmap_source src_mmap = mio::make_mmap_source(filename, 0, mio::map_entire_file, error); if (error) { const auto\u0026amp; errmsg = error.message(); SPDLOG_ERROR(\"Failed to mmap file: {} - {}\", filename, errmsg); return ds; } int32_t sid{}; // sample id size_t offset{}; const size_t num_samples = (src_mmap.size() / sample_size); ds.ds_.reserve(num_samples); for (size_t i = 0; i \u0026lt; num_samples; i++) { auto* ptr1 = reinterpret_cast\u0026lt;const double*\u0026gt;(src_mmap.data() + offset); const auto x = ptr1[0], y = ptr1[1]; offset += sizeof(double) * 2; auto* ptr2 = reinterpret_cast\u0026lt;const int32_t*\u0026gt;(src_mmap.data() + offset); const auto ch_id = ptr2[0]; offset += sizeof(int32_t); sample_data_t sample{}; sample.cid_ = ch_id - 1, sample.sid_ = sid++; sample.sample_data_ = {x, y}; ds.ds_.emplace_back(std::move(sample)); } return ds;}} // namespace2. 使用 fstream 读取文件namespace {sample_dataset_t loadDataXYFromFileFS(const std::string\u0026amp; filename) { sample_dataset_t ds; std::ifstream ifs(filename, std::ios::in | std::ios::binary); const auto sample_size = sizeof(double) * 2 + sizeof(int32_t); int32_t sid{}; if (!ifs.is_open()) { SPDLOG_WARN(\"Failed to open file: {}\", filename); return ds; } ifs.seekg(0, std::ios::end); const auto file_size = ifs.tellg(); const auto num_samples = file_size / sample_size; ifs.seekg(0, std::ios::beg); ds.ds_.reserve(num_samples); SPDLOG_INFO(\"Predicate loading {} samples from file: {}\", num_samples, filename); while (ifs.good()) { double x, y; int32_t ch{}; ifs.read((char*)(\u0026amp;x), sizeof(double)); ifs.read((char*)(\u0026amp;y), sizeof(double)); ifs.read((char*)(\u0026amp;ch), sizeof(int32_t)); if (ifs.good()) { sample_data_t sample_data{}; sample_data.sample_data_ = {x, y}; sample_data.cid_ = ch - 1; sample_data.sid_ = sid++; ds.ds_.emplace_back(std::move(sample_data)); } } ifs.close(); return ds;}} // namespace3. 性能对比读取及解析文件耗时对比(10次)，使用mmap时间约为fstream的1/4。[2024-08-27 23:57:14.722] [info] 1. Loaded 1503894 samples in 45.58 ms[2024-08-27 23:57:14.784] [info] 1. Loaded 1503894 samples in 49.41 ms[2024-08-27 23:57:14.842] [info] 1. Loaded 1503894 samples in 46.25 ms[2024-08-27 23:57:14.898] [info] 1. Loaded 1503894 samples in 46.14 ms[2024-08-27 23:57:14.955] [info] 1. Loaded 1503894 samples in 45.98 ms[2024-08-27 23:57:15.007] [info] 1. Loaded 1503894 samples in 42.34 ms[2024-08-27 23:57:15.063] [info] 1. Loaded 1503894 samples in 45.53 ms[2024-08-27 23:57:15.121] [info] 1. Loaded 1503894 samples in 48.29 ms[2024-08-27 23:57:15.178] [info] 1. Loaded 1503894 samples in 45.35 ms[2024-08-27 23:57:15.233] [info] 1. Loaded 1503894 samples in 44.41 ms[2024-08-27 23:57:15.483] [info] 2. Loaded 1503894 samples in 238.79 ms[2024-08-27 23:57:15.743] [info] 2. Loaded 1503894 samples in 250.12 ms[2024-08-27 23:57:16.011] [info] 2. Loaded 1503894 samples in 255.81 ms[2024-08-27 23:57:16.261] [info] 2. Loaded 1503894 samples in 238.10 ms[2024-08-27 23:57:16.502] [info] 2. Loaded 1503894 samples in 230.86 ms[2024-08-27 23:57:16.749] [info] 2. Loaded 1503894 samples in 235.96 ms[2024-08-27 23:57:17.007] [info] 2. Loaded 1503894 samples in 246.65 ms[2024-08-27 23:57:17.251] [info] 2. Loaded 1503894 samples in 231.93 ms[2024-08-27 23:57:17.499] [info] 2. Loaded 1503894 samples in 237.23 ms[2024-08-27 23:57:17.751] [info] 2. Loaded 1503894 samples in 240.08 ms4. 资料 github – mio" },

  { "title": "Range-v3 用法积累，及资料", "url": "/blog/2024/c++-range-v3/", "categories": "Cpp",
  "tags": "Cpp", "date": "2024-08-27", "content":
  "1. view – 返回子集 sub_range#include \u0026lt;ranges\u0026gt;using dataset_slice_t = std::ranges::subrange\u0026lt;std::vector\u0026lt;rias::data_type::sample_data_t\u0026gt;::iterator\u0026gt;;dataset_slice_t sub_range(it_min, it_max);2. filter – 根据条件过滤出子集Block blk;auto ss = sample_dataset_-\u0026gt;ds_ | std::views::filter([\u0026amp;](const sample_data_t\u0026amp; ds) { return ds.dbg_blk_id_ == blkId; });std::ranges::for_each(ss, [\u0026amp;](const sample_data_t\u0026amp; ds) { blk.signals_.push_back(ds); }); github – range-v3 User Manual – range-v3 pdf – A GentleIntroductiontoRangesv3 Range-v3 practical examples" },

  { "title": "总结：内存访问优化", "url": "/blog/2024/cpp-perf-repos/", "categories": "Cpp",
  "tags": "Cpp, Performance, Memory", "date": "2024-08-27", "content":
  "1. 虚拟内存分配1.1 mmapmmap用于建立文件映射，或者匿名映射。当用于文件映射时，mmap将文件内容缓存进内核空间的page cache里面，然后将用户的一段虚拟内存空间直接映射到page cache。用户通过访问这段虚拟内存，直接读写内核空间上的page cache，避免buffer拷贝开销及用户态的切换。用mmap用户建立匿名映射时，将用户空间的一段虚拟内存空间直接映射到某段物理内存上，这段虚拟内存称为匿名页。匿名映射用于malloc操作（大于128KB）。mmap文件知识点: 通常情况下(除了MAP_POPULATE)，mmap创建时，只是在用户空间分配一段地址空间(VMA)，只有访问地址空间时，才会分配物理地址空间(Page fault中断分配内存)，并更新映射到VMA，建立映射关系。 mmap映射的物理内存，可以跨进程共享，但需要进程之间加锁访问(写操作)。如果多个进程写同一个mmap映射的物理内存，会触发Copy On Write(COW)，内核重新分配一个新的物理内存，并复制原有物理内存的内容。 通过msync()将内存写回硬盘，munmap()释放内存。MAP_POPULATE标志位: 建立页表，这将使得内核进行一些预读(实测没有性能提升)。使用方式： 使用open + 选项O_RDONLY | O_DIRECT打开文件; 以及使用mmap + MAP_POPULATE选项，在打开文件时建立页表。MAP_LOCKED标志位: 锁定映射内存，阻止被换出。类似于mlock()。更多I/O相关: about IO performance1.2 malloc / free在现代操作系统中，malloc的作用是分配虚拟内存空间，并不实际分配物理内存。当分配的虚拟内存空间第一次被访问时，才会真正的分配物理内存（OS的写时分配行为）。malloc行为： 首先尝试在进程空间内存池中查找有没有可以重用的内存空间，如果没有才会进行系统调用。 如果申请的内存小于128KB，会通过调用brk()函数申请内存：根据申请内存大小，将堆顶指针向上移动，并返回新申请的内存地址给用户；当free掉brk()分配的内存时，并不会将物理内存缓存归还给操作系统，而是放到malloc的内存池中，待下次再次申请时直接使用。 申请的内存大于128KB，会通过调用mmap()函数申请内存：通过匿名映射获取虚拟内存；当free掉mmap()分配的内存时，会将物理内存缓存归还给操作系统。1.3 new / deletenew / delete操作时，在调用malloc/free基础上，对non-trival对象，调用其构造/析构函数；对于trival对象，不需要调用构造/析构函数，直接分配/释放内存。2. 用户态 malloc 用户态内存分配：intel TBB malloc, tcmalloc，Vulkan Memory Allocator等。（ 经测试，microsoft mimalloc适配性不是很好，使用过程中会出错；intel TBB malloc overhead似乎比较大） 内存池。（见下面资料链接） 对象池。2.1 参考资料 Vulkan Memory Allocator Vulkan Memory Allocator – docs 游戏架构设计：内存管理 游戏架构设计：高性能并行编程2.2 内存池仓库 github – memory github – poolSTL3. Linux Huge page设置Huge Page，减少内存访问需要的的缺页中断，提高内存访问效率。以及减少TLB未命中导致的性能下降 – 未命中TLB则需要逐级查询page table。Linux使用Huge Page有两种方式：3.1 Transparent Huge Page主流Linux kernel发布版本都是默认支持THP的（TRANSPARENT_HUGEPAGE）。但是在用户态使能，需要开启设置khugepaged: kernel – Transparent Hugepage Support$ cat /sys/kernel/mm/transparent_hugepage/enabledalways [madvise] never设置khugepaged:# run as rootecho \"always\" \u0026gt;! /sys/kernel/mm/transparent_hugepage/enabled选项: always: 开启THP，内核尝试将连续的内存页合并成一个THP页。用户层不需要任何操作。 madvise: 开启THP，内核尝试将连续的内存页合并成一个THP页。用户层可以使用madvise()系统调用，将内存标记为THP。对x86-64系统，THP页大小为2MB。查看Huge Page的页大小:cat /sys/kernel/mm/transparent_hugepage/hpage_pmd_size使用THP，需要在分配内存时，需要将分配的内存对齐，比如2MB大页对齐:#include \u0026lt;iostream\u0026gt;#include \u0026lt;sys/mman.h\u0026gt;// ... definition of is_huge() and is_thp() ...constexpr size_t HPAGE_SIZE = 2 * 1024 * 1024;int main() { auto size = 4 * HPAGE_SIZE; void *mem = aligned_alloc(HPAGE_SIZE, size); madvise(mem, size, MADV_HUGEPAGE); // Make sure the page is present static_cast\u0026lt;char *\u0026gt;(mem)[0] = 'x'; std::cout \u0026lt;\u0026lt; \"Is huge? \" \u0026lt;\u0026lt; is_huge(mem) \u0026lt;\u0026lt; \"\\n\"; std::cout \u0026lt;\u0026lt; \"Is THP? \" \u0026lt;\u0026lt; is_thp(mem) \u0026lt;\u0026lt; \"\\n\";}完整代码thp.cppg++ --std=c++17 thp.cpp -o thp3.2 HugeTLB可以通过仅仅链接libhugetlbfs，即可使用大页内存:# https://github.com/libhugetlbfs/libhugetlbfssudo apt-get install libhugetlbfs-dev libhugetlbfs-bin -ysudo ln -s /usr/bin/ld.hugetlbfs /usr/share/libhugetlbfs/ld# 分配1000个大页，一个page大小为2MB，总共2GB. 需要管理员权限echo 1000 \u0026gt; /proc/sys/vm/nr_hugepages# 确认大页是否可用cat /proc/meminfo | grep HugePages_链接libhugetlbfs:-B /usr/share/libhugetlbfs -Wl,--hugetlbfs-align -no-pie -Wl,--no-as-needed参考 ARM – Introduction to libhugetlbfs另外，使用HugeTLB分配一个10MB匿名映射:void *addr = mmap(0, 10*1024*1024, (PROT_READ | PROT_WRITE), (MAP_PRIVATE | MAP_ANONYMOUS | MAP_HUGETLB), 0, 0);参考: Linux HugeTLB: What is the advantage of the filesystem approach? Allocating Huge Pages on Linux" },

  { "title": "聚类算法（密度）：基于 nano-flann", "url": "/blog/2024/density-algrithm-knn/", "categories": "Algorithm",
  "tags": "Algorithm, Cpp", "date": "2024-08-26", "content":
  "0. DBSCAN 算法及 K-D 树介绍DBSCAN算法相关概念： 邻域半径 eps。 核心点，最少核心点 minPts。 直接密度可达。 密度可达。 密度相连。K-D树的时间复杂度：Kdtree 算法的构建时间复杂度为 O(nlogn)，搜索时间复杂度最好为 O($\\log_2 N$)，最坏为 O($N^{1-1/k}$)。 基于k-d 树的查询算法实现与二维可视化 KD-Tree详解: 从原理到编程实现1. 背景采集到的二维点云数据(samples)，生成K-D搜索树，使用广度优先搜索，聚合成block数据。后续的识别/分类算法，在block数据基础上进行。由于使用点云处理库PCL比较庞大，以及其依的FLANN基于C++14，使用C++17/20导致在自定义点云数据结构时，编译有些STL算法库被废弃，编译出错。故使用 nanoflann。2. 基于 nano-flann 的聚类算法实现2.1 自定义点云数据结构#pragma once// test\\test_flann\\flann_adaptor.h#include \u0026lt;vector\u0026gt;#include \"3rd_utils.h\"#include \"typedef.h\"struct PointCloud { using Point = rias::data_type::sample_data_t; using coord_t = double; //!\u0026lt; The type of each coordinate std::vector\u0026lt;Point\u0026gt; pts; // Must return the number of data points inline size_t kdtree_get_point_count() const { return pts.size(); } // Returns the dim'th component of the idx'th point in the class: // Since this is inlined and the \"dim\" argument is typically an immediate // value, the // \"if/else's\" are actually solved at compile time. inline double kdtree_get_pt(const size_t idx, const size_t dim) const { CHECK2(dim \u0026lt; 2, \"Invalid dimension: \" \u0026lt;\u0026lt; dim); if (dim == 0) return pts[idx].mapped_pos_.x_; else return pts[idx].mapped_pos_.y_; } // Optional bounding-box computation: return false to default to a standard // bbox computation loop. // Return true if the BBOX was already computed by the class and returned // in \"bb\" so it can be avoided to redo it again. Look at bb.size() to // find out the expected dimensionality (e.g. 2 or 3 for point clouds) template \u0026lt;class BBOX\u0026gt; bool kdtree_get_bbox(BBOX\u0026amp; /* bb */) const { return false; }};2.2 基于 K-D 树构建广度优先搜索算法#pragma once// test\\test_flann\\point_cloud_algorithm.h#include \u0026lt;queue\u0026gt;#include \u0026lt;vector\u0026gt;#pragma warning(push)#pragma warning(disable : 4267) // disable conversion warning#include \"nanoflann/nanoflann.hpp\"#pragma warning(pop)#include \"3rd_utils.h\"#include \"flann_adaptor.h\"#include \"typedef.h\"namespace rias::test {class BFSDensitySampleSearch { static constexpr int32_t kDIM = 2; using my_kd_tree_t = nanoflann::KDTreeSingleIndexAdaptor\u0026lt;nanoflann::L2_Simple_Adaptor\u0026lt;double, PointCloud\u0026gt;, PointCloud, kDIM\u0026gt;; using sample_t = rias::data_type::sample_data_t; public: BFSDensitySampleSearch(const std::vector\u0026lt;sample_t\u0026gt;\u0026amp; pts) : pts_(pts) {} ~BFSDensitySampleSearch() = default; BFSDensitySampleSearch\u0026amp; setRadius(double radius) { radius_ = radius; return *this; } BFSDensitySampleSearch\u0026amp; setMinPts(double minPts) { minPts_ = minPts; return *this; } const std::vector\u0026lt;std::vector\u0026lt;size_t\u0026gt;\u0026gt;\u0026amp; clusters() const { return clusters_; } BFSDensitySampleSearch\u0026amp; search() { PointCloud cloud; cloud.pts = pts_; CHECK1(sizeof(double[kDIM]) == sizeof(data_type::point_double_t), \"sizeof(double[kDIM])\"); my_kd_tree_t index(kDIM /*dim*/, cloud, {10 /* max leaf */}); std::vector\u0026lt;double\u0026gt; densities(pts_.size(), double{}); for (size_t i = 0; i \u0026lt; pts_.size(); i++) { const auto* q_pt = (double*)(\u0026amp;pts_[i].mapped_pos_); std::vector\u0026lt;nanoflann::ResultItem\u0026lt;uint32_t, double\u0026gt;\u0026gt; ret_matches; const size_t n_matched = index.radiusSearch(q_pt, radius_, ret_matches); densities[i] = static_cast\u0026lt;double\u0026gt;(n_matched); } std::vector\u0026lt;bool\u0026gt; visited(pts_.size(), false); for (size_t i = 0; i \u0026lt; pts_.size(); i++) { if (visited[i] /*|| densities[i] \u0026lt; minPts_*/) { // 0. skip visited or low density points continue; } std::vector\u0026lt;size_t\u0026gt; cluster; std::queue\u0026lt;size_t\u0026gt; q; q.push(i); // 1. add self ot queue at first visited[i] = true; while (!q.empty()) { const auto idx = q.front(); q.pop(); cluster.push_back(idx); // 2. elements in the same queue are in the same cluster const auto* q_pt = (double*)(\u0026amp;pts_[idx].mapped_pos_); std::vector\u0026lt;nanoflann::ResultItem\u0026lt;uint32_t, double\u0026gt;\u0026gt; ret_matches; const size_t n_matched = index.radiusSearch(q_pt, radius_, ret_matches); for (size_t j = 0; j \u0026lt; n_matched; j++) { // 3. add matched points to the same queue const auto n_idx = ret_matches[j].first; // const auto n_dist = ret_matches[j].second; if (!visited[n_idx] \u0026amp;\u0026amp; densities[n_idx] \u0026gt;= minPts_) { visited[n_idx] = true; q.push(n_idx); } } } // 4. a cluster done (one or more points) clusters_.push_back(cluster); } return *this; } private: const std::vector\u0026lt;sample_t\u0026gt;\u0026amp; pts_; double radius_{1.0}, minPts_{1.0}; std::vector\u0026lt;std::vector\u0026lt;size_t\u0026gt;\u0026gt; clusters_;};} // namespace rias::test2.3 测试代码{ const auto t0 = std::chrono::high_resolution_clock::now(); auto dss = BFSDensitySampleSearch(sample_ds-\u0026gt;ds_).setRadius(3.0).setMinPts(1.0); auto clusters = dss.search().clusters(); const auto t1 = std::chrono::high_resolution_clock::now(); const auto dur = std::chrono::duration\u0026lt;double, std::milli\u0026gt;(t1 - t0).count(); spdlog::info(\"search in {} points. found {} clusters. time: {:.3f} seconds\\n\", sample_ds-\u0026gt;ds_.size(), clusters.size(), dur / 1000.0); /*for (const auto\u0026amp; cluster : clusters) { //spdlog::info(\"{}\", fmt::join(cluster, \", \")); spdlog::info(\"cluster: {}\", cluster); }*/}2.4 测试[2024-08-26 21:47:35.648] [warning] [dataxy_loader.cc:146] under flow 870 samples in dataset[2024-08-26 21:47:35.649] [info] load data from file: 0.37393689999999996 seconds[2024-08-26 21:47:37.740] [info] search in 1503894 points. found 438931 clusters. time: 2.133 secondsTODO: 测试小数据集下的性能对比。2.5 参考 nanoflann 空间数据结构(四叉树/八叉树/BVH树/BSP树/k-d树) 数据结构-k-d树 nanoflann库使用笔记3. 基于templated的聚类算法实现(线性搜索 O2时间复杂度)template \u0026lt;typename ElemType, typename LinerFuncType, typename AdjFuncType\u0026gt;class BFSLinerMerge { public: BFSLinerMerge(const std::vector\u0026lt;ElemType\u0026gt;\u0026amp; blocks, LinerFuncType linerCondFunc, AdjFuncType adjCondFunc) : blocks_(blocks), liner_cond_func_(linerCondFunc), adj_cond_func_(adjCondFunc) {}; BFSLinerMerge\u0026amp; search() { if (!clusters_.empty()) clusters_.clear(); std::vector\u0026lt;bool\u0026gt; visited(blocks_.size(), false); clusters_.reserve(blocks_.size()); for (size_t i = 0; i \u0026lt; blocks_.size(); i++) { if (visited[i]) continue; std::queue\u0026lt;size_t\u0026gt; q; q.push(i); visited[i] = true; std::vector\u0026lt;Block*\u0026gt; cluster; while (!q.empty()) { size_t idx = q.front(); q.pop(); const auto this_blk = blocks_[idx]; cluster.push_back(blocks_[idx]); for (size_t j = 0; j \u0026lt; blocks_.size(); j++) { if (visited[j]) continue; const auto\u0026amp; blkj = blocks_[j]; if (liner_cond_func_(this_blk, blkj) \u0026amp;\u0026amp; adj_cond_func_(this_blk, blkj)) { q.push(j); visited[j] = true; } } } clusters_.push_back(cluster); }; return *this; } std::vector\u0026lt;std::vector\u0026lt;ElemType\u0026gt;\u0026gt;\u0026amp; clusters() { return clusters_; } private: const std::vector\u0026lt;ElemType\u0026gt;\u0026amp; blocks_; LinerFuncType liner_cond_func_; AdjFuncType adj_cond_func_; std::vector\u0026lt;std::vector\u0026lt;ElemType\u0026gt;\u0026gt; clusters_;};4. 聚类算法资料收集 几种常用的基于密度的聚类算法 DBSCAN密度聚类算法 K紧邻法(KNN)原理小结" },

  { "title": "Ubuntu 安装 OpenCV", "url": "/blog/2024/ubuntu-install-opencv/", "categories": "Cpp",
  "tags": "Cpp, OpenCV", "date": "2024-08-23", "content":
  "依赖项安装sudo apt-get updatesudo apt-get install -y build-essential cmake pkg-configsudo apt-get install -y libjpeg-dev libpng-dev libtiff-devsudo apt-get install -y libavcodec-dev libavformat-dev libswscale-dev libv4l-devsudo apt-get install -y libxvidcore-dev libx264-devsudo apt-get install -y libgtk2.0-dev libgtk-3-devsudo apt-get install -y libatlas-base-dev gfortransudo apt-get install -y python3-dev python3-numpysudo apt-get install -y libtbb2 libtbb-dev libdc1394-22-devsudo apt-get install -y libopencv-dev下载 OpenCV 源码git clone https://github.com/opencv/opencv.gitgit clone https://github.com/opencv/opencv_contrib.git# opencv opencv_contrib 检出同一个版本编译 OpenCVcmake配置过程中，需要下载依赖包。设置代理：export http_proxy=http://192.168.9.165:10809export https_proxy=http://192.168.9.165:10809 编译选项OPENCV_EXTRA_MODULES_PATH设置extra模块路径：~/tmp/opencv/opencv_contrib/modules。 OpenCV还依赖其他第三方库：TBB、Eigen、VTK 等。可以先安装上。参考 cmake编译opencv指南" },

  { "title": "CMake 检查系统和编译器", "url": "/blog/2024/cmake-check-system-andcompiler/", "categories": "CMake",
  "tags": "CMake", "date": "2024-08-22", "content":
  "判断操作系统IF (CMAKE_SYSTEM_NAME MATCHES \"Linux\")ELSEIF (CMAKE_SYSTEM_NAME MATCHES \"Windows\")ELSEIF (CMAKE_SYSTEM_NAME MATCHES \"FreeBSD\")ELSE ()MESSAGE(STATUS \"other platform: ${CMAKE_SYSTEM_NAME}\")ENDIF (CMAKE_SYSTEM_NAME MATCHES \"Linux\")判断编译器if (\"${CMAKE_CXX_COMPILER_ID}\" STREQUAL \"Clang\")# using Clangelseif (\"${CMAKE_CXX_COMPILER_ID}\" STREQUAL \"GNU\")# using GCCelseif (\"${CMAKE_CXX_COMPILER_ID}\" STREQUAL \"Intel\")# using Intel C++elseif (\"${CMAKE_CXX_COMPILER_ID}\" STREQUAL \"MSVC\")# using Visual Studio C++endif()" },

  { "title": "备份：Ubuntu Bash Alias, 以及 bash 显示 git status", "url": "/blog/2024/ubuntu-bash-alias/", "categories": "Bash",
  "tags": "Bash, Linux", "date": "2024-08-21", "content":
  "bash alias#################### bash aliasalias g='git status -sb'alias ll='ls -alF'alias la='ls -A'alias l='ls -ltrhA'alias gl='ls|grep --color'# alias .='cd ../'# alias ..='cd ../..'alias ..='cd ..'alias ...='cd ..; cd ..'alias ....='cd ..; cd ..; cd ..'alias c='clear'alias r='reset'bash 显示 git status# Show git branch nameforce_color_prompt=yescolor_prompt=yesparse_git_branch() { git branch 2\u0026gt; /dev/null | sed -e '/^[^*]/d' -e 's/* \\(.*\\)/(\\1)/'}if [ \"$color_prompt\" = yes ]; then PS1='${debian_chroot:+($debian_chroot)}\\[\\033[01;32m\\]\\u@\\h\\[\\033[00m\\]:\\[\\033[01;34m\\]\\w\\[\\033[01;31m\\]$(parse_git_branch)\\[\\033[00m\\]\\$ 'else PS1='${debian_chroot:+($debian_chroot)}\\u@\\h:\\w$(parse_git_branch)\\$ 'fiunset color_prompt force_color_prompt参考：How do I show the git branch with colours in Bash prompt?Bash增强：自动完成 ble.sh Setup ble.sh for PowerShell-like history completion in bash" },

  { "title": "随机一致性抽样算法（RANSAC）", "url": "/blog/2024/RANSAC/", "categories": "Algorithm",
  "tags": "Cpp, Algorithm", "date": "2024-08-16", "content":
  "1. RANSAC 算法过程最小二乘法拟合只进行一次迭代，计算所有离散点平均值，得到最终拟合直线或曲线。RANSAC通过多次迭代，寻找拟合直线或曲线的最佳（最近）权重的点。第N次拟合，得到第N次迭代的内点（集合）。第N+1次迭代，得到第N+1次的内点（集合），如果第N+1次迭代计算的内点其权重大于第N次迭代的内点的权重，则更新最佳内点为第N+1次的内点（集合）。其入参有： 点集； 拟合误差delta，作为判断内点的阈值； 迭代次数loopNum；其计算流程为： 随机选取一个点作为起点，随机选取另一个点作为终点； 从集合中查找所有误差在delta范围内的点集合：作为本次内点（集合）：this_inlier，计算其权重，与保存的权重比较best_weight，如果本次权重较好，则更新best_inlier为本次内点（集合），且更新best_weight为本次权重； 标记本次循环所访问过的点：包括终点在内的this_inlier为visited状态，后面的迭代中不再访问这些点； 重复步骤2-4，直到达到设定的迭代次数loopNum；Demo计算结果如下：2. 算法改进实际应用中，发现不能直接应用RANSAC算法。一个原因离散点（点云）数量过大，实际应用不能拟合期望的直线。另外就是其迭代次数在大数据量下，计算量过大。改进措施包括： 通过其他分块算法，初步得到拟合线段所在区域； 根据行业应用特点，限定起点、限定密度、限定斜率； 综合使用其中一个或多个，特别是初步筛选及限定起点；3. 实现TBD4. 更多资料 Ransac 随机一致性采样" },

  { "title": "Intel TBB malloc 使用 (windows)", "url": "/blog/2024/intel-tbb-malloc-usage/", "categories": "Cpp",
  "tags": "Cpp, Performance", "date": "2024-08-13", "content":
  "CMake 查找 Intel TBBfind_package( TBB COMPONENTS tbb tbbmalloc tbbmalloc_proxy REQUIRED)if(TBB_FOUND) message(STATUS \"TBB version: ${TBB_VERSION}\")endif()链接引入 TBB 库if (MSVC) target_link_libraries(${target_lib} PRIVATE TBB::tbb TBB::tbbmalloc TBB::tbbmalloc_proxy)endif()cpp 代码中引入 TBB 符号，防止库链接被优化掉#if defined(WITH_TBB_MALLOC) \u0026amp;\u0026amp; defined(_MSC_VER) //\u0026amp;\u0026amp; defined(NDEBUG)#include \"oneapi/tbb/scalable_allocator.h\"#include \"oneapi/tbb/tbbmalloc_proxy.h\"#pragma comment(lib, \"tbbmalloc_proxy.lib\")#pragma comment(linker, \"/include:__TBB_malloc_proxy\")#endif查看内存分配函数替换是否出错#if defined(WITH_TBB_MALLOC) \u0026amp;\u0026amp; defined(_MSC_VER) //\u0026amp;\u0026amp; defined(NDEBUG) { // https://emfomenk.github.io/versions/latest/elements/oneTBB/source/memory_allocation/c_interface_to_scalable_allocator.html const auto mode_ret = scalable_allocation_mode(TBBMALLOC_USE_HUGE_PAGES, 1); if (mode_ret == TBBMALLOC_NO_EFFECT) { spdlog::warn(\"huge pages not supported by OS\"); } char** func_replacement_log; int func_replacement_status = TBB_malloc_replacement_log(\u0026amp;func_replacement_log); if (func_replacement_status != 0) { printf(\"tbbmalloc_proxy cannot replace memory allocation routines\\n\"); for (char** log_string = func_replacement_log; *log_string != 0; log_string++) { spdlog::warn(\"{}\", *log_string); } } }#endif资料 How to Use oneTBB for Efficient Memory Allocation in C++ Applications" },

  { "title": "复习：std::function 用法笔记", "url": "/blog/2024/std-function/", "categories": "Cpp",
  "tags": "Cpp", "date": "2024-08-13", "content":
  "std::function 可以将函数,函数对象（仿函数）,lambda表达式包装成一个对象。std::function对象本身可以作为函数参数，并且是可复制的（复制构造、赋值）。1. 封装函数指针int add(int a, int b) { return a + b; }int main() { std::function\u0026lt;int(int, int)\u0026gt; f = add; int result = f(1, 2); std::cout \u0026lt;\u0026lt; result \u0026lt;\u0026lt; std::endl; return 0;}2. 封装函数对象（仿函数）struct Adder { int operator()(int a, int b) { return a + b; }};int main() { Adder adder; std::function\u0026lt;int(int, int)\u0026gt; f = adder; int result = f(1, 2); std::cout \u0026lt;\u0026lt; result \u0026lt;\u0026lt; std::endl; return 0;}3. 封装 lambda 表达式int main() { // using func_t = int(int, int); std::function\u0026lt;int(int, int)\u0026gt; f = [](int a, int b) { return a + b; }; int result = f(1, 2); std::cout \u0026lt;\u0026lt; result \u0026lt;\u0026lt; std::endl; return 0;}4. 与 std::bind 结合使用int add(int a, int b) { return a + b; }int main() { std::function\u0026lt;int(int)\u0026gt; f = std::bind(add, 1, std::placeholders::_1); int result = f(2); std::cout \u0026lt;\u0026lt; result \u0026lt;\u0026lt; std::endl; return 0;}5. lambda 表达式lambda表达式是一个匿名函数对象，即编译器会创建一个仿函数( 调用时，调用 operator()(....) )，并将外部捕获的变量，添加到该匿名对象中。这些是lambda的额外开销。注：在O2编译时，仿函数也会被优化掉，直接编译成跳转代码段。int main(){ // Lambda \u0026amp; auto int member=10; auto endGame = [=](int a, int b){ return a+b+member;}; endGame(4,5); return 0;}展开成：int main(){ int member = 10; class __lambda_6_18 { int member; public: inline /*constexpr */ int operator()(int a, int b) const { return a + b + member; } public: __lambda_6_18(int _member) : member{_member} {} }; __lambda_6_18 endGame = __lambda_6_18{member}; endGame.operator()(4, 5); return 0;}参考 What is a lambda expression, and when should I use one?" },

  { "title": "整理：内存一致模型", "url": "/blog/2024/%E6%95%B4%E7%90%86-%E5%A4%84%E7%90%86%E5%99%A8%E5%86%85%E5%AD%98%E6%A8%A1%E5%9E%8B/", "categories": "CPU",
  "tags": "Cpp, CPU", "date": "2024-08-11", "content":
  "内存一致性模型（Memory Consistency Model）是并发编程中的一个重要概念，它定义了在多线程环境下，内存操作的可见性和顺序性规则。内存一致性模型还涉及到Cache一致性（Cache Coherence）和CPU指令的乱序执行（Out-of-Order Execution）。理解内存一致性模型对于编写正确、高效的并发程序至关重要。1. CPU Cache 内部结构一个core内部结构： cache store buffer invalidate queue结构如下图所示：1.1. Cache一致性协议 MESIMESI是CPU内部多个core同步通讯协议，保证多个core中的cache的数据一致性。MESI这四个字母分别代表了每一个cache line可能处于的四种状态：Modified、Exclusive、Shared 和 Invalid。通过给cache line设置状态位，以及CPU core（也可能有内存控制器参与）之间的消息同步逻辑，让多个core中的cache数据保持一致性。在没有store buffer, invalidate queue之前，MESI可以保证不需要memory fence指令也可以保证数据的一致性。1.2. False sharingFalse sharing的原因是两个CPU访问的变量，在内存中的位置，同时落入一个cache line范围内，根据MESI协议，一个CPU写操作，将导致另一个CPU的读写操作之前，需要进行memory及两个CPU的cache line同步操作。通常发生在两个线程操作同一个数据结构体的时候。#define CACHE_ALIGN_SIZE 64#define CACHE_ALIGNED __attribute__((aligned(CACHE_ALIGN_SIZE)))struct aligned_value { int64_t val;} CACHE_ALIGNED; // Note: aligning the struct to a cache line sizealigned_value aligned_data[2] CACHE_ALIGNED;// sizeof(aligned_value) == 1281.3. 现代CPU上MESI的局限由于MESI同步协议导致处理器之间同步的代价很高，现代处理器再每个core里面增加两个异步队列: store buffer和invalidate queue来减少CPU的空闲等待。这两个异步队列，导致MESI协议失效。 store buffer: CPU将write/store操作数据放入store buffer，cache负责flush操作。 invalidate queue: cache收到Invalidate消息，不是马上执行，而是放入invalidate queue，等待CPU空闲时，再执行。 注意: Store Buffer / Load Buffer 在硬件中并不是简单的 FIFO 队列，而是具有调度逻辑的结构（类似 CAM — Content Addressable Memory），支持乱序发射、乱序完成。Intel 官方手册称之为 “Memory Ordering Buffer”，ARM 架构称之为 “Load-Store Queue (LSQ)”。但即使内部结构支持乱序，x86 TSO 模型仍然保证 Store Buffer 按程序顺序提交到 cache（FIFO 语义），这是 x86 强内存模型的关键。其原因是：针对发起方CPU，其认为自己的store / invalidate操作已经完成，但是由于数据/消息是放在store buffer / invalidate queue中，所以可能还没来得及被其他CPU看到，导致数据不一致。其中的一个解决办法是：发起方的store buffer被清空，接收方的invalidate queue被处理掉。在此之后，MESI协议可以正常工作。1.4. 补充：CPU指令流水线1. Fetch \u0026lt;-- ✅ 顺序取指，从指令缓存中取出下一条指令2. Decode \u0026lt;-- 顺序解码，分析操作数和目的寄存器3. Rename \u0026lt;-- 寄存器重命名4. Dispatch \u0026lt;-- 投递到调度窗口，等待执行条件满足5. Execute \u0026lt;-- ✅ 乱序执行（由调度器决定），实际在执行单元上运行指令6. Writeback \u0026lt;-- 写结果到 ROB7. Commit \u0026lt;-- ✅ 按程序顺序提交（retire），更新寄存器状态或进行内存写入2. memory barrier2.1. 概念及理论 同步点：针对同一个原子变量的load操作与store操作，分别构成一个同步点。其概念有三要素：(1)：load/store操作，(2)：针对同一个原子变量，(3)：以及在不同线程中； synchronize-with 关系，该概念包含两个含义：(1)：同一个同步点，(2)：读取的值是另一个同步点写入的值； happens-before 关系；memory fence定义的是同步点操作，即分别在store一方插入一个write barrier指令，在load一方插入一个read barrier指令。因此，memory barrier需要成对出现，否则达不到同步效果。 所谓同步点只是语义上的概念，实际实现中，memory barrier指令并不直接作用于某个特定的变量，而是通过屏障指令的执行来保证之前/之后的所有内存操作（针对任何变量）满足特定的顺序和可见性要求。2.2. 详细解释由于多核处理器 CPU 之间独立的L1/L2 cache，会出现cache line不一致的问题，为了解决这个问题，有相关协议模型，比如 MESI 协议来保证 cache 数据一致，同时由于 CPU 对 MESI 进行的异步优化，对写和读分别引入了「store buffer」和「invalid queue」，很可能导致后面的指令查不到前面指令的执行结果（各个指令的执行顺序非代码执行顺序），这种现象很多时候被称作「CPU乱序执行」。为了解决乱序问题（也可以理解为可见性问题，修改完没有及时同步到其他的CPU），又引出了「内存屏障」的概念；内存屏障可以分为三种类型：写屏障，读屏障以及全能屏障（包含了读写屏障），屏障可以简单理解为：在操作数据的时候，往数据插入一条特殊的指令。只要遇到这条指令，那前面的操作都得「完成」。 写屏障指令(write barrier, or sfence)，等待之前的写操作完成，并把该指令「之前」存在于「store Buffer」中的所有写指令刷入cache。就可以让CPU修改的数据马上暴露给其他CPU(MESI)，达到「写操作」可见性的效果。 读屏障指令(read barrier, or lfence)，会把该指令「之前」存在于「invalid queue」中的所有的指令都处理掉。通过这种方式就可以确保当前CPU的缓存状态是准确的，达到「读操作」一定是读取最新的效果。 由于不同CPU架构的缓存体系不一样、缓存一致性协议不一样、重排序的策略不一样、所提供的内存屏障指令也有差异，所以一些语言c++/java/go/rust 都有实现自己的内存模型, 比如golang大牛Russ Cox写的内存模型系列文章 Memory Models 值得深入了解。2.3. x86上面的fence实操演示ARM架构CPU有Store Buffer、Invalidate Queue，是一个松散内存一致性模型。x86架构只有Store Buffer，是一个强内存一致性模型。在x86架构下，对StoreLoad操作进行重排(乱序)。其余几种保持顺序：StoreStore, LoadLoad, LoadStore，即不需要设置fence指令也可以保持CPU之间的内存一致性。禁止编译器重排：X = 1;asm volatile(\"\" ::: \"memory\"); // Prevent compiler reorderingr1 = Y;禁止编译器及CPU重排：X = 1;asm volatile(\"mfence\" ::: \"memory\"); // Prevent compiler and CPU reorderingr1 = Y;详细知识参考： CPU 缓存一致性与内存屏障 Cache一致性和内存一致性 Acquire and Release Fences 從硬體觀點了解 memory barrier 的實作和效果 CPU架构和MESI缓存一致性-\u0026gt;内存模型一致性-\u0026gt;内存屏障和原子操作-\u0026gt;内存序-\u0026gt;C++内存序3. C++11 内存一致性模型定义 内存序 语义 x86-64 实现 ARM64 实现 relaxed 只保证原子性，不保证顺序 普通 MOV 普通 LDR/STR acquire (load) Acquire操作之后的读/写不能重排到此操作之前，Acquire操作之后的读操作能看到Release线程的Release操作之前的写入 普通 MOV + 编译器屏障 LDAR release (store) Release操作之前的读/写不能重排到此操作之后，Release操作之前的写操作对Acquire线程可见 普通 MOV + 编译器屏障 STLR acq_rel (RMW) 同时具有 acquire 和 release 语义 LOCK 前缀指令（隐含全屏障） LDAXR/STLXR seq_cst 全局唯一顺序一致性，所有线程看到相同的 seq_cst 操作顺序 store: XCHG 或 MOV+MFENCE; load: MOV LDAR/STLR (+ 额外屏障) 关键认知：C++ 内存序是语言层面的抽象，其硬件实现因架构而异。 Release/Aquire语义要保证happens-before关系，所以Release同时约束前面的读以及写操作不能重排到Release store之后。Acquire同时约束后面的读以及写操作不能重排到Acquire load之前。acq_rel则同时添加了 acquire 和 release 语义，适用于 read-modify-write 操作（如 fetch_add、compare_exchange）。相当于：load(memory_order_acquire);modify();store(memory_order_release);acq_rel 约束本线程中前面的读写操作重排到此操作之前，后面的读写操作重排到此操作之后。且可以看见其他线程在 release 操作之前的写入，另外，其他线程在 acquire 操作之后的读取也能看到本线程Release之前的写入。seq_cst 提供所有线程的一致性，用于多生产-多消费场景，即多线程、多个atomic变量，保证全局顺序一致性。 在 x86 TSO （Total Store Order）模型下，硬件已经保证了 LoadLoad、LoadStore、StoreStore 顺序，只允许 StoreLoad 重排。因此 acquire 和 release 在 x86 上不需要任何硬件屏障指令，只需编译器屏障（asm volatile(\"\" ::: \"memory\")）来阻止编译器重排。 在 ARM 弱内存模型下，所有四种重排都可能发生，因此需要显式的硬件指令（LDAR/STLR/DMB）。 seq_cst 在 x86 上需要额外处理 StoreLoad 重排：编译器在 store 端使用 XCHG（隐含 LOCK）或 MOV + MFENCE，load 端使用普通 MOV。 参见：https://www.felixcloutier.com/x86/sfence，https://www.felixcloutier.com/x86/lfence，https://www.felixcloutier.com/x86/mfence。3.1. release/acquire 如何实现局部同步release/acquire 提供的是成对的、局部的同步关系：当线程 B 的 acquire load 读到线程 A 的 release store 写入的值时，线程 A 在 release store 之前的所有写入，对线程 B 在 acquire load 之后的所有读取都可见。概念层面：Thread A (producer): Thread B (consumer): data = 42; while (!(p = flag.load(acquire))); str = \"Hello\"; // 以下读取能看到 Thread A 在 release 之前的所有写入 flag.store(true, release); ──→ assert(data == 42); // ✅ 保证成立 ↑ assert(*p == \"Hello\"); // ✅ 保证成立 | └── release 保证：之前的写入（data, str） 不会被重排到此 store 之后示例代码：#include \u0026lt;thread\u0026gt;#include \u0026lt;atomic\u0026gt;#include \u0026lt;cassert\u0026gt;#include \u0026lt;string\u0026gt;std::atomic\u0026lt;std::string*\u0026gt; ptr{nullptr};int data{42};void producer() { std::string* p = new std::string(\"Hello\"); data = 42; ptr.store(p, std::memory_order_release);}void consumer() { std::string* p2; while (nullptr == (p2 = ptr.load(std::memory_order_acquire))); assert(*p2 == \"Hello\"); // never fires assert(data == 42); // never fires}int main() { std::thread t1(producer); std::thread t2(consumer); t1.join(); t2.join(); return 0;}x86-64 编译器实际生成的汇编（GCC/Clang）：;; release store (x86-64);; ptr.store(p, std::memory_order_release); (ptr 是 std::atomic\u0026lt;std::string*\u0026gt;)mov QWORD PTR ptr[rip], rax ; 普通 MOV 指令！ ; x86 TSO 已保证 StoreStore 顺序 ; 编译器只需确保不重排（编译器屏障）;; acquire load (x86-64);; std::string* p = ptr.load(std::memory_order_acquire);mov rax, QWORD PTR ptr[rip] ; 普通 MOV 指令！ ; x86 TSO 已保证 LoadLoad、LoadStore 顺序ARM64 编译器生成的汇编（GCC/Clang）：;; release store (ARM64);; ptr.store(p, std::memory_order_release); (ptr 是 std::atomic\u0026lt;std::string*\u0026gt;)stlr x0, [x1] ; Store-Release 指令（64位指针用 x 寄存器） ; 硬件保证：之前的读写不会重排到此 store 之后;; acquire load (ARM64);; std::string* p = ptr.load(std::memory_order_acquire);ldar x0, [x1] ; Load-Acquire 指令（64位指针用 x 寄存器） ; 硬件保证：之后的读写不会重排到此 load 之前从硬件角度理解 release/acquire 的作用：Release Store 的效果： ┌──────────────────────────────┐ │ 之前的所有 store 已经进入 │ │ Store Buffer（x86: FIFO │ │ 保证顺序提交到 cache） │ ├──────────────────────────────┤ │ Release Store 本身 │ ← 保证在之前的 store 之后可见 └──────────────────────────────┘ x86: Store Buffer 是 FIFO，天然保证 StoreStore 顺序 ARM: STLR 指令包含隐式屏障，阻止重排Acquire Load 的效果： ┌──────────────────────────────┐ │ Acquire Load 本身 │ ← 此时可能触发 cache miss， │ │ 从其他 CPU 的 cache 获取最新值 ├──────────────────────────────┤ │ 之后的 load/store │ ← 保证在 acquire load 之后执行 │ 能看到 release 之前的写入 │ └──────────────────────────────┘ x86: 硬件已保证 LoadLoad/LoadStore 顺序，Invalidate Queue 中的 失效消息在 load 前被处理（x86 不真正延迟处理 invalidation） ARM: LDAR 指令包含隐式屏障，且强制处理 Invalidate Queue3.2. memory_order_seq_cst 如何实现全局同步seq_cst 比 release/acquire 更强：它提供一个全局唯一的操作顺序（single total order），所有线程看到的 seq_cst 操作顺序完全一致。x86-64 编译器实际生成的汇编（GCC/Clang）：;; seq_cst store (x86-64);; ptr.store(p, std::memory_order_seq_cst); (ptr 是 std::atomic\u0026lt;std::string*\u0026gt;)xchg rax, QWORD PTR ptr[rip] ; XCHG 隐含 LOCK 前缀（64位指针用 QWORD/rax） ; 效果： ; 1. 排空 Store Buffer（之前的 store 全部提交） ; 2. 原子地写入新值 ; 3. 等待本次写入被其他 CPU 的 cache 确认（ACK） ; 4. 隐含全屏障（相当于 MFENCE）;; 或者等价写法：mov QWORD PTR ptr[rip], rax ; 普通 storemfence ; 排空 Store Buffer + 阻止 StoreLoad 重排s;; seq_cst load (x86-64);; std::string* p = ptr.load(std::memory_order_seq_cst);mov rax, QWORD PTR ptr[rip] ; 普通 MOV 指令！（64位指针用 QWORD/rax） ; ⚠️ 注意：x86 上 seq_cst load 不需要 MFENCE ; 因为屏障已经放在 store 端，而 x86 TSO 保证 ; load 不会越过之前的 load（LoadLoad 有序）ARM64 编译器生成的汇编（GCC/Clang）：;; seq_cst store (ARM64);; ptr.store(p, std::memory_order_seq_cst); (ptr 是 std::atomic\u0026lt;std::string*\u0026gt;)stlr x0, [x1] ; Store-Release（64位指针用 x 寄存器，与 release 相同指令） ; ARMv8 中 STLR+LDAR 组合天然提供 seq_cst 语义;; seq_cst load (ARM64);; std::string* p = ptr.load(std::memory_order_seq_cst);ldar x0, [x1] ; Load-Acquire（64位指针用 x 寄存器，与 acquire 相同指令）为什么 x86 的 seq_cst 把屏障放在 store 端而不是 load 端？在 x86 TSO 下，唯一允许的重排是 StoreLoad（后面的 load 可能在前面的 store 提交到 cache 之前就执行了）。因此只需要在 store 之后 / load 之前插入屏障。编译器选择在 store 端处理，使用 XCHG（隐含 LOCK）或 MOV + MFENCE，这样 load 端可以保持为普通 MOV 指令，性能更好（因为 load 通常比 store 更频繁）。关键差异：Release Store（局部同步，成对使用）：-----------------------------------------CPU 0: store data=42 ; 进入 Store Buffer store flag=true ; release store，也进入 Store Buffer ; x86 FIFO 保证 data 先于 flag 提交 ↓ cache coherence (MESI)CPU 1: load flag (acquire) ; 如果读到 true， load data ; 则保证读到 42Seq_cst Store（全局同步，强制唯一顺序）：-----------------------------------------CPU 0: XCHG [x], 1 ; 隐含 LOCK 前缀 ↓ 1. 排空 Store Buffer（之前所有 store 提交到 cache） 2. 锁定 cache line，原子写入 3. 通过 MESI 发送 Invalidate 消息 4. 等待所有 CPU 的 Invalidate ACK ↓ 此时所有 CPU 都知道 x 已经被修改 ↓CPU 1,2,3: 后续的 seq_cst load 必须看到一致的全局顺序3.3. release/acquire 与 seq_cst 的本质区别：IRIW 问题release/acquire 只提供成对的局部同步，不保证全局顺序。经典的 IRIW（Independent Reads of Independent Writes） 问题展示了这一区别：std::atomic\u0026lt;int\u0026gt; x{0}, y{0};// Thread 1: 只写 xvoid thread1() { x.store(1, memory_order); }// Thread 2: 只写 yvoid thread2() { y.store(1, memory_order); }// Thread 3: 先读 x，再读 yvoid thread3() { int r1 = x.load(memory_order); // 读到 1 int r2 = y.load(memory_order); // 读到 0}// Thread 4: 先读 y，再读 xvoid thread4() { int r3 = y.load(memory_order); // 读到 1 int r4 = x.load(memory_order); // 读到 0} memory_order r1=1, r2=0, r3=1, r4=0 是否可能？ 原因 acquire/release ✅ 可能 没有全局顺序要求；Thread 3 和 Thread 4 可能各自看到不同的写入顺序 seq_cst ❌ 不可能 全局唯一顺序：如果 Thread 3 看到 x=1 在 y=1 之前，Thread 4 也必须看到相同的顺序 这就是 seq_cst 的额外开销所换来的保证：所有线程对所有 seq_cst 操作看到完全相同的执行顺序。4. 原子操作的硬件实现上面讨论的 memory order 解决的是可见性和顺序问题。而原子操作本身要解决的是不可分割性（atomicity） 问题——确保一个操作在观察者看来要么完全完成，要么完全没有发生。4.1. 原子性的基础：自然对齐的 load/store在大多数现代 CPU 上，自然对齐的基本类型（≤ 机器字长）的 load 和 store 天然就是原子的：x86-64: - 对齐的 1/2/4/8 字节 load/store 是原子的 - 即使是普通的 MOV 指令，只要地址对齐，就是原子操作 - 这就是为什么 relaxed load/store 在 x86 上就是普通 MOVARM64: - 对齐的 1/2/4/8 字节 LDR/STR 是原子的因此，std::atomic\u0026lt;int\u0026gt; 的 load() 和 store() 本质上不需要特殊的硬件支持来保证原子性——对齐的 MOV/LDR/STR 天然就是原子的。std::atomic 在这里主要提供的是编译器屏障和内存序语义。std::atomic（即使 relaxed）与普通变量的区别即使使用最弱的 memory_order_relaxed，std::atomic 仍然比普通变量多做了关键的几件事： 特性 普通变量 atomic（relaxed） 操作原子性 不保证（可能被拆分为多条指令） 保证不可分割 编译器优化 可合并、删除、重排 禁止合并/删除（每次操作都会发出指令） cache coherence 可能一直停留在寄存器中 触发 MESI 的 invalidate，保证跨核可见性 顺序保证 无 不保证顺序（relaxed），但保证最终可见 // 普通变量：编译器可能优化掉前一条int x;x = 42;x = 43; // 编译器可能直接只保留 x = 43// atomic：每条 store 都必须执行std::atomic\u0026lt;int\u0026gt; x;x.store(42, std::memory_order_relaxed); // 不会被优化掉x.store(43, std::memory_order_relaxed); // 也不会被优化掉 关键理解：atomic 的”原子性”和”内存序”是两个独立的概念。relaxed 只放弃了顺序保证，但原子性、编译器屏障（禁止优化）、cache coherence 触发一个都不少。4.2. Read-Modify-Write (RMW) 操作真正需要特殊硬件支持的是 RMW 操作：fetch_add、compare_exchange、exchange 等。这些操作需要在”读取-计算-写入”这个过程中保证没有其他 CPU 插入修改。x86 实现：LOCK 前缀;; std::atomic\u0026lt;int\u0026gt; x; x.fetch_add(1, relaxed);lock add DWORD PTR [rdi], 1 ; LOCK 前缀保证原子 Read-Modify-Write;; x.compare_exchange_strong(expected, desired);lock cmpxchg DWORD PTR [rdi], esi ; LOCK + CMPXCHG = 原子 CAS;; x.exchange(val);xchg DWORD PTR [rdi], esi ; XCHG 隐含 LOCK（不需要显式写） 以 lock cmpxchg 为例，它在硬件中不是拆成多条独立微指令，而是作为一组融合微操作（fused micro-op） 执行：CPU 在执行期间锁定资源，保证”比较-交换”整个过程不可被打断。具体流程：(1) 发出 LOCK# 信号或锁住 cache line → (2) 通过 MESI 抢占独占权限 → (3) 等待其他核 invalidation ACK → (4) 比较并条件写入 → (5) 完成后更新 MESI 状态为 Modified。LOCK 前缀的硬件实现（现代 x86）：┌──────────────────────────────────────────────────┐│ 1. 检查目标地址是否在 L1 Cache 中 ││ ├─ 是 → Cache Line Locking（锁定 cache line） ││ │ - 将 cache line 设为 Modified 状态 ││ │ - 在操作期间拒绝其他 CPU 的访问请求 ││ │ - 操作完成后释放 ││ └─ 否 → Bus Locking（锁定总线，较慢） ││ - 发出 LOCK# 信号 ││ - 阻止其他 CPU 访问内存 ││ - 现代 CPU 极少走到这条路径 │├──────────────────────────────────────────────────┤│ 2. 额外效果： ││ - LOCK 前缀指令隐含全内存屏障（等同 MFENCE） ││ - 排空 Store Buffer ││ - 即使使用 memory_order_relaxed， ││ x86 上的 RMW 也有 MFENCE 效果 │└──────────────────────────────────────────────────┘ 重要: x86 上 LOCK 前缀指令始终隐含全屏障。这意味着 x.fetch_add(1, relaxed) 在 x86 上的性能和 x.fetch_add(1, seq_cst) 几乎相同——都会生成 lock add。relaxed 的性能优势主要体现在 ARM 等弱内存模型架构上。ARM 实现：LL/SC 或 LSE 原子指令;; fetch_add (ARM64, 传统 LL/SC 方式).retry: ldxr w1, [x0] ; Load-Exclusive：加载值并标记为\"独占\" add w2, w1, #1 ; 计算新值 stxr w3, w2, [x0] ; Store-Exclusive：仅当独占标记仍有效时写入 cbnz w3, .retry ; 如果 stxr 失败（其他 CPU 写了同一 cache line），重试;; fetch_add (ARM64, ARMv8.1 LSE 原子指令) ldadd w1, w0, [x0] ; 硬件原子 load-add，无需重试循环LL/SC（Load-Linked / Store-Conditional）机制：┌───────────────┐ ┌───────────────┐│ CPU 0 │ │ CPU 1 ││ │ │ ││ LDXR [addr] │ │ ││ ↓ 标记 addr │ │ ││ 为 exclusive │ │ ││ │ │ STR [addr] ││ │ │ ↓ 清除 CPU0 ││ │ │ 的 exclusive ││ │ │ 标记 ││ STXR [addr] │ │ ││ ↓ 发现标记 │ │ ││ 已被清除 │ │ ││ → 返回失败 │ │ ││ → 跳回 LDXR │ │ ││ 重试 │ │ │└───────────────┘ └───────────────┘4.3. 三种内存序对 RMW 的影响（以 fetch_add 为例）;; x86-64: 三种内存序生成的汇编完全相同！;; x.fetch_add(1, relaxed / acq_rel / seq_cst)lock add DWORD PTR [rdi], 1 ; 都是 lock add（LOCK 隐含全屏障）;; ARM64: 不同内存序生成不同汇编;; relaxed: ldxr w1, [x0] add w2, w1, #1 stxr w3, w2, [x0] ; 普通 LL/SC;; acquire-release: ldaxr w1, [x0] ; Load-Acquire-Exclusive add w2, w1, #1 stlxr w3, w2, [x0] ; Store-Release-Exclusive;; seq_cst: ldaxr w1, [x0] ; Load-Acquire-Exclusive add w2, w1, #1 stlxr w3, w2, [x0] ; Store-Release-Exclusive dmb ish ; 额外的全屏障（某些实现）4.4. 总结对比┌──────────────┬─────────────────┬─────────────────────────┬────────────────────┐│ 操作类型 │ 保证什么 │ x86-64 实现 │ ARM64 实现 │├──────────────┼─────────────────┼─────────────────────────┼────────────────────┤│ relaxed │ 原子性 │ MOV │ LDR/STR ││ load/store │ │ │ │├──────────────┼─────────────────┼─────────────────────────┼────────────────────┤│ acquire │ 原子性 + │ MOV（编译器屏障） │ LDAR ││ load │ 后续操作不前移 │ │ │├──────────────┼─────────────────┼─────────────────────────┼────────────────────┤│ release │ 原子性 + │ MOV（编译器屏障） │ STLR ││ store │ 之前操作不后移 │ │ │├──────────────┼─────────────────┼─────────────────────────┼────────────────────┤│ seq_cst │ 原子性 + │ store: XCHG/MOV+MFENCE │ LDAR/STLR ││ load/store │ 全局唯一顺序 │ load: MOV │ (+ 可能额外 DMB) │├──────────────┼─────────────────┼─────────────────────────┼────────────────────┤│ RMW │ 原子性 + │ LOCK 前缀指令 │ LDXR/STXR ││ (any order) │ 不可分割的 │ (始终隐含全屏障) │ 或 LSE 原子指令 ││ │ 读-改-写 │ │ (屏障取决于 order) │└──────────────┴─────────────────┴─────────────────────────┴────────────────────┘参考资料 内存模型 谈谈 C++ 中的内存顺序 (Memory Order) 程序员的自我修养（⑫）：C++ 的内存顺序·中更多资料 C++ Concurrency In Action 2ed 中文翻译 C++ Concurrency in Action, 2nd Edition 现代CPU性能分析与优化 – 现代CPU设计 现代CPU性能分析与优化 – README 现代CPU性能分析与优化 – pdf" },

  { "title": "总结：MMU -- 包括 TLB 以及 Table Walk Unit ，以及内存 Page Table", "url": "/blog/2024/mmu-tlb-table-walk-unit-page-table/", "categories": "CPU",
  "tags": "Linux, CPU", "date": "2024-08-08", "content":
  "1. MMU 结构以及工作过程大多数使用MMU的机器采用内存分页机制，虚拟地址空间以页(Page)为单位，相应的，物理地址空间也被划分为页帧(Frame)。页帧必须与页保持相同的大小，通常为4KB，对于大页，页帧可以是2MB或1GB。大页一般用于服务器，用于系统分配大量数据，减少缺页中断的发生。MMU通过页表(Page Table)将虚拟地址映射到物理地址，页表存储在主存中，由系统内核创建及管理。MMU由两部分组成：TLB(Translate Look-aside Buffer)，以及Table Walk Unit： TLB (Translation Lookaside Buffer)：缓存最近使用的 VA 到 PA 的映射； Table Walk Unit：如果TLB没有命中CPU发出的VA，则由Table Walk Unit根据位于物理内存中的页表(Page Table)完成VA到PA的查找。CPU访问内存的时候，将VA发给MMU，MMU先在TLB中查找是否有对应的PA，如果有，则直接返回对应的PA；如果没有，则由Table Walk Unit根据位于物理内存中的页表完成查找。在上述过程中，如果Table Walk Unit没有找到对应的PA，则向CPU发出Page fault中断，CPU处理缺页中断（具体见后面章节描述）。MMU的工作过程图示：CPU发出的VA由两部分组成：VPN(Virtual Page Number) + offset。对应的，转换之后的物理地址也有两部分：页框号PFN(Physical Frame Number) + offset。1.1 从VA到PA，MMU处理流程图2. 页表结构现代CPU一般采用四级页表结构。以32位地址空间的二级页表为例，CPU发出的虚拟地址被拆分分为：页目录(Page Directory)、页表(Page Table)、页内偏移(Page Offset)三级（以4k为例，即12位）。VPN的最高10位用于索引页目录，紧接的10位用于索引页表索引，最低12位为页内偏移地址。拆分结构如下图所示：2.1 MMU查找过程MMU先根据一级页表的物理地址和一级页表Index去一级页表中找PTE，PTE中的地址不再是最终的物理地址，而是二级页表的物理地址。根据二级页表物理地址和二级页表index去二级页表中找PTE，此时PTE中的地址才是真实的物理地址。根据此物理地址和offset找到最终的物理内存地址。使用二级页表的好处是如果一级页表中的某一个PTE没有命中，那这一PTE对应的整个二级页表就不存在。2.2 进程页表与 Address Space ID操作系统会为每个进程分配一个页表，该页表使用物理地址存储。当进程使用类似malloc等需要映射代码或数据的操作时，操作系统会在随后马上修改页表以加入新的 物理内存。每次切换进程都需要进行TLB清理。这样会导致切换的效率变低。为了解决问题，TLB引入了ASID(Address Space ID)。ASID的范围是0-255。ASID由操作系统分配，当前进程的ASID值被写在ASID寄存器(使用CP15 c3访问)。TLB在更新页表项时也会将ASID写入TLB。MMU在查找TLB时， 只会查找TLB中具有相同ASID值的TLB行。且在切换进程时，TLB中被设置了ASID的TLB行不会被清理掉，当下次切换回来的时候还在。所以ASID的出现使得切换进程时不需要清理TLB中的所有数据，可以大大减少切换开销。3. 页表项(PTE) 与 TLB 中的标志位在进程的虚拟内存空间中，每一个虚拟内存页在页表中都有一个PTE与之对应，在32位系统中，每个PTE占用4个字节大小，其中保存了虚拟内存页背后映射的物理内存页的起始地址，以及进程访问物理内存的一些权限标识位。由于内核将整个物理内存划分为一页一页的单位，每个物理内存页大小为 4K，所以物理内存页的起始地址都是按照 4K 对齐的，也就导致物理内存页的起始地址的后 12 位全部是 0，我们只需要在PTE中存储物理内存地址的高 20 位就可以了，剩下的低 12 位可以用来标记一些权限位。一些标志位的含义如下： P(0)：映射的物理内存是否在内存中。值为0表示可能被换出，需要从磁盘中读取。此时，其他bit位存放的是物理内存页在磁盘中的位置； R/W(1)：值为0表示该物理页只读，针对该页面的写操作触发page fault异常。比如，使用fork创建子进程之后，父子进程内存空间完全一样，页表中的内容也是一样的，父子进程中的PTE均指向同一物理内存。此时，内核将父子进程中的PTE均修改为只读的，并将父子进程共同映射的这个物理内存引用计数+1。当子进程需要进行写操作时，在内核的page fault异常处理时，发现这个物理页面的引用计数大于1，说明是多进程共享同一个物理内存，将进行写时拷贝(Copy On Write， COW)，内核为子进程重新分配一个物理页，复制原有物理页中的内容到新物理页，减少物理页面的引用计数，并修改子进程PTE中的R/W标志位。 PCD(2)：Page Cache Disable，表示PTE指向的物理内存页面中的内容，是否可以被缓存到Cache中。在SOC异构处理器共享同一块物理内存时，可以使用PCD创建禁止Cache的PTE。 PWT(3): Page Write-Through，表示PTE指向的CPU Cache中的内容，是直接写入物理内存(Write-Through)，还是在Cache-line被中的内容被替换时写入物理内存(Write-Back)。4. 缺页处理过程当调用malloc时，内核不会立即分配物理内存，仅在进程的虚拟地址空间中创建VMA(vm_area_struct)，记录该段虚拟地址的范围和权限。此时PTE尚未建立或P位为0。当进程首次访问该虚拟地址时，MMU触发Page Fault异常，CPU陷入内核态。内核根据异常地址查找进程的VMA，按以下三种情况处理：情况一：非法地址 —— 地址不属于任何VMA，内核向进程发送SIGSEGV（Segmentation Fault），进程被终止。情况二：首次访问（页面尚未分配物理内存） —— PTE为空，内核分配物理页框、清零、建立PTE映射（读操作可先映射到共享零页，写时再分配，即COW）。情况三：页面被换出到 Swap —— PTE的P位为0，其余位记录了Swap位置。内核分配新物理页框，从Swap读回数据，更新PTE并置P位为1。4.1 物理内存分配完成之后PTE映射建立后，内核刷新TLB使旧缓存失效，然后返回用户态重新执行触发异常的指令。MMU此次通过Table Walk Unit查到有效PTE，完成VA → PA转换并缓存到TLB，CPU正常访问物理内存。参考 知乎 – 图解MMU 简书 – ARM体系架构——MMU 一步一图带你构建 Linux 页表体系 —— 详解虚拟内存如何与物理内存进行映射 CPU入门扫盲篇之MMU内存管理单元——万字长文带你搞定MMU\u0026amp;TLB\u0026amp;TWU 一步一图带你深入理解Linux物理内存 五万字 | 深入理解Linux内存管理 How Does the Memory Management Unit (MMU) Work with the Unix/Linux Kernel? Linux Kernel Documentation – Page Tables – MMU, TLB, and Page Faults" },

  { "title": "自定义 operator new， placement new，以及释放内存", "url": "/blog/2024/operator-new-placement-new/", "categories": "Cpp",
  "tags": "Cpp", "date": "2024-07-29", "content":
  "1. new 操作符(new operator)new 操作符做两件事：分配内存 + 调用构造函数初始化。2. operator new通常声明如下：operator new 操作符的职责仅仅是分配内存，操作符返回一个未经处理（raw）的指针，未初始化的内存。void* operator new(size_t size);调用方式如下：void *rawMemory = operator new(sizeof(string));注意，显式调用 operator new 时，几乎没有意义，因为不能显式调用构造函数。2.1 对指定类型的 operator new, operator delete 进行重载#include \u0026lt;cstddef\u0026gt;#include \u0026lt;iostream\u0026gt;// class-specific allocation functionsstruct X { static void* operator new(std::size_t count) { std::cout \u0026lt;\u0026lt; \"custom new for size \" \u0026lt;\u0026lt; count \u0026lt;\u0026lt; '\\n'; return ::operator new(count); } static void* operator new[](std::size_t count) { std::cout \u0026lt;\u0026lt; \"custom new[] for size \" \u0026lt;\u0026lt; count \u0026lt;\u0026lt; '\\n'; return ::operator new[](count); } static void operator delete(void* ptr, bool b){ std::cout \u0026lt;\u0026lt; \"custom placement delete called, b = \" \u0026lt;\u0026lt; b \u0026lt;\u0026lt; '\\n'; ::operator delete(ptr); }};int main() { X* p1 = new X; delete p1; X* p2 = new X[10]; delete[] p2;}3. placement newplacement new 是一种特殊的 new 操作符，它允许在已分配的内存上调用构造函数。alignas(T) unsigned char buf[sizeof(T)];// Construct a “T” object, placing it directly into your pre-allocated storage at memory address “buf”.T* tptr = new(buf) T;tptr-\u0026gt;~T(); // 显式调用析构函数参考 operator new, operator new[] operator delete, operator delete[] new expression – Placement new" },

  { "title": "学术绘图工具 -- Engauge Digitizer", "url": "/blog/2024/Academic-plot-tools/", "categories": "Tools",
  "tags": "Tools", "date": "2024-07-17", "content":
  "Engauge Digitizer github repo 介紹 Engauge Digitizer 工具 曲线图转数据工具软件(Engauge Digitizer)提取文献中的数据WebPlotDigitizer WebPlotDigitizer其他资料 手把手教你从曲线图中提取原始数据 chart_digitizer" },

  { "title": "OpenCL 环境准备及资料", "url": "/blog/2024/openCL-Install-and-Resources/", "categories": "OpenCL",
  "tags": "OpenCL, Cpp", "date": "2024-07-06", "content":
  "0. GPU驱动相关sudo add-apt-repository ppa:oibaf/graphics-drivers1. OpenCL 环境准备1.1 查看 OpenCL 设备# 查看 GPU 设备lspci | grep -i vgasudo apt install clinfo使用 clinfo 命令查看 OpenCL 版本。C++ 程序需要定义OpenCL版本：target_compile_definitions(${target_name} PRIVATE CL_TARGET_OPENCL_VERSION=300)1.2 安装 OpenCL SDKsudo apt install libstb-dev libsfml-dev libglew-dev libglm-dev libtclap-dev ruby doxygen -y# 安装支持包：包括ICD Loader，SDK及头文件sudo apt install ocl-icd-opencl-devgit clone https://github.com/KhronosGroup/OpenCL-SDK.git --recursive# 编译 \u0026amp; 安装 OpenCL SDK ..... 额外安装：Intel OpenCL Runtime，用于支持在CPU上运行OpenCL程序（模拟GPU ??），需要先安装OpenCL Loader(Installable Client Driver Loader，ICD Loader)。参考 ubuntu安装OpenCL运行及编译环境。1.3 参考 Getting started with OpenCL on Ubuntu Linux CMake Build System Support Khronos Community – OpenCL Intel Community – OpenCL* for CPU2. 资源 官方 blog Exploiting Task Parallelism with OpenCL: A Case Study OpenCL – GPU 设备信息查询3. 安装 Intel OpenCL 支持3.1 安装 Intel OpenCL Runtime for GPU安装步骤，以及遇到问题的解决办法： 安装 Intel OpenCL Runtime：github – Intel OpenCL Runtime Intel ARC intel-i915-dkms dpkg error upgrading to HWE kernel 6.5 How To Deploy OpenCL™ Code on Intel® Hardware openCL on Ubuntu still can’t detect my intel graphics platform after I install some relevant drivers3.2 安装额外支持包安装如下包 (可能不是必须的)：sudo apt-get install xserver-xorg-video-intelsudo apt-get install mesa-utils3.3 使能 User 用户权限权限使能，并重启系统。(如果GPU已经使能，不添加权限只能使用sudo查看到GPU设备)sudo usermod -a -G render $USERsudo usermod -a -G video $USER3.4 验证clinfo -lPlatform #0: Intel(R) OpenCL `-- Device #0: Intel(R) Core(TM) i5-8260U CPU @ 1.60GHzPlatform #1: Intel(R) OpenCL Graphics `-- Device #0: Intel(R) UHD Graphics 6203.5 Intel OpenCL 调试Intel OpenCL 调试，需要作一些设置。安装驱动以及调试符号等，对GPU硬件有要求（如要求GPU gen 12以上）；以及不同GPU硬件，支持的Ubuntu等linux发行版也不同。 Get Started with Intel® SDK for OpenCL™ Applications 2020 on Linux* OS with Training Sample Get Started with Intel® Distribution for GDB* on Linux* OS Host4. 升级 OpenCL 驱动，支持 OpenCL 3.0# add PPAsudo add-apt-repository ppa:kisak/kisak-mesa# install MESA packagessudo apt updatesudo apt full-upgrade恢复默认MESA：sudo apt install ppa-purgesudo ppa-purge ppa:kisak/kisak-mesa How to Install The Latest Mesa Graphics Driver in Ubuntu 20.04 / 21.045. AMD Windows 平台安装vcpkg安装及OpenCL依赖包：# 使用命令行终端，不能使用powershell终端vcpkg --triplet x64-windows install sfml tclap glm编译OpenCL：git clone --recursive https://github.com/KhronosGroup/OpenCL-SDK.git# 需要指定 vcpkg.cmake 路径cmake -D CMAKE_TOOLCHAIN_FILE=D:\\dev_libs\\vcpkg\\scripts\\buildsystems\\vcpkg.cmake -D VCPKG_TARGET_TRIPLET=x64-windows -D BUILD_TESTING=OFF -D BUILD_DOCS=OFF -D BUILD_EXAMPLES=OFF -D BUILD_TESTS=OFF -D OPENCL_SDK_BUILD_SAMPLES=ON -D OPENCL_SDK_TEST_SAMPLES=OFF -D CMAKE_INSTALL_PREFIX=D:\\dev_libs\\opencl .." },

  { "title": "Redis常用命令总结", "url": "/blog/2024/redis-commands/", "categories": "Redis",
  "tags": "Redis", "date": "2024-07-05", "content":
  "1. redis monitor# 运行一下命令，进入monitor模式，可以实时查看redis的命令执行情况redis-cli monitormoniotor模式下，每执行一条redis命令，就会有类似如下输出：1720185799.917896 [0 127.0.0.1:43768] \"COMMAND\" \"DOCS\"1720185984.438276 [0 127.0.0.1:43768] \"set\" \"mykey\" \"hello\" \"EX\" \"60\"1720186045.464191 [0 127.0.0.1:43768] \"set\" \"mykey\" \"KEEPTTL\"1720186089.705980 [0 127.0.0.1:43768] \"set\" \"mykey\" \"hello\" 开启monitor之后，对性能有较大影响。2. redis 设置 key-value# 使用 redis-cli 命令进入redis命令行redis-cli127.0.0.1:6379\u0026gt; set mykey helloOK127.0.0.1:6379\u0026gt; get mykey\"hello\"参考 监控 redis 执行命令 redis 介绍和常用命令 Redis 常用命令及示例总结 如何查看 修改 Redis 密码" },

  { "title": "在 Markdown 中使用数学公式", "url": "/blog/2024/%E5%9C%A8Markdown%E4%B8%AD%E4%BD%BF%E7%94%A8%E6%95%B0%E5%AD%A6%E5%85%AC%E5%BC%8F/", "categories": "Markdown",
  "tags": "Markdown", "date": "2024-07-02", "content":
  "1. 在字符中添加空格有四种宽度的空格可以使用，如下表格： 语法 显示 \\, a b \\; a b \\quad a b \\qquad a b 一个示例如下：\\[\\begin{cases}(H_{y1} \u0026lt; y_i \u0026lt; H_{y2}) \\;and\\; (H_{y1} \u0026lt; y_o \u0026lt; H_{y2}) \\\\(50^\\circ \u0026lt; |K_i| \u0026lt; 90^\\circ) \\;and\\; (50^\\circ \u0026lt; |K_o| \u0026lt; 90^\\circ) \\\\\\sqrt{(x_i - x_o)^2 + (y_i - y_o)^2} \\leq L_{smin} \\\\k_i \\times k_o \u0026gt; 0 \\\\\\Delta x \u0026gt; \\Delta y \\\\x_i \u0026gt; x_o \u0026amp; \\text{后向探头数据} \\\\x_i \u0026lt; x_o \u0026amp; \\text{前向探头数据}\\end{cases}\\]更多数学公式的使用： 如何优雅地在Markdown中输入数学公式 Markdown 数学公式指导手册" },

  { "title": "Markdown 使用笔记", "url": "/blog/2024/Markdown%E8%AF%AD%E6%B3%95%E7%AC%94%E8%AE%B0/", "categories": "Markdown",
  "tags": "Markdown", "date": "2024-07-02", "content":
  "1. Markdown 转义字符列表使用反斜杠（”\\“）可以转义 Markdown 中的特殊字符，使其被当作普通文本显示。以下是 Markdown 中的特殊字符列表：\\ backslash` backtick* asterisk_ underscore{} curly braces[] square brackets() parentheses# hash mark+ plus sign- minus sign (hyphen). dot! exclamation mark理论上， “\\” 可以转义任何字符，如果该字符不是特殊字符也会原样输出。这里的理论上，其实说的是转义直接出现在正文中。2. 在 Markdown 中插入特殊符号2.1 任务列表符号 符号 用途 ✅ ☑ ✓ ✔ √ 完成、正确、同意 ❌ ☒ ✘ ✕ ✖ 未完成、错误、否决 ⏳ ⌛ 等待、进行中 2.2 状态与进度 类别 符号 说明 进度 ⏳ ⌛ ✅ 🚧 🔄 等待、完成、进行、循环 优先级 🔴 🟡 🟢 高、中、低 标记 📌 📍 🔖 重要、位置、书签 2.3 工具与操作 符号 用途 🔧 ⚙️ 🔩 工具、设置、配置 🔍 🔎 搜索、查看、浏览 📁 📂 📋 文件夹、文件、文档 🔗 ⛓️ 链接、连接、关联 2.4 提示与注意 符号 用途 📢 📣 通知、宣布、重要 ⚠️ 🚨 警告、注意、危险 💡 💭 建议、想法、思考 ℹ️ 📝 信息、说明、备注 🔍 📊 分析、详解、查看 2.5 美化装饰符号 符号 说明 🌹 🍀 🌙 🍂 🍃 🌷 自然、植物 💎 🔥 ⭐ 🍄 🏆 宝贵、热门、优秀 ↔️ ➡️ ⬅️ ⬆️ ⬇️ 方向箭头 ✨ 🌟 💫 闪耀、高亮" },

  { "title": "使用inotify监控文件目录中的文件变化（新建文件）", "url": "/blog/2024/monitor-files-change-in-dir-using-inotify/", "categories": "Cpp",
  "tags": "Cpp, Linux", "date": "2024-06-24", "content":
  "通过结合使用epoll和inotify 实现监控功能的同时，以超时的方式实现轮询，适合线程退出。int inotifyId = inotify_init();if (-1 == inotifyId) { SPDLOG_WARN(\"inotify_init failed\"); return;}int epfd = epoll_create(INOTIFY_FDS);if (-1 == epfd) { SPDLOG_WARN(\"epoll_create failed\"); return;}struct epoll_event ev;ev.data.fd = inotifyId;ev.events = EPOLLIN | EPOLLET;int ret = epoll_ctl(epfd, EPOLL_CTL_ADD, inotifyId, \u0026amp;ev);if (-1 == ret) { SPDLOG_WARN(\"epoll_ctl failed\"); return;}const char* pathName = \"\u0026lt;dir of files to monitor\u0026gt;\";const uint32_t watch_mask = (IN_CLOSE_WRITE | IN_MODIFY | IN_ATTRIB);int watchFd = inotify_add_watch(inotifyId, pathName, watch_mask);if (watchFd \u0026lt; 0) { SPDLOG_WARN(\"inotify_add_watch failed\"); return;}SPDLOG_INFO(\"start monitor path: {}\", pathName);// 循环监听事件{ /* Some systems cannot read integer variables if they are not properly aligned. On other systems, incorrect alignment may decrease performance. Hence, the buffer used for reading from the inotify file descriptor should have the same alignment as struct inotify_event. */ char buf[INOTIFY_BUF_LEN] = {} __attribute__ ((aligned(__alignof__(struct inotify_event)))); struct epoll_event events[20]; while (runningFlag) { int nfds = epoll_wait(epfd, events, 20, 100); if (nfds \u0026lt;= 0) { if (nfds \u0026lt; 0 \u0026amp;\u0026amp; errno != EINTR) { SPDLOG_WARN(\"epoll_wait error: {}\", strerror(errno)); } continue; } if (events[0].data.fd != inotifyId) { SPDLOG_WARN(\"epoll_wait error: events[0].data.fd!= inotifyId\"); continue; } int nread{}, length = read(inotifyId, buf, INOTIFY_BUF_LEN - 1); if (length \u0026lt; 0) { SPDLOG_WARN(\"read error: {}\", strerror(errno)); continue; } while (length \u0026gt; 0) { // 解析事件 } }}// 退出清理inotify_rm_watch(inotifyId, watchFd);close(epfd);close(inotifyId);参考 Linux C 使用 inotify 监控文件或目录变化 linux 文件监控之 inotify inotify(7) — Linux manual page Monitor file system activity with inotify" },

  { "title": "总结：使用 gperftools 进行性能分析", "url": "/blog/2024/gperf-tools/", "categories": "Cpp",
  "tags": "Cpp, Linux, Performance", "date": "2024-06-21", "content":
  "1. 安装 gperftools使用 gperftools Release页面 下载（不要使用git clone，且不要使用cmake编译，编译不生成pprof工具），编译命令：./configuremake \u0026amp;\u0026amp; make install安装graphviz：sudo apt-get install graphviz2. 将 gperftools 链接进待测试程序如何在CMake中查找gperftools的路径，参考笔记 CMake 编写FindPackage 模块: posts/2024-06-21-cmake-find_package.md，添加自定义CMake Find Package模块，并在CMakeLists.txt中添加find_package(gperftools)，即可找到gperftools的路径。编译脚本添加如下：option(ENABLE_PROFILER \"Enable google perftools\" ON)message(STATUS \"ENABLE_PROFILER: ${ENABLE_PROFILER}\")if(ENABLE_PROFILER) set(CMAKE_MODULE_PATH \"${PROJECT_SOURCE_DIR}/cmake;${CMAKE_MODULE_PATH}\") find_package(Gperftools REQUIRED) set(PROFILER_LIBS ${GPERFTOOLS_PROFILER_LIBRARY}) message(STATUS \"PROFILER_LIBS: ${PROFILER_LIBS}\") add_definitions(\"-DHAVE_PROFILER\")else() set(PROFILER_LIBS \"\")endif()使能frame-pointer：if(ENABLE_PROFILER) message(STATUS \"enable profiler\") target_compile_options(${target_test} PRIVATE -fno-omit-frame-pointer) target_link_options(${target_test} PRIVATE -fno-omit-frame-pointer) # set(CMAKE_CXX_FLAGS \"${CMAKE_CXX_FLAGS} -fno-omit-frame-pointer\") # set(CMAKE_LINKER_FLAGS \"${CMAKE_LINKER_FLAGS} -fno-omit-frame-pointer\")endif()3. 使用 gperftools 分析程序性能运行被测试程序：CPUPROFILE=test_flow_benchmark.prof CPUPROFILE_FREQUENCY=500 ./test --gtest_filter=FlowBenchMarkTest10/FlowBenchMarkTest.*# 生成pdf报告pprof --pdf ./server server.prof \u0026gt; perf.pdf4. 使用 GNU gprof 分析程序性能target_compile_options(${target_test} PRIVATE -pg -g) target_link_options(${target_test} PRIVATE -pg -g)gprof 不支持多线程应用，多线程下只能采集主线程性能数据。多线程需要重写pthread_create()。参考：Linux性能优化gprof使用 。参考 gperftools gperftools文档 – CPU profiler 使用 gperftools 分析程序性能 Profiler – 链接选项 Linux下使用gperftools gperftools 的安装与使用" },

  { "title": "CMake 编写FindPackage 模块", "url": "/blog/2024/cmake-find_package/", "categories": "CMake",
  "tags": "Cpp, CMake", "date": "2024-06-21", "content":
  "例如编写CMake查找模块，名称为 Findgperftools.cmake，内容如下：# Try to find gperftools# Once done, this will define## gperftools_FOUND - system has Profiler# GPERFTOOLS_INCLUDE_DIR - the Profiler include directories# Tcmalloc_INCLUDE_DIR - where to find Tcmalloc.h# GPERFTOOLS_TCMALLOC_LIBRARY - link it to use tcmalloc# GPERFTOOLS_TCMALLOC_MINIMAL_LIBRARY - link it to use tcmalloc_minimal# GPERFTOOLS_PROFILER_LIBRARY - link it to use Profiler# TCMALLOC_VERSION_STRING# TCMALLOC_VERSION_MAJOR# TCMALLOC_VERSION_MINOR# TCMALLOC_VERSION_PATCHfind_path(GPERFTOOLS_INCLUDE_DIR gperftools/profiler.h HINTS $ENV{GPERF_ROOT}/include)find_path(Tcmalloc_INCLUDE_DIR gperftools/tcmalloc.h HINTS $ENV{GPERF_ROOT}/include)if(Tcmalloc_INCLUDE_DIR AND EXISTS \"${Tcmalloc_INCLUDE_DIR}/gperftools/tcmalloc.h\") foreach(ver \"MAJOR\" \"MINOR\" \"PATCH\") file(STRINGS \"${Tcmalloc_INCLUDE_DIR}/gperftools/tcmalloc.h\" TC_VER_${ver}_LINE REGEX \"^#define[ \\t]+TC_VERSION_${ver}[ \\t]+[^ \\t]+$\") string(REGEX REPLACE \"^#define[ \\t]+TC_VERSION_${ver}[ \\t]+(\\\".)?([0-9]*)\\\"?$\" \"\\\\2\" TCMALLOC_VERSION_${ver} \"${TC_VER_${ver}_LINE}\") unset(TC_VER_${ver}_LINE) endforeach() set(TCMALLOC_VERSION_STRING \"${TCMALLOC_VERSION_MAJOR}.${TCMALLOC_VERSION_MINOR}\") if(NOT TCMALLOC_VERSION_PATCH STREQUAL \"\") set(TCMALLOC_VERSION_STRING \"${TCMALLOC_VERSION_STRING}.${TCMALLOC_VERSION_PATCH}\") endif()endif()foreach(component tcmalloc tcmalloc_minimal profiler) string(TOUPPER ${component} COMPONENT) find_library(GPERFTOOLS_${COMPONENT}_LIBRARY ${component} HINTS $ENV{GPERF_ROOT}/lib) list(APPEND GPERFTOOLS_LIBRARIES GPERFTOOLS_${COMPONENT}_LIBRARY)endforeach()set(_gperftools_FIND_REQUIRED_VARS \"GPERFTOOLS_INCLUDE_DIR\")if(gperftools_FIND_COMPONENTS) foreach(component ${gperftools_FIND_COMPONENTS}) string(TOUPPER ${component} COMPONENT) list(APPEND _gperftools_FIND_REQUIRED_VARS \"GPERFTOOLS_${COMPONENT}_LIBRARY\") endforeach()else() list(APPEND _gperftools_FIND_REQUIRED_VARS \"GPERFTOOLS_LIBRARIES\")endif()include(FindPackageHandleStandardArgs)find_package_handle_standard_args(gperftools FOUND_VAR gperftools_FOUND REQUIRED_VARS ${_gperftools_FIND_REQUIRED_VARS} VERSION_VAR TCMALLOC_VERSION_STRING)mark_as_advanced(${GPERFTOOLS_LIBRARIES} GPERFTOOLS_INCLUDE_DIR)if(gperftools_FOUND) foreach(component tcmalloc tcmalloc_minimal profiler) if(NOT (TARGET gperftools::${component})) string(TOUPPER ${component} COMPONENT) add_library(gperftools::${component} UNKNOWN IMPORTED) set_target_properties(gperftools::${component} PROPERTIES INTERFACE_INCLUDE_DIRECTORIES \"${GPERFTOOLS_INCLUDE_DIR}\" IMPORTED_LINK_INTERFACE_LANGUAGES \"CXX\" IMPORTED_LOCATION \"${GPERFTOOLS_${COMPONENT}_LIBRARY}\") endif() endforeach() foreach(component tcmalloc tcmalloc_minimal) if(NOT (TARGET gperftools::${component})) set_target_properties(gperftools::${component} PROPERTIES INTERFACE_COMPILE_OPTIONS \"-fno-builtin-malloc -fno-builtin-calloc -fno-builtin-realloc -fno-builtin-free\") endif() endforeach()endif()在CMakeLists.txt中，如下使用：option(ENABLE_PROFILER \"Enable google perftools\" ON)message(STATUS \"ENABLE_PROFILER: ${ENABLE_PROFILER}\")if(ENABLE_PROFILER) set(CMAKE_MODULE_PATH \"${PROJECT_SOURCE_DIR}/cmake;${CMAKE_MODULE_PATH}\") find_package(Gperftools REQUIRED) set(PROFILER_LIBS ${GPERFTOOLS_PROFILER_LIBRARY}) message(STATUS \"PROFILER_LIBS: ${PROFILER_LIBS}\") add_definitions(\"-DHAVE_PROFILER\")else() set(PROFILER_LIBS \"\")endif()参考 使用find_package引入外部依赖包 CMakeLists.txt" },

  { "title": "简单线程安全队列", "url": "/blog/2024/simple-thread-safe-queue/", "categories": "Cpp",
  "tags": "Cpp", "date": "2024-06-20", "content":
  "#include \u0026lt;iostream\u0026gt;#include \u0026lt;queue\u0026gt;#include \u0026lt;mutex\u0026gt;template \u0026lt;typename T\u0026gt;class ThreadSafeQueue {public: void push(const T\u0026amp; item) { std::lock_guard\u0026lt;std::mutex\u0026gt; lock(mutex_); queue_.push(item); } bool try_pop(T\u0026amp; item) { std::lock_guard\u0026lt;std::mutex\u0026gt; lock(mutex_); if (queue_.empty()) { return false; } item = queue_.front(); queue_.pop(); return true; } bool empty() const { std::lock_guard\u0026lt;std::mutex\u0026gt; lock(mutex_); return queue_.empty(); }private: std::queue\u0026lt;T\u0026gt; queue_; mutable std::mutex mutex_;};" },

  { "title": "使能 C++ 程序的核心转储", "url": "/blog/2024/enable-core-dump-for-cpp/", "categories": "Cpp",
  "tags": "Cpp", "date": "2024-06-19", "content":
  "1. 使能 core dump# 查看是否使能 core dump, -a 显示所有设置sudo ulimit -c# 使能 core dump，不限制core dump文件大小sudo ulimit -c unlimited# 限制 core dump 文件大小为 2Gsudo ulimit -c 4194304# 关闭 core dumpsudo ulimit -c 02. 修改 core dump 文件位置临时修改 core dump 文件位置：# 修改 core dump 文件位置为 /tmp/corefile，以及格式echo /tmp/corefile/core-%e-%p-%t \u0026gt; /proc/sys/kernel/core_pattern永久修改 core dump 文件位置：# /etc/sysctl.confkernel.core_pattern = /tmp/corefile/core-%e-%p-%tkernel.core_uses_pid = 0# 生效sysctl –p /etc/sysctl.confcore dump文件格式：%p - insert pid into filename 添加pid(进程id)%u - insert current uid into filename 添加当前uid(用户id)%g - insert current gid into filename 添加当前gid(用户组id)%s - insert signal that caused the coredump into the filename 添加导致产生core的信号%t - insert UNIX time that the coredump occurred into filename 添加core文件生成时的unix时间%h - insert hostname where the coredump happened into filename 添加主机名%e - insert coredumping executable name into filename 添加导致产生core的命令名3. core dump 管理工具：coredumpctl# 查看 core dump 列表sudo coredumpctl list# 显示指定 core dump 文件信息sudo coredumpctl info core_filename# 分析 core dump 文件sudo coredumpctl gdb core_filename4. 使用 gdb 调试 core dump 文件编译 core dump 文件时，需要加上 -g 参数。使用file命令查看 core dump 文件是由哪个可执行文件产生的：4.1 调试 core dump 文件gdb program core（gdb + 可执行文件 +core文件）5. 参考 coredump文件生成，以及GDB工具使用" },

  { "title": "Python venv 环境搭建及 VSCode 环境配置", "url": "/blog/2024/VSCode%E8%BF%9B%E8%A1%8CPython%E5%BC%80%E5%8F%91%E7%8E%AF%E5%A2%83%E8%AE%BE%E7%BD%AE/", "categories": "Python",
  "tags": "VSCode, Python", "date": "2024-06-10", "content":
  "1. Python venv 环境搭建1.1. venv安装pip install virtualenv# 设置永久国内 pip 源pip config set global.index-url https://pypi.tuna.tsinghua.edu.cn/simple/pip config set install.trusted-host pypi.tuna.tsinghua.edu.cn1.2. 创建venv环境# python -m venv \u0026lt;目录\u0026gt;python -m venv .venv1.3. 激活venv环境Linux/Mac：source ./.venv/bin/activateWindows：.\\.venv\\Scripts\\activate# 退出虚拟环境.\\.venv\\Scripts\\deactivate.bat1.4. 安装依赖# pip install -r requirements.txt# pip install pygame2. vscode 环境配置2.1. settings.json配置需要安装插件： Python Pylance Yapf\"terminal.integrated.env.linux\": { \"PYTHONPATH\": \"${workspaceFolder}/python;${env:PYTHONPATH}\"},\"python.envFile\": \"${workspaceFolder}/.env\",\"python.analysis.extraPaths\": [ \"${workspaceFolder}/python\"],\"[python]\": { \"diffEditor.ignoreTrimWhitespace\": false, \"editor.defaultFormatter\": \"eeyore.yapf\", \"editor.formatOnSaveMode\": \"file\", \"editor.formatOnSave\": true, \"editor.indentSize\": 2, \"editor.wordBasedSuggestions\": \"off\", \"files.trimTrailingWhitespace\": true,},\"python.languageServer\": \"Pylance\",\"yapf.args\": [\"--style\", \"{based_on_style: pep8, indent_width: 2}\"],2.2. extensions.json配置{ \"recommendations\": [\"eeyore.yapf\", \"dangmai.workspace-default-settings\", \"ms-python.flake8\", \"ms-python.isort\", \"ms-python.python\"], \"unwantedRecommendations\": [\"ms-python.black-formatter\", \"ms-python.pylint\"]}附件见 python_prj/.vscode.zip 。" },

  { "title": "nginx 用户的权限配置", "url": "/blog/2024/nginx-autho-config/", "categories": "Nginx",
  "tags": "Nginx", "date": "2024-06-06", "content":
  "# nginx 解决 403 错误问题sudo setenforce Permissivesudo setsebool -P httpd_can_network_connect onsudo chcon -Rt httpd_sys_content_t /usr/share/nginx/html/nginx不能启动，提示端口不能绑定 8090 端口查看服务状态信息，提示如下信息：nginx: [emerg] bind() to 0.0.0.0:80 failed (13: permission denied)# 查看是否在http_port_t类型下sudo semanage port -l | grep http_port_t# 如果不在，添加到http_port_t类型下sudo semanage port -a -t http_port_t -p tcp 8090学习资料 nginx 安装配置及使用 启动权限拒绝问题" },

  { "title": "VSCode Remote SSH 使用局域网代理", "url": "/blog/2024/VSCode-remotessh-%E5%B1%80%E5%9F%9F%E7%BD%91%E4%BB%A3%E7%90%86/", "categories": "Bash",
  "tags": "Bash, VSCode", "date": "2024-06-06", "content":
  "配置路径：远程[SSH:xxxx] -\u0026gt; 应用程序 -\u0026gt; Proxy。填入代理地址以及端口。另外，去掉Proxy Strict SSL复选框的选中，即将Copilot配置为忽略证书错误。" },

  { "title": "VSCode Remote SSH 免密登陆", "url": "/blog/2024/VSCode-remotessh-%E5%85%8D%E5%AF%86%E7%99%BB%E9%99%86/", "categories": "Bash",
  "tags": "Bash, VSCode", "date": "2024-06-06", "content":
  "cat id_rsa.pub \u0026gt;\u0026gt; ~/.ssh/authorized_keyschmod 600 ~/.ssh/authorized_keyschmod 700 ~/.sshservice sshd restart" },

  { "title": "OpenGL 流水线（待消化吸收）", "url": "/blog/2024/opengl-pipeline/", "categories": "OpenGL",
  "tags": "Cpp, OpenGL", "date": "2024-05-30", "content":
  "1. 流水线概念2. 流水线例子参考 GLSL教程（一）图形流水线" },

  { "title": "OCC boolean operations", "url": "/blog/2024/OCCT-boolean-ops/", "categories": "OCCT",
  "tags": "OCCT", "date": "2024-05-22", "content":
  "OpenCASCADE 布尔运算简介 Boolean Operations OpenCascade功能及模块简介个人理解：由于OCCT采用BRep（边界表示）表示方式，在做boolean操作时，计算出诸如edge/edge，face/face的交界面/点等，并存储供后面使用。1. 干涉检测计算出Objects与Tools的各种干涉类型：如vertex与vertex/edge/face/solid之间的干涉数据，edge与edge/face/solid之间的干涉数据，face与face/solid之间的干涉数据。主要是edge与edge之间的干涉检查。存储的干涉数据以P_curves，Pave形式存储。P_Curves是指比如一条边与一个面相交，需要在面上构建一条曲线来表示这条相交边。P_Curves存储边与面的相交信息（如交点、切线）。Pave是表示边界点的数据结构，出了位置外，还包含了曲率、方向等信息。在boolean过程中，需要用Pave存储相交信息。Pave是构建P_Curves（以及其他数据）的基础。2. General Fuse Algorithm (GFA)步骤如下：init -\u0026gt; calc vertex/vertex interferences -\u0026gt; calc vertex/edge interferences -\u0026gt; update Pave block -\u0026gt; calc edge/edge interferences -\u0026gt; calc vertex/face interferences -\u0026gt; calc edge/face interferences -\u0026gt; build split edges（构建分割边）-\u0026gt; calc face/face interferences -\u0026gt; build Section Edges（构建截面边）-\u0026gt; build P_Curves -\u0026gt;Process Degenerated Edges（处理退花边）。2.1 分割边其中，构建分割边，是指： 识别形状之间的交线 根据交线在Objects的形状上构建新的边 将这些新构建的边作为分割边 根据分割边，将Objects的形状划分为若干个子形状构建好分割边后,后续的布尔运算就可以基于这些分割后的子形状进行计算,得到最终的结果形状。2.2 截面边" },

  { "title": "CAM资料及笔记（待学习）", "url": "/blog/2024/CAM%E8%B5%84%E6%96%99%E5%8F%8A%E7%AC%94%E8%AE%B0/", "categories": "CAM",
  "tags": "CAM, OCCT", "date": "2024-05-22", "content":
  "1. CAM路径规划TODO1.1. CAM路径规划资料 数控加工中路径规划与速度插补综述：在线地址https://arxiv.org/pdf/2303.01368。 OpenCAMLib：开源CAM库，提供3D加工路径生成，适用于CNC车铣。 中科院副研究员袁春明：自由曲面刀具路径规划的现状和新进展 基于运动学约束的五轴机床路径规划与最优速度规划：论文，暂时未下载" },

  { "title": "Windows下安装 MSMPI", "url": "/blog/2024/windows-install-msmpi/", "categories": "MPI",
  "tags": "MPI", "date": "2024-05-21", "content":
  "下载安装包及环境变量设置下载安装msmpi以及msmpisdk：MSMPI github releases。注意安装路径不要有空格及中文。命令行设置：set MSMPI输出信息如下图所示：建立第一个测试程序hello_mpifind_package(MPI REQUIRED)# list(APPEND myMPI_INC_DIR $ENV{MSMPI_INC})# list(APPEND myMPI_LIBS $ENV{MSMPI_LIB64})message(STATUS \"MPI_FOUND=${MPI_FOUND}\")message(STATUS \"MPI_CXX_INCLUDE_DIRS=${MPI_CXX_INCLUDE_DIRS}\")message(STATUS \"MPI_LIBRARIES=${MPI_LIBRARIES}\")#include \u0026lt;mpi.h\u0026gt;#include \u0026lt;stdio.h\u0026gt;int main(int argc, char** argv) { // Initialize the MPI environment MPI_Init(\u0026amp;argc, \u0026amp;argv); // Get the number of processes ssociated with the communicator int world_size{}; MPI_Comm_size(MPI_COMM_WORLD, \u0026amp;world_size); // Get the rank of the calling process int world_rank{}; MPI_Comm_rank(MPI_COMM_WORLD, \u0026amp;world_rank); // Get the name of the processor char processor_name[MPI_MAX_PROCESSOR_NAME]{}; int name_len; MPI_Get_processor_name(processor_name, \u0026amp;name_len); printf(\"Hello world from process %s with rank %d out of %d processors\\n\", processor_name, world_rank, world_size); // Finalize: Any resources allocated for MPI can be freed MPI_Finalize();}教程参考 MPI Tutorial Open MPI 入门笔记" },

  { "title": "基于VTK的3D软件", "url": "/blog/2024/vtk-based-3d-software/", "categories": "VTK",
  "tags": "VTK", "date": "2024-05-17", "content":
  "1. VisIt VisIt。 github。2. Inviwo它支持微软的Windows、苹果的Mac OS和开源的Linux，核心模块采用C++开发，支持Python 3.9版本以上的调用。底层的图形渲染采用 OpenGL 3.3 以上的核心模式，图形用户界面已经支持Qt6。 Inviwo。 github。3. Voreen website source4. MegaMol website github5. ParaView website github6. ttk website github7. PolyscopePolyscope是一个年轻、有趣的可视化软件包。 website github8. GLVis website github9. libigllibigl本身是一个轻量化的C++计算几何处理库，但是最终的处理结果会以可视化的方式呈现出来，因此笔者把它视为可视化软件。很多大学和知名机构都使用libigl。 website github10. morphologica它是一个非常年轻的数据可视化工具包。它的定位有点类似之前介绍的VTK，是以C++头文件库方式提供的、底层渲染使用现代OpenGL、适用于各类数值模拟的可视化开发工具包。目前orphologica可以结合GLFW、Qt和wxWidgets等图形用户界面库进行开发，从而可以提供比较友好的窗口管理和用户界面。 doc github11. F3D F3D github引用 科学可视化软件介绍 – VTK" },

  { "title": "OpenGL/Vulkan 学习网站（记录）", "url": "/blog/2024/opengl-vulkan%E5%AD%A6%E4%B9%A0%E8%B5%84%E6%96%99/", "categories": "OpenGL",
  "tags": "Vulkan, OpenGL", "date": "2024-05-17", "content":
  "opengl-tutorial (cn) LearnOpenGL-CN官方教程 OpenGL Tutorials TyphoonLabs’ OpenGL Shading Language tutorials TyphoonLabs’ OpenGL Shader Designer IDE有关3D计算机图形学的基础知识 Learning Modern 3D Graphics Programming github – Ray Tracing in One Weekend Book Series Physically Based Rendering:From Theory To ImplementationVulkan 学习资源SDK以及资源： LunarG Vulkan SDK Vulkan GPU 资源 Getting Started with the Linux Tarball Vulkan SDK学习资料： Vulkan Tutorial。大神Overv写的，非常棒的Vulkan入门教程。 VulkanTutorialCN。Vulkan Tutorial的中文翻译版。 VulkanTutorialCN pdf。Vulkan Tutorial的中文翻译版PDF。 VkGPUDrivenCNGuide。 Vulkan GPU Driven 渲染中文指南。 VkGPUDrivenCNGuide pdf。Vulkan GPU Driven 渲染中文指南PDF。 vulkan_tutorial_doc_cn – 图形管线基础 Vulkan教程 blog – Vulkan学习笔记 yangwc – Vulkan学习笔记" },

  { "title": "为什么要使用 std::enable_shared_from_this，以及使用场景", "url": "/blog/2024/%E4%B8%BA%E4%BB%80%E4%B9%88%E4%BD%BF%E7%94%A8enable_shared_from_this/", "categories": "Cpp",
  "tags": "Cpp", "date": "2024-05-16", "content":
  "为什么要使用 std::enable_shared_from_this 当使用原始指针构造或者初始化一个shared_ptr时，将会创建一个新的控制块。为了确保一个对象仅由一个共享的控制块管理，必须通过复制已存在的shared_ptr对象来创建一个新的shared_ptr实例。 但是在某些情况下，shared_ptr管理的对象需要为自己获取shared_ptr，类似于下面这样尝试从自身指针创建shared_ptr的方式是行不通的： struct Egg { std::shared_ptr\u0026lt;Egg\u0026gt; get_self_ptr() { return std::shared_ptr\u0026lt;Egg\u0026gt;(this); }};void spam() { auto sp1 = std::make_shared\u0026lt;Egg\u0026gt;(); auto sp2 = sp1-\u0026gt;get_self_ptr(); // undefined behavior // sp1 and sp2 have two different control blocks managing same Egg} 类似下面这样，在类内部持有自身的shared_ptr对象导致其释放不了：#include \u0026lt;iostream\u0026gt;#include \u0026lt;memory\u0026gt;#include \u0026lt;boost/core/ignore_unused.hpp\u0026gt;struct Immortal { std::shared_ptr\u0026lt;Immortal\u0026gt; self; ~Immortal() { std::cout \u0026lt;\u0026lt; \"dtor of Immortal called. self.use_count() = \" \u0026lt;\u0026lt; self.use_count() \u0026lt;\u0026lt; std::endl; // self.reset(); // (1) }};int main(int argc, char* argv[]) { boost::ignore_unused(argc, argv); auto immortal = new Immortal(); immortal-\u0026gt;self = std::shared_ptr\u0026lt;Immortal\u0026gt;(immortal); std::cout \u0026lt;\u0026lt; \"main exit. immortal-\u0026gt;self.use_count() = \" \u0026lt;\u0026lt; immortal-\u0026gt;self.use_count() \u0026lt;\u0026lt; std::endl; // delete immortal; (2) return 0;}log 输出：main exit. immortal-\u0026gt;self.use_count() = 1如果恢复(1)，(2)两处代码，导致同样多次析构出错。测试代码 test_enable_shared_from_this。 为了解决这个问题，我们就需要用到std::enable_shared_from_this。public 继承std::enable_shared_from_this 的类中可以通过调用shared_from_this()方法来获取自身的shared_ptr。struct Thing;void some_api(const std::shared_ptr\u0026lt;Thing\u0026gt;\u0026amp; tp);struct Thing : public std::enable_shared_from_this\u0026lt;Thing\u0026gt; { void method() { some_api(shared_from_this()); }};std::enable_shared_from_this 的使用场景创建异步任务，并且需要把自身的shared_ptr传给异步任务，此时适用于使用 std::enable_shared_from_this。不使用 std::enable_shared_from_this，使用 std::weak_ptr 替代的解决办法weak_ptr是一种弱引用，它不会影响受管理对象的生命周期，但是在需要时可以用来获取强引用。 如果一个对象持有自身的weak_ptr，那么在需要的时候，就可以获取自身的shared_ptr：class Naive {public: static std::shared_ptr\u0026lt;Naive\u0026gt; create() { auto sp = std::shared_ptr\u0026lt;Naive\u0026gt;(new Naive); sp-\u0026gt;weak_self_ = sp; return sp; } auto async_method() { return std::async(std::launch::async, [self = weak_self_.lock()]() { self-\u0026gt;do_something(); }); } void do_something() { std::this_thread::sleep_for(std::chrono::seconds(1)); }private: Naive() = default; Naive(const Naive\u0026amp;) = delete; const Naive\u0026amp; operator=(const Naive\u0026amp;) = delete; std::weak_ptr\u0026lt;Naive\u0026gt; weak_self_;};void test() { std::future\u0026lt;void\u0026gt; ft; { auto pn = Naive::create(); ft = pn-\u0026gt;async_method(); } ft.get();}保证对象的线程安全在多线程中，访问shared_ptr指向的对象时，一个方法是加锁，另一个是使用原子变量：std::atomic\u0026lt;std::shared_ptr\u0026lt;int\u0026gt;\u0026gt; atomic_shared;参考 C++：深入理解 std::enable_shared_from_this 与 shared_from_this 从现代C++实现的百行线程池来了解一下线程池相关技术" },

  { "title": "std::shared_ptr 线程安全及性能考量", "url": "/blog/2024/shared_ptr_thread_safe/", "categories": "Cpp",
  "tags": "Cpp", "date": "2024-05-16", "content":
  "1. 线程安全根据cppreference的描述，std::shared_ptr线程安全如下（机器翻译）： 如果是每个线程都拥有自己的std::shared_ptr对象，则针对线程自己的std::shared_ptr对象，其所有成员函数都是线程安全的； 如果多个线程共享同一个std::shared_ptr对象，其const成员函数的访问是线程安全的，但其非const成员函数的访问需要同步。 多线程访问同一个std::shared_ptr对象时，使用std::atomic\u0026lt;std::shared_ptr\u0026lt;T\u0026gt;\u0026gt;可以防止数据竞争。原文：All member functions (including copy constructor and copy assignment) can be called by multiple threads on different shared_ptr objects without additional synchronization even if these objects are copies and share ownership of the same object. If multiple threads of execution access the same shared_ptr object without synchronization and any of those accesses uses a non-const member function of shared_ptr then a data race will occur; the std::atomic can be used to prevent the data race.总结： 多线程拷贝同一个std::shared_ptr对象，是线程安全的(引用计数是线程安全的)。 多线程访问std::shared_ptr指向的同一个内存对象时，访问const成员函数是线程安全的。 其余情况，需要使用同步。2. 性能考量 使用std::make_shared，std::make_shared将被指向对象的内存分配与控制块的内存分配合并为一次分配； std::make_shared的性能接近new；但std::shared_ptr\u0026lt;T\u0026gt;(new T)耗时较明显； 在x86平台，std::shared_ptr的访问时间应该接近T*；但ARM平台应该会有明显的性能差异；使用原始指针并赋值给shared_ptr，shared_ptr管理对象与控制块内存布局:而使用std::make_shared，shared_ptr管理对象与控制块内存布局:3. 更多资料 合集 - C++系列(18) C++: weak_ptr到底有什么用？ 合集 - C++系列(18) C++ 高效使用智能指针的8个建议 合集 - C++系列(18) C++: 16个基础的C++代码性能优化实例" },

  { "title": "系统设计资料，以及23种设计模式彩图", "url": "/blog/2024/design-patterns-23/", "categories": "Cpp",
  "tags": "DesignPattern", "date": "2024-05-16", "content":
  "linux-insides 中文翻译 system-design-primer参考 万字详解 GoF 23 种设计模式（多图、思维导图、模式对比），让你一文全面理解 Design Patterns In Modern C++ 中文版翻译 -第24章：访问者模式" },

  { "title": "C++实现一个简洁的 LRU 缓存", "url": "/blog/2024/cpp%E5%AE%9E%E7%8E%B0%E4%B8%80%E4%B8%AALRU/", "categories": "Cpp",
  "tags": "Cpp", "date": "2024-05-16", "content":
  "template \u0026lt;typename T, typename KeyT = int\u0026gt;struct lru_cache_t { size_t sz_; std::list\u0026lt;std::pair\u0026lt;KeyT, T\u0026gt;\u0026gt; cache_; using ListIt = typename std::list\u0026lt;std::pair\u0026lt;KeyT, T\u0026gt;\u0026gt;::iterator; std::unordered_map\u0026lt;KeyT, ListIt\u0026gt; hash_; lru_cache_t(size_t sz) : sz_(sz) {} bool full() const { return (cache_.size() == sz_); } template \u0026lt;typename F\u0026gt; bool lookup_update(KeyT key, F slow_get_page) { auto hit = hash_.find(key); if (hit == hash_.end()) { if (full()) { hash_.erase(cache_.back().first); cache_.pop_back(); } cache_.emplace_front(key, slow_get_page(key)); hash_.emplace(key, cache_.begin()); return false; } auto eltit = hit-\u0026gt;second; cache_.splice(cache_.begin(), cache_, eltit); return true; }}; 实现关键点：使用 std::list 存储避免 resize，使用 std::unordered_map 实现快速查找。使用 splice 方法将命中的元素移动到链表头部。int slow_get_page_int(int key) { return key; }int main() { int hits = 0; int n; size_t m; std::cin \u0026gt;\u0026gt; m \u0026gt;\u0026gt; n; assert(std::cin.good()); caches::lru_cache_t\u0026lt;int\u0026gt; c{m}; for (int i = 0; i \u0026lt; n; ++i) { int q; std::cin \u0026gt;\u0026gt; q; assert(std::cin.good()); if (c.lookup_update(q, slow_get_page_int)) hits += 1; } std::cout \u0026lt;\u0026lt; hits \u0026lt;\u0026lt; std::endl;}资料： LRU and LFU in C++" },

  { "title": "POD、trivial copyable，standard layout", "url": "/blog/2024/cpp11-pod-type/", "categories": "Cpp",
  "tags": "Cpp", "date": "2024-05-16", "content":
  "1. trival copyabletrival copyable 是指内存布局可以使用 memcpy 进行内存拷贝的类型，但不定义与C类型兼容。C++11 要求如下： 构造函数，复制构造函数，移动构造函数（move ctor），析构函数，均为默认。 赋值运算符，移动赋值运算符，均为默认。 没有虚函数和虚基类。1.1 说明 默认的构造等函数，包含可以使用 =default，如果包含自定义的构造函数，且提供 =default 的构造函数，其依然是 trival copyable。 可以使用初始化列表初始化成员，其依然是 trival copyable。 如果包含了虚函数 / 虚基类 / 虚析构函数，因为会生成 vptr，所以不是 trival copyable。 可以包含 static 成员，且可以是 None trivial copyable。 可以包含 private等访问控制符。 C++11 中，我们使用模版类 std::is_trivially_copyable\u0026lt;T\u0026gt;::value 来判断一个类型是否为平凡类型。1.2 示例// empty classes are trivialstruct Trivial1 {};// all special members are implicitstruct Trivial2 { int x;};struct Trivial3 : Trivial2 { // base class is trivial Trivial3() = default; // not a user-provided ctor int y;};struct Trivial4 {public: int a;private: // no restrictions on access modifiers int b;};struct Trivial5 { Trivial1 a; Trivial2 b; Trivial3 c; Trivial4 d;};struct Trivial6 { Trivial2 a[23];};struct Trivial7 { Trivial6 c; void f(); // it's okay to have non-virtual functions};struct NonTrivial1 : Trivial3 { virtual void f(); // virtual members make non-trivial ctors};struct Trivial8 { int x; static NonTrivial1 y; // no restrictions on static members};struct Trivial9 { Trivial9() = default; // not user-provided // a regular constructor is okay because we still have default ctor Trivial9(int x) : x(x) {}; int x;};struct NonTrivial2 { NonTrivial2() : z(42) {} // user-provided ctor int z;};struct NonTrivial3 { NonTrivial3(); // user-provided ctor int w;};NonTrivial3::NonTrivial3() = default; // defaulted but not on first declaration // still counts as user-providedstruct NonTrivial5 { virtual ~NonTrivial5(); // virtual destructors are not trivial};2. standard layoutstandard layout 是指定义数据结构成员的内存布局规则，使其可以与其他语言交互，如可以使用 extern \"C\" 进行类型声明。 没有虚函数，没有虚基类。 所有非静态成员都有相同的访问控制。 非静态成员不能是引用类型。 如果有基类，第一个非静态成员类型不能是其基类类型。 如果有基类，要求基类也是 standard layout。 继承树中，只能有一个类有非静态数据成员。2.1 示例// empty classes have standard-layoutstruct StandardLayout1 {};struct StandardLayout2 { int x;};struct StandardLayout3 {private: // both are private, so it's ok int x; int y;};struct StandardLayout4 : StandardLayout1 { int x; int y; void f(); // perfectly fine to have non-virtual functions};struct StandardLayout5 : StandardLayout1 { int x; StandardLayout1 y; // can have members of base type if they're not the first};struct StandardLayout6 : StandardLayout1, StandardLayout5 { // can use multiple inheritance as long only // one class in the hierarchy has non-static data members};struct StandardLayout7 { int x; int y; StandardLayout7(int x, int y) : x(x), y(y) {} // user-provided ctors are ok};struct StandardLayout8 {public: StandardLayout8(int x) : x(x) {} // user-provided ctors are ok// ok to have non-static data members and other members with different accessprivate: int x;};struct StandardLayout9 { int x; static NonStandardLayout1 y; // no restrictions on static members};struct NonStandardLayout1 { virtual f(); // cannot have virtual functions};struct NonStandardLayout2 { NonStandardLayout1 X; // has non-standard-layout member};struct NonStandardLayout3 : StandardLayout1 { StandardLayout1 x; // first member cannot be of the same type as base};struct NonStandardLayout4 : StandardLayout3 { int z; // more than one class has non-static data members};struct NonStandardLayout5 : NonStandardLayout3 {}; // has a non-standard-layout base class参考资料： What are aggregates and trivial types/PODs, and how/why are they special? – What changes for C++11?附加：类对象的生命周期经历#include \u0026lt;memory\u0026gt;template \u0026lt;typename T\u0026gt;void life_of_an_object{ std::allocator\u0026lt;T\u0026gt; alloc; // 1. 通过 allocator 抑或是 malloc 等其他方式分配出空间 T * p = alloc.allocate(1); // 2. 通过 placement new，(在需要的时候) 动态地构造对象 new (p) T(); // 这里是默认构造，也可能是带参数的构造方式如 new (p) T(构造参数...); // 3. 通过显式调用析构函数，(在需要的时候) 动态地销毁对象 p-\u0026gt;~T(); // 4. 通过分配函数的对应的解分配手段，解分配空间 alloc.deallocate(p, 1);} c++ trivial/pod 是什么意思？" },

  { "title": "CMake 创建自定义目标：在构建前复制文件", "url": "/blog/2024/cmake-create-custom-target-copy-files-before-build/", "categories": "CMake",
  "tags": "CMake", "date": "2024-05-14", "content":
  "创建自定义目标add_custom_target(TargetCopy3rdPartyLibs ALL COMMAND ${CMAKE_COMMAND} -E make_directory \"${CMAKE_RUNTIME_OUTPUT_DIRECTORY}\" COMMAND cp_3rd_libs.bat \"${PROJECT_SOURCE_DIR}/3rd_libs\" \"${CMAKE_RUNTIME_OUTPUT_DIRECTORY}\" WORKING_DIRECTORY \"${PROJECT_SOURCE_DIR}/src/cmake\" COMMENT \"Copying 3rd party libraries to bin directory\")引用自定义目标在编译APP，LIB的CMakeLists.txt文件中，添加：set_target_properties(${target_app} PROPERTIES DEPENDS Copy3rdPartyLibs)" },

  { "title": "OpenCascade拓扑与几何的关系", "url": "/blog/2024/Topology-and-Geometry-in-Open-CASCADE/", "categories": "OCCT",
  "tags": "OCCT", "date": "2024-05-12", "content":
  "1. OPENCASCADE 中的拓扑结构（TopoDS包）抽象结构是以TopoDS包的C++类来实现的。下面是一个继承图，取自Doxygen生成的文档。TopoDS_Shape 是通过值来操作的，包含3个字段–location、orientation 和一个 myTShape 句柄 （TopoDS_TShape类型），见下图（只包含最重要的字段）：myTShape 和 Location 被用来在各种形状之间共享数据，从而节省大量的内存。例如，两个连接在一起的面中间的边，这个边具有相同的位置（location）和 myTShape，但有不同的朝向（orientation）（在其中一个面是正向，在另一个面中是反向）。2. OPENCASCADE 中的边界表示法（BRep包）边界表示（Boundary Representation）也称为BRep表示，它是几何造型中最成熟、无二义的表示法。实体的边界通常是由面的并集表示，而每个面又由它所在曲面的定义加上其边界来表示，面的边界是边的并集，而边又是由点来表示。如下图1.1所示，曲面的汇合处形成曲线，而曲线的汇合处形成点。所以点、线、面是描述一个形状所需要的基本组成单元。几何信息（在BRep包中）通过继承TopoDS_Shape类来实现，只有3种类型的拓扑对象有几何特征–顶点（vertex）、边(edge)和面(face)（见下图）：2.1 Brep 进一步说明几何信息（BRep）描述了形状的形状、位置、方向、边界等信息，而拓扑信息（TopoDS）描述了形状的拓扑结构，如顶点、边、面的连接关系。以BRep Edge为例说明：边有几种几何表示方式（参考上图，BRep_TEdge之下有节点List of Edge Representions）: 三维空间的曲线C(t)，用类Geom_Curve实现，这是基础表示方式。 曲线P(t)为二维空间中的参数曲线，用于描述位于曲面上的曲线。这些通常被称为pcurves，用类Geom2d_Curve实现。 多边形（Polygonal representation）用一组三维点来描述，用类Poly_Polygon3D实现。 多边形（Polygonal representation），也可以用一组三维点的编号来描述，在类Poly_PlygonOnTriangulation实现。3. OpenCascade 中拓扑（Topo）与几何（BRep）的关系边界表示的一个重要特点是描述形状的信息包括几何信息（geometry）和拓朴（topology）信息两个方面： 拓朴信息描述形状上的顶点、边、面的连接关系，它形成物体边界表示的“骨架”。 形状的几何信息犹如附着在“骨架”上的肌肉。在OpenCascade中，形状的几何信息包含曲线和曲面的参数解析表示Geom_Curve/Geom_Surface。这样我们就可以用平面方程和柱面方程来描述曲面，用直线或圆弧方程来描述曲线。这时会出现一个问题，即代数表达式只能定义无边界的几何体。除了单个点、圆以及球体，经典的解析几何仅能表示无限延伸的曲线和曲面。为了解决这个问题，边界表示法按下述方法明确地定义曲线或曲面的边界： 曲线的边界由位于曲线上的一对点来确定； 曲面的边界由位于曲面上的一组曲线来确定；通过这个方法，就可以定义一段曲线或一片曲面。这时，不同几何元素之间的关系的组织问题就出现了，为此我们将记录如下信息： 哪些点界定哪些曲线； 哪些曲线界定哪些曲面；这些关于谁关联谁的信息，就是几何造型系统经常提到的拓朴。在边界表示法中，理论上表示一个物理模型只需要三个拓朴体（顶点TopoDS_Vertex、边TopoDS_Edge和面TopoDS_Face），但在实际应用中，为了提高计算机处理的速度或提供高级的操作功能，还要引入其他一些概念，如环TopoDS_Wire、壳TopoDS_Shell、复合体TopoDS_Compound等。z. 原文及参考z.1 ROMAN LYGIN TOPOLOGY AND GEOMETRY IN OPEN CASCADE. PART 1 TOPOLOGY AND GEOMETRY IN OPEN CASCADE. PART 2 TOPOLOGY AND GEOMETRY IN OPEN CASCADE. PART 3 TOPOLOGY AND GEOMETRY IN OPEN CASCADE. PART 4 TOPOLOGY AND GEOMETRY IN OPEN CASCADE. PART 5在occ dev上对应的讨论帖子： Blog: Topology and Geometry in Open CASCADEz.2 国内 eryar 等同文章系列 Topology Shapes of OpenCascade BRep Topology and Geometry in OpenCascade-Vertex Topology and Geometry in OpenCascade-Edge Topology and Geometry in OpenCascade-Face Topology and Geometry in OpenCascade-Topology Topology and Geometry in OpenCascade-Adapterseryar 其他相关文章及代码： Surface Normal Vector in OpenCascade Surface Normal Averaging Parametric Curves and Surfaces PCurve - Curve on Surfacez.3 ROMAN LYGIN 原文的中文翻译 【OpenCascade】拓扑与几何的关系z.4 其他资料 OpenCasCade官方开发文档翻译(1)–occt整体介绍 Victoria Rudakova" },

  { "title": "使用 Starship 定制Bash提示符", "url": "/blog/2024/%E7%BB%88%E7%AB%AF%E7%BE%8E%E5%8C%96/", "categories": "Bash",
  "tags": "Bash", "date": "2024-05-12", "content":
  "安装 Starship 并使能：curl -sS https://starship.rs/install.sh | sh# Bash导入eval \"$(starship init bash)\"添加配置文件，路径~/.config/starship.toml。备份文件starship.toml。官方文档：Starship 配置。" },

  { "title": "使用 Starship 定制Bash提示符", "url": "/blog/2024/zed%E4%BD%BF%E7%94%A8%E7%AC%94%E8%AE%B0/", "categories": "zed",
  "tags": "zed", "date": "2024-05-12", "content":
  "本笔记记录自zed 1.6.3（202606）开始的使用笔记。1. 使用zed进行远程开发类似VSCode远程开发环境搭建，在首次连接到远程服务器的时候，remote ssh会连接到远程，在远程服务器上下载remote server。在zed的菜单File -\u0026gt; Open Remote，选择Remote SSH连接到远程服务器，这个菜单下面，使用Connect SSH Server菜单项，输入远程服务器地址，诸如user@remote-server。另外，zed可以检测到目录下.ssh/config文件，这个文件存储远程服务器地址及用户名，这个跟VSCode的远程开发配置相同。由于远程下载的时候，没有使用VPN，导致下载失败。所以写了一个Python脚本，用于在本地下载remote server，然后上传到远程服务器，并按zed规定的目录名称格式解压zed server，脚本：setup_remote_zed_server.py。脚本来自博客：Zed 手动/离线 安装 zed_server, 解决远程开发一直卡下载问。" },

  { "title": "VSCode 改用 MathJax 渲染数学公式", "url": "/blog/2024/VSCode%E6%94%B9%E7%94%A8MathJax%E6%B8%B2%E6%9F%93%E6%95%B0%E5%AD%A6%E5%85%AC%E5%BC%8F/", "categories": "VSCode",
  "tags": "VSCode, MathJax", "date": "2024-05-12", "content":
  "首先安装Markdown Preview Enhanced插件，并做如下配置： Ctrl+Shift+P → Markdown Preview Enhanced: Open Config Script (Global)将如下内容添加到配置文件config.js中：{ \"mathjaxConfig\": { \"tex\": { \"inlineMath\": [[\"$\", \"$\"], [\"\\\\(\", \"\\\\)\"]], \"displayMath\": [[\"$$\", \"$$\"], [\"\\\\[\", \"\\\\]\"]], \"tags\": \"ams\" } }}" },

  { "title": "我的VSCode插件清单", "url": "/blog/2024/VSCode%E6%8F%92%E4%BB%B6%E6%B8%85%E5%8D%95/", "categories": "VSCode",
  "tags": "VSCode", "date": "2024-05-12", "content":
  "1. 全量清单extensions.json文件内容如下所示：{ \"recommendations\": [ \"bierner.markdown-mermaid\", \"davidanson.vscode-markdownlint\", \"devtbi.vscode-cppinsights\", \"donjayamanne.python-extension-pack\", \"dotjoshjohnson.xml\", \"github.copilot\", \"github.copilot-chat\", \"ibm.output-colorizer\", \"josetr.cmake-language-support-vscode\", \"kevinrose.vsc-python-indent\", \"llvm-vs-code-extensions.vscode-clangd\", \"mechatroner.rainbow-csv\", \"meta.pyrefly\", \"ms-ceintl.vscode-language-pack-zh-hans\", \"ms-dotnettools.vscode-dotnet-runtime\", \"ms-python.black-formatter\", \"ms-python.debugpy\", \"ms-python.isort\", \"ms-python.python\", \"ms-python.vscode-pylance\", \"ms-python.vscode-python-envs\", \"ms-vscode-remote.remote-ssh\", \"ms-vscode-remote.remote-ssh-edit\", \"ms-vscode-remote.vscode-remote-extensionpack\", \"ms-vscode.cmake-tools\", \"ms-vscode.cpptools\", \"ms-vscode.cpptools-extension-pack\", \"ms-vscode.cpptools-themes\", \"ms-vscode.remote-explorer\", \"repreng.csv\", \"theqtcompany.qt\", \"theqtcompany.qt-core\", \"theqtcompany.qt-cpp\", \"theqtcompany.qt-cpp-pack\", \"theqtcompany.qt-qml\", \"theqtcompany.qt-ui\", \"vscode-icons-team.vscode-icons\", \"yzhang.markdown-all-in-one\" ]}插件分类总结： 分类 插件数 主要工具 C++/Clangd 4 clangd, cpptools, cppinsights Qt 开发 6 Qt官方全套扩展 Python 7 Python, Pylance, Black等 CMake 2 cmake-tools, cmake-language-support GitHub/Copilot 2 Copilot + Chat 远程开发 4 Remote SSH/WSL/Container Markdown 3 markdown-all-in-one等 其他工具 10 CSV, XML, Icons等 命令行安装extensions.json：code --list-extensions | ForEach-Object { code --install-extension $_ }2. 写 Markdown 笔记插件{ \"recommendations\": [\"yzhang.markdown-all-in-one\"]}3. 设置多个开发环境VSCode左下角齿轮图标 -\u0026gt; 配置文件（默认） -\u0026gt; 配置文件，打开配置文件页面，创建新的配置文件，完成之后，选择该配置文件，然后再重新安装需要的插件即可。" },

  { "title": "VSCode 开发，PowerShell 高效开发完整配置", "url": "/blog/2024/VSCode-PowerShell%E9%AB%98%E6%95%88%E5%BC%80%E5%8F%91%E5%AE%8C%E6%95%B4%E9%85%8D%E7%BD%AE/", "categories": "PowerShell",
  "tags": "PowerShell, VSCode, posh-git, oh-my-posh", "date": "2024-05-12", "content":
  "1. 配置 VSCode 使用 PowerShell 7VSCode 快捷键打开用户配置文件 (Json)：Ctrl+Shift+P，输入 open user settings，选择打开 settings.json 文件，添加以下内容：\"terminal.integrated.profiles.windows\": { \"PowerShell7\": { \"path\": \"C:\\\\Program Files\\\\PowerShell\\\\7\\\\pwsh.exe\", \"args\": [], \"icon\": \"terminal-powershell\" }, \"Command Prompt\": { \"path\": [ \"${env:windir}\\\\Sysnative\\\\cmd.exe\", \"${env:windir}\\\\System32\\\\cmd.exe\" ], \"args\": [], \"icon\": \"terminal-cmd\" }, \"Git Bash\": { \"source\": \"Git Bash\" }},\"terminal.integrated.defaultProfile.windows\": \"PowerShell7\",2. PowerShell 执行策略设置Windows 默认不允许自动运行脚本，需要以管理员身份执行：Set-ExecutionPolicy RemoteSigned -Scope CurrentUser -Confirm3. 安装 posh-gitposh-git 为 PowerShell 提供 Git 状态信息（branch、暂存区、工作区变更数量等），并支持 Git 命令 Tab 补全。# 更新 PowerShellGet（如有必要）Install-Module PowerShellGet -Force# 全新安装PowerShellGet\\Install-Module posh-git -Scope CurrentUser -Force# 升级已有版本PowerShellGet\\Update-Module posh-git3.1. posh-git 状态符号说明Prompt 格式为 [{branch} S +A ~B -C | +E ~F -G !H]： S：与远端的同步状态（≡ 同步、↑n 领先、↓n 落后、× 远端分支已删除） | 左侧：暂存区（Index）变更（+ 新增、~ 修改、- 删除、! 冲突） | 右侧：工作区（Working Tree）未暂存变更3.2. 配置 posh-git 到 Profile安装完成后，执行一次以下命令（仅执行一次，它会自动将 Import-Module posh-git 写入 profile 文件）：# 仅对当前 host（推荐）Add-PoshGitToProfile# 对所有 host（Console、ISE 等）Add-PoshGitToProfile -AllHosts4. 安装配置 oh-my-poshoh-my-posh 提供更美观的 PowerShell Prompt 主题。4.1. 安装# 安装 PSReadLine（命令行增强，支持历史预测等）Install-Module -Name PSReadLine -Scope CurrentUser# 安装 oh-my-poshwinget install JanDeDobbeleer.OhMyPosh -s winget4.2. 配置 Profile打开 profile 文件：notepad $PROFILE添加以下内容（注意替换 Administrator 为你的用户名）：# oh-my-posh 主题（选择一个）# hotstick.minimal.omp.json -- 单行，Prompt 简短# agnoster.minimal.omp.json -- 经典单行风格oh-my-posh init pwsh --config 'C:\\Users\\Administrator\\AppData\\Local\\Programs\\oh-my-posh\\themes\\hotstick.minimal.omp.json' | Invoke-Expression可用主题参考：ohmyposh.dev/docs/themes4.3. 乱码处理（字体设置）oh-my-posh Prompt 使用特殊符号，需安装支持 Box-drawing 的字体，推荐 Fira Code 或 Cascadia Code。安装字体后，在 Windows Terminal 的默认值 → 外观 → 字体中选择对应字体。5. PSReadLine 高效命令行设置以下配置让 PowerShell 命令历史记录更易用，加入 $PROFILE 文件中：# 设置预测文本来源为历史记录Set-PSReadLineOption -PredictionSource History# Tab 键显示可选菜单（类似 zsh）Set-PSReadLineKeyHandler -Key Tab -Function MenuComplete# 上下方向键：按已输入前缀搜索历史记录，并将光标移到行尾Set-PSReadLineKeyHandler -Key UpArrow -ScriptBlock { [Microsoft.PowerShell.PSConsoleReadLine]::HistorySearchBackward() [Microsoft.PowerShell.PSConsoleReadLine]::EndOfLine()}Set-PSReadLineKeyHandler -Key DownArrow -ScriptBlock { [Microsoft.PowerShell.PSConsoleReadLine]::HistorySearchForward() [Microsoft.PowerShell.PSConsoleReadLine]::EndOfLine()} 说明：使用 ScriptBlock 方式（而非直接 -Function HistorySearchBackward）的好处是：搜索后光标自动跳到行尾，直接可以继续补充参数。6. Git 快捷命令设置在 $PROFILE 中添加以下别名函数，需放在 Import-Module posh-git 之前（posh-git 会识别已定义的 git 别名用于 Tab 补全）：# 美化的 git log 图形视图function gl { git log --color --graph --pretty=format:'%Cred%h%Creset -%C(yellow)%d%Creset %s %Cgreen(%cr) %C(bold blue)\u0026lt;%an\u0026gt;%Creset' --abbrev-commit}# 简洁 git statusfunction g { git status -sb }7. 清屏快捷键（VSCode 终端）VSCode 中 cls 命令仅滚动缓冲区，建议用快捷键清屏：Ctrl+Shift+P → 搜索 Terminal: Clear → 绑定快捷键 Ctrl+L。8. 设置终端编码为 UTF-8# 在 $PROFILE 中添加，防止中文乱码[Console]::OutputEncoding = [System.Text.Encoding]::UTF8$OutputEncoding = [System.Text.Encoding]::UTF8 同时需要将终端字体改为支持 Box-drawing 的字体（如 Cascadia Code、Fira Code）。点阵字体（如新宋体）不支持。附：完整 Profile 示例# git 别名函数需在 Import-Module posh-git 之前定义，以便 posh-git Tab 补全识别# git log 图形视图function gl { git log --color --graph --pretty=format:'%Cred%h%Creset -%C(yellow)%d%Creset %s %Cgreen(%cr) %C(bold blue)\u0026lt;%an\u0026gt;%Creset' --abbrev-commit }# git status 简洁视图function g { git status -sb }Import-Module posh-git# 若同时使用 oh-my-posh，需开启此变量以让 oh-my-posh 读取 posh-git 状态$env:POSH_GIT_ENABLED = $trueImport-Module PSReadLine# 设置预测文本来源为历史记录Set-PSReadLineOption -PredictionSource History# Tab 键显示可选补全菜单Set-PSReadLineKeyHandler -Key Tab -Function MenuComplete# 上下方向键：按已输入前缀搜索历史记录，并将光标移至行尾Set-PSReadLineKeyHandler -Key UpArrow -ScriptBlock { [Microsoft.PowerShell.PSConsoleReadLine]::HistorySearchBackward() [Microsoft.PowerShell.PSConsoleReadLine]::EndOfLine()}Set-PSReadLineKeyHandler -Key DownArrow -ScriptBlock { [Microsoft.PowerShell.PSConsoleReadLine]::HistorySearchForward() [Microsoft.PowerShell.PSConsoleReadLine]::EndOfLine()}# 设置 PowerShell 输出编码为 UTF-8，还需要将字体改为支持 Box-drawing 的字体# 比如 Cascadia Code、Consolas。新宋体这种点阵字体不行。[Console]::OutputEncoding = [System.Text.Encoding]::UTF8$OutputEncoding = [System.Text.Encoding]::UTF8参考资料 posh-git GitHub oh-my-posh 主题列表 Fira Code 字体 VSCode 终端高级功能 Windows 安装 posh-git PowerShell 7 升级笔记 Windows Terminal 更新后 oh-my-posh 报错的问题分析与解决" },

  { "title": "Qt 笔记及开源库收集", "url": "/blog/2024/qt-notes/", "categories": "Qt",
  "tags": "Qt", "date": "2024-05-11", "content":
  "1. 控件的 eventFilter如下代码，将为edtCleanCount部件安装一个事件过滤器，当用户按下Up或Down键时，并且参数obj为edtCleanCount时，会自动修改edtCleanCount的值。也可以继承QLineEdit类，重写keyPressEvent方法，实现相同的功能。// eventFilter 声明// bool eventFilter(QObject* obj, QEvent* event) override;// initvoid MyDialog::init() { // 设置只允许edtCleanCount输入数字 QValidator* validator = new QRegExpValidator(QRegExp(\"[0-9]+\"), this); ui-\u0026gt;edtCleanCount-\u0026gt;setValidator(validator); //在编辑器部件安装事件过滤器 ui-\u0026gt;edtCleanCount-\u0026gt;installEventFilter(this);}bool MyDialog::eventFilter(QObject* obj, QEvent* event) { if (obj == ui-\u0026gt;edtCleanCount \u0026amp;\u0026amp; event-\u0026gt;type() == QEvent::KeyPress) { QKeyEvent* key_event = static_cast\u0026lt;QKeyEvent*\u0026gt;(event); const auto key = key_event-\u0026gt;key(); if (key == Qt::Key_Up || key == Qt::Key_Down) { auto flag = (key == Qt::Key_Up) ? 1 : -1; auto clean_count = ui-\u0026gt;edtCleanCount-\u0026gt;text().toInt() + flag; clean_count = std::max(clean_count, 0); // 最小值为0 ui-\u0026gt;edtCleanCount-\u0026gt;setText(QString::number(clean_count)); return true; } } return QWidget::eventFilter(obj, event);}2. 清除Layout及其所有子部件#include \u0026lt;QtWidgets\u0026gt;inline void clearLayout(QLayout* layout) { while (QLayoutItem* item = layout-\u0026gt;takeAt(0)) { if (QWidget* widget = item-\u0026gt;widget()) { widget-\u0026gt;deleteLater(); } if (QLayout* childLayout = item-\u0026gt;layout()) { clearLayout(childLayout); } if (QSpacerItem* spaerItem = item-\u0026gt;spacerItem()) { layout-\u0026gt;removeItem(spaerItem); } delete item; }}QT 开源库收集 QaterialHotReload Qt Advanced Docking System github – SARibbon github – ElaWidgetTools github – HuskarUI Qt Material Icons Library QtSerialMonitor其他资源收集 IconScout 图标资源" },

  { "title": "Linux系统及编译相关笔记", "url": "/blog/2024/linux-and-compile-notes/", "categories": "Linux",
  "tags": "Linux, GCC", "date": "2024-05-11", "content":
  "1. 加载多个DSO中存在同名符号APP加载的多个DSO，如果存在重复的符号，则只加载第一个遇到的符号。规则适合于包括函数，全局变量。加载顺序，由链接顺序，以及环境变量定义lib路径优先级，如RPATH，LD_LIBRARY_PATH。路径搜索优先级为：RPATH \u0026gt; LD_LIBRARY_PATH \u0026gt; ld.so.cache。1.1. RPATH编译期间，设定RPATH，如cmake脚步设置如下：set_target_properties(${target_name} PROPERTIES LINK_FLAGS \"-Wl,-rpath='$ORIGIN' \")如果存在间接依赖，则也优先使用RPATH设定的路径进行搜索。1.1.1. ORIGIN 占位符$ORIGIN是一个特殊的占位符，代表可执行文件或库文件自身的目录，当设置为$ORIGIN时，它告诉动态链接器在可执行文件或库所在的同一目录下查找依赖的库。1.2. 调试lib加载–LD_DEBUG设置环境变量LD_DEBUG使能系统级加载信息：LD_DEBUG=libs ./test_app1.3. 参考链接 Linux运行时动态库搜索路径优先级 gitee – 测试代码2. Linux命令：查看APP/DSO的编译器信息查看架构：file bin/jouav_cluster_msg_simu# readelf -h bin/jouav_cluster_msg_simu | grep Machine查看编译器版本：readelf -p .comment bin/jouav_cluster_msg_simu3. 其他文章 linux更换题目(可执行文件)libc版本问题" },

  { "title": "如何使用 github pages 以及 Chirpy Theme 创建博客", "url": "/blog/2024/how-to-create-blog-using-github-pages-and-Chirpy-theme/", "categories": "Computer",
  "tags": "Website", "date": "2024-05-11", "content":
  "使用Chirpy Theme在GitHub Pages部署Jekyll 使用Jekyll + Github Pages搭建静态网站 Text and Typography" },

  { "title": "Ubuntu 24.04 分区，以及更换 kernel 6.18", "url": "/blog/2024/Ubuntu%E5%AE%89%E8%A3%85%E6%96%B0%E7%89%88%E6%9C%ACkernel/", "categories": "Linux",
  "tags": "Linux", "date": "2024-05-11", "content":
  "Ubuntu 24.04 分区建议更换 kernel 6.18# 查看已经有的 kernel 版本# sudo apt-cache search linux-headers | grep 6.1sudo add-apt-repository ppa:cappelikan/ppa -ysudo apt update \u0026amp;\u0026amp; sudo apt install mainline -y安装新版的 kernel 6.18，执行命令后，工具会自动开始内核相关文件的下载：sudo mainline install 6.18# 安装完成后，重启系统sudo reboot# 重启后，查看当前内核版本uname -r实际上，所有可用的kernel版本都可以从Ubuntu官方仓库看到：https://kernel.ubuntu.com/mainline/。安装完成并重启之后，可以锁定当前版本：sudo apt-mark hold $(dpkg -l | grep -E \"linux-(headers|image|unsigned|modules|modules-extra)\" | grep 6.12.3 | awk '{print $2}')删除旧版本 kernel（需要先锁定当前新版本，不然新版本也会删除）：# 查看已经安装的 kernel 版本dpkg -l | grep -E \"linux-(headers|image|modules-extra)+\" | grep -v 6.12.3 | awk '{print $2}'# 删除旧版本 kerneldpkg -l | grep -E \"linux-(headers|image|modules-extra)+\" | grep -v 6.12.3 | awk '{print $2}' | xargs -I {} sudo apt remove -y {}参考链接： Ubuntu Linux 内核版本升级指南：mainline" },

  { "title": "C++ 17 新功能： std::visit 和 std::variant 配合使用 (待更新删除冗余描述)", "url": "/blog/2024/combind-usage-of-std-visit-and-std-variant/", "categories": "Cpp",
  "tags": "Cpp", "date": "2024-05-09", "content":
  "1. std::variant （变体）在很多编程场景中，我们经常遇到需要处理多种类型的情况。传统上，这可以通过多种方式来实现，例如使用 union 或 void* 指针，甚至使用一系列的 if-else 语句和类型转换。但这些方法通常都有各种缺点，如类型不安全、容易出错或难以维护。std::variant 为这一问题提供了一个现代、类型安全的解决方案。它允许你在一个单一的变量中存储多种不同的类型，并能在运行时安全地访问它们，并能获取他们的类型信息。可以把它看作是一个可以存储多种类型中的任一种的类型安全的容器。下面是一个基本用法的例子：#include \u0026lt;variant\u0026gt;#include \u0026lt;iostream\u0026gt;int main() { std::variant\u0026lt;int, double, std::string\u0026gt; v1 = 42; std::variant\u0026lt;int, double, std::string\u0026gt; v2 = 3.14; std::variant\u0026lt;int, double, std::string\u0026gt; v3 = \"hello\"; // 访问存储的值（不安全，需确保类型正确） std::cout \u0026lt;\u0026lt; std::get\u0026lt;int\u0026gt;(v1) \u0026lt;\u0026lt; std::endl; // 安全地访问存储的值 auto pval = std::get_if\u0026lt;int\u0026gt;(\u0026amp;v1); if (pval) { std::cout \u0026lt;\u0026lt; *pval \u0026lt;\u0026lt; std::endl; } return 0;}1.1 std::variant 的局限尽管 std::variant 非常强大，但它并不是万能的。它的一个主要限制是，虽然它可以存储多种类型，但在任何给定时间点，它只能存储其中一种。这意味着，如果你想存储多种类型，你需要使用 std::visit 函数来访问它们。当然，你可以使用 std::holds_alternative 或 std::get_if 进行手动检查，但这样做的代码通常既繁琐又容易出错。std::variant\u0026lt;int, double, std::string\u0026gt; v = 42;if (std::holds_alternative\u0026lt;int\u0026gt;(v)) { int value = std::get\u0026lt;int\u0026gt;(v); // 安全} else if (std::holds_alternative\u0026lt;double\u0026gt;(v)) { double value = std::get\u0026lt;double\u0026gt;(v); // 运行时错误！}创建一个四行三列的MD表格 方法 优点 缺点 std::holds_alternative 简单、直观 不能提取值 std::get 可以直接提取值 类型错误会抛出异常 std::get_if 可以检查和提取值，不会抛出异常 返回指针，需要额外的空指针检查 2. std::visit （访问）函数原型如下：template\u0026lt;class Visitor, class... Variants\u0026gt;constexpr visit(Visitor\u0026amp;\u0026amp; vis, Variants\u0026amp;\u0026amp;... vars);其底层工作原理： 访问存储的值：当 std::visit 被调用时，它首先需要确定 std::variant 当前存储的具体类型。这是通过检查内部的类型标记完成的。 函数模板实例化：std::visit 接受一个可调用对象和一个或多个 std::variant 对象。这个可调用对象通常是一个重载的函数对象或 lambda 表达式，其具有多个重载以处理不同的类型。编译器会为这些重载生成函数模板实例。 类型恢复和函数调用：一旦确定了 std::variant 中的类型，std::visit 通过生成的模板代码来“恢复”此类型，并调用与该类型匹配的函数重载。如果有多个 std::variant 参数，std::visit 将处理所有组合的可能性，并调用适当的重载。 编译时多态：这一切都在编译时发生。编译器生成适用于所有可能的类型组合的代码。因此，std::visit 实现了一种编译时的多态，而不是运行时多态（如虚函数）。 综上所述，std::visit 的核心在于它能够在编译时处理多态性，允许编译器生成处理 std::variant 中所有可能类型的代码。这种方法确保了类型安全，并允许进行高效的代码优化。3. 如何优雅地使用 std::visit3.1 使用泛型 lambda 表达式std::visit 允许你传入一个可调用对象（callable object），通常是一个 lambda 表达式。现代 C++ 提供了一种特殊的 lambda 表达式，称为泛型 lambda 表达式（generic lambda）。泛型 lambda 是一个使用 auto 关键字作为参数类型的 lambda 表达式。这意味着 lambda 可以接受任何类型的参数，并在函数体内进行处理。auto generic_lambda = [](auto x) { // do something with x};这种灵活性在处理 std::variant 时尤为有用，因为你可能需要根据多种可能的类型来编写逻辑。3.2 使用 if constexpr 和类型萃取编程就像是一场高级的拼图游戏。你需要一种机制来判断哪块拼图适用于当前的情况。在 std::visit 的上下文中，这通常是通过 if constexpr 和类型萃取（type traits）来完成的。if constexprif constexpr 是 C++17 引入的一种编译时 if 语句，它允许在编译时进行条件判断。这意味着编译器会根据条件来优化生成的代码，这通常会带来更高的性能。使用 if constexpr，你可以在一个统一的代码块中处理多种类型，而无需使用多个繁琐的 if-else 语句。这不仅让代码看起来更简洁，而且更易于维护。类型萃取：认识你的类型类型萃取（Type Traits）是 C++11 引入的一组模板，用于在编译时获取类型的属性。例如，std::is_same_v\u0026lt;T1, T2\u0026gt; 可以告诉你 T1 和 T2 是否是同一种类型。通过结合 if constexpr 和类型萃取，你可以写出高度灵活且类型安全的代码。这也是 std::visit 能发挥最大威力的地方。综合应用：泛型 lambda 与类型判断std::variant\u0026lt;int, double, std::string\u0026gt; v = \"hello\";std::visit([](auto\u0026amp;\u0026amp; arg) { using T = std::decay_t\u0026lt;decltype(arg)\u0026gt;; if constexpr (std::is_same_v\u0026lt;T, int\u0026gt;) { std::cout \u0026lt;\u0026lt; \"int: \" \u0026lt;\u0026lt; arg \u0026lt;\u0026lt; std::endl; } else if constexpr (std::is_same_v\u0026lt;T, double\u0026gt;) { std::cout \u0026lt;\u0026lt; \"double: \" \u0026lt;\u0026lt; arg \u0026lt;\u0026lt; std::endl; } else { static_assert(std::is_same_v\u0026lt;T, std::string\u0026gt;); std::cout \u0026lt;\u0026lt; \"string: \" \u0026lt;\u0026lt; arg \u0026lt;\u0026lt; std::endl; }}, v);这里，我们使用了泛型 lambda 来接受任何类型的 arg，然后用 if constexpr 和类型萃取来确定 arg 的实际类型，并据此执行相应的操作。4. 实例一：std::visit 和访问者 模式一个简单的 std::visit 使用示例。在这个例子中，我将使用 std::variant 来存储不同类型的数据，并展示如何使用 std::visit 以类型安全的方式访问和处理这些数据。#include \u0026lt;iostream\u0026gt;#include \u0026lt;variant\u0026gt;#include \u0026lt;string\u0026gt;#include \u0026lt;functional\u0026gt;// 定义 variant 类型using MyVariant = std::variant\u0026lt;int, double, std::string\u0026gt;;// 访问者函数对象struct VariantVisitor { void operator()(int i) const { std::cout \u0026lt;\u0026lt; \"处理 int: \" \u0026lt;\u0026lt; i \u0026lt;\u0026lt; std::endl; } void operator()(double d) const { std::cout \u0026lt;\u0026lt; \"处理 double: \" \u0026lt;\u0026lt; d \u0026lt;\u0026lt; std::endl; } void operator()(const std::string\u0026amp; s) const { std::cout \u0026lt;\u0026lt; \"处理 string: \" \u0026lt;\u0026lt; s \u0026lt;\u0026lt; std::endl; }};int main() { MyVariant v1 = 10; // v1 存储 int MyVariant v2 = 3.14; // v2 存储 double MyVariant v3 = \"hello\"; // v3 存储 string std::visit(VariantVisitor(), v1); // 输出: 处理 int: 10 std::visit(VariantVisitor(), v2); // 输出: 处理 double: 3.14 std::visit(VariantVisitor(), v3); // 输出: 处理 string: hello return 0;}在这个例子中： 我们定义了一个 std::variant 类型 MyVariant，它可以存储 int、double 或 std::string。 VariantVisitor 是一个重载了 operator() 的结构体，对每种可能的类型提供了一个处理方法。 在 main 函数中，我们创建了三个 MyVariant 实例，分别存储不同的类型。 使用 std::visit 调用 VariantVisitor 实例，它会自动选择并调用与 variant 当前存储的类型相匹配的重载函数。5. 实例二：进一步研究 std::visit 与访问者模式的兼容如果您想要在 operator() 中添加额外的参数，std::visit 本身不会直接支持这种用法，因为 std::visit 期望的可调用对象的参数必须与传递给它的 std::variant 类型匹配。不过，您可以通过一些技巧来实现这个功能。一种常用的方法是使用 lambda 表达式或绑定器（如 std::bind）来封装您的访问者对象和额外的参数。这里有一个简单的示例说明如何做到这一点：#include \u0026lt;variant\u0026gt;#include \u0026lt;iostream\u0026gt;#include \u0026lt;functional\u0026gt;struct MyVisitor { void operator()(int i, const std::string\u0026amp; extra) const { std::cout \u0026lt;\u0026lt; \"Int: \" \u0026lt;\u0026lt; i \u0026lt;\u0026lt; \", Extra: \" \u0026lt;\u0026lt; extra \u0026lt;\u0026lt; '\\n'; } void operator()(float f, const std::string\u0026amp; extra) const { std::cout \u0026lt;\u0026lt; \"Float: \" \u0026lt;\u0026lt; f \u0026lt;\u0026lt; \", Extra: \" \u0026lt;\u0026lt; extra \u0026lt;\u0026lt; '\\n'; } void operator()(const std::string\u0026amp; s, const std::string\u0026amp; extra) const { std::cout \u0026lt;\u0026lt; \"String: \" \u0026lt;\u0026lt; s \u0026lt;\u0026lt; \", Extra: \" \u0026lt;\u0026lt; extra \u0026lt;\u0026lt; '\\n'; }};int main() { std::variant\u0026lt;int, float, std::string\u0026gt; v; std::string extraInfo = \"Some extra information\"; v = 12; std::visit([\u0026amp;](auto\u0026amp;\u0026amp; arg){ MyVisitor{}(arg, extraInfo); }, v); v = 3.14f; std::visit([\u0026amp;](auto\u0026amp;\u0026amp; arg){ MyVisitor{}(arg, extraInfo); }, v); v = \"Hello World\"; std::visit([\u0026amp;](auto\u0026amp;\u0026amp; arg){ MyVisitor{}(arg, extraInfo); }, v);}在 C++ 中，[\u0026amp;](auto\u0026amp;\u0026amp; arg){ MyVisitor{}(arg, extraInfo); } 是一个 lambda 表达式，用于创建一个匿名函数。这个特定的 lambda 表达式用于 std::visit 调用中，以便将 std::variant 的值和额外的参数一起传递给 MyVisitor 类的 operator()。 (auto\u0026amp;\u0026amp; arg) 参数列表：这表示 lambda 接受一个名为 arg 的参数，auto\u0026amp;\u0026amp; 是一个通用引用，它可以接受任何类型的参数。在 std::visit 的上下文中，这个参数将是 std::variant 中当前存储的值。 函数体：{ MyVisitor{}(arg, extraInfo); } 是 lambda 表达式的函数体。在这里，它创建了 MyVisitor 类的一个临时实例，并调用其 operator()，传递两个参数：arg（从 std::variant 中得到的值），和 extraInfo（从外部作用域捕获的额外信息）。综合起来，当这个 lambda 表达式被 std::visit 调用时，它会根据 std::variant 当前存储的类型将相应的值作为 arg 传递给 MyVisitor 的 operator()，同时携带一个额外的参数 extraInfo。这允许 MyVisitor 的方法根据当前的 variant 类型和额外的信息执行相应的操作。6. 使用 std::visit 的优缺点6.1 优点 代码简洁 类型安全 扩展性std::visit 的优点是扩展性（Extensibility）。如果 std::variant 添加了新的类型，你只需要更新 std::visit 的访问器函数，而无需改动其他代码。7. 参考 C++17之std::visit的具体使用 【C++ 17 新功能 std::visit 】深入解析 C++17 中的 std::visit：从原理到实践 std::variant 与 std::visit Visiting a std::variant with the Overload Pattern" },

  { "title": "VTK 笔记", "url": "/blog/2023/vtk-notes/", "categories": "VTK",
  "tags": "VTK", "date": "2023-07-09", "content":
  "VTK可视化基本流程/* 二维等值线提取与可视化代码片段 */// 等值线 FiltervtkContourFilter *contourFilter = vtkContourFilter::New();contourFilter-\u0026gt;SetValue(0, ui-\u0026gt;selectedValueLabel-\u0026gt;text().toDouble());contourFilter-\u0026gt;SetInputData(polyData);// 将几何数据转换为可被渲染引擎绘制的可视化表达vtkPolyDataMapper *contourMapper = vtkPolyDataMapper::New();contourMapper-\u0026gt;SetInputConnection(contourFilter-\u0026gt;GetOutputPort());contourMapper-\u0026gt;ScalarVisibilityOff();// 需要被渲染绘制的对象contourActor = vtkActor::New();contourActor-\u0026gt;SetMapper(contourMapper);contourActor-\u0026gt;GetProperty()-\u0026gt;SetColor(1.0, 0.0, 0.0);contourActor-\u0026gt;GetProperty()-\u0026gt;SetLineWidth(2.0);// 添加到渲染器renderer-\u0026gt;AddActor(contourActor);VTK 主要 class VTK 简明教程 – 主要 classVTK开发精要：数据与管线机制（济南友泉软件有限公司） VTK开发精要：数据与管线机制FilterVTK中通过管线机制来实现组合各种算法处理数据。每一种算法是一个Filter，多个Filter连接在一起形成VTK管线。$\\color{#FF0000}{每个Filter可以分为两个组成部分：一个是算法对象，继承自vtkAlgrithm，主要负责处理输入的数据和信息；另一个是执行对象，继承自vtkExecute()，负责通知算法对象何时运行以及传递需要处理的数据和信息。}$Filter类继承自vtkAlgrithm及其子类，实例化时，内部会生成一个默认的Executive()对象，用于管理执行管线。Filter的输入数据与信息存储在输入端口中。一个Filter可能有0个输入端口（例如 Reader 对象）；也可能有一个或多个输入端口（例如，vtkGlyph3D 类需要两个输入端口，每个输入端口可以建立多个连接）。一个Filter可能有1个或多个输出端口，每个输出端口对应一个逻辑输出。例如vtkExtractVectorComponents类，从一个三维向量数据中提取每个分量数据，该Filter需要一个输入端口接受向量数据，三个输出端口用于输出三个分量数据，端口号分别为 0, 1, 2。Filter之间通过端口（Port）建立连接（Connection）。例如一个标准的连接代码如下：Filter2-\u0026gt;SetInputConnection( Filter1-\u0026gt;GetOutputPort() );该句代码将Filter1的输出端口与Filter2的输入端口建立连接，连接中只涉及了一个输入端口和一个输出端口。VTK中还有许多Filter可能需要多个输入，例如vtkGlyph3D，该类需要两个输入数据并生成一个输出数据。因此这里需要建立两个连接，相应的函数分别为SetInputConnection()和SetSourceConnection()，其中，SetInputConnection()输入的是几何点集数据，对应输入端口0，SetSourceConnection()输入的是Glyph图形数据，对应输入端口1。vtkGlyph3D中输入的两个数据具有不同的意义，因此建立了两个不同的输入端口。另外，对一个Filter的多个输入数据具有相同意义时，则只需要建立一个输入端口，并使用AddInputConnection()来添加新的连接。例如vtkAppendFilter类实现数据的合并，其多个输入数据具有相同意义，而不像vtkGlyph3D的两个输入表示不同的对象，因此其连接建立如下：apeend = vtkAppendFilter::New();append-\u0026gt;AddInputConnection( foo-\u0026gt;GetOutputPort );append-\u0026gt;AddInputConnection( bar-\u0026gt;GetOutputPort );管线的接口是通过逻辑端口（Logical Port）而不是数据流实现的，因此在形成连接的过程中不需要知道实际的数据类型，而是在执行时进行数据类型检查，以决定管线是否执行。" },

  { "title": "Linux 系统console使用命令清屏", "url": "/blog/2023/clean-screen-and-reset/", "categories": "Linux",
  "tags": "Bash, Linux", "date": "2023-07-02", "content":
  "echo -en \"\\e[H\\e[J\\e[3J\"reference clear command in Konsole" },

  { "title": "vtkUnstructuredGrid 显示 HDF5 数据", "url": "/blog/2023/show-hdf5-data-using-vtkUnstructuredGrid/", "categories": "VTK",
  "tags": "VTK", "date": "2023-06-22", "content":
  "Reference VTK examples添加 field (dataset) 数据第一步，创建VTK的field数据vtkSmartPointer\u0026lt;vtkDoubleArray\u0026gt; fieldDataArray = vtkSmartPointer\u0026lt;vtkDoubleArray\u0026gt;::New();fieldDataArray-\u0026gt;SetNumberOfComponents(1); // assuming scalar data// Assuming \"dataset\" is your std::map\u0026lt;std::string, std::vector\u0026lt;double\u0026gt;\u0026gt;for (const auto\u0026amp; pair : dataset) { const std::string\u0026amp; fieldName = pair.first; const std::vector\u0026lt;double\u0026gt;\u0026amp; fieldValues = pair.second; fieldDataArray-\u0026gt;SetName(fieldName.c_str()); for (double value : fieldValues) { fieldDataArray-\u0026gt;InsertNextValue(value); } // Assuming \"grid\" is your vtkUnstructuredGrid object grid-\u0026gt;GetPointData()-\u0026gt;AddArray(fieldDataArray);} 根据dataset的属性field是scalar数据还是tensor数据，设置SetNumberOfComponents的参数； 根据dataset的属性field是scalar数据还是tensor数据，选择InsertNectValue或者InsertNextTuple； 根据dataset的属性location type属性是vetex还是element，grid选取GetPointData或者GetCellData；第二步，添加VTK field数据到mappervtkSmartPointer\u0026lt;vtkDataSetMapper\u0026gt; mapper = vtkSmartPointer\u0026lt;vtkDataSetMapper\u0026gt;::New();mapper-\u0026gt;SetInputData(grid);mapper-\u0026gt;SetScalarModeToUsePointData(); // or SetScalarModeToUseCellData()mapper-\u0026gt;SelectColorArray(\"fieldName\"); // replace \"fieldName\" with the name of the field you want to use for coloringvtkSmartPointer\u0026lt;vtkLookupTable\u0026gt; lut = vtkSmartPointer\u0026lt;vtkLookupTable\u0026gt;::New();lut-\u0026gt;SetRange(minValue, maxValue); // set range according to your datamapper-\u0026gt;SetLookupTable(lut);vtkSmartPointer\u0026lt;vtkActor\u0026gt; actor = vtkSmartPointer\u0026lt;vtkActor\u0026gt;::New();actor-\u0026gt;SetMapper(mapper);vtkSmartPointer\u0026lt;vtkScalarBarActor\u0026gt; scalarBar = vtkSmartPointer\u0026lt;vtkScalarBarActor\u0026gt;::New();scalarBar-\u0026gt;SetLookupTable(mapper-\u0026gt;GetLookupTable());vtkSmartPointer\u0026lt;vtkRenderer\u0026gt; renderer = vtkSmartPointer\u0026lt;vtkRenderer\u0026gt;::New();renderer-\u0026gt;AddActor(actor);renderer-\u0026gt;AddActor2D(scalarBar); 根据dataset的属性location type属性是vetex还是element，mapper调用SetScalarModeToUsePointData或者SetScalarModeToUseCellData；颜色转换的另一种方式使用vtkColorTranslationFunction转换：// Create a vtkColorTransferFunction to map scalar values to colorsvtkSmartPointer\u0026lt;vtkColorTransferFunction\u0026gt; colorTransferFunction = vtkSmartPointer\u0026lt;vtkColorTransferFunction\u0026gt;::New();colorTransferFunction-\u0026gt;AddRGBPoint(minScalarValue, r, g, b); // Add as many points as needed// Set the color transfer function for the mappermapper-\u0026gt;SetLookupTable(colorTransferFunction);" },

  { "title": "Linux 系统下编译 gcc 9.2", "url": "/blog/2023/compile-gdb-commands/", "categories": "GCC",
  "tags": "GCC", "date": "2023-06-17", "content":
  "操作下载 gcc源码包，解压生成gcc-9.2.0：http://ftp.gnu.org/gnu/gcc/gcc-9.2.0/gcc-9.2.0.tar.gz在四个依赖包在以下文件中有描述：./contrib/download_prerequisites gmp-6.1.0.tar.bz2 mpfr-3.1.4.tar.bz2 mpc-1.0.3.tar.gz isl-0.18.tar.bz2下载下来之后，放到gcc_9.2.0目录下，并执行上面的这个download_prerequisites脚本。之后即可开始编译gcc9.2.0 。Build commandsCC=\"$CC\"CXX=\"$CXX\"CFLAGS=\"$OPT_FLAGS\" CXXFLAGS=\"`echo \" $OPT_FLAGS \" | sed 's/ -Wall / /g'`\"../configure --prefix=/meda_home/ai0157/opt --enable-bootstrap --enable-shared --enable-threads=posix --enable-checking=release --with-system-zlib --enable-__cxa_atexit --disable-libunwind-exceptions --enable-linker-build-id --enable-languages=c,c++,lto --disable-vtable-verify --with-default-libstdcxx-abi=new --enable-libstdcxx-debug --without-included-gettext --enable-plugin --disable-initfini-array --disable-libgcj --enable-plugin --disable-multilib --with-tune=generic --build=x86_64-unknown-linux-gnu --target=x86_64-unknown-linux-gnu --host=x86_64-unknown-linux-gnuReferences build_gcc_9.sh linux gcc-9.2.0 源码编译 Ubuntu环境下LLVM 15.0 完全编译 附windows编译LLVM master" },

  { "title": "OCCT 使用 BVH 树加速 bounding box 查找遍历", "url": "/blog/2023/OCCT-BHV-tree-bounding-box/", "categories": "OCCT",
  "tags": "OCCT", "date": "2023-06-14", "content":
  "Demo codeDemo code 1#include \u0026lt;BVH_Tree.hxx\u0026gt;#include \u0026lt;TopoDS_Shape.hxx\u0026gt;#include \u0026lt;vector\u0026gt;// Assuming you have a vector of TopoDS_Shape objects called 'shapes'// Create a BVH treeHandle(BVH_Tree\u0026lt;Standard_Real, 3\u0026gt;) bvhTree = new BVH_Tree\u0026lt;Standard_Real, 3\u0026gt;();// Insert the shapes into the BVH treefor (const auto\u0026amp; shape : shapes) { // Compute the bounding box for the shape Bnd_Box box; BRepBndLib::Add(shape, box); // Insert the bounding box into the BVH tree bvhTree-\u0026gt;Insert(box, shape);}// a query on the BVH tree to find possible interferencebvhTree-\u0026gt;Select([](const TopoDS_Shape\u0026amp; shape1, const TopoDS_Shape\u0026amp; shape2) { // Check for interference between shape1 and shape2 // If there is interference, process it accordingly});Demo code 2#include \u0026lt;SpatialHash.hxx\u0026gt;#include \u0026lt;TopoDS_Shape.hxx\u0026gt;int main(){ // Create a spatial hash from the shapes. SpatialHash hash(shapes.size()); for (TopoDS_Shape\u0026amp; shape : shapes) { hash.Add(shape); } // Find possibly interfering objects. for (TopoDS_Shape\u0026amp; shape : shapes) { std::vector\u0026lt;TopoDS_Shape\u0026gt; interferingObjects = hash.Find(shape); for (TopoDS_Shape\u0026amp; otherShape : interferingObjects) { // Check if the shapes interfere. } } return 0;}References 几何体数据结构学习（6）BVH树" },

  { "title": "OCC TopExp_Explorer 用法", "url": "/blog/2023/OCCT-explore-api-usage/", "categories": "OCCT",
  "tags": "OCCT", "date": "2023-06-12", "content":
  "#include \u0026lt;iostream\u0026gt;#include \u0026lt;TopoDS_Shape.hxx\u0026gt;#include \u0026lt;TopExp_Explorer.hxx\u0026gt;using namespace std;int main(){ // Create a TopoDS_Shape object. TopoDS_Shape shape; // Create a TopExp_Explorer object and explore the shape. TopExp_Explorer exp(shape, TopAbs_FACE); // While there are more faces, print the face's name. while (exp.More()) { cout \u0026lt;\u0026lt; exp.Current().Name() \u0026lt;\u0026lt; endl; exp.Next(); } return 0;}" },

  { "title": "OCC boundding box 以及 distance", "url": "/blog/2023/OCCT-bounding-box/", "categories": "OCCT",
  "tags": "OCCT", "date": "2023-06-12", "content":
  "bounding box 方式检测#include \u0026lt;Bnd_Box.hxx\u0026gt;#include \u0026lt;BRepBndLib.hxx\u0026gt;TopoDS_Shape shape1, shape2; // Assume these shapes are already definedBnd_Box boundingBox1, boundingBox2;BRepBndLib::Add(shape1, boundingBox1);BRepBndLib::Add(shape2, boundingBox2);bool isInterfering = !boundingBox1.IsOut(boundingBox2);distance 方式检测#include \u0026lt;BRepExtrema_DistShapeShape.hxx\u0026gt;BRepExtrema_DistShapeShape distShapeShape(shape1, shape2);Standard_Real minDistance = distShapeShape.Value();bool isInterfering = (minDistance \u0026lt;= Precision::Confusion()); BRepExtrema_DistShapeShape 比 bounding box 方式费时； 如果bouding box方式检测出来出现干涉（interference, overlap），则可以使用distance检测方式确认；" },

  { "title": "Ubuntu 及 Windows 系统下安装 Qt5", "url": "/blog/2023/install-qt5-on-ubuntu-windows/", "categories": "Qt",
  "tags": "Qt", "date": "2023-06-10", "content":
  "1. Ubuntu 安装 Qt5 命令安装Qt5：sudo apt install qtbase5-dev qtchooser qt5-qmake qtbase5-dev-tools qtmultimedia5-dev qttools5-dev qttools5-dev-tools qtcreator libqt5svg5-dev libqt5charts5 libqt5charts5-dev qtdeclarative5-dev libqt5xmlpatterns5-dev libqt5x11extras5-dev cmake-qt-guisudo apt-get -y install build-essential cmake gcc git lib32ncurses-dev lib32z1 libfox-1.6-dev libsdl1.2-dev software-properties-common wget zip python3-pip-whl python3-pil libgtest-dev python3-pip python3-tk python3-setuptools clang-14 python3-clang-14 libusb-1.0-0-dev stlink-tools openocd npm pv libncurses5:i386 libpython2.7:i386 libclang-14-dev python-is-python32. cmake find_package for Qt5find_package( Qt5 COMPONENTS Core Gui Widgets Network Svg Charts REQUIRED)3. 在线安装包 Qt5同时适用于windows/linux平台：下载在线安装包：Qt downloads安装命令：.\\qt-unified-windows-x64-online.exe --mirror https://mirrors.ustc.edu.cn/qtproject升级/维护时使用国内镜像：# 可以创建快捷方式MaintenanceTool.exe --mirror https://mirror.nju.edu.cn/qt4. 额外链接 gist Which dev packages are needed to build a QtQuick application on Ubuntu 20.04?" },

  { "title": "从Makefile创建compile_commands.json", "url": "/blog/2023/Create_CompileDB_from_Makefile/", "categories": "Bash",
  "tags": "Bash", "date": "2023-06-05", "content":
  "操作步骤# https://github.com/nickdiego/compiledbpip install compiledb使用方法：compiledb make" },

  { "title": "Ubuntu 编译 Qt + VTK + OCCT + samples", "url": "/blog/2023/OCCT-build-with-samples/", "categories": "OCCT",
  "tags": "OCCT", "date": "2023-05-25", "content":
  "1. 卸载安装的Qt，编译安装Qt实践发现使用安装的Qt，编译带samples的OCCT在CMake阶段就出错（可能是Ubuntu系统中环境有问题，或者是没有完整设置Qt相关变量）。使用手动编译安装的Qt，并设置好相关环境变量可正常编译带samples的OCCT。手动编译安装Qt见：VTK（1）：ubuntu 22.04 源码编译安装 Qt5.15.6 。2. 编译 VTKVTK （2）：ubuntu 22.04 编译 VTK 9.2 rc23. 编译 OCCT 及 samples顶层CMakeLists.txt中有变量BUILD_SAMPLES_QT控制是否编译samples（包括qt相关samples）。cmake -D3RDPARTY_QT_DIR=/usr/local/Qt-5.15.6 -DCMAKE_BUILD_TYPE=Debug -DBUILD_SAMPLES_QT=ON ..make -j4. windows 安装 Qt下载在线安装包：Qt downloads安装命令：.\\qt-unified-windows-x64-online.exe --mirror https://mirrors.ustc.edu.cn/qtproject" },

  { "title": "VTK学习资源收集", "url": "/blog/2023/VTK-official-examples/", "categories": "VTK",
  "tags": "VTK, Cpp", "date": "2023-05-24", "content":
  "学习文档 VTK textbook – online VTK Book Figure Examples VTK textbook – pdf VTK Userguide – pdf示例关于CXX的示例说明，在官方examples代码的相关README里面有简单介绍，路径：src/Cxx.md github HDF5 \u0026amp; TDR代码 JosefWeinbub C++示例代码列表 Cxx 演示vtkUnstructureGrid及mesh显示(EdgeVisibilityOn) UGridColor a mesh by dotting a vector from the origin to each point with a specified vector SimpleElevationFilterVector field – 显示矢量场箭头 VectorField VectorFieldNonZeroExtractionVTK blog VTK入门范例2" },

  { "title": "改用zsh以及oh-my-zsh", "url": "/blog/2023/%E5%AE%89%E8%A3%85zsh%E4%BB%A5%E5%8F%8A%E4%BD%BF%E7%94%A8oh-my-zsh/", "categories": "Shell",
  "tags": "Shell, zsh, oh-my-zsh", "date": "2023-05-24", "content":
  "安装zsh以及oh-my-zsh：sh -c \"$(curl -fsSL https://raw.githubusercontent.com/ohmyzsh/ohmyzsh/master/tools/install.sh)\"安装autosuggestions插件（根据历史命令自动补全）：git clone https://github.com/zsh-users/zsh-autosuggestions ${ZSH_CUSTOM:-~/.oh-my-zsh/custom}/plugins/zsh-autosuggestions在.zshrc中启用插件，plugins列表中，添加zsh-autosuggestions。例如：plugins=(git zsh-autosuggestions)添加alias：# 实用别名alias ls='ls --color=auto'alias l='ls -lh'alias ll='ls -lAh'alias la='ls -A'alias grep='grep --color=auto'# Git 别名alias g='git'alias ga='git add'alias gc='git commit'alias gp='git push'alias gl='git pull'alias gs='git status'alias gd='git diff'# 目录导航别名alias ..='cd ..'alias ...='cd ../..'alias c='clear'alias h='history'一些定制化的设置：############################ helpers: add to PATH / LD_LIBRARY_PATH only once_path_append() { [[ \":$PATH:\" != *\":$1:\"* ]] \u0026amp;\u0026amp; export PATH=\"$PATH:$1\"; }_ldpath_append() { [[ \":$LD_LIBRARY_PATH:\" != *\":$1:\"* ]] \u0026amp;\u0026amp; export LD_LIBRARY_PATH=\"${LD_LIBRARY_PATH}:$1\"; }############################ oneAPIif [[ -z \"$SETVARS_COMPLETED\" ]]; then source /opt/intel/oneapi/setvars.shfi############################ dotnet 6_path_append \"$HOME/dotnet6\"[[ -z \"$DOTNET_ROOT\" ]] \u0026amp;\u0026amp; export DOTNET_ROOT=$HOME/dotnet6############################ cargo_path_append \"$HOME/.cargo/bin\"############################ go[[ -z \"$GOPATH\" ]] \u0026amp;\u0026amp; export GOPATH=$HOME/go[[ -z \"$GOPROXY\" ]] \u0026amp;\u0026amp; export GOPROXY=https://goproxy.cn_path_append \"$GOPATH/bin\"############################ fastdds[[ -z \"$FASTDDS_ROOT\" ]] \u0026amp;\u0026amp; export FASTDDS_ROOT=$HOME/dev_path_append \"$FASTDDS_ROOT/bin\"_ldpath_append \"$FASTDDS_ROOT/lib\"A. 资料 Oh My Zsh 官方网站 Oh My Zsh GitHub 仓库 Zsh插件加载性能优化：优化zsh脚本的启动及加载速度，2026年1月10日 用 Starship + 原生插件打造极速终端博客：使用 Starship 以及 oh-my-zsh 的原生插件打造极速终端 Fish Shell 官方网站" },

  { "title": "kill信号不同分类的影响", "url": "/blog/2023/kill-signals-of-linux-system/", "categories": "Linux",
  "tags": "Bash, Linux", "date": "2023-05-24", "content":
  "与kill -KILL不同的是，kill -INT -PID 将通知被结束进程，等同于Ctrl+C。例如如果结束一个script，该script中同步启动了一个APP，使用kill -INT -\u0026lt;PIDofScript\u0026gt;可以同时将这个APP结束掉，kill -KILL则不行。 分类 信号 程序不可捕获、阻塞或忽略的信号 SIGKILL, SIGSTOP 不能恢复至默认动作的信号 SIGILL, SIGTRAP 默认会导致进程流产的信号 SIGABRT, SIGBUS, SIGFPE, SIGILL, SIGIOT, SIGQUIT, SIGSEGV, SIGTRAP, SIGXCPU, SIGXFSZ 默认会导致进程退出的信号 SIGALRM, SIGHUP, SIGINT, SIGKILL, SIGPIPE, SIGPOLL, SIGPROF, SIGSYS, SIGTERM, SIGUSR1, SIGUSR2, SIGVTALRM 默认会导致进程停止的信号 SIGSTOP, SIGTSTP, SIGTTIN, SIGTTOU 默认进程忽略的信号 SIGCHLD, SIGPWR, SIGURG, SIGWINCH Linux 下的KILL函数的用法 - 拂 晓 - 博客园 (cnblogs.com)如果kill不能结束掉，则尝试使用pkill：pkill -TERM -P \u0026lt;PID of script\u0026gt;" },

  { "title": "Bash参数特殊变量符号", "url": "/blog/2023/bash-args-and-special-symbols/", "categories": "Bash",
  "tags": "Bash, Shell", "date": "2023-05-24", "content":
  "Shell 特殊参数解释首先来看几个特殊变量：$0, $#, $*, $@, $?, $$, $_#!/bin/bashecho $0 # 当前脚本的文件名（间接运行时还包括绝对路径）。echo $n # 传递给脚本或函数的参数。n 是一个数字，表示第几个参数。例如，第一个参数是 $1 。echo $# # 传递给脚本或函数的参数个数。echo $* # 传递给脚本或函数的所有参数。echo $@ # 传递给脚本或函数的所有参数。被双引号 (\" \") 包含时，与 $* 不同，下面将会讲到。echo $? # 上个命令的退出状态，或函数的返回值。echo $$ # 当前 Shell 进程 ID。对于 Shell 脚本，就是这些脚本所在的进程 ID。echo $_ # 上一个命令的最后一个参数echo $! # 后台运行的最后一个进程的 ID 号执行结果如下：$ ./test.sh test test1 test2 test3 test4./test.sh # $0 # $n5 # $#test test1 test2 test3 test4 # $*test test1 test2 test3 test4 # $@0 # $?12305 # $$12305 # $_ # $!参数中的双引号$* 和 $@ 都表示传递给函数或脚本的所有参数，不被双引号 (“”) 包含时，都以 $1\"\"$2\" … \"$n 的形式输出所有参数。但是当它们被双引号 (“”) 包含时，\"$*\" 会将所有的参数作为一个整体，以 $1 $2 … $n 的形式输出所有参数；$@ 会将各个参数分开，以 \"$1\"\"$2\" … \"$n\" 的形式输出所有参数。传递参数方式一采用 $0, $1, $2 … 等方式获取脚本命令行传入的参数，值得注意的是，$0 获取到的是脚本路径以及脚本名，后面按顺序获取参数，当参数超过10个时(包括10个)，需要使用${10}, ${11}….才能获取到参数。#!/bin/bashecho \"脚本$0\"echo \"第一个参数$1\"echo \"第二个参数$2\"运行结果：$ ./test.sh 1 2#shell中将会输出：脚本./test.sh第一个参数1第二个参数2" },

  { "title": "Ubuntu安装Intel GPU驱动", "url": "/blog/2023/Ubuntu%E5%AE%89%E8%A3%85Intel-GPU%E9%A9%B1%E5%8A%A8/", "categories": "Linux, GPU",
  "tags": "ubuntu, Intel, GPU", "date": "2023-05-24", "content":
  "安装命令：sudo apt install intel-media-va-driver vainfo mesa-utils# 添加用户组sudo usermod -aG render,video hxf0223安装完成后，使用vainfo命令检查是否正确安装：# tty需要加上--display drm参数，否则会提示无法打开显示设备(无桌面会话)vainfo --display drm --device /dev/dri/renderD128顺便安装上FFmpeg：sudo apt install libv4l-dev v4l-utilssudo apt install mesa-common-devsudo apt install ffmpeg libavfilter-dev libavdevice-dev libavutil-dev libavformat-dev libswresample-dev libswscale-dev FFmpeg tutorials：github上的FFmpeg教程，包含了很多使用FFmpeg的示例代码，可以参考学习。" },

  { "title": "Ubuntu 安装 OCCT", "url": "/blog/2023/Ubuntu-install-OCCT/", "categories": "OCCT",
  "tags": "OCCT", "date": "2023-05-24", "content":
  "OCCT 依赖库 Introduction – Requirements下载OpenCascadeOCC 7.7编译有问题，使能VTK时编译错误，下载最新版的OCC修复该问题：# https://dev.opencascade.org/resources/git_repositorygit clone https://git.dev.opencascade.org/repos/occt.git occt编译安装命令1. 编译安装第三方库 注意要 tcl-dev, tk-dev, tcllib, tklib，除非自己编译安装，此时要设置TCL/TK相关的路径，比较麻烦。sudo apt install doxygen doxygen-gui graphviz graphviz-doc libx11-xcb-devsudo apt-get install tcllib tklib tcl-dev tk-dev libfreetype-dev libx11-dev libgl1-mesa-dev libfreeimage-devsudo apt-get install rapidjson-dev libdraco-devtcl 8.6编译安装命令(tk 8.6相同的编译配置命令)：# https://www.tcl.tk/software/tcltk/8.6.htmlcd tcl8.6.13/unix./configure --enable-gcc --enable-shared --enable-threads --enable-64bitmake \u0026amp;\u0026amp; sudo make installfreeType编译安装命令：# https://freetype.org/download.htmlCFLAGS='-m64 -fPIC' CPPFLAGS='-m64 -fPIC' ./configuremake \u0026amp;\u0026amp; sudo make installFreeImage编译安装命令：# https://freeimage.sourceforge.io/download.html# 修改 Makefile.fip：增加 CXXFLAGS += -std=c++11make -f Makefile.fip# 拷贝头文件及 .so, .a 到相应目录编译安装fdk-acc及ffmpeg:# install dependenciessudo apt install libavformat-dev libavcodec-dev libswresample-dev libswscale-dev libavutil-dev libsdl1.2-dev libx264-dev nasm# tdk-acc./configure \u0026amp;\u0026amp; make \u0026amp;\u0026amp; sudo make install# ffmpeg./configure --enable-gpl --enable-nonfree --enable-libfdk-aac --enable-libx264 --enable-filter=delogo --enable-shared --enable-pthreadsmake \u0026amp;\u0026amp; sudo make install2. 编译安装 OCCTcd opencascade-7.7.0mdkir build \u0026amp;\u0026amp; cd build# 配置OCC编译，使能VTK, ITK等功能ccmake ..make -j6sudo make install几次配置截图之后，配置如下3. 安装或编译 Qt5 安装 Qt5sudo apt-get install qtbase5-dev qtchooser qt5-qmake qtbase5-dev-toolssudo apt-get install qtcreator 编译 Qt5编译及设置环境变量见 VTK（1）：ubuntu 22.04 源码编译安装 Qt5.15.6国内学习博客 OpenCascade基本框架介绍-昨夜星辰 (hustlei.github.io) OpenCASCADE入门指南 - opencascade - 博客园 (cnblogs.com)http 链接：2.OpenCASCADE - eryar - C++博客 (cppblog.com) http://cppblog.com/eryar/category/17808.html?Show=All引用资料 Build 3rd-parties - Open CASCADE Technology Documentation Build OCCT - Open CASCADE Technology Documentation Download - Open CASCADE Technology 3rd party Components TCL/TK 8.6" },

  { "title": "OCCT projects on github", "url": "/blog/2023/OCCT-sample-projects/", "categories": "OCCT",
  "tags": "OCCT", "date": "2023-05-24", "content":
  "OCCT samples需要安装 OpenGL： 基于Ubuntu搭建OpenGL开发环境sudo apt-get install libglfw3-devsudo apt install mesa-utils以下两个项目均来自 github gkv311 occt-samples-qopenglwidget occt-helloOCCT 文章Boolean一些http链接：Open Cascade中的布尔操作 http://cppblog.com/eryar/archive/2013/01/17/197357.html OpenCASCADE 布尔运算简介 OpenCasCade拓扑几何的布尔运算 Boolean Operations OpenCascade Modeling Algorithms Boolean OperationsFillet autoCAD doc FilletMesh Mesh" },

  { "title": "Office笔记", "url": "/blog/2023/%E5%8A%9E%E5%85%AC%E8%BD%AF%E4%BB%B6-Office%E7%AC%94%E8%AE%B0/", "categories": "Computer",
  "tags": "Computer, Office", "date": "2023-05-05", "content":
  "1. 自定义页脚格式1：第x页/共y页首先双击页面底部进入页脚编辑，输入：第页/共页然后鼠标放在需要编辑的位置，比如”第”与”页”之间，鼠标操作如下：“共”与”页”之间同样操作，从对话框的列表中选择NUMPAGES即可。最终样式如下：格式2：第x页/共y页，且总页数从正文数起假设前面4页是封面、目录等。首先依然输入”“第页/共页”“，鼠标依旧放在需要编辑的位置（以”共”与”页”之间为例），作如下操作： 按下快捷键Ctrl + F9，会出现一对大括号{}。这个步骤为输入最外层括号。 输入等于号：=。 再次按下快捷键Ctrl + F9，会出现一对大括号{}，在大括号内输入NUMPAGES，表示总页数。 最后在等于号后输入减去4的表达式：-4。同样的可以输入”PAGE”。一个样例如下：在编辑的过程中，使用快捷键Alt + F9可以切换显示域代码和域结果，使用快捷键F9可以更新域代码的结果。要注意的是，不能使用直接输入/粘贴的大括号，而应该使用word中function功能的快捷键创建大括号。快捷键如下：1.1. 参考 How to make NUMPAGES show 1 less page? The Function and Shortcut Keys that manipulate fields" },

  { "title": "Ubuntu 系统安装 LLVM 套件 （可选择版本）", "url": "/blog/2023/install-versioned-llvm-ubuntu/", "categories": "LLVM",
  "tags": "LLVM, Clang", "date": "2023-05-05", "content":
  "wget https://mirrors.tuna.tsinghua.edu.cn/llvm-apt/llvm.shchmod u+x llvm.shsudo ./llvm.sh 18 -m https://mirrors.tuna.tsinghua.edu.cn/llvm-apt # install llvm, clang 18# sudo ./llvm.sh 18 all -m https://mirrors.tuna.tsinghua.edu.cn/llvm-apt也可以使用apt命令安装：sudo apt install clangd-21 clang-format-21 clang-21设置默认版本：sudo update-alternatives --install /usr/bin/clang clang /usr/bin/clang-21 210sudo update-alternatives --install /usr/bin/clang++ clang++ /usr/bin/clang++-21 210sudo update-alternatives --install /usr/bin/clangd clangd /usr/bin/clangd-21 210sudo update-alternatives --install /usr/bin/clang-format clang-format /usr/bin/clang-format-21 210查看当前版本：update-alternatives --display clangupdate-alternatives --display clang++完整版安装以及设置：LLVM_VERSION=21sudo ./llvm.sh $LLVM_VERSION all -m https://mirrors.tuna.tsinghua.edu.cn/llvm-aptsudo update-alternatives --install /usr/bin/clang clang /usr/bin/clang-$LLVM_VERSION $LLVM_VERSION --slave /usr/bin/clang++ clang++ /usr/bin/clang++-$LLVM_VERSIONsudo update-alternatives --install /usr/bin/llvm-config llvm-config /usr/bin/llvm-config-$LLVM_VERSION $LLVM_VERSION --slave /usr/bin/llvm-ar llvm-ar /usr/bin/llvm-ar-$LLVM_VERSION --slave /usr/bin/llvm-as llvm-as /usr/bin/llvm-as-$LLVM_VERSION --slave /usr/bin/llvm-link llvm-link /usr/bin/llvm-link-$LLVM_VERSION --slave /usr/bin/llvm-nm llvm-nm /usr/bin/llvm-nm-$LLVM_VERSION --slave /usr/bin/llvm-objdump llvm-objdump /usr/bin/llvm-objdump-$LLVM_VERSION --slave /usr/bin/llvm-ranlib llvm-ranlib /usr/bin/llvm-ranlib-$LLVM_VERSION How to install Clang 17 or 18 in Ubuntu 22.04 20.04LLVM下载镜像（包含Windows/Ubuntu安装包）：清华大学开源软件镜像站。" },

  { "title": "通过函数指针地址找到函数", "url": "/blog/2023/get-function-via-address/", "categories": "GDB",
  "tags": "GDB, Linux, Cpp", "date": "2023-05-05", "content":
  "# 通过地址找到函数声明info symbol 0x7f0db14cf57e# 通过地址找到函数在哪一行info line *0x7f0db14cf57e# 查看加载了哪些共享库info sharedlibrary参考 16 Examining the Symbol Table" },

  { "title": "写给大家看的设计模式", "url": "/blog/2023/design-modes-for-all/", "categories": "Cpp",
  "tags": "DesignPattern", "date": "2023-05-05", "content":
  "写给大家看的设计模式 图解TensorFlow源码1. 抽象工厂模式 被创建抽象类 Iuser； 抽象工厂类 IFactory，定义创建接口 IUser* createUser()。当需要新增被创建类型时，需要新增被创建类型及对应工厂类型；// 代表数据库中User表中的一条记录class RecUser { int32_t uid; std::string uname;};// 代表数据库中Department表中一条记录class RecDepartment { int32_t did; std::string dname;};// 1. 用户表操作抽象接口class IUser {public: virtual void insert(const RecUser *user) = 0; virtual RecUser *getUserRecord(int32_t id) = 0; virtual ~IUser() = default;};// 2.1 User表操作实现类：SQL Serverclass SqlServerUser : public IUser {public: void insert(const RecUser *user) override { std::cout \u0026lt;\u0026lt; \"Insert a user record into SQL Server table.\" \u0026lt;\u0026lt; std::endl; } RecUser *getUserRecord(int32_t id) override { std::cout \u0026lt;\u0026lt; \"Get a record from SQL Server table.\" \u0026lt;\u0026lt; std::endl; return nullptr; } ~SqlServerUser() override { std::cout \u0026lt;\u0026lt; \"dtor of SqlServerUser.\" \u0026lt;\u0026lt; std::endl; }};// 2.2 User表操作实现类：access 表class AccessUser : public IUser {public: void insert(const RecUser *user) override { std::cout \u0026lt;\u0026lt; \"Insert a user record into access table.\" \u0026lt;\u0026lt; std::endl; } RecUser *getUserRecord(int32_t id) override { std::cout \u0026lt;\u0026lt; \"Get a record from access table.\" \u0026lt;\u0026lt; std::endl; return nullptr; } ~AccessUser() override { std::cout \u0026lt;\u0026lt; \"dtor of AccessUser. \" \u0026lt;\u0026lt; std::endl; }};// 3. 抽象工厂类class IFactory {public: virtual IUser *createUser() = 0;};// 4.1 SQL Server 具体工厂类class SqlFactory : public IFactory {public: IUser *createUser() override { return new SqlServerUser; }};// 4.2 access 具体工厂类class AccessFactory : public IFactory {public: IUser *createUser() override { return new AccessUser; }};TEST_CASE(\"design mode: abstract factory\") { RecUser user_rec; // 创建类型1实例：SQL表操作 std::shared_ptr\u0026lt;IFactory\u0026gt; factory(new SqlFactory); std::shared_ptr\u0026lt;IUser\u0026gt; user_table(factory-\u0026gt;createUser()); user_table-\u0026gt;insert(\u0026amp;user_rec); user_table-\u0026gt;getUserRecord(1); // 创建类型2实例：access表操作 factory.reset(new AccessFactory); user_table.reset(factory-\u0026gt;createUser()); user_table-\u0026gt;insert(\u0026amp;user_rec); user_table-\u0026gt;getUserRecord(1);}1.1 简化的抽象工厂模式struct HotDrinkFactory { virtual unique_ptr\u0026lt;HotDrink\u0026gt; make() const = 0;};struct CoffeeFactory : public HotDrinkFactory { unique_ptr\u0026lt;HotDrink\u0026gt; make() const override { return make_unique\u0026lt;Coffee\u0026gt;(); }};class DrinkFactory { std::map\u0026lt;std::string, std::unique_str\u0026lt;HotDrinkFactory\u0026gt;\u0026gt; hot_factories_;public: DrinkFactory() { hot_factories_[\"coffee\"] = make_unique\u0026lt;CoffeeFactory\u0026gt;(); hot_factories_[\"tea\"] = make_unique\u0026lt;TeaFactory\u0026gt;(); } unique_ptr\u0026lt;HotDrink\u0026gt; make_drink(const string\u0026amp; name) { auto drink = hot_factories_[name]-\u0026gt;make(); drink-\u0026gt;prepare(200); // oops! return drink; }};2. Builder模式class PersonAddressBuilder : public PersonBuilderBase { typedef PersonAddressBuilder self;public: explicit PersonAddressBuilder(Person\u0026amp; person) : PersonBuilderBase{ person } {} self\u0026amp; at(std::string street_address) { person.street_address = street_address; return *this; } self\u0026amp; with_postcode(std::string post_code) { /*...*/ return *this; } self\u0026amp; in(std::string city) { /*...*/ return *this; }};// 使用，赋值属性后返回引用可以持续进行属性赋值Person p = Person::create() .lives().at(\"123 London Road\") .with_postcode(\"SW1 1GB\") .in(\"London\") .works().at(\"PragmaSoft\") .as_a(\"Consultant\") .earning(10e6);3. 享元模式 (flyweight)享元(flyweight)模式，是一种结构型设计模式，用于减少内存使用。典型由三个部分组成： flyweight – 享元：存储共享的状态，不随时间变化的存储对象，可以被共享。 享元工厂：创建并管理享元。 客户端：使用享元工厂获取享元并操作。#include \u0026lt;iostream\u0026gt;#include \u0026lt;string\u0026gt;#include \u0026lt;unordered_map\u0026gt;#include \u0026lt;memory\u0026gt;// 享元类class Character {public: virtual void display(int width, int height, int pointSize) = 0;};// 具体享元类class ConcreteCharacter : public Character {public: ConcreteCharacter(char symbol) : symbol_(symbol) {} void display(int width, int height, int pointSize) override { std::cout \u0026lt;\u0026lt; \"Character: \" \u0026lt;\u0026lt; symbol_ \u0026lt;\u0026lt; \", Width: \" \u0026lt;\u0026lt; width \u0026lt;\u0026lt; \", Height: \" \u0026lt;\u0026lt; height \u0026lt;\u0026lt; \", Point Size: \" \u0026lt;\u0026lt; pointSize \u0026lt;\u0026lt; std::endl; }private: char symbol_;};// 享元工厂类class CharacterFactory {public: std::shared_ptr\u0026lt;Character\u0026gt; getCharacter(char symbol) { if (characters_.find(symbol) == characters_.end()) { characters_[symbol] = std::make_shared\u0026lt;ConcreteCharacter\u0026gt;(symbol); } return characters_[symbol]; }private: std::unordered_map\u0026lt;char, std::shared_ptr\u0026lt;Character\u0026gt;\u0026gt; characters_;};// 客户端代码int main() { CharacterFactory factory; std::shared_ptr\u0026lt;Character\u0026gt; characterA = factory.getCharacter('A'); std::shared_ptr\u0026lt;Character\u0026gt; characterB = factory.getCharacter('B'); std::shared_ptr\u0026lt;Character\u0026gt; characterC = factory.getCharacter('A'); // 共享已有的'A'对象 characterA-\u0026gt;display(10, 20, 12); characterB-\u0026gt;display(15, 25, 14); characterC-\u0026gt;display(10, 20, 12); // 共享的'A'对象 return 0;}3.1 Boost.FlyweightBoost.Flyweight可以用来创建和管理共享对象。struct User2{ boost::flyweight\u0026lt;std::string\u0026gt; first_name, last_name; //类似一个缓存 User2(const std::string \u0026amp;first_name, const std::string \u0026amp;last_name) : first_name(first_name), last_name(last_name) {}};void boost_flyweight(){ User2 user1{\"John\", \"Smith\"}; User2 user2{\"Jane\", \"Smith\"}; std::cout \u0026lt;\u0026lt; user1.first_name \u0026lt;\u0026lt; endl; std::cout \u0026lt;\u0026lt; std::boolalpha //std::boolalpha 可以把 bool 变成 true/false 字符串 \u0026lt;\u0026lt; (\u0026amp;user1.first_name.get() == \u0026amp;user2.first_name.get()) \u0026lt;\u0026lt; std::endl; std::cout \u0026lt;\u0026lt; std::boolalpha \u0026lt;\u0026lt; (\u0026amp;user1.last_name.get() == \u0026amp;user2.last_name.get()) \u0026lt;\u0026lt; std::endl;}" },

  { "title": "Python与C++混合调试", "url": "/blog/2023/debug-combind-Python-and-C++/", "categories": "GDB",
  "tags": "GDB", "date": "2023-05-05", "content":
  "vscode启动python与C++混合调试时，gdb需要管理员权限。Remote attach using non-root account would fail directly1. 取消限制sudo sysctl -w kernel.yama.ptrace_scope=02. 取消限制，永久有效设置pkexec的权限。新建文件，并重启Ubuntu系统：# /usr/share/polkit-1/actions/com.ubuntu.pkexec.gdb.policy\u0026lt;?xml version=\"1.0\" encoding=\"UTF-8\"?\u0026gt;\u0026lt;!DOCTYPE policyconfig PUBLIC \"-//freedesktop//DTD PolicyKit Policy Configuration 1.0//EN\" \"http://www.freedesktop.org/standards/PolicyKit/1/policyconfig.dtd\"\u0026gt;\u0026lt;policyconfig\u0026gt; \u0026lt;action id=\"com.ubuntu.pkexec.gdb-settings\"\u0026gt; \u0026lt;icon_name\u0026gt;gdb-settings\u0026lt;/icon_name\u0026gt; \u0026lt;defaults\u0026gt; \u0026lt;allow_any\u0026gt;yes\u0026lt;/allow_any\u0026gt; \u0026lt;allow_inactive\u0026gt;yes\u0026lt;/allow_inactive\u0026gt; \u0026lt;allow_active\u0026gt;yes\u0026lt;/allow_active\u0026gt; \u0026lt;/defaults\u0026gt; \u0026lt;annotate key=\"org.freedesktop.policykit.exec.path\"\u0026gt;/usr/bin/gdb\u0026lt;/annotate\u0026gt; \u0026lt;annotate key=\"org.freedesktop.policykit.exec.allow_gui\"\u0026gt;true\u0026lt;/annotate\u0026gt; \u0026lt;/action\u0026gt;\u0026lt;/policyconfig\u0026gt;3. 如何启动调试 按照 launch.json 中定义的两个启动配置，先启动 “Python: Current File”，并设置好断点，暂停执行。 再启动 “GDB Attach proc 0”。在启动 attach gdb的过程中，需要在列出来的进程中手动选择刚才启动的python调试进程。 使用vscode调试时，由于命令行较长，且子进程较多。此时如果需要找对应的两个进程ID，方法如下# 搜索被调试的python文件ps -ef | grep relay00_graph.py显示结果如下：hxf0223 3877 2970 0 19:13 pts/5 00:00:00 /home/hxf0223/anaconda3/bin/python /home/hxf0223/.vscode-server/extensions/ms-python.python-2022.18.2/pythonFiles/lib/python/debugpy/adapter/../../debugpy/launcher 40627 -- /home/hxf0223/work/tvm_study/python/relay00_graph.pyhxf0223 3882 3877 3 19:13 pts/5 00:00:03 /home/hxf0223/anaconda3/bin/python /home/hxf0223/.vscode-server/extensions/ms-python.python-2022.18.2/pythonFiles/lib/python/debugpy/adapter/../../debugpy/launcher/../../debugpy --connect 127.0.0.1:39707 --configure-qt none --adapter-access-token 138af13a7101628b9d16326d38bf3a390f664ea5d7afcdbebbbc4a67c10708ec /home/hxf0223/work/tvm_study/python/relay00_graph.pyhxf0223 3945 3486 0 19:15 pts/7 00:00:00 grep --color=auto relay00_graph.p其中，带有 “–connect 127.0.0.1:39707”的进程即为gdb进程，即libtvm.so等在这个进程中加载。参考 Debugging Mixed Python C++ code in Visual Studio Code gdb dashboard – gdbinit" },

  { "title": "收藏：C++ 代码仓库，CPU \u0026 性能资料", "url": "/blog/2023/cpu_and_perf_study_docs_collections/", "categories": "Cpp",
  "tags": "Cpp, Performance, CPU", "date": "2023-05-05", "content":
  "1. oneTBBIntel并行库 oneTBB.包含tbb malloc.2. 高性能 json 库glaze3. 100行头文件代码实现的线程池ThreadPool4. 事件驱动 Sigslot Signals eventpp5. 学习资料 MoreEffectiveC++笔记 现代CPU性能分析与优化 – CPU架构，性能分析 C++17 the complete guide Learn LLVM 17" },

  { "title": "C++ 学习资源 及 代码片段积累", "url": "/blog/2023/cpp-code-snippet/", "categories": "Cpp",
  "tags": "Cpp", "date": "2023-05-05", "content":
  "1. C++ 学习资源 Fluent C++ Modern C++ foonathan::​blog() C++ Stories Sutter’s Mill 这些资源帮助你深入学习C++2. C++ 代码片段2.1. 获取数组长度#include \u0026lt;type_traits\u0026gt;struct StructDef { int32_t arr[32];};StructDef sd;const size_t len = std::extent\u0026lt;decltype(sd.arr)\u0026gt;::value;std::vector\u0026lt;int32_t\u0026gt; vec(sd.arr, sd.arr + len);2.2. CHECK#include \u0026lt;iostream\u0026gt;#include \u0026lt;cstdlib\u0026gt;#if defined DEBUG || defined _DEBUG#define CHECK2(condition, message) \\ (!(condition)) ? (std::cerr \u0026lt;\u0026lt; \"Assertion failed: (\" \u0026lt;\u0026lt; #condition \u0026lt;\u0026lt; \"), \" \\ \u0026lt;\u0026lt; \"function \" \u0026lt;\u0026lt; __FUNCTION__ \u0026lt;\u0026lt; \", file \" \u0026lt;\u0026lt; __FILE__ \u0026lt;\u0026lt; \", line \" \u0026lt;\u0026lt; __LINE__ \u0026lt;\u0026lt; \".\" \u0026lt;\u0026lt; std::endl \\ \u0026lt;\u0026lt; message \u0026lt;\u0026lt; std::endl, \\ abort(), 0) \\ : 1#else#define CHECK2(condition, message) (void)0#endif2.3. 获取线程 ID通过 pthread_self 及 std::this_thread::getid 函数获取的线程 ID，跟使用 top, htop 命令呈现的线程 ID 不对应。通过如下代码获取跟 top, htop 命令一致的 TID：#include \u0026lt;syscall.h\u0026gt;pid_t gettid(void) { return (pid_t)syscall(SYS_gettid);}// 设置所在线程名称const char* thd_name = \"demo_thread\";prctl(PR_SET_NAME, reinterpret_cast\u0026lt;unsigned long\u0026gt;(thd_name ), 0, 0, 0);2.4. 获取本地IP地址列表#include \u0026lt;arpa/inet.h\u0026gt;#include \u0026lt;sys/socket.h\u0026gt;#include \u0026lt;netdb.h\u0026gt;#include \u0026lt;ifaddrs.h\u0026gt;#include \u0026lt;stdio.h\u0026gt;#include \u0026lt;stdlib.h\u0026gt;#include \u0026lt;unistd.h\u0026gt;#include \u0026lt;string.h\u0026gt;#include \u0026lt;vector\u0026gt;#include \u0026lt;string\u0026gt;std::vector\u0026lt;std::string\u0026gt; get_local_ip_v2() { std::vector\u0026lt;std::string\u0026gt; addr_list; struct ifaddrs *ifaddr, *ifa; char host[NI_MAXHOST]; if (getifaddrs(\u0026amp;ifaddr) == -1) { perror(\"getifaddrs\"); exit(EXIT_FAILURE); } for (ifa = ifaddr; ifa != nullptr; ifa = ifa-\u0026gt;ifa_next) { if (ifa-\u0026gt;ifa_addr == nullptr) { continue; } const auto s = getnameinfo(ifa-\u0026gt;ifa_addr, sizeof(struct sockaddr_in), host, NI_MAXHOST, nullptr, 0, NI_NUMERICHOST); if (ifa-\u0026gt;ifa_addr-\u0026gt;sa_family == AF_INET) { // strcmp(ifa-\u0026gt;ifa_name,\"wlan0\") == 0 if (s != 0) { printf(\"getnameinfo failed: %s\\n\", gai_strerror(s)); break; } // printf(\"\\tInterface : \u0026lt;%s\u0026gt;\\n\", ifa-\u0026gt;ifa_name); // printf(\"\\t Address : \u0026lt;%s\u0026gt;\\n\", host); addr_list.push_back(host); } } freeifaddrs(ifaddr); return addr_list;}3. 其他3.1 doxygen 注释// Doxygen documentation generator set \"doxdocgen.c.triggerSequence\": \"/\", // 触发自动注释的生成 \"doxdocgen.c.commentPrefix\": \" * \", // 注释行的前缀 \"doxdocgen.c.firstLine\": \"/**\", // 注释行的首行 \"doxdocgen.c.lastLine\": \"*/\", // 注释行的尾行 // Smart text snippet for factory methods/functions. \"doxdocgen.c.factoryMethodText\": \"Create a {name} object\", // Smart text snippet for getters. \"doxdocgen.c.getterText\": \"Get the {name} object\", // Smart text snippet for setters. \"doxdocgen.c.setterText\": \"Set the {name} object\", // Smart text snippet for constructors. \"doxdocgen.cpp.ctorText\": \"Construct a new {name} object\", // Smart text snippet for destructors. \"doxdocgen.cpp.dtorText\": \"Destroy the {name} object\", // The template of the template parameter Doxygen line(s) that are generated. If empty it won't get generated at all. \"doxdocgen.cpp.tparamTemplate\": \"@tparam {param} \", // 文件注释：版权信息模板 \"doxdocgen.file.copyrightTag\": [ \"@copyright Copyright (c) {year} 深圳市奥瓦机器人有限公司\" ], // 文件注释：自定义模块，这里我添加一个修改日志 // \"doxdocgen.file.customTag\": [ // \"@par 修改日志:\", // \"\u0026lt;table\u0026gt;\", // \"\u0026lt;tr\u0026gt;\u0026lt;th\u0026gt;Date \u0026lt;th\u0026gt;Version \u0026lt;th\u0026gt;Author \u0026lt;th\u0026gt;Description\", // \"\u0026lt;tr\u0026gt;\u0026lt;td\u0026gt;{date} \u0026lt;td\u0026gt;1.0.1 \u0026lt;td\u0026gt;zhoulq \u0026lt;td\u0026gt;内容\", // \"\u0026lt;/table\u0026gt;\", // ], // 文件注释的组成及其排序 \"doxdocgen.file.fileOrder\": [ \"file\", // @file \"brief\", // @brief 简介 \"author\", // 作者 \"version\", // 版本 \"date\", // 日期 \"empty\", // 空行 \"copyright\",// 版权 \"empty\", \"custom\" // 自定义 ], // 下面时设置上面标签tag的具体信息 \"doxdocgen.file.fileTemplate\": \"@file {name}\", \"doxdocgen.file.versionTag\": \"@version 1.0.1\", \"doxdocgen.generic.authorEmail\": \"liqun_zhou91@163.com\", \"doxdocgen.generic.authorName\": \"zhoulq\", \"doxdocgen.generic.authorTag\": \"@author {author} ({email})\", // 日期格式与模板 \"doxdocgen.generic.dateFormat\": \"YYYY-MM-DD\", \"doxdocgen.generic.dateTemplate\": \"@date {date}\", // 根据自动生成的注释模板（目前主要体现在函数注释上） \"doxdocgen.generic.order\": [ \"brief\", \"tparam\", \"param\", \"return\", \"author\", \"date\" ], \"doxdocgen.generic.paramTemplate\": \"@param{indent:8}{param}{indent:8}\", \"doxdocgen.generic.returnTemplate\": \"@return {type} \", \"doxdocgen.generic.splitCasingSmartText\": true, \"doxdocgen.generic.includeTypeAtReturn\": true, // return 中包含类型信息 \"doxdocgen.generic.boolReturnsTrueFalse\": false, // bool 返回值拆分成 true 和 false 两种情况 \"doxdocgen.generic.linesToGet\": 20, // 回车后最多向下多少行去找函数声明 \"doxdocgen.generic.useGitUserName\": false, // {author} 是都根据 git config --get user.name 替换 \"doxdocgen.generic.useGitUserEmail\": false, //declarations or definitions anymore. VSCode用Doxygen自定义代码注释配置 doxygen注释规范" }

]