CANN ops-nn 算子解析SigmoidCrossEntropyWithLogitsV2 与 aclnnBinaryCrossEntropyWithLogits 二分交叉熵损失详解【免费下载链接】ops-nn本项目是CANN提供的神经网络类计算算子库实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-nnSigmoidCrossEntropyWithLogitsV2 是 CANN ops-nn 神经网络算子库中用于计算「Sigmoid 二分类交叉熵BCEWithLogits」融合损失的算子功能上与 PyTorch 的torch.nn.BCEWithLogitsLoss对齐适用于单标签与多标签二分类任务的损失计算与训练优化。本文以 算子 README 为主体结合该算子目录下的 aclnn 接口文档、C 示例、op_api / op_host / op_kernel 源码与单元测试系统讲解其数学原理、参数语义、支持平台、aclnn 两段式调用流程与底层实现机制帮助读者在 CANN 环境下正确、高效地完成二分类交叉熵损失的计算。一、算子概述与产品支持情况SigmoidCrossEntropyWithLogitsV2 位于 loss/sigmoid_cross_entropy_with_logits_v2 目录是 ops-nn 损失loss模块中的成员。其接口功能是计算输入 logitsself与标签target之间的 BCELoss 损失即对 logits 先做 Sigmoid 激活再与标签计算二分类交叉熵并把权重weight、正类权重pos_weight与 reduction 归约策略一并内建在损失公式中。从算子定义protobuf注释可知该算子的设计目标是与 PyTorch 的BCEWithLogitsLoss保持第三方框架兼容详见 sigmoid_cross_entropy_with_logits_v2_proto.h。产品支持矩阵产品是否支持Ascend 950PR / Ascend 950DT√Atlas A3 训练系列产品 / Atlas A3 推理系列产品√Atlas A2 训练系列产品 / Atlas A2 推理系列产品√Atlas 200I/500 A2 推理产品×Atlas 推理系列产品√Atlas 训练系列产品√需要特别注意的是Atlas 推理系列产品与 Atlas 训练系列产品上数据类型不支持 BFLOAT16其余平台支持 FLOAT16 / FLOAT / BFLOAT16。从源码侧可以印证这一支持矩阵算子的 AICore 配置在 sigmoid_cross_entropy_with_logits_v2_def.cpp 中仅注册了ascend950与ascend350两个平台并开启了动态编译、动态 rank 与动态 shape 支持而在 aclnn_binary_cross_entropy_with_logits.cpp 中GetDtypeSupportList()依据 SoC 版本区分了两套数据类型支持列表ASCEND910_DTYPE_SUPPORT_LIST仅DT_FLOAT、DT_FLOAT16对应 Atlas 训练/推理系列不支持 BF16与 README 描述一致ASCEND910B_DTYPE_SUPPORT_LISTDT_FLOAT、DT_FLOAT16、DT_BF16对应 A2/A3 及以上平台。二、数学原理单标签与多标签 BCELoss 公式单标签场景定义逐元素损失向量$$ \ell(self, target) L {l_{1},..., l_{n}}^{T} $$其中每个元素的损失为$$ \ell_{n} -weight_{n}[target_{n} \cdot log(\sigma(self_{n})) (1 - target_{n}) \cdot log(1 - \sigma(self_{n}))] $$这里 $\sigma(\cdot)$ 为 Sigmoid 函数$\sigma(self_{n}) \frac{1}{1 e^{-self_{n}}}$即对连接层logits输出先做 Sigmoid 归一化到 (0, 1) 区间再与二值标签 target 计算交叉熵weight为逐样本/逐元素的损失权重。reduction 归约当指定 reduction 时对损失向量做归约$$ \ell(self, target) \begin{cases} L, if\ reduction none\ mean(L), if\ reduction mean\ sum(L), if\ reduction sum\ \end{cases} $$none不做任何操作直接返回逐元素损失mean对损失取平均值sum对损失求和。多标签多类别场景当 self 与 target 为多维张量、每个样本同时对应多个二分类标签时对第 c 个类别计算损失$$ \ell_c(self, target) L_c {l_{1,c},..., l_{n,c}}^{T} $$$$ \ell_{n,c} -weight_{n,c}[pos_weight_{n,c} \cdot target_{n,c} \cdot log(\sigma(self_{n,c})) (1 - target_{n,c}) \cdot log(1 - \sigma(self_{n,c}))] $$与单标签公式相比多标签场景引入了pos_weight各类的正类权重当target 1时损失被pos_weight放大从而缓解正负样本不均衡问题当target 0时pos_weight不参与计算。这与 PyTorchBCEWithLogitsLoss的pos_weight语义一致。值得说明的是从 op_api 的实现看见下文“实现机制”一节aclnn 接口内部会先以reductionnone调用融合算子得到逐元素损失再通过 ReduceMean / ReduceSum 完成 mean / sum 归约因此上述公式中的 none 分支是融合算子的核心计算路径。三、参数说明以下参数表来自 算子 README 并补充了 aclnn 接口文档 aclnnBinaryCrossEntropyWithLogits.md 中的细化信息参数名输入/输出描述使用说明数据类型数据格式维度(shape)非连续Tensorself输入连接层输出logits-FLOAT16、FLOAT、BFLOAT16ND1-8 维√target输入label 标签值-与 self 一致ND与 self 一致√weightOptional输入二分交叉熵权重shape 需要能够 broadcast 到 target与 self 一致ND1-8 维√posWeightOptional输入各类的正类权重shape 需要能够 broadcast 到 target与 self 一致ND1-8 维√reduction输入输出结果计算方式aclnn 直调接口支持 0(none) / 1(mean) / 2(sum)INT64---out输出输出误差reduction none 时 shape 与 self 一致其他情况 shape 为 [1]与 target 一致ND与 self 一致√几个关键点需要重点理解reduction 的取值语义在 aclnn_binary_cross_entropy_with_logits.cpp 中enum Reduction { None 0, Mean 1, Sum 2 }定义了三种取值且REDUCTION_MAX_NUM 2CheckReduction()会对超出 [0, 2] 范围的取值报ACLNN_ERR_PARAM_INVALID错误码 161002。README 中“直调算子当前仅支持 none 值”的描述指的是底层融合算子的reduction属性见 tiling 源码 中REDUCTION_MODE_KEY {{none, 0}}tiling 只处理 none 分支而aclnn 接口层完整支持 0/1/2 三种归约mean 与 sum 由接口内部的 Reduce 算子补齐。weight / pos_weight 的广播要求两者的 shape 必须能够 broadcast 到 target 的 shape广播规则详见 broadcast_relationship.md不满足时报ACLNN_ERR_PARAM_INVALID。数据类型一致性target、weight、posWeight、out 的数据类型均要求与 self 保持一致aclnn 层还会校验 out 与 target 类型一致见CheckDtypeValid()out 仅在reduction none时与 self 同 shapemean/sum 时为标量 shape [1]。空 tensor 特例当 target 为空 tensor 时代码不做 broadcast 校验与 GPU 行为保持一致见 aclnn_binary_cross_entropy_with_logits.cpp 中CheckShape()的实现。四、约束说明算子的约束说明为“无”即对输入 shape 无额外格式约束ND 格式、1-8 维。确定性计算aclnnBinaryCrossEntropyWithLogits默认为确定性实现见 aclnnBinaryCrossEntropyWithLogits.md 的“约束说明”。平台数据类型差异Atlas 推理/训练系列产品不支持 BFLOAT16规划模型时需注意精度选择。五、调用方式aclnn 两段式接口详解SigmoidCrossEntropyWithLogitsV2 算子通过 aclnn 接口对外暴露接口名为aclnnBinaryCrossEntropyWithLogits。与 CANN 算子库其他算子一致该接口采用两段式two-phase调用模型参见 two_phase_api.md先调用 GetWorkspaceSize 接口完成参数校验、流程构建与 workspace 大小计算再调用执行接口完成计算。5.1 函数原型aclnnStatus aclnnBinaryCrossEntropyWithLogitsGetWorkspaceSize( const aclTensor *self, const aclTensor *target, const aclTensor *weightOptional, const aclTensor *posWeightOptional, int64_t reduction, aclTensor *out, uint64_t *workspaceSize, aclOpExecutor **executor)aclnnStatus aclnnBinaryCrossEntropyWithLogits( void *workspace, uint64_t workspaceSize, aclOpExecutor *executor, const aclrtStream stream)5.2 第一段接口参数说明GetWorkspaceSize参数名输入/输出描述selfaclTensor*输入连接层输出logitsFLOAT16 / FLOAT / BFLOAT16ND 格式1-8 维支持非连续 TensortargetaclTensor*输入label 标签值与 self 保持类型与 shape 一致weightOptionalaclTensor*输入二分交叉熵权重可选需能 broadcast 到 targetposWeightOptionalaclTensor*输入各类正类权重可选需能 broadcast 到 targetreductionint64_t输入0(none) / 1(mean) / 2(sum)outaclTensor*输出输出误差none 时与 self 同 shape否则为标量 shapeworkspaceSizeuint64_t*输出返回需要在 Device 侧申请的 workspace 大小executoraclOpExecutor**输出返回 op 执行器包含算子计算流程5.3 第一段接口的入参校验与返回码第一段接口会完成入参校验出现以下场景时报错返回值错误码描述ACLNN_ERR_PARAM_NULLPTR161001传入的 self 或 out 为空指针ACLNN_ERR_PARAM_INVALID161002self、target、weightOptional 和 posWeightOptional 的数据类型和数据格式不在支持范围内ACLNN_ERR_PARAM_INVALID161002self 和 target 维度不一致ACLNN_ERR_PARAM_INVALID161002weightOptional、posWeightOptional 不能扩展成 self/target 形状更完整的返回码含义可参见 aclnn_return_code.md。从源码看上述校验对应 aclnn_binary_cross_entropy_with_logits.cpp 中CheckParams()的四步检查CheckNotNull()空指针、CheckDtypeValid()数据类型与输出类型匹配、CheckReduction()reduction 取值在 [0,2]、CheckShape()维度 ≤ 8、self 与 target 同 shape、权重可广播。5.4 第二段接口参数说明执行接口参数名输入/输出描述workspace输入在 Device 侧申请的 workspace 内存地址workspaceSize输入workspace 大小由第一段接口计算得到executor输入op 执行器包含算子计算流程stream输入指定执行任务的 Stream5.5 完整调用示例C以下示例来自 examples/test_aclnn_sigmoid_cross_entropy_with_logits_v2.cpp完整覆盖了 acl 初始化、Tensor 构造、两段式调用、结果回拷与资源释放的 7 个步骤可直接作为开发模板#include iostream #include vector #include acl/acl.h #include aclnnop/aclnn_binary_cross_entropy_with_logits.h #define CHECK_RET(cond, return_expr) \ do { \ if (!(cond)) { \ return_expr; \ } \ } while (0) #define LOG_PRINT(message, ...) \ do { \ printf(message, ##__VA_ARGS__); \ } while (0) int64_t GetShapeSize(const std::vectorint64_t shape) { int64_t shapeSize 1; for (auto i : shape) { shapeSize * i; } return shapeSize; } int Init(int32_t deviceId, aclrtStream* stream) { // 固定写法资源初始化 auto ret aclInit(nullptr); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclInit failed. ERROR: %d\n, ret); return ret); ret aclrtSetDevice(deviceId); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclrtSetDevice failed. ERROR: %d\n, ret); return ret); ret aclrtCreateStream(stream); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclrtCreateStream failed. ERROR: %d\n, ret); return ret); return 0; } template typename T int CreateAclTensor(const std::vectorT hostData, const std::vectorint64_t shape, void** deviceAddr, aclDataType dataType, aclTensor** tensor) { auto size GetShapeSize(shape) * sizeof(T); // 调用aclrtMalloc申请Device侧内存 auto ret aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclrtMalloc failed. ERROR: %d\n, ret); return ret); // 调用aclrtMemcpy将Host侧数据拷贝到Device侧内存上 ret aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclrtMemcpy failed. ERROR: %d\n, ret); return ret); // 计算连续tensor的strides std::vectorint64_t strides(shape.size(), 1); for (int64_t i shape.size() - 2; i 0; i--) { strides[i] shape[i 1] * strides[i 1]; } // 调用aclCreateTensor接口创建aclTensor *tensor aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND, shape.data(), shape.size(), *deviceAddr); return 0; } int main() { // 1.固定写法device/stream初始化参考acl API手册 int32_t deviceId 0; aclrtStream stream; auto ret Init(deviceId, stream); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(Init acl failed. ERROR: %d\n, ret); return ret); // 2. 构造输入与输出 std::vectorint64_t inputShape {4, 2}; std::vectorint64_t targetShape {4, 2}; std::vectorint64_t weightShape {4, 2}; std::vectorint64_t posWeightShape {4, 2}; std::vectorint64_t outShape {4, 2}; void* inputDeviceAddr nullptr; void* targetDeviceAddr nullptr; void* weightDeviceAddr nullptr; void* posWeightDeviceAddr nullptr; void* outDeviceAddr nullptr; aclTensor* input nullptr; aclTensor* target nullptr; aclTensor* weight nullptr; aclTensor* posWeight nullptr; aclTensor* out nullptr; std::vectorfloat inputHostData {0.1, 0.1, 0.2, 0.2, 0.3, 0.3, 0.4, 0.4}; std::vectorfloat targetHostData {0.2, 0.2, 0.1, 0.1, 0.2, 0.2, 0.1, 0.1}; std::vectorfloat weightHostData {0.5, 0.5, 0.5, 0.5, 0.5, 0.5, 0.5, 0.5}; std::vectorfloat posWeightHostData {0.5, 0.5, 0.5, 0.5, 0.5, 0.5, 0.5, 0.5}; std::vectorfloat outHostData {0, 0, 0, 0, 0, 0, 0, 0}; ret CreateAclTensor(inputHostData, inputShape, inputDeviceAddr, aclDataType::ACL_FLOAT, input); CHECK_RET(ret ACL_SUCCESS, return ret); ret CreateAclTensor(targetHostData, targetShape, targetDeviceAddr, aclDataType::ACL_FLOAT, target); CHECK_RET(ret ACL_SUCCESS, return ret); ret CreateAclTensor(weightHostData, weightShape, weightDeviceAddr, aclDataType::ACL_FLOAT, weight); CHECK_RET(ret ACL_SUCCESS, return ret); ret CreateAclTensor(posWeightHostData, posWeightShape, posWeightDeviceAddr, aclDataType::ACL_FLOAT, posWeight); CHECK_RET(ret ACL_SUCCESS, return ret); ret CreateAclTensor(outHostData, outShape, outDeviceAddr, aclDataType::ACL_FLOAT, out); CHECK_RET(ret ACL_SUCCESS, return ret); int64_t reduction 0; // 0none, 1mean, 2sum uint64_t workspaceSize 0; aclOpExecutor* executor; // 3. 调用aclnnBinaryCrossEntropyWithLogits第一段接口 ret aclnnBinaryCrossEntropyWithLogitsGetWorkspaceSize(input, target, weight, posWeight, reduction, out, workspaceSize, executor); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclnnBinaryCrossEntropyWithLogitsGetWorkspaceSize failed. ERROR: %d\n, ret); return ret); // 根据第一段接口计算出的workspaceSize申请device内存 void* workspaceAddr nullptr; if (workspaceSize 0) { ret aclrtMalloc(workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(allocate workspace failed. ERROR: %d\n, ret); return ret); } // 调用aclnnBinaryCrossEntropyWithLogits第二段接口 ret aclnnBinaryCrossEntropyWithLogits(workspaceAddr, workspaceSize, executor, stream); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclnnBinaryCrossEntropyWithLogits failed. ERROR: %d\n, ret); return ret); // 4.固定写法同步等待任务执行结束 ret aclrtSynchronizeStream(stream); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclrtSynchronizeStream failed. ERROR: %d\n, ret); return ret); // 5. 获取输出的值将Device侧内存上的结果拷贝至Host侧 auto size GetShapeSize(outShape); std::vectorfloat resultData(size, 0); ret aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr, size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(copy result from device to host failed. ERROR: %d\n, ret); return ret); for (int64_t i 0; i size; i) { LOG_PRINT(result[%ld] is: %f\n, i, resultData[i]); } // 6. 释放aclTensor aclDestroyTensor(input); aclDestroyTensor(target); aclDestroyTensor(weight); aclDestroyTensor(posWeight); aclDestroyTensor(out); // 7. 释放device资源 aclrtFree(inputDeviceAddr); aclrtFree(targetDeviceAddr); aclrtFree(weightDeviceAddr); aclrtFree(posWeightDeviceAddr); aclrtFree(outDeviceAddr); if (workspaceSize 0) { aclrtFree(workspaceAddr); } aclrtDestroyStream(stream); aclrtResetDevice(deviceId); aclFinalize(); return 0; }示例的编译与运行流程可参考 compile_and_run_sample.md。若不需要 weight 与 pos_weight可将对应参数传nullptr接口层会自动以OnesLike构造全 1 张量补齐见下节。六、源码级实现机制6.1 算子定义层op_graph / op_hostGE 算子注册sigmoid_cross_entropy_with_logits_v2_proto.h 中通过REG_OP(SigmoidCrossEntropyWithLogitsV2)注册算子输入为 predict / target / 可选 weight / 可选 pos_weight输出为 loss属性reduction为字符串类型、默认值为 mean数据类型限定为DT_FLOAT16 / DT_FLOAT / DT_BF16。OpDef 实现sigmoid_cross_entropy_with_logits_v2_def.cpp 补充了平台配置AICore 仅配置ascend950与ascend350并开启DynamicCompileStaticFlag、DynamicRankSupportFlag、DynamicShapeSupportFlag即支持动态 shape / 动态 rank 的编译与执行。二进制配置ascend950 的 binary 配置 与 ascend350 的 binary 配置 列出了 5 种输入/输出 dtype 组合float16→float16、float16→float32、float32→float32、bfloat16→bfloat16、bfloat16→float32shape 均为-2动态维度ND 格式。这也解释了为什么 loss 输出 dtype 可以与 predict 不同如 half 输入、float 输出。6.2 aclnn 接口层op_api融合 归约的两段式构图aclnn_binary_cross_entropy_with_logits.cpp 是整个 aclnn 接口的核心实现其计算流程可以概括为参数校验CheckParams()完成空指针、数据类型、reduction、shape 与广播四类检查任一失败即返回对应错误码。空 tensor 处理HandleEmptyTensor()针对空输入做特化——reduction 为 none 时直接返回空 tensormean 时输出填充NAN的张量sum 时输出填充0的张量与 GPU 行为对齐。连续性转换与类型提升BinaryCrossEntropyWithLogitsStub()先对 self / target / weight / posWeight 执行Contiguous转为连续张量再按PromoteType提升公共类型确保混合精度输入如 float16 float32可以正常计算。可选参数补齐posWeight 缺省时在非 regbase 场景下自动用OnesLike构造全 1 张量等价于pos_weight 1不引入额外权重。融合算子调用以reduction none调用l0op::SigmoidCrossEntropyWithLogitsV2一次性完成 Sigmoid、交叉熵与权重计算融合算子见 binary_cross_entropy_with_logits.cpp其中输出类型会在 half 输入时自动提升为 float再通过INFER_SHAPE推导形状。归约若 reduction 为 mean / sumHandleMeanAndSumReductionOut()构造全维度 axes分别调用l0op::ReduceMean/l0op::ReduceSum完成归约none 分支则校验输出 shape 与逐元素损失 shape 一致。结果回写将结果 Cast 到 out 的目标 dtype再通过ViewCopy拷贝到输出张量支持非连续 out。这种“融合算子算逐元素损失 Reduce 算子做归约”的设计既避免了逐元素阶段的多轮 kernel 调度又复用了 Reduce 算子保证 mean/sum 的数值一致性。6.3 内核层op_kernel基于 BroadcastSch 的模板化 DAGsigmoid_cross_entropy_with_logits_v2.cpp 定义了带 4 个模板参数的__global__ __aicore__内核入口schMode调度模式、Reduction、HasWeight、HasPosWeight通过if constexpr在编译期分派到 4 种组合同时有 weight 与 pos_weightSigmoidCEWithLogitsV2HasTwoWeight仅有 weightSigmoidCEWithLogitsV2WeightOnly仅有 pos_weightSigmoidCEWithLogitsV2PosWeightOnly均无SigmoidCEWithLogitsV2。计算 DAG 定义位于 sigmoid_cross_entropy_with_logits_v2_dag.h并通过BroadcastSch调度模板来自atvoss/broadcast/broadcast_sch.h处理广播场景——这与 weight / pos_weight 需 broadcast 到 target 的语义一致。tiling 侧sigmoid_cross_entropy_with_logits_v2_tiling.cpp只注册了reduction none的 tiling keyREDUCTION_MODE_KEY {{none, 0}}再次印证融合内核本身只负责 none 分支的逐元素计算mean / sum 归约完全由 aclnn 层的 Reduce 算子完成。6.4 测试与验证仓库为算子提供了多层测试op_api 单元测试test_aclnn_binary_cross_entropy_with_logits.cpp 覆盖了非法 dtype如ACL_DOUBLE报ACLNN_ERR_PARAM_INVALID、默认 mean 归约、sum 归约等用例并做了精度校验TestPrecision。tiling 单元测试test_sigmoid_cross_entropy_with_logits_v2_tiling.cpp 验证 tiling 参数计算。系统测试STatk_aclnnBinaryCrossEntropyWithLogits.json 与对应的 executor 脚本 组成 ATK 系统测试套件用于真实 NPU 环境下的端到端验证。七、典型应用场景与使用建议应用场景单标签二分类如 CTR 预估、垃圾邮件检测等target 为 0/1 标量或与 logits 同 shape 的张量多标签分类如多标签图像标注logits 每个通道对应一个独立二分类配合pos_weight缓解正负样本不均衡训练框架对接算子与 PyTorchBCEWithLogitsLoss语义对齐适合作为自定义反向训练链路中的前向损失算子。使用建议reduction 与 out shape 配套reduction none时 out 必须与 self 同 shapemean/sum时 out 应为标量 shape维度 0接口层会对 shape 不符给出告警并可能报错。权重广播weight / pos_weight 维度可以比 target 少但必须可广播需要逐类别加权时将 pos_weight 构造成与类别维对齐的 shape 即可。精度与平台匹配Atlas 训练/推理系列产品不支持 BF16half 输入时输出会提升为 float计算中间精度更高最终 Cast 回 out 的 dtype。确定性该接口为确定性实现同一输入多次执行结果一致便于调试与精度对比。空输入场景空 tensor 输入时 mean 返回 NaN、sum 返回 0、none 返回空行为与 GPU 对齐无需在业务侧做额外保护。参考资料与延伸阅读算子说明loss/sigmoid_cross_entropy_with_logits_v2/README.mdaclnn 接口文档loss/sigmoid_cross_entropy_with_logits_v2/docs/aclnnBinaryCrossEntropyWithLogits.md调用示例loss/sigmoid_cross_entropy_with_logits_v2/examples/test_aclnn_sigmoid_cross_entropy_with_logits_v2.cppaclnn 接口实现loss/sigmoid_cross_entropy_with_logits_v2/op_api/aclnn_binary_cross_entropy_with_logits.cpp融合算子封装loss/sigmoid_cross_entropy_with_logits_v2/op_api/binary_cross_entropy_with_logits.cpp算子定义loss/sigmoid_cross_entropy_with_logits_v2/op_host/sigmoid_cross_entropy_with_logits_v2_def.cppGE 原型loss/sigmoid_cross_entropy_with_logits_v2/op_graph/sigmoid_cross_entropy_with_logits_v2_proto.h内核实现loss/sigmoid_cross_entropy_with_logits_v2/op_kernel/sigmoid_cross_entropy_with_logits_v2.cpp两段式接口说明docs/zh/context/two_phase_api.md广播规则docs/zh/context/broadcast_relationship.mdaclnn 返回码docs/zh/context/aclnn_return_code.md编译运行样例说明docs/zh/context/compile_and_run_sample.md【免费下载链接】ops-nn本项目是CANN提供的神经网络类计算算子库实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-nn创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考