每一个进入投资交易市场的交易员，其成长经历必然是不尽相同的。当一个投资人初次推开投资交易市场的大门时，因为对市场缺乏足够的认识，必然会在市场存在的众多流派中选择一个。

有被巴菲特的价值投资所吸引的，有受到行业大咖影响首先研究基本面的，有受到某些日内交易员影响研究日内走势规律的，还有被资深投资人影响研究技术分析或者量化的……

那么，切入角度的不同会影响交易生涯的走向吗？其中到底有没有一个大致的成长框架呢？

归纳总结出以下几个阶段：

第一阶段：无知

投资者进行市场的理由五花八门，但是究其根本，都是受到了利益的诱惑。交易高手们入市的理由也是各不相同。

比如，迈克尔.马科斯。一个在10年时间内使其公司的账户的净值增长了2500倍的期货交易员，他就是因为听说某位交易员朋友可以一周让资金翻两倍，才对投机市场产生了浓厚的兴趣。

保罗.都德.琼斯，一个创造了在3年半的时间内仅亏损一个月记录的私募基金经理。是因为看到了某位投资大师的文章后，震撼不已，认为投资交易行业是太阳下最光辉的职业，毅然决然的踏入了交易行业。

初次之外，还有的是因为穷苦潦倒渴望财务自由，还有的是受到了长辈经历的影响，还有的是向往资本市场的“华丽光鲜”，渴望改良现状…

但这个阶段，由于刚进入市场，而且带着强烈的“自信”。基本都将以亏损而结束。

无知者无畏，是这个阶段的典型特点。

 

 

在第一阶段对交易员产生了强大的伤害后，除了少量崩溃放弃的交易员之外，大多数投资者将进入了第二阶段。

他们开始重视市场，开始明白如果想要赚钱，需要研究熟悉的东西非常多。

基本面分析，技术分析，资金管理，止盈止损，合约规则，市场信息，他们开始大量的学习，大量的研究，疯狂的吸取知识和信息。

在国外众多被采访的交易员中，很多人曾经谈起过自己的积累阶段。有疯狂的读书的，有偏执的迷恋过技术指标的，有一天超过14个小时都在研究盘面的。

积累阶段是一个做加法的过程。慢慢的学习，慢慢的摸索，慢慢的堆积经历，在这个过程中，将伴随着大量的挫折。

比如，无畏交易大师拉蒙，巴托斯。其在1975年-1980年中经历了无数次的失败后，才成功的摸索到了一套适合他的交易方法。

正如埃德，塞克塔，一个在16年内将净值增加了2500倍的天才交易员所说：求胜意志强烈的人，一定会寻求各种方法来满足其求胜的欲望。

第三阶段：突破

量变引起质变，厚积而薄发。当一个投资人积累的知识信息足够多，交易经历中的方方面面会不停的产生共鸣与冲突，他会不停的去掉糟粕，留下精华。

正如2006年罗宾斯世界杯期货交易大赛的冠军，同时也是2005年和2007年的亚军的凯文.戴维在一次采访中说：在近20年的交易生涯中，我学到了很多的知识。包括技术分析，基本分析，战略规划，资金管理以及交易心理学。然而，在读过了大量的交易丛书，参加了无数的研讨会，花了无数夜晚研究交易思想之后，我才终于顿悟到，如无必要，勿增实体。在交易中，简单才是最好。

从此，他的策略越来越简单。最新的某一套策略甚至仅以最近两天的收盘价作为交易依据。

这个阶段是一个顿悟的过程，由横向的知识收集变为纵向的深度思考。

 

第四阶段：系统

这个阶段，交易员进入了系统的构建阶段。他们开始尝试将复杂多变的信息依据融合成为一个系统化的入场方式。开始尝试将所有对于出场的理解融合成为一套完整的出场规则，他们开始进行系统的组装。

他们关注的重点不在是交易走势的研判，而是对具体的走势的处理过程。

李.哥特斯，世界级的机械交易系统设计大师和交易大师，他在职业生涯的关键时期开发了volpat策略。

机构间交易的重量级人物汤姆德马克，他创造了汤姆德马克TD指标，并形成了系统化交易。

传奇交易员理查德，丹尼斯构建了各种交易系统，最经典的海龟交易法则广为流传。

第五阶段：认知突破

这个阶段，交易员基本都明白了唯有适合自己，适合交易的交易系统化体系才是交易成功的关键。他们会发现，无穷种入场方式都不过是在寻求一个试错依据，他们会发现，无穷种出场都是为了截断亏损，让利润奔跑。他们会发现，无穷种资金管理方法都是为了控制风险。

这个阶段，能否持续的运行一套系统，并坚持下去，考验的是绝对的交易认知。包括，对取舍的认知，对系统的认知，对盈亏本源的认知。一个交易员交易认知中的漏洞，一定会在这个阶段出现，并且很有可能会对他的资金和心理造成重创。

能否辨别出运气成分和逻辑漏洞，并且作出应有的恰当的评估，不停的丰富并且完善自己的交易认知，是这个阶段最重要的突破点。

 

第六阶段：信仰

如果一个交易员能够长期，持续，稳定的运行一套系统。其必然是渡过了上个阶段的认知冲击期。这说明他对他所使用的系统的各个环节都无比坚信，任何诱惑，任何怀疑都无法动摇其对系统的信任。

他可以淡然的看着各方明星崛起，他可以冷静的面对系统长时期的不利，他还可以平静的看着利益无处不在，到处充满诱惑的市场而稳如磐石，不为所动。因为他明白，他的专属系统，只要坚持下去，便可在未来登上金字塔的顶端。他会对自己的系统坚信不疑，任何其他的事情都无法对其产生冲击。

大师们的看法和建议，其实就是围绕三个方面来阐述的，资金管理，方法和心理。

正如著有《趋势的性质》一书，世界上第一批外包资金的交易大师，拉蒙巴罗斯所说：成功的交易并非来自于某一条神奇的方法论，成功的交易=交易心理*有效的风险管理*书面交易规则，和上述三个方面不谋而合。

其实说的就是，成功的交易=交易心理*资金管理*交易规则。我们也可以把资金管理和交易规则统称为交易系统，而交易心理的顶级形态就是如信仰般坚定。

以上便是我们提炼总结的投资者必经的六个阶段，那么读者你处于哪个阶段？

汇通网9月8日讯——德国商业银行固定收入(Fix Income, Commodity and Currency)技术分析部门主管琼斯(Karen Jones)周一(9月8日)指出，英镑兑美元反弹恐难持久，前景仍然看跌。

他指出,英镑兑美元(稍早)已经跌至了6月低点1.5172，目前正从该点位反弹。目前，日线图上的TD指标(德马克指标)已经完美地转好，这将使得该货币对在进一步的下跌之前小幅反弹。

分析师指出，此轮涨势可能勉力收复200日均线(当前处于1.5354)，但理想状态下将会被1.5425/27(8月7日低点)水平位所阻。因此他表示，将继续看跌英镑兑美元的走势前景，预计其将会跌至1.5088这一61.8%回撤位，然后是1.4895这一78.6%回撤位。

内容版权归汇通财经品牌所有，转载本文请务必标明："文章来源于汇通网"。违者必究！

更多热帖

作者 | 刘垚

编辑 | 尔悦

小 T 导读：在使用或者实现分布式数据库（Distributed Database）时，会面临把一个表的数据按照一定的策略分散到各个数据库节点上的情况，随之而来的是多节点数据查询复杂性的问题，例如 Join 和子查询。本文将会为你解读分布式数据库下子查询和 Join 等复杂 SQL 如何实现，来帮助你更好地解决上述问题。

首先简单讲一下 SQL 的执行过程：

SQL ==> Parser ==> Translate & Semantic Check ==> Optimizer ==> Coordinator ==> Executer

• Parser 产生的是语法树，即 Abstract Syntax Tree；
• Translate & Semantic Check，这一步会从 Catalog 读取元数据，用元数据完善语法树，便于 Optimizer 使用。例如：常见的 select * from tableA，一般会在这一步把“*”换成 tableA 的列；
• Optimizer 产生的是优化之后的逻辑执行计划，即 Optimized Logical Plan，执行计划是个有向无环图，即 DAG；
• Coordinator 负责分发逻辑执行计划给各个节点去计算；
• Executer 会把逻辑执行计划转成物理执行计划，即 Physical Plan。

开源的数据库有很多，我们可以结合一些主流数据库的源代码来理解子查询和 Join 的实现方式，比如关系型数据库 ：Impala、Presto、ClickHouse，时序数据库（Time- Series Database）： TDengine 等。下面从子查询和 Join 两部分进行分析。

子查询部分

逻辑执行计划有多种 Node，分别对应着 SQL 中的各种计算，包括 Scan Node、Join Node、Aggregate Node、Sort Node、Project Node 等等，相应的物理执行计划的算子为 Scan Operator 、Join Operator、Aggregate Operator、Sort Operator、Project Operator 等等。而数据库一般没有计算子查询的算子，这是因为将抽象语法树转成逻辑执行计划之后，就已经没有子查询的概念了，其运行逻辑是数据算子之间自下而上逐层传递，并逐层计算，并不特别计算子查询。下面讲一下分布式数据库针对子查询的一些相关处理。

首先，分布式数据库的优化器会将子查询扁平化处理，这种方式一般分为两种，一种是直接在语法树（AST）上做子查询扁平化（Subquery Flatten），另外一种是在生成逻辑执行计划时进行扁平化。这两种方式本质上大同小异，都要保证语义的等价性。但也并不是所有的子查询都能扁平化，有如下几种特殊情况：

• 子查询和父查询都有聚集函数
• 子查询有聚集函数，并且父查询有分组计算（Group By）
• 子查询有聚集函数，并且用子查询聚集函数的结果关联（Join）父查询的表
• 父查询有聚集函数，并且子查询有分组计算（Group By）
• 子查询有 Limit（限制返回结果的行数），并且父查询有过滤条件（Where）或者分组计算、排序（Order By）
• 其他

基于 AST 进行子查询扁平化时，需要先遍历语法数据，并按规则进行判断，进而去除不必要的子查询。对于生成逻辑执行计划时的子查询扁平化，在生成 Plan Node 时需要先去除冗余的 Node，举个例子，SQL：select colA from (select * from tA) group by colA;

一般来说，逻辑执行计划会有多个子计划，通常在需要网络传输时才会产生子计划，需要注意的是子计划和子查询之间并没有必然的联系，即有子查询不一定对应一个子计划。

Join 部分

首先，分布式数据库会对 Join 进行优化，包括 Join 消除（例如基于主键外键去除不必要的 Join）、外连接消除（Outer Join 转成 Inner Join）、Join Order 优化（基于数据的统计信息，用动态规划算法、贪心算法或遗传算法等优化 Table 的 Join 顺序）等等。

再讲一下 Join 的三种基本算法：Hash Join（必须要有等值连接条件，例如 t1.colA = t2.colB）、Merge Join（左表和右表的数据都是有序的，按连接条件中的列有序）、Nestloop Join（含有非等值连接条件并且数据无序）。在实际当中，会把三种算法进行混合使用，这是因为 Join 条件可以同时包含等值连接和非等值连接，例如 t1.colA = t2.colB AND t1.colC > t2.colC

Hash Join

在进行 Join Order 优化时，优化器会调整左表和右表的顺序，一般把小表放右边，大表放左边，并且选择 Join 模式：Shuffle Join（按照关联条件，同时 shuffle 左表和右表，然后再计算 Join） 或 Boradcast Join（把右表广播到左表所在的节点，注意左表不动，然后再计算 Join）。一般是基于代价去选择 Join Order 优化，但考虑到统计信息可能会存在误差，因此很多数据库可以通过 Hint、Query Option 等方式，由用户来指定 Join 顺序、Join 模式等。

Hash Join 是目前最常用的 Join 算法，大部分数据库都实现了 Hash Join。这种算法会先读取右表，并把右表的数据放入 Hash Map 里，如果存不下就会放入外存。通常情况下，各个数据库都会实现自己的 Hash Map，很少直接使用 STL 或 Boost 等第三方库中的 Hash Map，原因主要有两点：

• 定制化 Hash Map 会提升 Join 计算速度。
• 定制化 Hash Map 能更准确地控制内存使用，当内存不足时，会使用外存，定制化 Hash Map 可以根据 Join 算法，优化 Swap 机制，减少 Swap 的数据量。Hash Map 的结构如下：

右表可能含有重复的数据，所以会有 Duplicate Node。这里的重复数据是指 Join Key（Join 条件对应的列）的数据重复，并且其他列不重复，所以要分别缓存。注意上述图中，是通过 Hash 算法解决 Hash 冲突的问题，即不会把不同的 Join Key 放在同一个桶中。当然，现实操作中也有把不同的 Join Key 放在同一个桶中的情况，那需要遍历 List 才能确定查找的 Join Key 是否存在。

Merge Join

Merge Join 一般是在左表和右表的数据是有序的情况下使用。例如时序数据库 TDengine，数据按时间戳列有序，那么用时间戳列做 Join 时，TDengine database 会用 Merge Join 来计算，这样的一个好处是处理速度非常快，并且占用内存非常小。

Nestloop Join

这种 Join 算法速度非常慢，但对于全功能数据库而言是不可缺少的。使用这种算法时，可以结合索引来提速。

总结而言，Hash Join 使用最广，适用于很多数据分析的场景，并且大部分数据库都支持；Merge Join 一般是在左右表数据有序时才会使用，不需要缓存数据，所以使用内存非常少，计算速度是三种 Join 算法中最快的；Nestloop Join 性能很差，分布式数据库一般很少使用，有些分布式数据库就不支持，可以通过索引来加速 Nestloop Join。

写在最后

上面我们对子查询和 Join 两种复杂 SQL 的实现方式做了具体解读，大家可以结合一些开源数据库的源代码来理解，像 TDengine 的源代码都可以在 GitHub 上看到，如果你对时序数据库的复杂 SQL 实现有兴趣，这就是一个不错的观摩对象。也欢迎大家在下方评论区进行交流。

---

点击了解更多 TDengine Database 的具体细节。

撰文｜姚迟、郑泽康

本文将以开发一个 leaky_relu（准确说是 leaky_relu_yzh op，因为 master 分支的 leaky_relu 组合了其它知识点）为例介绍如何在 OneFlow 中新增算子（https://github.com/Oneflow-Inc/oneflow/pull/8350）。

 

1

背景

op 与 kernel

op 与 kernel 是两个有关联的概念。op 是逻辑上的算子，包含 OneFlow Compiler 在构建计算图时所需要的必要信息，如输入、输出形状，哪些张量需要自动求导等信息。有了 op 中的信息，OneFlow Compiler 就可以构建计算图并依据计算图做资源申请、构建等操作（如根据张量的输入输出大小申请内存）， 但是 op 中不包含具体的处理数据的逻辑。

在真正需要处理数据时，OneFlow Runtime 会启动 kernel 完成计算，所以 kernel 中包含了具体处理数据的逻辑。对于一个逻辑上的 op，OneFlow Runtime 会根据数据类型、硬件设备（比如是 CPU 还是 CUDA）的具体情况，选择启动不同的 kernel。

OneFlow 中的系统 op 与 user op

在 OneFlow 系统中存在两类算子（op）：系统 op 和 user op。

系统 op 定义在：oneflow/core/operator/ 目录， 对应的 kernel 实现在：oneflow/core/kernel 目录。系统 op 是对构图、流水等系统性能较为关键的一些 op。

除极少数 op 属于系统 op 外，大多数 op 都是 user op，这些 user op 和用户模型业务逻辑相关。OneFlow user op 的定义及 kernel 实现分别在 oneflow/user/ops 和 oneflow/user/kernels 目录下。

目前 OneFlow 已实现了丰富的算子库，但是当已有的算子库无法满足搭建模型的需求时，就需要新增算子。本文介绍的新增算子指的是新增 user op。

ODS 与 TableGen

TableGen（https://llvm.org/docs/TableGen/index.html） 是一个代码生成工具，简单而言，它读取并解析一个 .td 格式（语法接近 C++ 模板）的文件，然后交给 TableGen 后端

（https://llvm.org/docs/TableGen/BackEnds.html）生成另外格式的语言。

MLIR 基于 TableGen 制定了一套算子定义规范ODS（https://mlir.llvm.org/docs/OpDefinitions/）以及对应的后端 OpDefinitionsGen（https://github.com/llvm/llvm-project/blob/main/mlir/tools/mlir-tblgen/OpDefinitionsGen.cpp。）

OneFlow 在 ODS 的基础上，实现了 TableGen OneFlow 后端（https://github.com/Oneflow-Inc/oneflow/tree/master/tools/oneflow-tblgen），并使用它来定义 OneFlow user op。

因此，OneFlow 的 user op 定义写在 OneFlowUserOps.td 文件中。

 

2

开发 op

在 OneFlow 中开发一个新的 user op，主要分为以下4步：

1. 定义 op
2. 实现 kernel 计算逻辑
3. 导出 functional 接口
4. 实现用于求导的反向逻辑

定义 op

定义 op 指的是，对 op 的名称，op 的输入、输出数据类型和 op 的属性进行声明。OneFlow 遵循 MLIR 的 ODS（Operation Definition Specification）（
https://mlir.llvm.org/docs/OpDefinitions/） 实现了自己的 MLIR OneFlow Dialect。在算子定义方面，这样做的好处是，各种推导函数和序列化/反序列化的接口都可以委托给 ODS，降低了人工手写出错的概率，后续优化、格式转化等流程可以更灵活。

定义一个 OneFlow user op，主要包括 5 个部分，分别是：

• op class
• 输入 input
• 输出 output
• 属性 attrs
• 导出并实现推导接口

op class

可以在
oneflow/ir/include/OneFlow/OneFlowUserOps.td 查看 op 定义的源码。

以 def 关键字开头定义一个 op，该 op 继承 OneFlow_BaseOp，同时指定 OneFlow_BaseOp 的模版参数。模版参数依次为 op type name、Trait （
https://mlir.llvm.org/docs/Traits/）列表。

def OneFlow_LeakyReluYZHOp : OneFlow_BaseOp<"leaky_relu_yzh", [NoSideEffect, DeclareOpInterfaceMethods<UserOpCompatibleInterface>]> { //... }

其中 "leaky_relu_yzh" 是指定的 op type name。每个 op 都需要指定一个全局唯一的 op type name 作为全局标识符。

第二个模板参数是一个 list（[...]），其中的每一项都是一个 Trait，OneFlow 中常用的有：

• NoSideEffect 表示该算子无副作用（即不会改变内存、网络、管道、磁盘等的系统状态），这个特性可以指导某些优化操作
• NoGrad 表示该算子在数学上没有梯度（不可导）
• CpuOnly 表示该算子只支持在 CPU 设备上执行
• SupportNonContiguous 表示该算子是否支持 NonContiguous 张量（关于 Contiguous Tensor 的概念，可以参考 PyTorch Internals 中的相关内容 ）

输入 input 与输出 output

通过重写 input 域来定义 op 的输入，比如

// 一个输入 x let input = (ins OneFlow_Tensor:$x );

定义了一个输入张量 x。输入的格式为 输入类型:$name。

输入类型目前包括：

• OneFlow_Tensor
• Variadic<OneFlow_Tensor>：指可变 tensor，比如 concat op，支持 concat 可变个数的 tensor。
• Optional<OneFlow_Tensor>：表示这个 tensor 是可选的，既可以有也可以没有，比如 conv op 中的 add_output。

一个 op 也可以定义多个输入，比如：

// 两个输入：a, b let input = (ins OneFlow_Tensor:$a, OneFlow_Tensor:$b );

通过重写 output 域来定义 op 的输出，比如下面定义了 2 个输出张量：

let output = (outs OneFlow_Tensor:$out0, OneFlow_Tensor:$out1 );

属性 attrs

通过重写 attrs 域定义 op 的属性，比如定义 dropout （
https://oneflow.readthedocs.io/en/master/functional.html#
oneflow.nn.functional.dropout）中的 rate 属性：

let attrs = (ins DefaultValuedAttr<F32Attr, "0.">:$rate );

它表示名为 $rate 的类型是 F32Attr，默认值是 0.。这里也可以不指定默认值：

let attrs = (ins F32Attr:$rate );

I32Attr、F32Attr、BoolAttr、StrAttr、I32ArrayAttr 等常见基础数据类型定义在 OpBase.td

（https://github.com/llvm/llvm-project/blob/main/mlir/include/mlir/IR/OpBase.td#L1077-L1086）中。

OneFlow 自定义数据类型，如 ShapeAttr、DTArrayAttr 等定义在 OneFlowBase.td

（https://github.com/Oneflow-Inc/oneflow/blob/master/oneflow/ir/include/OneFlow/OneFlowBase.td#L27-L35）中。

导出并实现推导接口

还有一些其它域，用于指定是否生成对应的接口。这些接口往往是构建计算图过程中的推导接口。

比如 shape 推导（根据输入的 shape 推导输出的推导）、data type 推导、SBP 推导等。

OneFlow-TableGen 仅负责生成这些函数的接口，开发者需要在其自动生成的 cpp 文件中实现这些接口。默认情况不会生成下列任何接口，开发者需要显式指定需要生成哪些接口。

let has_check_fn = 1; // 生成属性检查接口 let has_logical_tensor_desc_infer_fn = 1; // 生成 logical shape 推导接口 let has_physical_tensor_desc_infer_fn = 1; // 生成 physical shape 推导接口 let has_get_sbp_fn = 1; // 生成 get sbp 接口 let has_sbp_signature_infer_fn = 1; // 生成 sbp signature 推导接口，未来会移除，推荐使用 has_nd_sbp_infer_fn let has_data_type_infer_fn = 1; // 生成 data type 推导接口 let has_device_and_stream_infer_fn = 1; // 生成 device 推导接口 let has_input_arg_modify_fn = 1; // 生成输入 modify 接口，比如设置 is_mutable、requires_grad（用于Lazy）等 let has_output_arg_modify_fn = 1; // 生成输出 modify 接口，比如设置 is_mutable、requires_grad（用于Lazy）等 let has_output_blob_time_shape_infer_fn = 1; // 生成输出 time shape 推导接口 let has_nd_sbp_infer_fn = 1; // 生成 nd sbp 推导接口

一般常用的是下面几个：

let has_logical_tensor_desc_infer_fn = 1; let has_physical_tensor_desc_infer_fn = 1; let has_data_type_infer_fn = 1; let has_get_sbp_fn = 1;

了解完上面这些概念和用法后，可以开始修改
oneflow/ir/include/OneFlow/OneFlowUserOps.td文件。

leaky_relu_yzh op 完整的定义见 这里（
https://github.com/Oneflow-Inc/oneflow/blob/7ab4b0f08c86a6f8af08b44daa510725942288fb/oneflow/ir/include/OneFlow/OneFlowUserOps.td#L8418-L8433）。

在 OneFlowUserOps.td 中新增Op定义之后，重新 make 后会自动在 build 目录下的 oneflow/core/framework/ 目录下生成文件以下几个文件：

• op_generated.h：由解析 .td 文件生成的 op C++ 类
• op_generated.cpp：由解析 .td 文件生成的 op 注册代码（包含调用 REGISTER_USER_OP 宏的代码）

之后需要做的就是在 oneflow/user/ops （https://github.com/Oneflow-Inc/oneflow/tree/master/oneflow/user/ops）目录下新加一个 cpp 文件，用于实现 op 的接口。

比如 leaky_relu_yzh 对应的文在 oneflow/user/ops/leaky_relu_yzh_op.cpp（https://github.com/Oneflow-Inc/oneflow/blob/7ab4b0f08c86a6f8af08b44daa510725942288fb/oneflow/user/ops/leaky_relu_yzh_op.cpp#L21-L79），实现了推导逻辑张量、推导物理张量、推导 SBP 信息以及推导输出数据类型各接口。

实现 Kernel 逻辑

op 的计算支持多种设备（如 CPU、GPU、DCU 等），所以要分别实现计算逻辑。

相关代码：

• Leaky ReLU CPU Kernel
• （https://github.com/Oneflow-Inc/oneflow/blob/7ab4b0f08c86a6f8af08b44daa510725942288fb/oneflow/user/kernels/leaky_relu_yzh_kernel.cpp）
• Leaky ReLU GPU KernelCPU
• （https://github.com/Oneflow-Inc/oneflow/blob/7ab4b0f08c86a6f8af08b44daa510725942288fb/oneflow/user/kernels/leaky_relu_yzh_kernel.cu）

计算逻辑

template<typename T> class CpuLeakyReluYZHKernel final : public user_op::OpKernel { public: CpuLeakyReluYZHKernel() = default; ~CpuLeakyReluYZHKernel() = default; private: void Compute(user_op::KernelComputeContext* ctx) const override const user_op::Tensor* x = ctx->Tensor4ArgNameAndIndex("x", 0); user_op::Tensor* y = ctx->Tensor4ArgNameAndIndex("y", 0); const int32_t elem_cnt = x->shape().elem_cnt(); const T* x_ptr = x->dptr<T>(); T* y_ptr = y->mut_dptr<T>(); const auto alpha = ctx->Attr<float>("alpha"); FOR_RANGE(int32_t, i, 0, elem_cnt) { y_ptr[i] = x_ptr[i] > 0 ? x_ptr[i] : alpha * x_ptr[i]; } } bool AlwaysComputeWhenAllOutputsEmpty() const override { return false; } };

在 OneFlow 中实现 kernel， 必须定义一个继承自
oneflow::user_op::OpKernel 的类，并重写其中的虚函数。在以上代码中，重写了 Compute 与
AlwaysComputeWhenAllOutputsEmpty 两个虚函数，它们的意义分别是：

• Compute 必须重写，在其中实现具体的运算逻辑
• AlwaysComputeWhenAllOutputsEmpty 必须重写，对于绝大多数 op 而言，直接返回 false 即可。对于极少数内部需要维护状态，即使输出为空也需要调用 kernel 进行计算的 op 而言，应该返回 true

Compute 方法中通过调用
user_op::KernelComputeContext* ctx 中的接口，可以获取输入张量、输出张量、attr 具体的数据，再按照算子的算法逻辑对它们进行处理。以下是对
CpuLeakyReluKernel::Compute 处理逻辑的解读：

• 首先取得 "x"，"y" 两个 Tensor。传入Tensor4ArgNameAndIndex的字符串要和之前在OneFlowUserOps.td设置的名称一致
• 获取 x 的元素个数，以便后续用于 for 循环进行计算
• 获取属性 alpha
• 进入次数为 elem_cnt 的 for 循环，将结果写入

注册 Kernel

实现 kernel 类后，需要调用 REGISTER_USER_KERNEL 注册。

#define REGISTER_CPU_LEAKY_RELU_YZH_KERNEL(dtype) \ REGISTER_USER_KERNEL("leaky_relu_yzh") \ .SetCreateFn<CpuLeakyReluYZHKernel<dtype>>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCPU) \ && (user_op::HobDataType("y", 0) == GetDataType<dtype>::value));

这里会调用REGISTER_USER_KERNEL宏，包括以下信息：

1. op type name：为哪个 op 注册 kernel
2. SetCreateFn<T>()：该模板方法的模板参数 T，就是我们实现的 kernel 类，OneFlow Runtime 将使用它创建 kernel 对象。
3. SetIsMatchedHob：因为一个 op 可能有多个 kernel，要想根据物理设备及数据格式的不同而选择不同的 kernel 进行计算，就需要调用 SetIsMatchedHob 进行设置。该方法接受一个表达式，表达式为 true 时，OneFlow 将调用该 kernel 完成计算。以上代码的匹配逻辑是：当硬件设备为 cpu，且 y 的数据类型和 dtype 一致时，选择调用注册的 kernel 类（CpuLeakyReluYZHKernel<dtype>）。

GPU 计算逻辑

CUDA 编程基础知识入门可以参考：

• 视频：CUDA 的由来（https://www.bilibili.com/video/BV1Mb4y1p7BG）
• 视频：CUDA 的入门小程序（https://www.bilibili.com/video/BV1bF411s76k）
• 视频：线程层级（https://www.bilibili.com/video/BV1MZ4y127Sq）

不过以上的视频都无法替代自己认真学习官方资料：CUDA C Programming Guide（https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html）

了解了 CUDA 的基础知识，就不难理解 leaky_relu CUDA 版本的实现。

首先定义了 leaky_relu 前向运算的 CUDA 核函数

template<typename T> __global__ void LeakyReluForwardGpu(const int n, const float alpha, const T* x, T* y) { CUDA_1D_KERNEL_LOOP(i, n) { y[i] = x[i] > 0 ? x[i] : x[i] * alpha; } }

其中调用了宏 CUDA_1D_KERNEL_LOOP （https://github.com/Oneflow-Inc/oneflow/blob/master/oneflow/core/device/cuda_util.h#L91-L94）进行运算

在 Compute 函数中，调用了 RUN_CUDA_KERNEL (也是定义在 cuda_util.h 这个文件中)这个宏启动核函数。

对应的 GPU kernel 类的实现见：

https://github.com/Oneflow-Inc/oneflow/blob/7ab4b0f08c86a6f8af08b44daa510725942288fb/oneflow/user/kernels/leaky_relu_yzh_kernel.cu#L32-L49

其中用到了启动 kernel 的宏 RUN_CUDA_KERNEL，它的定义是:

#define RUN_CUDA_KERNEL(func, device_ctx_ptr, thread_num, ...) \ func<<<SMBlocksNum4ThreadsNum(thread_num), kCudaThreadsNumPerBlock, 0, \ (device_ctx_ptr)->cuda_stream()>>>(__VA_ARGS__)

1. 第一个参数是核函数名字
2. 第二个参数是 device context，后续获取对应的 cuda_stream
3. 第三个参数是要启动的线程数量，会根据线程数量来计算所需的 Block 数目。

因为 leaky relu 是 elementwise 运算，各个元素互不影响，所以我们启动了 elem_cnt 个线程。

后续的注册与 CPU 版本类似，这里不再赘述。直接参考以下代码即可：

https://github.com/Oneflow-Inc/oneflow/blob/7ab4b0f08c86a6f8af08b44daa510725942288fb/oneflow/user/kernels/leaky_relu_yzh_kernel.cu#L51-L62

可以看到不同设备类的 Compute 中大部分代码是重复的。一种更优的代码组织方式是用一个 .cpp 文件完成 kernel 和注册的逻辑，.cu 文件编写 GPU Kernel 函数和 GPU 模板特化的代码，.h 文件用于定义和编写注册宏。可参考 dim_gather_kernel_*

（https://github.com/Oneflow-Inc/oneflow/tree/master/oneflow/user/kernels）中的代码。

OneFlow 为了适配多种设备，还提供了 Primitive 组件，可以参考：Primitive PR（
https://github.com/Oneflow-Inc/oneflow/pull/6234）

导出 functional 接口

关于 functional 接口层的详细介绍在这里：
https://github.com/Oneflow-Inc/oneflow/wiki/Functional-Interface

概括而言，functional 层起到了“上接 Python，下联 C++”的作用：

┌─────────────┐ │ Module │ │ (Python) │ ├─────────────┤ │ │ │ Functional │ ├─────────────┤ │ │ │ Op/Kernels │ │ (C++) │ └─────────────┘

因此，在上文定义 op 和注册 kernel 后，需要为算子导出 functional 接口，才能使用户通过 Python 代码调用该算子。

导出 functional 接口分为以下几个步骤：

1. 实现对应的 functor 并注册
2. 在 oneflow/core/functional/functional_api.yaml 中添加接口描述

实现对应的 functor 并注册

对于 leaky_relu_yzh op，在 activation_functor.cpp

（https://github.com/Oneflow-Inc/oneflow/blob/7ab4b0f08c86a6f8af08b44daa510725942288fb/oneflow/core/functional/impl/activation_functor.cpp#L391-L421） 中，对其进行定义：

class LeakyReluYZHFunctor { public: LeakyReluYZHFunctor() { op_ = CHECK_JUST(one::OpBuilder("leaky_relu_yzh").Input("x").Output("y").Build()); } Maybe<Tensor> operator()(const std::shared_ptr<one::Tensor>& x, const float& alpha) const { MutableAttrMap attrs; JUST(attrs.SetAttr<float>("alpha", alpha)); return OpInterpUtil::Dispatch<one::Tensor>(*op_, {x}, attrs); } private: std::shared_ptr<OpExpr> op_; };

• 在构造函数里，构造了 leaky_relu 这个op
• 实现 operator() 重载运算符，通过 Dispatch 调用构造好的 op，并分别传入输入，属性

类似的我们也给 LeakyReluGrad 导出 functional 接口，以便后续编写求导逻辑使用。

最后我们需要注册到 Functional Library：

https://github.com/Oneflow-Inc/oneflow/blob/7ab4b0f08c86a6f8af08b44daa510725942288fb/oneflow/core/functional/impl/activation_functor.cpp#L610-L611

m.add_functor<impl::LeakyReluYZHFunctor>("LeakyReluYZH"); // 注意最后字符串中的名字在后续的 functional_api.yaml 中会用到

通过 m.add_functor 注册后的 functor，可以在 C++ 层使用，如通过 functional::LeakyRelu 就可以调用 LeakyReluFunctor。

在 functional_api.yaml 中添加接口描述

functional 通过解析 yaml 配置文件，在 build 过程中自动帮我们生成接口。

在functional_api.yaml

（https://github.com/Oneflow-Inc/oneflow/blob/master/oneflow/core/functional/functional_api.yaml）文件中，编写相关配置。

https://github.com/Oneflow-Inc/oneflow/pull/8350/files#diff-4b35c1dcdbae81b75439ba570bc149554ca85b83757430613fcb612ae25972afR572-R579

- name: "leaky_relu_yzh" signature: "Tensor (Tensor x, Float alpha) => LeakyReluYZH" bind_python: True

• 其中 name 表示导出到 Python 接口后函数的名字，比如导出后在 Python 下使用就是

flow._C.leaky_relu_yzh(...)

• signature 用于描述接口原型及 C++ 代码的对应关系。=> 左边的为原型；=> 右边为对应的 Functional Library 中的名字。这里LeakyRelu 和前面导出时指定的字符串是一致的。
• bind_python，表示这个接口是否需要绑定 Python 接口 。比如下面的 leaky_relu_grad，我们不会在 Python 层用到（但会在 C++ 层求导使用），所以设置为 False。

完成以上工作后，新增的算子已经支持正向运算，编译好代码便可以进行如下简单的测试：

import oneflow as flow import numpy as np x_tensor = flow.Tensor(np.random.randn(3, 3)) out = flow._C.leaky_relu_yzh(x_tensor, alpha=0.2)

但是，还需要注册反向，才能支持反向传播。我们也先将反向需要的 LeakyReluGrad 导出为 functional 接口。

- name: "leaky_relu_yzh_grad" signature: "Tensor (Tensor x, Tensor dy, Float alpha) => LeakyReluYZHGrad" bind_python: False实现用于求导的反向逻辑

反向传播的本质就是高数中的链式法则，只不过 Autodiff 将链式法则变得模块化、易复用。

可以先阅读 CSC321 Lecture 10: Automatic Differentiation（https://www.cs.toronto.edu/~rgrosse/courses/csc321_2018/slides/lec10.pdf ）了解 autodiff 的基本概念。

从逻辑上而言，一个算子在反向过程中能够求导数，一般需要以下信息：

• 正向过程中的输入、输出
• 正向过程的 attr
• 反向过程中上一层（正向过程中的下一层）传递过来的正向输出的梯度

未来 Graph 模式和 Eager 模式下的反向逻辑会合并，但目前还是需要分别注册。

为 Eager 模式注册反向

求导部分在
oneflow/core/autograd/gradient_funcs/activation.cpp

（https://github.com/Oneflow-Inc/oneflow/pull/8350/files#diff-36aeebf7fd5a8b88bd5af87774e7b0b4f76fed42cfb75044d6eebdfe65628347R213-R253）完成

主要有以下几部分：

• LeakyReluYZHCaptureState ：用于存储数据的结构体

这是一个简单的结构体，继承自 AutoGradCaptureState，用于存储 op 的属性，以便于后续求导。

struct LeakyReluYZHCaptureState : public AutoGradCaptureState { bool requires_grad; // 输入x是否需要梯度 float alpha=0.0; // 输入的参数alpha };

• LeakyReluYZH 类：继承自 OpExprGradFunction 的类。需要重写三个函数：Init、Capture、Apply。

class LeakyReluYZH : public OpExprGradFunction<LeakyReluYZHCaptureState> { public: Maybe<void> Init(const OpExpr& op) override { //... } Maybe<void> Capture(LeakyReluYZHCaptureState* ctx, const TensorTuple& inputs, const TensorTuple& outputs, const AttrMap& attrs) const override { //... } Maybe<void> Apply(const LeakyReluYZHCaptureState* ctx, const TensorTuple& out_grads, TensorTuple* in_grads) const override { //... } };

• Init：做的是一些初始化的工作，可以根据前向 op 的配置，来初始化属性。

Maybe<void> Init(const OpExpr& op) override { const auto* fw_op_expr = dynamic_cast<const UserOpExpr*>(&op); CHECK_NOTNULL_OR_RETURN(fw_op_expr); base_attrs_ = MakeAttrMapFromUserOpConf(fw_op_expr->proto()); return Maybe<void>::Ok(); }

• Capture：用于捕捉前向的 Tensor，属性，用于后续的求导。

以 LeakyReluYZH 为例子，我们需要得到：a) 输入的 Tensor，当 Tensor 数值大于 0，梯度为 1，当小于 0，梯度为 alpha b) alpha的数值

Maybe<void> Capture(LeakyReluYZHCaptureState* ctx, const TensorTuple& inputs, const TensorTuple& outputs, const AttrMap& attrs) const override { CHECK_EQ_OR_RETURN(inputs.size(), 1); // 判断输入个数是否为1 ctx->requires_grad = inputs.at(0)->requires_grad(); // 判断输入是否需要梯度 if (!ctx->requires_grad) { return Maybe<void>::Ok(); } // 如果不需要梯度，也就不需要求导了，直接返回 Maybe<void>::Ok() ComposedAttrMap composed_attrs(attrs, base_attrs_); ctx->alpha = JUST(composed_attrs.GetAttr<float>("alpha")); // 获取 alpha，并存入 ctx->alpha 中 ctx->SaveTensorForBackward(inputs.at(0)); // 调用 SaveTensorForBackward 方法，保存输入的 Tensor return Maybe<void>::Ok(); }

• Apply：实际计算梯度的函数，我们可以拿到先前的 Tensor，并调用 functional 接口下注册的 LeakyReluGrad，求得梯度，并返回

Maybe<void> Apply(const LeakyReluYZHCaptureState* ctx, const TensorTuple& out_grads, TensorTuple* in_grads) const override { CHECK_EQ_OR_RETURN(out_grads.size(), 1); // 检查梯度 Tensor 个数是否为 1 in_grads->resize(1); // 因为输入只有一个，所以我们 resize(1) if (ctx->requires_grad) { const auto& x = ctx->SavedTensors().at(0); // 调用 SavedTensors 接口，拿到先前通过 SaveTensorForBackward 接口保存的 Tensor in_grads->at(0) = JUST(functional::LeakyReluYZHGrad(x, out_grads.at(0), ctx->alpha)); // 拿到 x，dy，alpha，传入给 LeakyReluYZHGrad 计算，并将梯度返回给 in_grads->at(0) } return Maybe<void>::Ok(); }

最后一步是注册，第一个参数是 op type name，第二个参数是继承自 OpExprGradFunction 的类。

REGISTER_OP_EXPR_GRAD_FUNCTION("leaky_relu_yzh", LeakyReluYZH); // 第二个参数是用于求导的类名

为 Graph 模式注册反向

为 Graph 模式注册 leaky_relu_yzh op 的反向代码在：

https://github.com/Oneflow-Inc/oneflow/pull/8350/files#diff-ef94ddb8f5c25689f2c6fab7a9675f16c95a22018a8c01647b4398ce2fb85a8bR81-R97

REGISTER_USER_OP_GRAD("leaky_relu_yzh") .SetBackwardOpConfGenFn([](user_op::BackwardOpConfContext* ctx) -> Maybe<void> { // 根据前向的 op type name，拼凑出一个 leaky_relu_yzh_grad_op_name （leaky_relu_yzh_grad） const std::string leaky_relu_yzh_grad_op_name = ctx->FwOp().op_name() + "_grad"; ctx->DefineOp(leaky_relu_yzh_grad_op_name, [&ctx](user_op::BackwardOpBuilder& builder) { // 构建一个 op（op type name 为 leaky_relu_yzh_grad） // 把前向输出 y 的梯度，作为 leaky_relu_yzh_grad 的输入 dy // 把前向的 x 作为 leaky_relu_yzh_grad 的输入 x // 输出为 dx // attr alpha 同前向一样 return builder.OpTypeName("leaky_relu_yzh_grad") .InputBind("dy", ctx->FwOp().output_grad("y", 0)) .InputBind("x", ctx->FwOp().input("x", 0)) .Attr("alpha", ctx->FwOp().attr<float>("alpha")) .Output("dx") .Build(); }); // 把 leaky_relu_yzh_grad_op_name 算子的输出 dx 的结果 // 绑定到前向输入 x 的反向梯度上 // 即： // leaky_relu_yzh 的输入 x 的梯度 = leaky_relu_yzh_grad 的输出 dx ctx->FwOp().InputGradBind(user_op::OpArg("x", 0), [&ctx, &leaky_relu_yzh_grad_op_name]() -> const std::string& { return ctx->GetOp(leaky_relu_yzh_grad_op_name).output("dx", 0); }); return Maybe<void>::Ok(); });

3

测试与文档

本文覆盖的内容完成后，只是“跑通”算子，还需要进一步完善，包括为算子添加测试和 API 文档，这些将在后续的文章中介绍。

 

欢迎下载体验 OneFlow v0.8.0 最新版本：
https://github.com/Oneflow-Inc/oneflow/

上一篇：短线是金客的博客(短线是金书籍)

下一篇：股票(科达机电黎智清)