luoganttcc-CSDN博客

原创摩西十戒

第一条：“我是耶和华－你的神，曾将你从埃及地为奴之家领出来，除了我之外，你不可有别的神。”第二条：“不可为自己雕刻偶像，也不可做什么形象仿佛上天、下地，和地底下、水中的百物。不可跪拜那些像，也不可事奉它，因为我耶和华－你的神是忌邪的神。恨我的，我必追讨他的罪，自父及子，直到三四代；爱我、守我戒命的，我必向他们发慈爱，直到千代。”第三条：“不可妄称耶和华－你神的名；因为妄称耶和华名的，耶和...

2019-01-11 17:57:25 3441 1

本文实现了一个教学版的FlashAttention算法，通过分块计算和增量融合来优化注意力机制的内存访问效率。代码展示了如何将Q、K、V矩阵分块处理（Bc和Br为块大小），并逐步更新输出O、最大值m和归一化因子l。与标准注意力（standard_attention）相比，FlashAttention在保持数值精度的同时（误差<1e-5），显著减少了内存访问开销。该实现可根据SRAM大小（SRAM_size_M）自动调整分块策略，适用于大矩阵的高效计算。

2026-04-28 14:44:10 23

原创 ptx 代码解释

这段PTX代码实现了CUDA中计算线程坐标(row,col)的核心逻辑。它分为x和y两个方向计算，每个方向4条指令：首先将blockDim、blockIdx、threadIdx值加载到寄存器，然后通过乘加指令(mad.lo.s32)计算全局坐标。x方向计算列号(col=blockIdx.xblockDim.x+threadIdx.x)，y方向计算行号(row=blockIdx.yblockDim.y+threadIdx.y)。这种分段计算方式体现了PTX作为低级中间表示的机械特性，与CUDA的高层抽象形成对

2026-04-22 11:25:10 346

原创木棱堡策略是不是对应创业建立护城河

《木棱堡策略：商业护城河的野蛮生长法则》摘要莫斯科公国用木棱堡扩张的军事智慧，揭示了创业公司构建护城河的核心逻辑：先以MVP（木堡）快速验证需求，通过"种城战术"形成产品矩阵网络，构建初期防御体系。关键在于同步发展"防守（木墙+火枪手）"与"进攻（哥萨克骑兵）"的双重能力——用护城河守住基本盘，以增长飞轮开拓市场。最终要将临时性的木堡升级为石制堡垒，把先发优势转化为结构性壁垒。这套在恶劣环境中野蛮生长的生存哲学，正是创业者在巨头环伺下杀出血路的实

2026-04-21 15:52:54 444

原创蒙古人当年是如何控制莫斯科的

杀死第一人格：蒙古的铁骑和屠城，彻底摧毁了基辅罗斯那种开放、浪漫、具有一定民主色彩的城邦文明（第一人格），让罗斯人明白，在绝对的暴力面前，道德和尊严一文不值。植入黑暗基因：蒙古人把莫斯科大公当成猎犬来养。莫斯科为了活下去，必须比其他公国更残忍、更狡诈、更善于揣摩主子的心意。莫斯科的中央集权、特务统治（后来的特辖军）、对暴力的崇拜，无一不是从蒙古主子那里学来的“黑暗生命力”。反噬主子：最讽刺的结局是，蒙古人为了方便统治，赋予了莫斯科代征贡赋和替天行道的特权；

2026-04-21 15:48:05 317

原创为什么像 FlashAttention / GEMM 这种高性能 kernel，会“故意吃很多 register”，甚至降低 occupancy，但性能反而更高

高性能计算中，减少HBM访问比提高occupancy更重要。FlashAttention和GEMM通过增加寄存器使用，将中间数据保留在SRAM（寄存器/shared memory）中，显著减少了昂贵的HBM访问次数。虽然这会降低occupancy，但HBM访问的延迟（约400周期）远高于SRAM（约1周期），当数据需要多次使用时，SRAM缓存策略能带来10倍以上的性能提升。这种优化将计算从memory-bound转变为compute-bound，使得减少HBM访问的收益远高于增加并发线程的收益。这正是高性能

2026-04-21 14:45:53 158

原创华为昇腾（Ascend）等芯片，同样存在“寄存器 / 片上存储资源限制并发”的问题

华为昇腾（Ascend）等芯片，同样存在“寄存器 / 片上存储资源限制并发”的问题，本质等价于 CUDA 的 register + occupancy。不是 CUDA 才有这个问题，而是所有高性能并行架构的共性：片上 SRAM（register/UB）会限制并发，从而影响整体性能。👉 为什么 FlashAttention / GEMM 会“故意吃满 register”，但性能反而更高。昇腾没有叫 occupancy，但“能同时跑多少 task”就是 occupancy 的等价物。

2026-04-21 14:41:50 293

原创 cuda的 . occupancy + register

本文分析了CUDA编程中寄存器(register)占用与occupancy(占用率)的关系。寄存器作为线程私有高速工作区，其数量会直接影响SM上同时活跃的warp数量，从而决定occupancy高低。occupancy定义为SM上实际活跃warp数与理论最大warp数的比值，高occupancy有利于隐藏访存延迟。文中通过具体数值示例说明：当寄存器需求增加时，SM能驻留的warp数会减少，导致occupancy下降，进而影响性能优化。NVIDIA官方文档也指出，occupancy是衡量GPU计算资源利用率的

2026-04-21 14:32:36 497

原创 bank conflict 举例说明

本文通过具体计算解释了为什么在CUDA共享内存中，32×32的二维数组按列访问会导致bank冲突，而32×33的数组不会。对于32×32数组，相邻行的线性地址差为32，模32后bank相同，导致所有线程访问同一bank。而32×33数组的相邻行地址差为33，模32后bank编号依次递增1，使线程访问分散到不同bank，避免了冲突。关键结论是：32的倍数会导致bank对齐冲突，而33通过错位设计实现了访问分散。

2026-04-21 10:56:11 433

原创 CUDA grid/block 到矩阵映射示例（矩阵加法）

摘要：本文详细讲解了CUDA中grid/block到矩阵索引的映射关系，以8×8矩阵加法为例。配置采用4×4的block和2×2的grid，完整覆盖矩阵。核心公式展示了blockIdx和threadIdx如何组合计算矩阵坐标，并通过区域划分图直观展示映射关系。文章提供了完整的CUDA kernel代码示例，实现C[row][col]=A[row][col]+B[row][col]的并行计算，包含边界检查。最后指出这种映射关系直接影响内存访问效率、warp执行等性能关键因素，是理解GPU并行计算的基础。

2026-04-15 11:22:21 418

原创 global=block×blockSize+local 线性展开

本文系统阐述了高性能计算中的索引映射核心公式 global = block × blockSize + local 的数学本质与工程应用。该公式本质是一种多进制展开，将分块坐标（高位）与块内坐标（低位）线性组合为连续地址，其中blockSize作为权重（stride）决定访问步长。文中详细对比了该公式与进制系统、CUDA线程索引、矩阵内存布局的统一性，指出其在高维张量、编译器优化等场景的普适性。特别强调索引展开的连续性对性能的关键影响：合理权重（如stride=1）可实现内存合并访问，而错误权重会导致带宽利

2026-04-15 09:41:07 384

原创 CUDA Grid/Block 在 A100 上的任务切分与执行

本文分析了CUDA任务在A100 GPU上的切分与执行过程。首先通过示例代码计算总任务规模：24个block，每个block128线程，共3072线程，分为96个warp。重点阐述了block内warp的切分方式，在block(4,8,4)配置下，每个z平面正好对应一个warp（4×8=32线程）。文章详细展示了线性thread id的计算公式及其展开方式，说明warp沿x→y连续扩展，z维负责切分warp。最后指出A100执行时以warp为调度单位，并分析该配置的性能特点：结构清晰但总warp数偏少。理解

2026-04-14 17:16:04 320

原创单卡性能 4090 比A100 强

摘要： NVIDIA RTX 4090 与 A100 的性能对比需分场景讨论。在单卡、计算密集型任务（如中小模型推理、图像生成）中，4090凭借更多SM单元、更高频率和更强单卡算力（FLOPs）表现更优。然而，在内存/带宽敏感（如大模型推理、长序列）、多卡协同或训练场景下，A100凭借HBM显存、NVLink及高系统吞吐能力显著领先。核心差异在于：4090擅长单点爆发，但受限于显存和通信；A100则强在持续供给与多卡扩展，体现“系统吞吐＞峰值算力”的工程逻辑。

2026-04-14 15:34:40 275

原创 # NVIDIA RTX 4090 vs A100：SM / Core 对比

NVIDIA RTX 4090与A100的核心对比显示：4090采用Ada架构，拥有128个SM和16384个CUDA核心，是A100(108SM/6912核心)的2.3倍，理论算力更强。但A100凭借HBM2e显存(1.5-2TB/s带宽)、NVLink和多卡扩展能力，在大规模AI任务中表现更优。关键差异在于4090侧重算力堆叠，而A100实现系统级平衡（算力+带宽+通信），特别适合数据中心场景。最终结论：4090适合计算密集型任务，A100则在大模型训练等场景中更具优势。

2026-04-14 15:19:18 267

原创 dim3 grid_size(2, 3, 4)； dim3 block_size(4, 8, 4）算例

本文分析了CUDA中3D网格和块配置的执行模型，重点讨论了warp对齐问题。以一个典型配置dim3 block_size(4,8,4)为例，每个block包含128个线程（4×8×4），正好划分为4个完整的32线程warp。特别指出每个z层（tz）对应一个warp，这种结构完全符合warp=32的要求且没有资源浪费。文章还解释了线程线性化计算方式、warp划分逻辑，并强调虽然这种配置结构规整，但实际性能还需考虑内存访问连续性、共享内存使用等因素。最后提供了一个完整的CUDA示例代码，展示如何打印3D线程布局

2026-04-13 22:40:04 105

原创给出一个算例就是大矩阵相乘用share memery 和不用时间的差别

本文通过具体算例分析矩阵乘法的GPU实现。以1024×1024矩阵乘法为例，总计算量约10^9次乘加操作。首先介绍不使用共享内存的naive版本，每个线程独立计算一个输出元素C[i][j]。后续将对比不同优化方法的效果，帮助读者理解矩阵乘法在GPU上的实现原理和性能优化思路。

2026-04-13 22:12:49 309

原创给你一张 GPU（比如 A100）参数，让你自己算出 flash attention “最优 tile size”

本文介绍了在A100 GPU上优化FlashAttention tile size的工程方法。核心结论是：常见配置为BLOCK_M/BLOCK_N=64/128，BLOCK_K=32/64/128，总线程数128或256。这些参数由四个关键资源决定：共享内存容量、寄存器压力、warp数量和load/compute重叠能力。文章提供了A100 SM的关键资源数据：每个SM最多2048线程、64 warps、65536个32-bit寄存器和约160KB共享内存。优化过程需要平衡这些资源限制，计算每个block占

2026-04-13 21:58:43 260

原创 share memery 就是 sram 吗

文章摘要：共享内存（shared memory）本质上是GPU芯片上的SRAM（静态随机存储器），属于高速片上存储。在NVIDIA架构中，它作为SM（流式多处理器）内部的可编程存储区域，具有比HBM（高带宽内存）更快的访问速度。关键区别在于：共享内存是片上SRAM，延迟极低；而HBM是片外存储，带宽高但延迟较大。这种设计使得共享内存特别适合需要频繁数据交换的并行计算任务。

2026-04-13 21:38:15 300

原创一个 warp 同时运行 32 个thread 就是同时运行 32 core

GPU中的warp（32个线程）执行机制解析摘要：本文澄清了关于GPU中warp执行的常见误解。一个warp包含32个线程，这些线程会被映射到一组执行单元（core pipeline）并行执行，但并非固定占用32个物理核心。在理想情况下，当SM（流式多处理器）拥有足够计算单元时，32个线程可以同时执行，近似相当于32个核心的并行计算能力。然而实际执行中，warp的线程调度还受限于硬件资源分配和指令流水线等因素，不能简单等同于32个物理核心的固定占用。正确理解warp的执行机制对于优化GPU程序性能至关重

2026-04-13 21:32:35 298

原创 gridDim 最好是sm 的整数吗

GPU调度常见误区：gridDim不需要是SM数量的整数倍核心结论：误区纠正：gridDim（网格维度）不必是SM（流多处理器）数量的整数倍关键要求：gridDim应远大于SM数量才能充分利用GPU算力常见误解来源：直觉认为"1个SM对应1个block"是最优分配实际GPU采用动态调度机制，每个SM会并行处理多个block 调度机制说明：现代GPU采用warp调度器，SM会同时处理多个block 需要足够多的block（远大于SM数）才能保持所有SM持续满载典型的经验值是

2026-04-13 21:24:44 354

原创为什么财产都是留给儿子

文章指出，女性择偶时重视对方物质条件的倾向，客观上导致了重男轻女现象的延续。通过两个家庭的对比分析发现：将资产集中给儿子的家庭（甲）比平分资产的家庭（乙）更容易延续后代——甲家庭的儿子因资产多更易婚配，女儿因资产少也更容易"上嫁"；而乙家庭子女的婚配优势都较弱。这种婚恋选择机制使得重男轻女的家庭在进化中更具优势，形成了道德伦理与生存现实的矛盾。作者认为，正是女性自身对物质条件的重视，无意中维系了重男轻女的传统。

2026-04-13 10:44:34 238

原创看书看到大脑疲劳以后，应如何快速恢复大脑状态

【大脑疲劳恢复指南】当长时间阅读导致大脑疲劳时，可采取以下科学恢复方法：1）短暂闭目养神5分钟，配合4-7-8呼吸法（吸气4秒-屏息7秒-呼气8秒）；2）进行10分钟低强度运动如散步或拉伸，促进脑部供氧；3）补充含Omega-3的食物（如坚果）和水分；4）切换不同感官活动，如听轻音乐；5）采用番茄工作法（25分钟专注+5分钟休息）。建议单次用脑不超过90分钟，保持充足睡眠是预防疲劳的关键。

2026-04-12 17:42:52 380

原创详细介绍一下 IPU CGRA dataflow ASIC

本文对比分析了三种新型计算架构：IPU、CGRA和Dataflow ASIC。IPU（如Graphcore）采用图计算执行方式，适合稀疏/不规则计算；CGRA通过可重构计算阵列提供高灵活性，适合自定义算子；Dataflow ASIC则为固定高性能推理场景提供极致性能。三者各具优势：IPU=软件图执行灵活性，CGRA=硬件可编程性，Dataflow ASIC=专用高性能。选择取决于具体应用场景对灵活性和性能的需求。

2026-04-09 14:40:58 365

原创是不是低秩表达三阶 attention 的本质还是二阶

摘要：低秩分解后的三阶表达式本质上是多个二阶结构的组合，而非真正的三阶结构。关键在于它并非退化为单一二阶，而是转化为一组可叠加的二阶形式。通过公式分析，原始三阶张量$f(x,y,z)=\sum_{i,j,k}T_{ijk}x_iy_jz_k$在低秩分解后表现为二阶结构的线性组合，揭示了高阶模型的底层逻辑。这一区别对理解模型复杂度与表达能力具有重要意义。

2026-04-09 12:46:29 308

原创如果用高阶attention O (n3 ) 复杂度，需要多少算力

三阶attention的算力需求远超二阶，从O(n²)提升到O(n³)，导致计算成本呈灾难性增长。普通Transformer的attention计算基于点积S_ij=q_i⊤k_j，而三阶attention会引入更复杂的计算模式。当序列长度增加时，三阶attention的计算资源消耗将变得难以承受，使多数实际应用场景无法运行。这种指数级增长的计算复杂度是三阶attention难以推广的根本原因。

2026-04-09 11:32:04 337

原创华为的 npu 架构如何进行 flash attention

SRAM与HBM/DRAM的成本差异显著。SRAM每GB成本约为HBM的10-100倍，甚至在某些工艺下每bit成本高达DRAM的50-150倍。这是由于SRAM采用6晶体管结构，面积大（100-150 F²/bit），而DRAM/HBM采用1晶体管+电容结构，面积小（6-10 F²/bit）且需要刷新。这种物理密度差异导致SRAM成本远高于HBM，对理解FlashAttention等技术的设计考量至关重要。

2026-04-09 11:22:18 371

原创 flashattention 是适配英伟达的 gpgpu ,如果计算架构改了，flashattention也会改

FlashAttention的核心思想是优化内存IO调度策略，通过减少HBM与SRAM之间的数据传输次数来提升性能。其本质是解决"Data Movement ≫ Compute"的问题，而非特定算法。当前实现主要针对NVIDIA GPU架构特性，但其思想具有通用性。具体实现需要根据不同硬件架构进行调整，以最大化减少IO瓶颈。

2026-04-09 09:55:32 329

原创为何我读书越多，内心却越感到空虚

书里告诉你努力可能只是系统的剥削，道德只是统治阶级的工具，甚至连你引以为傲的自我意识，不过是潜意识和环境的随机产物。你挖出了自己的一只眼睛扔进泉水，看到了世界的荒诞和虚无。他失去了一只眼睛换来了全知，然后彻底陷入了无法逆转的虚无和焦虑，因为他知道一切挣扎都是徒劳。你在书本里获取了上帝视角的快感，关上书，你还得去挤早高峰的地铁，还得面对老板的弱智绩效。但你忘了，神话里的奥丁在看到必然的毁灭后，没有原地等死，他还是拿起了永恒之枪。这世界上最残忍的事，就是你具备了看穿一切的眼力，却没具备改变这一切的能力。

2026-04-08 11:14:39 363

空空如也

空空如也