当前位置： 首页 > article >正文

triton原子操作研究

article 2026/4/6 7:01:50
背景使用Triton实现一个向量累加triton.jitdefreduction_kernel(input,output,N:int,BLOCK_SIZE:tl.constexpr,num_warps:tl.constexpr,):pidtl.program_id(0)idxtl.arange(0,BLOCK_SIZE)offsetBLOCK_SIZE*pididx maskoffsetN atl.load(inputoffset,maskmask,other0.0)tl.atomic_add(output,tl.sum(a))它的原子加法竟然直接使用不需要进行tid0的判断它是怎么做到一个block只加一次的参考对应的C版本应该是这样的templatesize_t BlockSize__global__voidreduce(constfloat*input,float*output,intN){intidxblockIdx.x*blockDim.xthreadIdx.x;inttidthreadIdx.x;intwarpIdtid/WARP_SIZE;intlaneIdtid%WARP_SIZE;__shared__floats_sum[BlockSize/WARP_SIZE];floatsum(idxN)?0:input[idx];sumwarp_reduce_sum(sum);if(laneId0){s_sum[warpId]sum;}__syncthreads();sums_sum[laneId];sumwarp_reduce_sumBlockSize/WARP_SIZE(sum);if(threadIdx.x0){atomicAdd(output,sum);}}SASS代码分析基本代码如下IMAD.MOV.U32 R1,RZ,RZ,c[0x0][0x28]// R9 threadIdx.xS2R R9,SR_TID.X ULDC.64UR4,c[0x0][0x118]BSSY B0,0x40f5ff1e0MOV R4,RZ// R3 blockIdx.xS2R R3,SR_CTAID.X// R0 R3 0x7f(127)LOP3.LUT R0,R9,0x7f,RZ,0xc0,!PT IMAD R3,R3,0x80,R0 ISETP.GE.AND P0,PT,R3,c[0x0][0x170],PT LEA R2,P1,R3,c[0x0][0x160],0x2LEA.HI.X.SX32 R3,R3,c[0x0][0x164],0x2,P1 P0 BRA0x40f5ff1d0LDG.E R4,[R2.64]BSYNC B0// shfl_xor_syncSHFL.BFLY PT,R3,R4,0x10,0x1f// P1 R9 0x1fLOP3.LUT P1,RZ,R9,0x1f,RZ,0xc0,!PT ISETP.GE.U32.AND P0,PT,R0,0x4,PT FADD R3,R3,R4 SHF.R.U32.HI R4,RZ,0x3,R9 SHFL.BFLY PT,R2,R3,0x8,0x1fLOP3.LUT R4,R4,0xc,RZ,0xc0,!PT FADD R2,R3,R2 SHFL.BFLY PT,R5,R2,0x4,0x1fFADD R5,R2,R5 SHFL.BFLY PT,R6,R5,0x2,0x1fFADD R6,R5,R6 SHFL.BFLY PT,R7,R6,0x1,0x1fFADD R7,R6,R7// P1 false才执行也就是 R9 0x1f 0即laneId 0!P1 STS[R4],R7 BAR.SYNC.DEFER_BLOCKING0x0// P1 P1 R0 0x20ISETP.LT.U32.AND P1,PT,R0,0x20,!P1 !P0 LDS R8,[R0.X4]LOP3.LUT P0,RZ,R9,0x3,RZ,0xc0,!PT ISETP.LT.U32.AND P0,PT,R0,0x4,!P0 SHFL.BFLY PT,R3,R8,0x2,0x1fFADD R3,R8,R3 SHFL.BFLY PT,R2,R3,0x1,0x1fFADD R5,R3,R2 P0 STS[R0.X4],R5 BAR.SYNC.DEFER_BLOCKING0x0// P1 为false 时直接退出!P1 EXIT LDS R5,[RZ]IMAD.MOV.U32 R2,RZ,RZ,c[0x0][0x168]MOV R3,c[0x0][0x16c]MEMBAR.ALL.GPU ERRBAR ATOMG.E.ADD.F32.FTZ.RN.STRONG.GPU PT,R2,[R2.64],R5 CCTL.IVALL从代码中可以看到它执行了7次蝶式交换也就是异或交换前5次是warp内部的累加然后后面2次是4128/324128/324128/32然后ISETP.LT.U32.AND P1,PT,R0,0x20,!P1 !P1 EXIT两个条件一起成立才是P1laneId0tid32不满足这个条件的直接退出在PTX中更加明确setp.eq.b32%p2,%r14,0;%p4 st.shared.b32[%r60],%r7;bar.sync0;$L__tmp16:setp.lt.u32%p7,%r13,32;and.pred%p5,%p2,%p7;ld.shared.b32%r9,[global_smem];mov.u32%r8,0x0;%p5 atom.global.gpu.acq_rel.add.f32%r8,[%rd20],%r9;最后一个原子加法是有谓词的所以triton确实不涉及向量的原子加法确实一个block只会执行一次并且合理推测只要不涉及向量化操作的都只会执行一次

相关文章：

triton原子操作研究

背景 使用Triton实现一个向量累加 triton.jit def reduction_kernel(input,output,N: int,BLOCK_SIZE: tl.constexpr,num_warps: tl.constexpr, ):pid tl.program_id(0)idx tl.arange(0, BLOCK_SIZE)offset BLOCK_SIZE * pid idxmask offset < Na tl.load(input offs…...

编程日记 2026/4/6 7:01:50

SEO_为什么你的SEO策略无效？常见原因与解决办法（372 ）

SEO策略无效的常见原因 在当今数字化时代&#xff0c;搜索引擎优化&#xff08;SEO&#xff09;是网站流量和业务增长的关键。不少企业在实施SEO策略后&#xff0c;却发现效果并不理想。为什么你的SEO策略无效&#xff1f;我们将从多个角度分析常见原因&#xff0c;并给出相应…...

编程日记 2026/4/6 6:59:49

Kandinsky-5.0-I2V-Lite-5s应用场景：游戏NPC立绘动态化+过场动画快速生成

Kandinsky-5.0-I2V-Lite-5s应用场景&#xff1a;游戏NPC立绘动态化过场动画快速生成 1. 游戏开发中的视觉挑战 在游戏开发过程中&#xff0c;NPC立绘动态化和过场动画制作一直是两个耗时费力的环节。传统方法需要美术团队逐帧绘制动画&#xff0c;或者使用复杂的3D建模工具&a…...

编程日记 2026/4/6 6:59:49

【算法精解】CEC2021竞赛亚军算法-MadDE框架及代码实现（Matlab）

本文核心内容&#xff1a;  MadDE算法主要框架及该算法创新点  Matlab代码实现&#xff08;可免费获取&#xff0c;包括代码及原文献&#xff09; 不少同学改进算法有时缺乏可落地思路&#xff0c;或从文献获得灵感却苦于写不出代码。为此&#xff0c;KAU 推出【算法精解】…...

编程日记 2026/4/6 6:59:49

Sony FCB-EV9500L LVDS图像闪烁问题分析

在基于高清图像采集与远距离传输的系统中&#xff0c;Sony FCB-EV9500L作为一款高性能一体化机芯模组&#xff0c;被广泛应用于安防监控、工业视觉及医疗设备等领域。在实际应用过程中&#xff0c;部分工程师反馈其在LVDS传输链路中出现图像闪烁问题。本文将围绕LVDS信号特性、…...

编程日记 2026/4/6 6:59:49

前后端分离网站系统｜SpringBoot+Vue+MyBatis+MySQL完整源码+部署教程

&#x1f4a1;实话实说&#xff1a;有自己的项目库存&#xff0c;不需要找别人拿货再加价&#xff0c;所以能给到超低价格。摘要 随着互联网技术的快速发展&#xff0c;传统的前后端耦合架构在开发效率和维护成本上逐渐显现出局限性&#xff0c;前后端分离架构因其灵活性、可扩…...

编程日记 2026/4/6 6:59:48

TensorRT-LLM与Triton Server部署实战：从环境配置到模型推理

1. 环境准备&#xff1a;从零搭建TensorRT-LLM与Triton Server基础环境 第一次接触TensorRT-LLM和Triton Server时&#xff0c;我花了整整三天时间在环境配置上踩坑。现在回想起来&#xff0c;大部分问题都源于对NVIDIA生态工具链的不熟悉。下面我会用最直白的语言&#xff0c;…...

编程日记 2026/4/6 6:57:48

保姆级教程：用ACE-Step一键生成多语言音乐，视频配乐不求人

保姆级教程&#xff1a;用ACE-Step一键生成多语言音乐&#xff0c;视频配乐不求人 你是不是也遇到过这样的烦恼&#xff1f;精心剪辑的视频&#xff0c;万事俱备&#xff0c;就差一段能完美烘托氛围的背景音乐。翻遍免费音乐库&#xff0c;要么风格不搭&#xff0c;要么听腻了…...

编程日记 2026/4/6 6:57:48

Guohua Diffusion国风绘画工具：5分钟快速部署，小白也能画水墨神兽

Guohua Diffusion国风绘画工具&#xff1a;5分钟快速部署&#xff0c;小白也能画水墨神兽 1. 工具简介&#xff1a;专为国风绘画而生的AI神器 Guohua Diffusion是一款专注于国风水墨画生成的本地AI绘画工具&#xff0c;基于原生Guohua-Diffusion模型开发。它最大的特点就是&q…...

编程日记 2026/4/6 6:57:48

DeOldify多用户并发测试：100+请求下服务稳定性与响应延迟实测

DeOldify多用户并发测试&#xff1a;100请求下服务稳定性与响应延迟实测 1. 引言&#xff1a;当AI上色服务遇到真实流量考验 想象一下&#xff0c;你搭建了一个很酷的AI图片上色服务&#xff0c;平时自己用着挺顺&#xff0c;处理一张老照片也就几秒钟。但突然有一天&#xf…...

编程日记 2026/4/6 6:57:48

OpenClaw电商运营助手：Qwen2.5-VL-7B批量生成商品图文详情

OpenClaw电商运营助手&#xff1a;Qwen2.5-VL-7B批量生成商品图文详情 1. 为什么需要自动化商品详情生成 每次大促前&#xff0c;运营团队最头疼的就是商品详情页的批量更新。去年双十一前&#xff0c;我手动处理了200多个SKU的图文优化&#xff0c;连续加班一周后&#xff0…...

编程日记 2026/4/6 6:57:47

避坑指南：ZYNQ lwIP Socket TCP服务器开发中，DHCP超时、内存泄漏和任务卡死的调试经验

ZYNQ lwIP TCP服务器开发实战&#xff1a;从实验室到工业环境的稳定性优化 在嵌入式网络开发中&#xff0c;ZYNQ平台结合lwIP协议栈的TCP服务器实现看似简单&#xff0c;但当代码从实验室走向真实工业环境时&#xff0c;开发者往往会遭遇一系列"幽灵问题"——DHCP获取…...

编程日记 2026/4/6 6:55:46

StructBERT文本相似度-中文-通用模型效果展示：电商商品描述语义聚类案例

StructBERT文本相似度-中文-通用模型效果展示&#xff1a;电商商品描述语义聚类案例 1. 项目概述 StructBERT中文文本相似度模型是一个基于百度深度学习技术的高精度语义理解工具&#xff0c;专门用于计算中文句子之间的语义相似度。这个模型能够理解中文语言的深层语义&…...

编程日记 2026/4/6 6:55:46

编写程序实现智能乐器音准检测偏差时，提示“需要调音”，新手也能调好音。

1. 实际应用场景描述场景&#xff1a;一名吉他初学者刚刚买回一把新吉他&#xff0c;或者在干燥天气后琴弦音准发生了偏移。他不知道电子调音表如何使用&#xff0c;也不具备绝对音感。本系统功能&#xff1a;用户拨动琴弦&#xff08;例如第 6 弦 E2&#xff09;&#xff0c;电…...

编程日记 2026/4/6 6:55:46

手机生成剧本杀软件2025推荐，创新剧情设计工具助力创作

手机生成剧本杀软件2025推荐&#xff0c;创新剧情设计工具助力创作随着剧本杀市场的蓬勃发展&#xff0c;越来越多的创作者和爱好者希望借助科技的力量来提升创作效率和质量。在2025年&#xff0c;一款名为量子探险AI剧本杀工坊的手机生成剧本杀软件脱颖而出&#xff0c;成为众…...

编程日记 2026/4/6 6:55:46

【从0开始学设计模式-6| 原型模式】

一个月没更新了&#xff0c;在找实习。。 其实还是懒了&#xff0c;其实每天花个半小时左右就能写一篇博客的。。。概念 原型模式(Prototype Pattern) 设计出来的目标就是&#xff1a;通过本体复制出与本体一样的分身&#xff08;分身具有本体一样特性&#xff09;定义&#xf…...

编程日记 2026/4/6 6:55:45

基于springboot+vue电子商务网站用户行为分析hx0901

文章目录详细视频演示技术介绍功能介绍核心代码系统效果图源码获取详细视频演示 文章底部名片&#xff0c;获取项目的完整演示视频&#xff0c;免费解答技术疑问 技术介绍 开发语言&#xff1a;Java 框架&#xff1a;ssm JDK版本&#xff1a;JDK1.8 服务器&#xff1a;tomca…...

编程日记 2026/4/6 6:53:45

OpenClaw定时任务管理：千问3.5-35B-A3B-FP8实现早间资讯自动推送

OpenClaw定时任务管理&#xff1a;千问3.5-35B-A3B-FP8实现早间资讯自动推送 1. 为什么需要自动化资讯推送 每天早上打开电脑第一件事&#xff0c;就是查看行业动态和技术新闻。但手动检索各大平台、整理关键信息要耗费20多分钟&#xff0c;经常打乱晨间工作节奏。直到发现Op…...

编程日记 2026/4/6 6:53:45

OpenClaw命令行增强：gemma-3-12b-it解释复杂指令并自动补全

OpenClaw命令行增强&#xff1a;gemma-3-12b-it解释复杂指令并自动补全 1. 为什么需要命令行增强工具 作为一个常年与终端打交道的开发者&#xff0c;我经常遇到这样的困境&#xff1a;记得某个命令的功能&#xff0c;却想不起具体参数&#xff1b;或者面对复杂的管道操作时&…...

编程日记 2026/4/6 6:53:45

智谱开源手机AI框架实测：一句话让Open-AutoGLM帮你搜索、购物、发微信

智谱开源手机AI框架实测&#xff1a;一句话让Open-AutoGLM帮你搜索、购物、发微信 1. 什么是Open-AutoGLM&#xff1f; Open-AutoGLM是智谱AI开源的手机端智能助理框架&#xff0c;它能像真人一样操作你的手机。想象一下&#xff0c;你只需要说"帮我订个外卖"&…...

编程日记 2026/4/6 6:53:45

Hudi 生产问题排障-乱序Upsert入湖数据丢失

一、背景与问题在大数据流式处理领域&#xff0c;乱序一直是一个无法越过的问题&#xff0c;如何正确处理乱序数据也是流式组件不断努力优化的方向&#xff0c;比如FLink提供的watermark机制&#xff08;forBoundedOutOfOrderness/allowedLateness/sideOutputLateData&#xff…...

编程日记 2026/4/6 6:53:44

深入解析Xilinx PCIe IP核示例工程的仿真与调试技巧

1. Xilinx PCIe IP核示例工程快速入门 第一次接触Xilinx PCIe IP核时&#xff0c;我完全被复杂的文件结构和专业术语搞懵了。后来发现&#xff0c;只要掌握几个关键点&#xff0c;就能快速上手这个强大的高速串行通信接口。PCIe&#xff08;Peripheral Component Interconnect …...

编程日记 2026/4/6 6:51:44

Kandinsky-5.0-I2V-Lite-5s多风格测试：卡通、写实、水墨画生成效果对比

Kandinsky-5.0-I2V-Lite-5s多风格测试&#xff1a;卡通、写实、水墨画生成效果对比 1. 开场&#xff1a;当静态艺术遇见动态魔法 想象一下&#xff0c;你珍藏的卡通插画突然活了过来&#xff0c;水墨画中的山水开始流动&#xff0c;写实照片里的场景有了生命。这正是Kandinsk…...

编程日记 2026/4/6 6:51:44

港大新作GS-SDF开源了！手把手教你用激光雷达+3DGS复现IROS2025论文效果（附避坑指南）

港大GS-SDF开源项目实战&#xff1a;从环境配置到效果复现全指南 当激光雷达遇上3D高斯溅射&#xff0c;会碰撞出怎样的火花&#xff1f;港大MARS实验室最新开源的GS-SDF项目给出了令人惊艳的答案。这个将LiDAR点云与神经符号距离场&#xff08;SDF&#xff09;相结合的创新方…...

编程日记 2026/4/6 6:51:44

从read()到硬盘：用strace和bpftrace动态追踪Linux内核文件读取的完整路径（附实战脚本）

从read()到硬盘&#xff1a;用strace和bpftrace动态追踪Linux内核文件读取的完整路径&#xff08;附实战脚本&#xff09; 当线上服务出现文件读取延迟时&#xff0c;大多数系统工程师的第一反应是检查磁盘I/O指标。但真正的挑战在于&#xff1a;如何准确定位从用户态系统调用到…...

编程日记 2026/4/6 6:51:42

5分钟部署Fun-ASR语音识别：支持中文、英文、日文等31种语言

5分钟部署Fun-ASR语音识别&#xff1a;支持中文、英文、日文等31种语言 1. 快速入门指南 1.1 学习目标 本文将带您快速完成Fun-ASR-MLT-Nano-2512多语言语音识别模型的部署与使用。通过本教程&#xff0c;您将掌握&#xff1a; 一键式Docker部署方法Web界面基本操作流程Pyt…...

编程日记 2026/4/6 6:51:37

PyG实战：用自定义MessagePassing为异构图构建一个简单的推荐系统消息传递层

PyG实战&#xff1a;构建异构图的推荐系统消息传递层 当我们在电商平台上浏览商品时&#xff0c;系统总能精准推荐我们可能感兴趣的内容。这背后往往隐藏着一个复杂的用户-商品交互网络&#xff0c;而图神经网络(GNN)正是处理这类异构关系的利器。今天&#xff0c;我们就来探索…...

编程日记 2026/4/6 6:49:35

YOLO26功能体验：官方镜像预置多种权重，开箱即用体验最新模型

YOLO26功能体验&#xff1a;官方镜像预置多种权重&#xff0c;开箱即用体验最新模型 1. 引言&#xff1a;告别环境配置&#xff0c;直接上手YOLO26 如果你对计算机视觉感兴趣&#xff0c;想试试最新的目标检测模型&#xff0c;那么YOLO26绝对值得关注。作为YOLO系列的最新成员…...

编程日记 2026/4/6 6:49:34

从零到一：手把手教你用cam_lidar_calibration标定自己的VLP-16与海康相机（附完整ROS Bag录制技巧）

从零到一&#xff1a;VLP-16激光雷达与海康相机联合标定实战指南 当激光雷达点云与相机图像在自动驾驶系统中完美对齐时&#xff0c;传感器融合的魔法才真正开始。作为机器人感知的核心环节&#xff0c;标定质量直接决定了后续目标检测、SLAM等模块的精度上限。本文将手把手带您…...

编程日记 2026/4/6 6:49:33

手把手教你用C语言解决Modbus TCP从站多主站连接的3个典型问题（含select使用避坑）

深度解析Modbus TCP从站多主站连接的三大实战难题与优化方案 在工业自动化领域&#xff0c;Modbus TCP协议因其简单可靠的特点被广泛应用于设备间通信。但当从站需要同时处理多个主站&#xff08;如SCADA系统、HMI人机界面和测试工具&#xff09;的连接请求时&#xff0c;开发者…...

编程日记 2026/4/6 6:49:33