GPU全栈博主

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

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

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

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

原创 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系统的并行计算能力，实现了高效的数据处理。

原创 Vulkan进阶系列11 - RenderPass 设置对渲染性能的影响

摘要：Vulkan中的RenderPass定义了渲染目标及其操作规则，包括颜色缓冲区和深度缓冲区的处理方式。文章以render_passes为例，分析不同操作对性能的影响：颜色缓冲区的LoadOp选项（Load、Clear、Don't care）影响显存读取带宽；深度缓冲区的StoreOp选项（Store、Don't care）影响写入带宽。此外，对比了LoadOp=Clear与手动vkCmdClearAttachments的效率差异，前者通常更高效。通过合理配置这些参数可以优化渲染性能。

原创 驱动开发系列70 - vkQueueSubmit实现

vkQueueSubmit 把一组命令送到指定的GPU队列中执行，它是 CPU -> GPU 的发车按钮。// 必填: VK_STRUCTURE_TYPE_SUBMIT_INFO// 例如挂 VkTimelineSemaphoreSubmitInfo// 提交前要等待的信号量数量// 等待哪些信号量// 对应每个 wait 的目标阶段掩码(长度=waitSemaphoreCount)// 要执行的主命令缓冲个数// 命令缓冲数组（必须已 end）

原创 linux内核 - 用journalctl查看系统日志(journal)

你问的 journal，一般是指提供的日志工具，主要通过命令 journalctl来查看和管理日志。

原创 linux内核 - 服务进程是内核的主要责任

用户调用 exec.内核必须通过虚拟文件系统解析二进制路径，使用底层文件系统驱动程序加载文件，分配和映射内存，通过安全模块验证访问，并计划进程执行。它仅在三种情况下进入执行状态：来自用户空间的系统调用、来自硬件设备的中断或计划执行系统工作的内部线程。内核的存在是为了服务用户进程。它是一个被动的、合作的系统。下面是四个图表：高级概述、系统调用路径、中断路径和内核线程路径，共同显示了内核如何仅在需要时运行，并且始终为该进程服务。内核之所以重要，不是因为它执行许多工作，而是它执行这些工作是为了服务于其他东西。

原创 驱动开发系列69 - GLSL编译器实现 - 符号表的定义

在编译器中，无论是C、C++ 还是 GLSL，都需要管理源代码里的标识符，比如（变量，函数，结构体，内置函数等），这个管理就是通过：符号表 来完成的。符号表有以下几个作用：1. 存储符号信息，每个符号都有对应的信息，拿变量来说，有名字，类型，作用域（局部/全局），数据地址，大小，引用次数，自定义析构函数等。2. 在 C、C++ 以及 GLSL 等语言中，作用域是分层次的，包括全局作用域、函数作用域以及代码块作用域等。符号表负责记录当前符号的作用域层级以及父子关系，从而实现快速、准确的符号查找和管理。

原创 vscode 中那些好用的C++插件

原创 linux内核 - 内核不是一个进程，而是一个系统。

内核代码在三个主要上下文中执行：（1） 通过用户进程发起的系统调用，（2） 通过硬件触发的中断处理程序，以及 （3） 在完全在内核空间中运行的长期内核线程中执行。引导加载程序加载后，内核开始在 start_kernel（） 中执行，在那里它初始化内存管理、设备接口和核心子系统。相反，它成为一个响应式执行层，仅在需要时调用——由用户进程、硬件事件或它自己的内部线程调用。相反，它始终存在（在启动时加载到内存中），并控制硬件和软件之间的所有交互。简而言之，内核不是系统内的一个进程，而是系统的核心。

原创 linux内核 - 文件系统相关的几个概念介绍

Ext2是Linux核心文件系统，设计优化了与Linux的互操作性，也可用于其他系统。其后续版本Ext3在保持兼容性的基础上增加了日志功能，提升了系统崩溃后的恢复能力。两者基本原理相同，但Ext3提供了更多实用选项。本章将重点介绍Ext3的日志机制及其优势。

原创 CUDA编程11 - CUDA异步执行介绍

CUDA Stream 是一条GPU命令队列，同一条GPU命令队列中的任务按顺序执行，比如下面这个调用会严格按顺序执行。不同Stream之间（stream0， stream1，...）的任务，可能并行执行（取决于GPU硬件）。CUDA Event 是插在 Stream 里的“标记点”， 主要用途是“计时”， “查询进度”，“同步” 等。// 一些 GPU 操作// GPU 还没执行到 stopcounter++;

原创 linux 内核 - 常见的文件系统介绍

目前 Linux 最常用，支持大文件、大分区、日志、延迟分配，性能和稳定性都很好。SUSE/openSUSE、Fedora、Ubuntu 都在推广。由 SGI 开发，擅长大文件和高并发场景（日志文件、数据库）。曾经以小文件性能著称，现在基本被 Btrfs/ext4 取代。针对 Flash/SSD 优化，移动设备和嵌入式系统用得多。支持快照、子卷、压缩、数据校验、RAID 等高级功能。IBM 开发，日志文件系统，资源占用小但使用较少。：早期标准文件系统，不支持日志。：增加了日志功能，提升可靠性。

原创 CUDA编程10 - CUDA Samples 在ubuntu 24.04 上编译安装, 解决Unsupported gpu architecture ‘compute_xxx‘的问题

本文介绍了cuda-samples项目的编译安装流程：1)克隆项目仓库；2)安装CMake、FreeImage等依赖库；3)查询GPU架构并修改多处CMakeLists.txt文件以适配当前硬件；4)创建build目录并使用CMake配置构建选项；5)使用make命令进行多线程编译；6)可选执行安装。重点说明了需要修改多处CMakeLists.txt文件来避免GPU架构不兼容问题，并给出了详细的文件列表。

原创 渲染引擎实践 - 如何拉取Unreal Engine 完整代码（所有分支和标签）

本文介绍了下载和编译Unreal Engine源码的完整步骤：1）注册/登录Epic Games和GitHub账号并绑定；2）访问Unreal私有仓库；3）提供三种源码下载方式（ZIP下载、Git克隆、Fork方式）；4）最后以Linux为例说明编译流程（运行Setup.sh、GenerateProjectFiles.sh和make命令）。整个过程涵盖了从账号准备到最终编译的完整指南。

原创 驱动开发系列68 - GLSL编译器实现 - 算数指令折叠及访存优化

指令优化技术分析与实现 本文介绍了两种指令优化技术：指令折叠和访存优化。指令折叠通过合并乘法指令(MUL)和乘加指令(MAD)，减少指令数量和寄存器占用，提升执行效率。文中给出了具体示例和C++实现代码，展示了如何将3条指令优化为2条。访存优化部分针对GPU内存访问性能问题，提出统一按字节计算偏移、分割大访问为多条burst等技术，特别处理了vec3/vec4等非标准长度向量类型，提供了C++实现方案。这些优化技术可有效减少指令数量、降低寄存器压力、提升内存访问效率。

原创 驱动开发系列67 - NVIDIA 开源GPU驱动open-gpu-kernel-modules分析-驱动初始化

NVIDIA驱动初始化流程涉及两个关键函数：nvidia_init_module和nvkms_init。nvidia_init_module是驱动加载总入口，负责初始化内存调试、procfs节点、GPU设备探测、PCI驱动加载、字符设备注册等核心功能，采用分层错误处理机制确保资源正确释放。nvkms_init专门初始化Kernel Mode Setting子系统，包括资源管理器分配、内核线程队列创建、定时器初始化、字符设备注册和配置文件读取，同样配备完善的错误处理流程。两个函数都包含详细的日志输出和错误检查

原创 linux 环境下 - 切换显卡（intel，nvidia）

使用nvidia-prime切换显卡的方法：首先克隆官方GitHub仓库，进入目录后执行安装命令，最后通过nvidia-prime-select指令选择NVIDIA或Intel显卡。该方案提供官方支持，适合需要稳定切换显卡的用户。

原创 linux环境下 - 如何干净地卸载掉nvidia驱动

本文介绍了彻底卸载NVIDIA驱动的详细步骤：1）对于apt安装的驱动，使用dpkg -l|grep nvidia查看后，执行sudo apt purge 'nvidia-*'卸载；2）对于.run安装包，运行sudo nvidia-uninstall卸载；3）清理步骤包括更新initramfs和grub，重启系统；4）删除残留内核模块，退出图形环境后执行modprobe -r移除模块；5）可选清理用户态驱动库；6）最后更新模块依赖并重启验证。通过完整执行这些步骤，可确保系统完全清除NVIDIA驱动及其相关

原创 linux内核 - I/O 端口和I/O内存

计算机系统中，RAM作为核心外设由内核管理，通过共享总线与其他内存映射设备通信（内存映射I/O）。32位系统中，CPU地址空间部分保留给外设（I/O内存），造成RAM访问空洞。解决方案包括内存映射I/O（统一寻址但地址受限）和端口I/O（专用指令访问，x86架构采用）。内核提供request_region()等函数管理I/O端口，支持in/out指令操作。两种方式各有优劣：内存映射简化编程但限制地址空间，端口I/O保留完整RAM访问但需专用指令。

原创 linux内核 - 内存映射介绍

内核内存有时需要进行 重映射（remap），可能是从内核空间映射到用户空间，也可能是从高端内存映射到低端内存（内核空间内部的映射）。在一个 4 GB 的系统上，内核剩下的 128 MB 地址空间 用于映射剩余的 3.2 GB 高端内存（high memory）。而对于高端内存（超过低端 896 MB 的部分），内核需要将请求的高端内存区域 临时映射到内核地址空间，前面提到的 128 MB 地址空间就是专门为此保留的。内核会将对该映射内存区域的访问，通过常规的指针解引用，转换为对应的文件操作。

原创 linux内核 - vmalloc 介绍

本文介绍了Linux内核中的vmalloc()分配器，它分配虚拟地址连续但物理内存不连续的页帧，适用于分配大块内存（如网络缓冲区）。相比kmalloc()，vmalloc()更慢但能处理更大的内存请求。文章还提供了vmalloc()的基本用法示例模块代码，展示如何分配、初始化和释放vmalloc内存，并输出调试信息。该分配器返回的地址仅适用于内核软件使用，不能用于DMA操作。

原创 linux内核 - slab 分配器

Slab分配器是Linux内核中用于管理小块内存分配的核心机制，主要解决伙伴系统分配小内存时产生的碎片问题。它通过将内存划分为特定大小的Slab（由连续页框组成）来管理对象，每个Slab分为相同大小的块。Cache则是由多个Slab组成的链表，专门存放同一类型的内核对象。Slab有三种状态：空（所有块空闲）、部分使用（含空闲和已使用块）和已满（所有块被占用）。这种设计显著提升了内存分配效率，同时减少了内存碎片。

原创 linux内核 - 内存分配接口介绍

接口概述。

原创 linux内核 - 内存分配机制介绍

Linux内核采用分层内存分配机制：页分配器以页为单位提供基础内存；其上的slab分配器将内存页分割为更小单元，支持kmalloc API；此外内核还提供直接分配虚拟内存的vmalloc接口。不同分配方式满足各类内存需求，开发者可根据用途选择最合适的机制。