GPU全栈博主

原创 一文带你掌握GPU编程：从入门到精通

本系列文章持续更新中...

原创 一文带你掌握C++编程：从入门到精通

本系列文章持续更新中...

原创 驱动开发系列84 – Mesa NIR 编译器各个 pass 功能介绍

本文介绍了NIR（Nouveau Intermediate Representation）编译器的九种关键优化技术：1）变量初始化拆分；2）返回语句结构化转换；3）函数内联；4）指针访问优化；5）变量SSA转换；6）死变量消除；7）拷贝传播；8）死代码消除；9）公共子表达式消除。这些优化技术通过简化控制流、消除冗余计算、优化内存访问等方式，显著提升中间代码的执行效率，涵盖了从变量处理到代码结构优化的完整编译流程。

原创 驱动开发系列83 – CPU到设备的可见性问题（memset、memcpy与内存屏障）

在普通应用开发中，memset 和 memcpy 可以安全地操作内存。但在驱动开发中（尤其是GPU驱动开发中），这种假设不成立。原因在于CPU写入不等于设备（GPU）立即可见。写入问题主要出现在以下场景：（1）写设备内存（MMIO和显存）。（2）写uncached/write-combine内存。（3） ARM64 等弱内存模型架构。这些场景下如果memset，memcpy使用不当，会出现GPU读取到旧数据，命令缓冲区执行异常以及偶现（极难复现）的bug。

原创 渲染引擎实践 - UnrealEngine 目录结构

原创 驱动开发系列82 – Mesa NIR 和 SPIRV 的Debug方法

备注：如果没有 spirv-dis，请安装 sudo apt install spirv-tools。然后用 spirv-dis 查看 spirv ，比如 spirv-dis dump01.spv。使用 ZINK_DEBUG环境变量，比如 ZINK_DEBUG=spirv。二：如何查看SPIRV？一：如何查看 NIR？

原创 C++26中的反射

反射（Reflection）是指程序能够检查、分析，甚至修改自身结构与行为的一种能力。在编程语言的发展过程中，反射逐渐成为实现元编程与自动化代码生成的重要工具。在 C++ 语言这边，反射能力正在不断增强。根据提案 P2996R5 中的表述：“我们不仅希望观察程序的结构，还希望简化基于这些观察结果的代码生成过程。这种结合有时被称为反射元编程（reflection metaprogramming）。“就反射和编译时元编程而言，本提案并非最终方案。

原创 驱动开发系列81 – Mesa NIR 数据结构介绍

在前面的文章中，我介绍了NIR的整体结构，今天将继续深入，重点讲解NIR的数据结构。理解这些数据结构至关重要，因为NIR的遍历和Pass优化都依赖于它们。掌握这些数据结构是全面理解NIR的基础。NIR的代码位置是在Mesa仓库下的 mesa/src/compiler/nir下面。NIR 的整体结构可以看作是一棵分层组织的 IR 树，从顶层的 nir_shader 一直到最底层的 nir_instr，层层递进。整体结构如下：1. nir_shader 是NIR的最顶层结构，代表一个完整的 Shader。

原创 Ubuntu 20.04 安装有道词典

安装：运行：

原创 驱动开发系列80 - Mesa SPIR-V 简介

本文通过一个输出固定红色的片段着色器示例，介绍了SPIR-V语言的基本结构。示例展示了GLSL代码编译生成的SPIR-V二进制格式，其中包含着色器版本声明、内存模型、入口点定义、类型声明和输出变量设置等关键部分。SPIR-V代码逐行对应GLSL源代码的功能实现，最终将红色值(1.0,0.0,0.0,1.0)存储到输出变量中。该示例为理解GLSL到SPIR-V的编译过程提供了入门参考，适合对着色器编译器开发感兴趣的读者学习。

原创 Ubuntu 20.04 Linux安装钉钉

安装钉钉：sudo dpkg -i ./com.alibabainc.dingtalk_8.1.0.6011301_amd64.deb。下载钉钉：https://www.dingtalk.com/download。打开钉钉：com.alibabainc.dingtalk。

原创 驱动开发系列79 - 搭建 Zink 调试环境

本文详细介绍了Mesa图形驱动调试环境的搭建方法。首先对比了两种获取Mesa源码的方式：发行版打包的稳定源码和上游最新开发源码，推荐新手使用前者。然后分别讲解了两种编译方法：通过apt构建发行版源码和手动配置编译上游源码，重点介绍了关键编译选项和依赖安装。最后说明了调试环境的配置方法，包括驱动路径设置和断点调试技巧。文章为开发者提供了从源码获取到编译调试的完整指南，特别适合需要进行Mesa驱动开发的读者参考。

原创 C++ SHA-1 算法实现原理

在密码学中，SHA-1（Secure Hash Algorithm 1，安全哈希算法1）是一种密码学哈希函数，它接受任意长度的输入数据，并生成一个 160 位（20 字节） 的哈希值，称为消息摘要，通常以 40 个十六进制字符表示。SHA-1 由**美国国家安全局（NSA）**设计，并被采纳为美国联邦信息处理标准（FIPS）。然而，SHA-1 已被密码学界成功破解。自 2005 年起，SHA-1 即被认为无法抵御具备充足计算资源的攻击者；到 2010 年，多个组织已明确建议停止使用该算法。

原创 Ubuntu 20.04 Linux安装微信

安装微信：sudo dpkg -i ./WeChatLinux_x86_64.deb。打开微信：wechat。

原创 驱动开发系列78 - 应用程序渲染与显示（二)

然而，这仅是冰山一角。对于驻留在系统内存中的缓冲区对象，应用程序会创建一个指向该缓冲区内存的文件描述符，并通过已建立连接的流式套接字以单条、低开销的消息形式将其传递给合成器。一种简单的实现方法是依次应用每个阶段的配置：先在 CRTC 上设置显示模式，再将所有缓冲区对象上传至图形内存，接着配置用于扫描输出的帧缓冲区和平面，最后启用编码器和连接器。合成器作为一种系统级服务，负责接收各个应用程序提交的输出缓冲区，并按照既定的布局策略将它们组合成最终的屏幕图像，常见的布局方式包括堆叠式和铺排式。

原创 驱动开发系列77 - 应用程序渲染与显示（一)

渲染过程中，所有图形数据都存放在缓冲区对象中，例如顶点数据、纹理数据以及最终的输出图像本身，因此图形渲染在很大程度上可以看作是一系列围绕图形内存的管理与计算操作。介于两者之间的还有具有DMA或共享图形内存的图形芯片，独立设备的图形地址重映射表（GART）内存，以及板载显卡中由固件或系统预留、专门用于图形处理的内存区域。因为可以映射影子缓冲区。Mesa 的软件渲染器采用的正是类似的方式：所有输入缓冲区对象都位于系统内存中，着色器指令由系统 CPU 执行，最终的渲染结果被写入一个作为输出的哑缓冲区对象。

原创 OpenGL进阶23 - 光线跟踪实现详述(Ray Tracing)

本文介绍了基于OpenGL的光线追踪基本原理。与传统光栅化不同，光线追踪通过逆向处理像素射线与几何体的相交关系，实现了反射、阴影和折射等效果。文章详细讲解了射线与球体的相交计算方法：通过二次方程求解射线参数t，判断交点位置。虽然OpenGL不直接支持光线追踪，但通过自定义着色器可实现简单递归光线追踪器，适用于球体和无限平面场景。这种方法突破了传统光栅化的限制，为复杂视觉效果提供了可能。

原创 C++ MD5 算法实现原理

MD5算法是一种广泛使用的哈希函数，可生成 128位哈希值。MD5 由Ronald Rivest于 1991 年设计，用于取代早期的哈希函数MD4。MD5 算法是把任意长度的字节流 → 通过固定的非线性函数 + 位运算 → 压缩成 128 bit 状态。MD5 可作为校验和算法，用于验证数据完整性并检测意外的数据损坏。历史上，它曾被广泛用作加密哈希函数，但随后人们发现其存在多种安全漏洞。因此，MD5 已不再适用于安全敏感场景。

原创 LLVM - 编译器优化（一）

使用我们之前所展示的基本代码生成方法，你可以构建一个能够生成完全可用、能正确运行代码的编译器。然而，如果你仔细检查编译器输出的代码，就会发现其中存在许多明显的低效之处。这是因为基本的代码生成策略是孤立地看待每一个程序元素的，为了把它们连接在一起，只能采用最保守的策略。在高级语言发展的早期，由于优化技术尚未普及，编译器生成的代码通常被认为不如人工编写的代码高效。而如今，现代编译器拥有大量优化技术，并且对底层体系结构有着非常深入的了解，因此编译后的代码通常（但并非总是）优于人工编写的代码。

原创 C++ 线性探测法哈希表

线性探测是计算机编程中用于解决哈希表冲突的一种方案。哈希表是一种数据结构，用于维护键值对集合，并查找与给定键关联的值。与二次探测和双重哈希一样，线性探测也是一种开放寻址方法。在这些方案中，哈希表的每个单元格存储一个键值对。当哈希函数将新键映射到哈希表中已被其他键占用的单元格时，线性探测会在表中寻找最近的空闲位置，并将新键插入其中。查找操作也以相同的方式进行，从哈希函数给出的位置开始，按顺序搜索表，直到找到匹配的键或空单元格为止。（引自维基百科）

原创 C++ 拉链法（Chaining）哈希表

拉链哈希表是一种使用链表解决哈希冲突的方法。当多个元素映射到同一哈希桶时，它们会被存储在链表中。文中介绍了拉链法的基本原理，区分了"桶"和"槽"的概念，并提供了C++实现代码。该实现包含插入、查找和显示功能，使用简单的取模哈希函数。测试数据展示了如何处理哈希冲突，将元素1、6、11都存储在同一个桶中。这种方法的优点是实现简单，但查找效率可能随链表长度增加而降低。

原创 驱动开发系列76 - Mesa NIR 通用中间语言表示

NIR 是Mesa 的通用中间表示，适用于所有图形API和硬件后端。本文介绍下NIR的设计架构。包括NIR核心数据结构，NIR优化过程，NIR降级和验证。

原创 驱动开发系列75 - Zink OpenGL 到 Vulkan 转换层解析

Zink 实现了 Gallium3D 驱动程序接口，并将其操作转换为 Vulkan API 调用。它不针对特定的 GPU 硬件；相反，它将操作委托给底层的 Vulkan 驱动程序实现。它在Mesa源码中的位置是: src/gallium/drivers/zink。

原创 linux内核 - 基于6.19内核的进程管理和调度介绍（一）

所有现代操作系统都能够同时运行多个进程——至少在用户看来是这样。如果系统只有一个处理器，那么在任意时刻实际上只能运行一个程序。在多处理器系统中，真正能够并行运行的进程数量取决于物理 CPU 的数量。内核和处理器通过以极快的速度在不同正在运行的应用程序之间反复切换，制造出一种多任务并行执行的假象。由于切换间隔非常短，用户几乎察觉不到中间这些极其短暂的停顿，从而感觉计算机仿佛在同时处理多项任务。

原创 Ubuntu 20.04 安装 搜狗输入法

4. 重启后选择sougoupinyin。2. 安装搜狗输入法。

原创 驱动开发系列74 - GPU中的I2C

I2C（内部集成电路总线）是一种只用两根线的串行通信总线，一根传数据（SDA），一根传时钟（SCL）。主设备通过 SCL 控制数据传输，SDA 可以双向传输数据，从设备也能通过拉低 SCL 来延长时钟。总线上的信号线都是开漏的，需要上拉电阻才能变高。I2C 常用于嵌入式系统连接 EEPROM、RTC、GPIO 扩展器、传感器等外设。

原创 C++语言特性32 - 三方比较(C++20)

在 C++20 中，引入了一个新的比较运算符<=>。表达式形如a <=> b，它的返回值不是简单的bool，而是一个“比较结果对象”。（比较类别 — comparison category object），根据a与b的大小/相等关系，返回不同结果。具体来说：如果a < b，那么如果a > b，那么如果a == b（或等价 / 相等 / 等价关系成立），那么。

原创 linux内核 - 内核如何看待文件、磁盘与设备

VFS 是整个流程的入口。它负责解析来自用户空间的文件访问请求，将路径名解析（pathname resolution）为内核内部使用的对象：dentry（目录项）、inode（索引节点）、file（文件对象）。接下来发生的并不是一次单一的 I/O 操作，而是一次在内核多个子系统之间的分层协作流程，每一层都有明确的职责划分。从用户意图到最终的 I/O 操作，这条路径要经过三个关键组件：VFS（虚拟文件系统层）、块 I/O 子系统（block layer）以及具体的设备驱动（device driver）。

原创 Linux 命令行快捷键

本文介绍了Linux终端操作的三大类快捷键：光标移动、命令修改和拷贝粘贴。通过Ctrl/Alt组合键可快速跳转行首尾、单词间移动；支持字符/单词交换和大小写转换；剪切环(kill-ring)机制实现高效文本操作。这些快捷键能显著提升命令行效率，是工程师专业能力的体现，特别在无GUI环境中尤为重要。掌握这些技巧可避免频繁使用方向键和鼠标，保持操作流畅性。

原创 驱动开发系列73 - clEnqueueNDRangeKernel实现

本文介绍了OpenCL API函数clEnqueueNDRangeKernel的使用方法和内部实现机制。该函数用于将kernel提交到命令队列执行，主要参数包括命令队列、kernel对象、工作维度、工作组大小等。文章通过1D/2D/3D三种维度的示例说明了参数设置方法。在实现层面，详细阐述了驱动程序的9个关键步骤：获取执行上下文、检查设备内存使用、验证命令队列有效性、检查kernel对象、参数设置验证、工作组大小合法性检查等。其中重点说明了系数寄存器的概念及其在GPU计算中的作用，以及当寄存器不足时的溢出处

原创 linux内核 - 内核架构概览

摘要：Linux内核采用单体内核架构，严格区分用户空间（受限权限）与内核空间（完全权限），通过系统调用实现安全交互。核心子系统包括进程调度器（CFS算法）、内存管理（虚拟地址隔离）、虚拟文件系统（统一接口）、设备驱动（硬件抽象）、网络协议栈（TCP/IP支持）等。内核通过可加载模块实现扩展性，利用MMU硬件强制内存保护确保进程隔离，为容器化技术提供基础。系统调用机制（如syscall指令）是用户程序访问内核服务的唯一安全通道，这种架构设计保障了系统的安全性、稳定性和多任务处理能力。

原创 CUDA编程14 - 如何并行执行kernel

本文介绍了如何在CUDA中使用流实现kernel并发执行。通过创建多个流，可以让不同kernel同时运行，提高GPU利用率。文中详细说明了检查设备并发能力、分配内存、创建流和事件、控制执行顺序等关键步骤，并展示了如何通过cudaEventRecord和cudaStreamWaitEvent来同步流，最后进行结果验证。该方法能有效提升GPU计算性能。

原创 CUDA编程13 - 测量每个Block的执行时间

GPU 程序性能不是靠 CPU 那样的“顺序执行”来衡量的，而是靠线程块（block）和多处理器（SM）利用率。每个 block 在 GPU 的不同多处理器上执行，顺序不确定。传统的 kernel 总体计时（比如cudaEvent计时整个 kernel）只能知道总时间，无法分析哪个 block 慢，为什么慢。

原创 linux内核 - 内核是一个分层的系统

Linux内核通过分层架构实现安全可控的执行环境，包括虚拟层、映射层、隔离层和受控层。它采用任务绑定上下文、间接访问和严格验证机制，避免全局状态，确保行为可预测。内核通过抽象接口（如VFS、块层）隐藏实现细节，使用函数指针动态路由操作。每个任务携带独立上下文（内存映射、凭证等），所有访问都经过验证。这种分层设计使内核能在并发、抢占和硬件故障下保持可靠性和安全性，同时通过资源限制（Cgroup）和策略强制（LSM）实现受控访问。

原创 驱动开发系列72 - GLSL编译器实现 - 指令选择（二）

摘要：比特级常量传播优化是一种编译器优化技术，通过分析位运算中的常量部分来消除无效计算。例如对于指令vec4 tmp = inColor & ivec4(0xFF00FF00)，可确定某些位恒为0，从而优化后续运算。该优化适用于位运算、移位和算术运算场景，能有效减少程序指令数量。

原创 linux内核 - 获取单调计时时间戳的方法

摘要：Linux内核中的ktime_get()函数提供高精度(纳秒级)单调时钟，用于获取系统时间。它内部保证一致性，无需禁用抢占。示例代码展示了如何通过内核模块调用ktime_get()获取时间戳，并处理不同内核版本下的输出格式差异。该函数适用于需要精确时间测量的场景，如策略时间间隔计算等。模块包含初始化和退出函数，以及必要的元信息声明。

原创 linux内核 - 获取内核日志时间戳的方法

本文介绍了Linux内核中sched_clock高精度计时器的使用注意事项。由于sched_clock在线程抢占时可能不安全，示例代码通过preempt_disable()/preempt_enable()保护时间戳获取过程。Clockns64()函数封装了这一安全机制，返回纳秒级时间戳。随附的模块代码展示了如何正确调用该函数，并包含基本的模块加载/卸载功能。该实现适用于需要高精度时间戳的内核开发场景。

原创 linux 内核 - 内核设计原则

系统调用、陷阱、中断和内部线程在不同的约束条件下进入内核。这些差异决定了哪些函数是有效的，哪些必须预先分配，以及如何返回控制。每次进入内核的调用都与当前线程的身份相关：内存空间、文件描述符、信号状态和特权级别。它是根据必须在并发、硬件交互和故障条件下保持的规则来组织的。这些规则定义了执行流、哪些代码可以被调用、以及哪些操作是允许的。内存的分配遵循特定的策略。各个子系统有不同的要求，内核通过分层分配器和受保护的 API 做出响应。用户空间通过定义好的接口进入内核，但这些接口并不定义内核的结构。

原创 驱动开发系列71 - GLSL编译器实现 - 指令选择（一）

GLSL变量分为Uniform和Varying两类，分别对应SGPR和VGPR寄存器。GPU指令分为标量(SOP)和向量(VOP)两类。本文主要探讨GLSL中间码(IR)如何转换为GPU指令的过程，分析不同变量类型与指令类型的对应关系及其在GPU架构中的处理方式。

原创 CUDA编程12 - 使用OpenMP控制多个GPU示例

本文介绍了如何在多GPU系统上使用OpenMP并行控制多个GPU处理数组数据。示例代码演示了：1）将CPU线程与GPU绑定，每个线程控制一个GPU；2）将数组均分到多个GPU并行处理；3）每个GPU执行简单的加法核函数；4）验证结果正确性。关键步骤包括：初始化数据、OpenMP线程分配、GPU内存管理、核函数调用和结果检查。该方案充分利用了多GPU系统的并行计算能力，实现了高效的数据处理。