add StatelessExponential for multinomial-ops-math-AtomGit

Mdocs/zh/op_list.md+10-0

	@@ -3166,6 +3166,16 @@
3166	<td>AI Core</td>	3166	<td>AI Core</td>
3167	<td>根据随机种子和随机便宜量和保活系数生成随机掩码mask。</td>	3167	<td>根据随机种子和随机便宜量和保活系数生成随机掩码mask。</td>
3168	</tr>	3168	</tr>
		3169	+ <tr>
		3170	+ <td>random</td>
		3171	+ <td><a href="../../random/stateless_exponential/README.md">stateless_exponential</a></td>
		3172	+ <td>√</td>
		3173	+ <td>√</td>
		3174	+ <td>√</td>
		3175	+ <td>×</td>
		3176	+ <td>AI Core</td>
		3177	+ <td>生成指数分布的随机数。</td>
		3178	+ </tr>
3169	<tr>	3179	<tr>
3170	<td>random</td>	3180	<td>random</td>
3171	<td><a href="../../random/stateless_random/README.md">stateless_random</a></td>	3181	<td><a href="../../random/stateless_random/README.md">stateless_random</a></td>

Mrandom/sim_thread_exponential/op_kernel/arch35/sim_thread_exponential_simt.h+7-10

	@@ -26,17 +26,15 @@ struct ExponentialTransform {
26	float lambda_;	26	float lambda_;
27	float halfEpsilon_;	27	float halfEpsilon_;
28		28
29	- __aicore__ ExponentialTransform(float lambda) : lambda_(lambda), halfEpsilon_(1.1920929e-07f / 2.0f)	29	+ __aicore__ ExponentialTransform(float lambda) : lambda_(lambda), halfEpsilon_(1.1920929e-07f / 2.0f) {}
30	- {}
31		30
32	- __simt_callee__ __aicore__ inline void operator()(	31	+ __simt_callee__ __aicore__ inline void operator()(__gm__ volatile T* outputGm, uint64_t li, const uint32_t* results,
33	- __gm__ volatile T* outputGm, uint64_t li, const uint32_t* results, uint32_t iStep,	32	+ uint32_t iStep, [[maybe_unused]] uint32_t unroll = 1)
34	- [[maybe_unused]] uint32_t unroll = 1)
35	{	33	{
36	float u = results[iStep] * RAND_2POW32_INV + RAND_2POW32_INV_HALF;	34	float u = results[iStep] * RAND_2POW32_INV + RAND_2POW32_INV_HALF;
37	float logVal = (u >= 1.0f - halfEpsilon_) ? -halfEpsilon_ : AscendC::Simt::Log(u);	35	float logVal = (u >= 1.0f - halfEpsilon_) ? -halfEpsilon_ : AscendC::Simt::Log(u);
38	float x = -1.0f / lambda_ * logVal;	36	float x = -1.0f / lambda_ * logVal;
39	- outputGm[li] = static_cast<T>(x);	37	+ const_cast<__gm__ T>(&outputGm[li]) = static_cast<T>(x);
40	}	38	}
41	};	39	};
42		40
	@@ -50,9 +48,8 @@ struct ExponentialLauncher {
50	: seed_(seed), lambda_(lambda), baseAddr_(baseAddr)	48	: seed_(seed), lambda_(lambda), baseAddr_(baseAddr)
51	{}	49	{}
52		50
53	- __aicore__ inline void operator()(	51	+ __aicore__ inline void operator()(const ExecutionPolicyKernel& policy, int64_t gmOffset, int64_t kernelOffset,
54	- const ExecutionPolicyKernel& policy, int64_t gmOffset, int64_t kernelOffset, int64_t numel, int64_t grid,	52	+ int64_t numel, int64_t grid, int64_t totalThreads)
55	- int64_t totalThreads)
56	{	53	{
57	__gm__ volatile T* gmPtr = (__gm__ volatile T*)baseAddr_ + gmOffset;	54	__gm__ volatile T* gmPtr = (__gm__ volatile T*)baseAddr_ + gmOffset;
58	ExponentialTransform<T> transform(lambda_);	55	ExponentialTransform<T> transform(lambda_);
	@@ -73,4 +70,4 @@ __aicore__ inline void Process(GM_ADDR self, const RandomUnifiedSimtTilingDataSt
73	}	70	}
74		71
75	} // namespace SimThreadExponential	72	} // namespace SimThreadExponential
76	-#endif	73	+#endif

Arandom/stateless_exponential/CMakeLists.txt+20-0

@@ -0,0 +1,20 @@
	1	+# ----------------------------------------------------------------------------
	2	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+# CANN Open Software License Agreement Version 2.0 (the "License").
	5	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+# See LICENSE in the root of the software repository for the full text of the License.
	9	+# ----------------------------------------------------------------------------
	10	+
	11	+# StatelessExponential only supports ascend950 (arch35 / SIMT). Kernel computation
	12	+# reuses sim_thread_exponential, so it must be listed as a build dependency.
	13	+set(SUPPORT_COMPUTE_UNIT "ascend950")
	14	+set(SUPPORT_TILING_DIR "arch35")
	15	+add_all_modules_sources(OPTYPE stateless_exponential
	16	+ ACLNNTYPE aclnn_exclude
	17	+ COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT}
	18	+ TILING_DIR ${SUPPORT_TILING_DIR}
	19	+ DEPENDENCIES random_common sim_thread_exponential
	20	+ DISABLE_IN_OPP TRUE)

Arandom/stateless_exponential/README.md+84-0

@@ -0,0 +1,84 @@
	1	+# StatelessExponential
	2	+
	C C chenjiao7月29日新增算子需同步刷新op_list.md
	3	+## 产品支持情况
	4	+
	5	+\| 产品 \| 是否支持 \|
	6	+\| :----------------------------------------------------------- \| :------: \|
	7	+\| <term>Ascend 950PR/Ascend 950DT</term> \| √ \|
	8	+\| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> \| × \|
	9	+\| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> \| × \|
	10	+\| <term>Atlas 200I/500 A2 推理产品</term> \| × \|
	11	+\| <term>Atlas 推理系列产品</term> \| × \|
	12	+\| <term>Atlas 训练系列产品</term> \| × \|
	13	+
	14	+## 功能说明
	15	+
	16	+- 算子功能：基于 Philox4x32-10 PRNG 生成服从参数为 `lambd` 的指数分布随机数，将输入张量 `self` 原地填充为服从 `Exp(lambd)` 分布的随机数。本算子是 `aclnnMultinomialTensor` 接口在 Ascend 950 场景下无放回采样路径使用的内部服务算子（生成指数分布随机扰动，供后续 `div + argmax/topk` 完成无放回多项采样）。
	17	+- 计算公式：
	18	+
	19	+ $$
	20	+ x = -\frac{1}{\lambda} \ln(u), \quad u \sim U(0, 1]
	21	+ $$
	22	+
	23	+ 其中 `u` 由 Philox4x32-10 PRNG 生成，`lambd > 0` 为指数分布速率参数。
	24	+
	25	+## 参数说明
	26	+
	27	+<table style="undefined;table-layout: fixed; width: 980px"><colgroup>
	28	+ <col style="width: 100px">
	29	+ <col style="width: 150px">
	30	+ <col style="width: 280px">
	31	+ <col style="width: 330px">
	32	+ <col style="width: 120px">
	33	+ </colgroup>
	34	+ <thead>
	35	+ <tr>
	36	+ <th>参数名</th>
	37	+ <th>输入/输出/属性</th>
	38	+ <th>描述</th>
	39	+ <th>数据类型</th>
	40	+ <th>数据格式</th>
	41	+ </tr></thead>
	42	+ <tbody>
	43	+ <tr>
	44	+ <td>self</td>
	45	+ <td>输入/输出</td>
	46	+ <td>待填充的张量（原地写入），决定随机数生成总量与输出 dtype；≤8 维。</td>
	47	+ <td>FLOAT16、BFLOAT16、FLOAT</td>
	48	+ <td>ND</td>
	49	+ </tr>
	50	+ <tr>
	51	+ <td>seed</td>
	52	+ <td>输入</td>
	53	+ <td>随机数生成器的种子，影响生成的随机数序列。</td>
	54	+ <td>INT64</td>
	55	+ <td>ND</td>
	56	+ </tr>
	57	+ <tr>
	58	+ <td>offset</td>
	59	+ <td>输入</td>
	60	+ <td>随机数生成器的偏移量，必须是 4 的倍数。</td>
	61	+ <td>INT64</td>
	62	+ <td>ND</td>
	63	+ </tr>
	64	+ <tr>
	65	+ <td>lambd</td>
	66	+ <td>属性</td>
	67	+ <td>指数分布速率参数 λ，必须大于 0。</td>
	68	+ <td>FLOAT</td>
	69	+ <td>-</td>
	70	+ </tr>
	71	+ </tbody></table>
	72	+
	73	+## 约束说明
	74	+
	75	+1. `self` 仅支持 FLOAT16/BFLOAT16/FLOAT，ND 格式，维度数 ≤8；非连续输入通过 AutoContiguous 转为连续。
	76	+2. `lambd` 必须 > 0。
	77	+3. `offset` 必须为 4 的倍数。
	78	+4. 仅支持 Ascend 950；不支持入图（op_graph）、不支持 L2 接口、不支持非连续 Tensor。
	79	+
	80	+## 调用说明
	81	+
	82	+\| 调用方式 \| 样例代码 \| 说明 \|
	83	+\| --------- \| ------------------------------------------------------------ \| ------------------------------------------------------------ \|
	84	+\| aclnn接口 \| [test_aclnn_multinomial_tensor](../stateless_sample_multinomial/examples/test_aclnn_multinomial_tensor.cpp) \| 通过[aclnnMultinomialTensor](../stateless_sample_multinomial/docs/aclnnMultinomialTensor.md)接口构建计算流程时，内部调用StatelessExponential服务算子。 \|

Arandom/stateless_exponential/op_api/stateless_exponential.cpp+46-0

@@ -0,0 +1,46 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file stateless_exponential.cpp
	13	+ * \brief Op API (l0) implementation for StatelessExponential.
	14	+ *
	15	+ * seed/offset are passed as INPUT tensors (not ATTRs) so they can flow from tensor-based
	16	+ * callers (e.g. aclnnMultinomialTensor). self is both INPUT and OUTPUT (in-place); the
	17	+ * caller is responsible for allocating the buffer to be filled.
	18	+ */
	19	+#include "stateless_exponential.h"
	20	+#include "op_api/aclnn_check.h"
	21	+#include "opdev/common_types.h"
	22	+#include "opdev/make_op_executor.h"
	23	+#include "opdev/op_def.h"
	24	+#include "opdev/op_dfx.h"
	25	+#include "opdev/op_executor.h"
	26	+#include "opdev/op_log.h"
	27	+
	28	+using namespace op;
	29	+namespace l0op {
	30	+
	31	+OP_TYPE_REGISTER(StatelessExponential);
	32	+
	33	+const aclTensor* StatelessExponential(const aclTensor* self, const aclTensor* seed, const aclTensor* offset,
	34	+ float lambd, aclOpExecutor* executor)
	35	+{
	36	+ L0_DFX(StatelessExponential, self, seed, offset, lambd);
	37	+
	38	+ // In-place: self is both the input and the output. seed/offset are value-dependent
	39	+ // scalar inputs consumed at tiling time.
	40	+ auto ret = ADD_TO_LAUNCHER_LIST_AICORE(StatelessExponential, OP_INPUT(self, seed, offset), OP_OUTPUT(self),
	41	+ OP_ATTR(lambd));
	42	+ CHECK_RET(ret == ACLNN_SUCCESS, nullptr);
	43	+ return self;
	44	+}
	45	+
	46	+} // namespace l0op

Arandom/stateless_exponential/op_api/stateless_exponential.h+42-0

@@ -0,0 +1,42 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file stateless_exponential.h
	13	+ * \brief Op API header for StatelessExponential
	14	+ */
	15	+#ifndef STATELESS_EXPONENTIAL_OP_API_H
	16	+#define STATELESS_EXPONENTIAL_OP_API_H
	17	+
	18	+#include "opdev/op_executor.h"
	19	+
	20	+namespace l0op {
	21	+
	22	+/**
	23	+ * @brief Fill `self` in-place with Exp(lambda) random numbers (Philox4x32-10).
	24	+ *
	25	+ * This is an in-place operator: the random numbers are written into the tensor passed as
	26	+ * `self`, and `self` itself is returned. To avoid overwriting an existing weights tensor,
	27	+ * the caller should pre-allocate a fresh tensor (e.g. executor->AllocTensor with self's
	28	+ * shape/dtype) and pass that as `self`.
	29	+ *
	30	+ * @param self Tensor to be filled in-place with Exp(lambda) samples (FP16/BF16/FP32, ND)
	31	+ * @param seed Seed tensor (INT64 scalar)
	32	+ * @param offset Offset tensor (INT64 scalar, must be a multiple of 4)
	33	+ * @param lambd Rate parameter lambda, must be > 0
	34	+ * @param executor Op executor
	35	+ * @return self (same tensor, now holding exponential random numbers)
	36	+ */
	37	+const aclTensor* StatelessExponential(const aclTensor* self, const aclTensor* seed, const aclTensor* offset,
	38	+ float lambd, aclOpExecutor* executor);
	39	+
	40	+} // namespace l0op
	41	+
	42	+#endif // STATELESS_EXPONENTIAL_OP_API_H

Arandom/stateless_exponential/op_host/arch35/stateless_exponential_tiling_arch35.cpp+145-0

@@ -0,0 +1,145 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file stateless_exponential_tiling_arch35.cpp
	13	+ * \brief Tiling implementation for StatelessExponential operator (ascend950 / SIMT).
	14	+ *
	15	+ * Hybrid of two references:
	16	+ * - sim_thread_exponential_tiling_arch35.cpp : exponential-specific config (lambd>0 check,
	17	+ * enableSplitBlocks, unrollFactor, prob=lambd, TilingKey routed by output dtype);
	18	+ * - stateless_normal_tiling_arch35.cpp : seed/offset are tensor inputs that must NOT
	19	+ * be read at tiling time (an offset that is a device-computed intermediate has no host
	20	+ * value here). Tiling stores 0 placeholders; the kernel reads real seed/offset from GM.
	21	+ */
	22	+#include "stateless_exponential_tiling_arch35.h"
	23	+#include <string>
	24	+#include "log/log.h"
	25	+#include "platform/platform_ascendc.h"
	26	+#include "register/op_def_registry.h"
	27	+#include "op_host/math_tiling_templates_registry.h"
	28	+#include "util/math_util.h"
	29	+#include "../../../random_common/op_host/arch35/random_tiling_base.h"
	30	+
	31	+namespace optiling {
	32	+static constexpr uint16_t INPUT_IDX_SELF = 0;
	33	+static constexpr uint16_t INPUT_IDX_SEED = 1;
	34	+static constexpr uint16_t INPUT_IDX_OFFSET = 2;
	35	+static constexpr uint16_t OUTPUT_IDX_SELF = 0;
	36	+
	37	+static constexpr int64_t DCACHE_SIZE = 128 * 1024;
	38	+static constexpr uint32_t NUM_FOUR = 4;
	39	+
	40	+static constexpr uint64_t TILING_KEY_FP16 = 1;
	41	+static constexpr uint64_t TILING_KEY_BF16 = 2;
	42	+static constexpr uint64_t TILING_KEY_FP32 = 3;
	43	+
	44	+OpTilingConfig StatelessExponentialTilingSimt::BuildOpConfig()
	45	+{
	46	+ OpTilingConfig config;
	47	+
	48	+ // self: FP16/BF16/FP32, any shape; seed/offset: INT64 scalar (shapeSize == 1).
	49	+ config.inputCheckRules = {{INPUT_IDX_SELF, {{ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT}, -1, {}, nullptr}},
	50	+ {INPUT_IDX_SEED, {{ge::DT_INT64}, 1, {}, nullptr}},
	51	+ {INPUT_IDX_OFFSET, {{ge::DT_INT64}, 1, {}, nullptr}}};
	52	+ config.outputCheckRules = {{OUTPUT_IDX_SELF, {{ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT}, -1, {}, nullptr}}};
	53	+
	54	+ // Output size == total elements of self (lambd is the only attr; count is derived).
	55	+ config.getOutputSize = [](gert::TilingContext* ctx, int64_t& size) {
	56	+ auto selfShape = ctx->GetInputShape(INPUT_IDX_SELF);
	57	+ OP_CHECK_NULL_WITH_CONTEXT(ctx, selfShape);
	58	+ size = selfShape->GetStorageShape().GetShapeSize();
	59	+ return ge::GRAPH_SUCCESS;
	60	+ };
	61	+
	62	+ // Avoid reading device-computed seed/offset during GetWorkspaceSize to prevent
	63	+ // host-side page faults. Store 0 placeholders; kernel reads real values from GM at runtime.
	64	+ config.getSeedAndOffset = [](gert::TilingContext* /ctx/, int64_t& seed, int64_t& offset) {
	65	+ seed = 0;
	66	+ offset = 0;
	67	+ return ge::GRAPH_SUCCESS;
	68	+ };
	69	+
	70	+ // lambd is attr index 0 (the only attribute) and must be > 0.
	71	+ config.attrCheckRules = {
	72	+ {0,
	73	+ [](gert::TilingContext* ctx) -> bool {
	74	+ auto attrs = ctx->GetAttrs();
	75	+ if (attrs == nullptr) {
	76	+ return false;
	77	+ }
	78	+ const auto* lambdAttr = attrs->GetAttrPointer<float>(0);
	79	+ if (lambdAttr == nullptr) {
	80	+ return false;
	81	+ }
	82	+ if (*lambdAttr <= 0.0f) {
	X X xuejinghui7月29日如果用户输入Nan，这里返回false，继续向下执行，是否符合预期？ zhangxiyan7 7月29日评论： 7月29日评论：
	83	+ std::string valueStr = std::to_string(*lambdAttr);
	84	+ std::string reasonMsg = "lambd must be greater than 0";
	85	+ OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(ctx->GetNodeName(), "attr lambd", valueStr.c_str(),
	86	+ reasonMsg.c_str());
	87	+ return false;
	88	+ }
	89	+ return true;
	90	+ }},
	91	+ };
	92	+
	93	+ config.kernelMode = RandomKernelMode::SIMT;
	94	+ config.DcacheSize = DCACHE_SIZE;
	95	+ config.isNeedSyncAll = false;
	96	+ config.unrollFactor = NUM_FOUR;
	97	+ config.enableSplitBlocks = true;
	98	+ return config;
	99	+}
	100	+
	101	+ge::graphStatus StatelessExponentialTilingSimt::UniqueProcess()
	102	+{
	103	+ // Store lambd into the generic `prob` field; the kernel reads tilingData->prob as lambda.
	104	+ auto attrs = context_->GetAttrs();
	105	+ OP_CHECK_NULL_WITH_CONTEXT(context_, attrs);
	106	+ const auto* lambdAttr = attrs->GetAttrPointer<float>(0);
	107	+ OP_CHECK_NULL_WITH_CONTEXT(context_, lambdAttr);
	108	+ simtTilingData_.prob = *lambdAttr;
	109	+
	110	+ // Route TilingKey by output dtype: 1=FP16, 2=BF16, 3=FP32 (matches kernel dispatch).
	111	+ auto outputDesc = context_->GetOutputDesc(OUTPUT_IDX_SELF);
	112	+ OP_CHECK_NULL_WITH_CONTEXT(context_, outputDesc);
	113	+ auto outputDtype = outputDesc->GetDataType();
	114	+ if (outputDtype == ge::DT_FLOAT16) {
	115	+ tilingKey_ = TILING_KEY_FP16;
	116	+ } else if (outputDtype == ge::DT_BF16) {
	117	+ tilingKey_ = TILING_KEY_BF16;
	118	+ } else {
	119	+ tilingKey_ = TILING_KEY_FP32;
	120	+ }
	121	+
	122	+ return ge::GRAPH_SUCCESS;
	123	+}
	124	+
	125	+} // namespace optiling
	126	+
	127	+namespace optiling {
	128	+
	129	+static ge::graphStatus Tiling4StatelessExponentialArch35(gert::TilingContext* context)
	130	+{
	131	+ OP_LOGD(context->GetNodeName(), "Tiling4StatelessExponential running tiling.");
	132	+ StatelessExponentialTilingSimt tilingObj(context);
	133	+ return tilingObj.DoTiling();
	134	+}
	135	+
	136	+ge::graphStatus TilingPrepare4StatelessExponentialArch35(gert::TilingParseContext* context)
	137	+{
	138	+ return RandomTilingParseArch35(context, "StatelessExponential");
	139	+}
	140	+
	141	+IMPL_OP_OPTILING(StatelessExponential)
	142	+ .Tiling(Tiling4StatelessExponentialArch35)
	143	+ .TilingParse<RandomOperatorCompileInfo>(TilingPrepare4StatelessExponentialArch35)
	144	+ .TilingInputsDataDependency({INPUT_IDX_SEED, INPUT_IDX_OFFSET});
	145	+} // namespace optiling

Arandom/stateless_exponential/op_host/arch35/stateless_exponential_tiling_arch35.h+38-0

@@ -0,0 +1,38 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file stateless_exponential_tiling_arch35.h
	13	+ * \brief Tiling class declaration for StatelessExponential operator (ascend950 / SIMT).
	14	+ */
	15	+#ifndef STATELESS_EXPONENTIAL_TILING_ARCH35_H
	16	+#define STATELESS_EXPONENTIAL_TILING_ARCH35_H
	17	+
	18	+#include "register/op_def_registry.h"
	19	+#include "register/op_impl_registry.h"
	20	+#include "../../../random_common/op_host/arch35/random_tiling_arch35.h"
	21	+
	22	+namespace optiling {
	23	+
	24	+// SIMT tiling framework: override BuildOpConfig() and UniqueProcess() only.
	25	+class StatelessExponentialTilingSimt : public RandomTilingArch35 {
	26	+public:
	27	+ explicit StatelessExponentialTilingSimt(gert::TilingContext* context) : RandomTilingArch35(context, BuildOpConfig())
	28	+ {}
	29	+
	30	+protected:
	31	+ ge::graphStatus UniqueProcess() override;
	32	+
	33	+private:
	34	+ static OpTilingConfig BuildOpConfig();
	35	+};
	36	+
	37	+} // namespace optiling
	38	+#endif // STATELESS_EXPONENTIAL_TILING_ARCH35_H

Arandom/stateless_exponential/op_host/config/ascend950/stateless_exponential_simplified_key.ini+13-0

@@ -0,0 +1,13 @@
	1	+; 该文件主要影响 opc 工具编译二进制kernel时， --simplified_key_mode 选项中填写的值，格式如下所示：
	2	+; [某算子]
	3	+; default=xx
	4	+; ascendxx=xx
	5	+; 其中，default为默认mode，ascendxx为可选mode，如果不同芯片有差异化要求时，需要配置；
	6	+; 1)如果没有配置：非ascendC算子继续按空处理，即opc编译命令中不添加 --simplified_key_mode 选项，AscendC算子按照 simplified_key_mode=0 处理
	7	+; 2)如果仅有default配置：各个版本按default配置
	8	+; 3)如果仅有某些平台的配置，没有default配置：对应平台的按照配置的值传递，非对应平台的：非AscendC算子继续按空处理，AscendC算子按照 simplified_key_mode=0 处理
	9	+; 4)如果default配置和平台配置都有：对应平台的使用平台的配置，非对应的平台的以default值配置。
	10	+; 5)对于自定义simplified key的情况，需要在binary_simplified_key_mode.ini 文件中显式配置为None，不传入 --simplified_key_mode 选项，由opc工具和FE框架自行判断使用何种模式
	11	+; 6)是否是AscendC算子，由 ops/built-in/tbe/op_info_cfg/parser/ascendc_config.json 中配置的算子名字和对于的平台决定
	12	+[StatelessExponential]
	13	+default=0

Arandom/stateless_exponential/op_host/stateless_exponential_def.cpp+56-0

@@ -0,0 +1,56 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file stateless_exponential_def.cpp
	13	+ * \brief Op definition for StatelessExponential (in-place Exp(lambda) via Philox4x32-10).
	14	+ */
	15	+
	16	+#include "register/op_def_registry.h"
	17	+
	18	+namespace ops {
	19	+class StatelessExponential : public OpDef {
	20	+public:
	21	+ explicit StatelessExponential(const char* name) : OpDef(name)
	22	+ {
	23	+ // self is both INPUT and OUTPUT (in-place). seed/offset are scalar tensors whose
	24	+ // values are consumed at tiling time via TilingInputsDataDependency.
	25	+ this->Input("self")
	26	+ .ParamType(REQUIRED)
	27	+ .DataType({ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT})
	28	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	29	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	30	+ .AutoContiguous();
	31	+ this->Input("seed")
	32	+ .ParamType(REQUIRED)
	33	+ .DataType({ge::DT_INT64, ge::DT_INT64, ge::DT_INT64})
	34	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	35	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	36	+ .ValueDepend(OPTIONAL);
	37	+ this->Input("offset")
	38	+ .ParamType(REQUIRED)
	39	+ .DataType({ge::DT_INT64, ge::DT_INT64, ge::DT_INT64})
	40	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	41	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	42	+ .ValueDepend(OPTIONAL);
	43	+ this->Output("self")
	44	+ .ParamType(REQUIRED)
	45	+ .DataType({ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT})
	46	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	47	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
	48	+
	49	+ this->Attr("lambd").Float(1.0);
	50	+
	51	+ this->AICore().AddConfig("ascend950");
	52	+ }
	53	+};
	54	+
	55	+OP_ADD(StatelessExponential);
	56	+} // namespace ops

Arandom/stateless_exponential/op_kernel/stateless_exponential.cpp+88-0

@@ -0,0 +1,88 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file stateless_exponential.cpp
	13	+ * \brief Kernel entry point for StatelessExponential.
	14	+ *
	15	+ * The per-element computation (Philox4x32-10 -> uniform -> -log/lambda -> cast) is identical
	16	+ * to SimThreadExponential, so this file reuses SimThreadExponential::ExponentialTransform<T>.
	17	+ *
	18	+ * Difference from SimThreadExponential: seed/offset are tensor INPUTS (not attrs). Callers
	19	+ * such as aclnnMultinomialTensor pass an offset that is a device-computed intermediate
	20	+ * (l0op::Add output) with no host value at tiling time, so the tiling layer cannot read
	21	+ * their values (host-gathering an unmapped device page faults). Following the stateless_normal
	22	+ * reference, the kernel reads the real seed/offset directly from GM at runtime and applies
	23	+ * them here. The tiling split-block kernelOffset carries only the per-block counter increments
	24	+ * (base offset 0), so the real offset from GM is added on top of kernelOffset (identical scheme
	25	+ * to stateless_normal). Values are the same as the original tiling-fed design intended; only
	26	+ * their source moves from a crashing host-gather to a valid device read.
	27	+ */
	28	+#include "../sim_thread_exponential/arch35/sim_thread_exponential_simt.h"
	29	+
	30	+namespace StatelessExponential {
	31	+using namespace AscendC;
	32	+using namespace RandomKernelBase;
	33	+
	34	+template <typename T>
	35	+struct StatelessExponentialLauncher {
	36	+ int64_t seed_;
	37	+ int64_t realOffset_;
	38	+ float lambda_;
	39	+ GM_ADDR baseAddr_;
	40	+
	41	+ __aicore__ StatelessExponentialLauncher(int64_t seed, int64_t realOffset, float lambda, GM_ADDR baseAddr)
	42	+ : seed_(seed), realOffset_(realOffset), lambda_(lambda), baseAddr_(baseAddr)
	43	+ {}
	44	+
	45	+ __aicore__ inline void operator()(const ExecutionPolicyKernel& policy, int64_t gmOffset, int64_t kernelOffset,
	46	+ int64_t numel, [[maybe_unused]] int64_t grid, int64_t totalThreads)
	47	+ {
	48	+ __gm__ volatile T* gmPtr = (__gm__ volatile T*)baseAddr_ + gmOffset;
	49	+ SimThreadExponential::ExponentialTransform<T> transform(lambda_);
	50	+ AscendC::Simt::VF_CALL<PhiloxSimtKernelDiscontinuous<T, SimThreadExponential::ExponentialTransform<T>>>(
	51	+ AscendC::Simt::Dim3(DEFAULT_SIMT_THREAD_NUM), gmPtr, realOffset_ + kernelOffset, seed_, numel, policy.magic,
	52	+ policy.shift, totalThreads, transform);
	53	+ }
	54	+};
	55	+
	56	+// self is both input and output (in-place). seed/offset are scalar INT64 GM tensors.
	57	+template <typename T>
	58	+__aicore__ inline void Process(GM_ADDR self, GM_ADDR seed, GM_ADDR offset,
	59	+ const RandomUnifiedSimtTilingDataStruct* __restrict tilingData)
	60	+{
	61	+ if (AscendC::GetBlockIdx() >= static_cast<uint32_t>(tilingData->usedCoreNum))
	62	+ return;
	63	+
	64	+ // Read real seed/offset from GM (tiling filled 0 placeholders; the Add producing offset
	65	+ // has already executed by kernel launch time).
	66	+ int64_t realSeed = (reinterpret_cast<__gm__ int64_t>(seed));
	67	+ int64_t realOffset = (reinterpret_cast<__gm__ int64_t>(offset));
	68	+
	69	+ StatelessExponentialLauncher<T> launcher(realSeed, realOffset, tilingData->prob, self);
	70	+ ProcessWithSplitBlocks(tilingData, launcher);
	71	+}
	72	+} // namespace StatelessExponential
	73	+
	74	+extern "C" __global__ __aicore__ void stateless_exponential(GM_ADDR self, GM_ADDR seed, GM_ADDR offset, GM_ADDR selfOut,
	75	+ GM_ADDR workspace, GM_ADDR tiling)
	76	+{
	77	+ REGISTER_TILING_DEFAULT(RandomUnifiedSimtTilingDataStruct);
	78	+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIV_1_0);
	79	+ GET_TILING_DATA_WITH_STRUCT(RandomUnifiedSimtTilingDataStruct, tilingData, tiling);
	80	+
	81	+ if (TILING_KEY_IS(3)) {
	82	+ StatelessExponential::Process<float>(self, seed, offset, &tilingData);
	83	+ } else if (TILING_KEY_IS(1)) {
	84	+ StatelessExponential::Process<half>(self, seed, offset, &tilingData);
	85	+ } else if (TILING_KEY_IS(2)) {
	86	+ StatelessExponential::Process<bfloat16_t>(self, seed, offset, &tilingData);
	87	+ }
	88	+}

Arandom/stateless_exponential/tests/ut/op_host/arch35/test_stateless_exponential_tiling_arch35.cpp+156-0

@@ -0,0 +1,156 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file test_stateless_exponential_tiling_arch35.cpp
	13	+ * \brief StatelessExponential tiling UT (ascend950 / SIMT)
	14	+ */
	15	+
	16	+#include <gtest/gtest.h>
	17	+#include <iostream>
	18	+#include <vector>
	19	+#include "tiling_case_executor.h"
	20	+#include "../../../../op_host/arch35/stateless_exponential_tiling_arch35.h"
	21	+
	22	+class StatelessExponentialTilingTest : public testing::Test {
	23	+protected:
	24	+ static void SetUpTestCase() { std::cout << "StatelessExponentialTilingTest SetUp" << std::endl; }
	25	+
	26	+ static void TearDownTestCase() { std::cout << "StatelessExponentialTilingTest TearDown" << std::endl; }
	27	+};
	28	+
	29	+// FP32 -> TilingKey = 3
	30	+TEST_F(StatelessExponentialTilingTest, one_dim_float)
	31	+{
	32	+ optiling::RandomOperatorCompileInfo compileInfo = {64, 196608};
	33	+ int64_t seedValue = 5;
	34	+ int64_t offsetValue = 4;
	35	+
	36	+ gert::TilingContextPara tilingContextPara("StatelessExponential",
	37	+ {
	38	+ {{{256}, {256}}, ge::DT_FLOAT, ge::FORMAT_ND},
	39	+ {{{1}, {1}}, ge::DT_INT64, ge::FORMAT_ND, true, &seedValue},
	40	+ {{{1}, {1}}, ge::DT_INT64, ge::FORMAT_ND, true, &offsetValue},
	41	+ },
	42	+ {
	43	+ {{{256}, {256}}, ge::DT_FLOAT, ge::FORMAT_ND},
	44	+ },
	45	+ {
	46	+ {"lambd", Ops::Math::AnyValue::CreateFrom<float>(1.0f)},
	47	+ },
	48	+ &compileInfo);
	49	+
	50	+ uint64_t expectTilingKey = 3;
	51	+ std::vector<size_t> expectWorkspaces = {0};
	52	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
	53	+}
	54	+
	55	+// FP16 -> TilingKey = 1
	56	+TEST_F(StatelessExponentialTilingTest, one_dim_float16)
	57	+{
	58	+ optiling::RandomOperatorCompileInfo compileInfo = {64, 196608};
	59	+ int64_t seedValue = 7;
	60	+ int64_t offsetValue = 8;
	61	+
	62	+ gert::TilingContextPara tilingContextPara("StatelessExponential",
	63	+ {
	64	+ {{{300}, {300}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	65	+ {{{1}, {1}}, ge::DT_INT64, ge::FORMAT_ND, true, &seedValue},
	66	+ {{{1}, {1}}, ge::DT_INT64, ge::FORMAT_ND, true, &offsetValue},
	67	+ },
	68	+ {
	69	+ {{{300}, {300}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	70	+ },
	71	+ {
	72	+ {"lambd", Ops::Math::AnyValue::CreateFrom<float>(1.0f)},
	73	+ },
	74	+ &compileInfo);
	75	+
	76	+ uint64_t expectTilingKey = 1;
	77	+ std::vector<size_t> expectWorkspaces = {0};
	78	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
	79	+}
	80	+
	81	+// BF16 -> TilingKey = 2
	82	+TEST_F(StatelessExponentialTilingTest, one_dim_bfloat16)
	83	+{
	84	+ optiling::RandomOperatorCompileInfo compileInfo = {64, 196608};
	85	+ int64_t seedValue = 99;
	86	+ int64_t offsetValue = 12;
	87	+
	88	+ gert::TilingContextPara tilingContextPara("StatelessExponential",
	89	+ {
	90	+ {{{1024}, {1024}}, ge::DT_BF16, ge::FORMAT_ND},
	91	+ {{{1}, {1}}, ge::DT_INT64, ge::FORMAT_ND, true, &seedValue},
	92	+ {{{1}, {1}}, ge::DT_INT64, ge::FORMAT_ND, true, &offsetValue},
	93	+ },
	94	+ {
	95	+ {{{1024}, {1024}}, ge::DT_BF16, ge::FORMAT_ND},
	96	+ },
	97	+ {
	98	+ {"lambd", Ops::Math::AnyValue::CreateFrom<float>(2.0f)},
	99	+ },
	100	+ &compileInfo);
	101	+
	102	+ uint64_t expectTilingKey = 2;
	103	+ std::vector<size_t> expectWorkspaces = {0};
	104	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
	105	+}
	106	+
	107	+// 2-D shape, FP32 -> TilingKey = 3, exercises multi-dim coalescing in splitBlocks
	108	+TEST_F(StatelessExponentialTilingTest, two_dim_float)
	109	+{
	110	+ optiling::RandomOperatorCompileInfo compileInfo = {64, 196608};
	111	+ int64_t seedValue = 1;
	112	+ int64_t offsetValue = 16;
	113	+
	114	+ gert::TilingContextPara tilingContextPara("StatelessExponential",
	115	+ {
	116	+ {{{4, 64}, {4, 64}}, ge::DT_FLOAT, ge::FORMAT_ND},
	117	+ {{{1}, {1}}, ge::DT_INT64, ge::FORMAT_ND, true, &seedValue},
	118	+ {{{1}, {1}}, ge::DT_INT64, ge::FORMAT_ND, true, &offsetValue},
	119	+ },
	120	+ {
	121	+ {{{4, 64}, {4, 64}}, ge::DT_FLOAT, ge::FORMAT_ND},
	122	+ },
	123	+ {
	124	+ {"lambd", Ops::Math::AnyValue::CreateFrom<float>(0.5f)},
	125	+ },
	126	+ &compileInfo);
	127	+
	128	+ uint64_t expectTilingKey = 3;
	129	+ std::vector<size_t> expectWorkspaces = {0};
	130	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
	131	+}
	132	+
	133	+// lambd == 0 -> attrCheckRules fails -> GRAPH_FAILED
	134	+TEST_F(StatelessExponentialTilingTest, lambd_zero_failed)
	135	+{
	136	+ optiling::RandomOperatorCompileInfo compileInfo = {64, 196608};
	137	+ int64_t seedValue = 5;
	138	+ int64_t offsetValue = 4;
	139	+
	140	+ gert::TilingContextPara tilingContextPara("StatelessExponential",
	141	+ {
	142	+ {{{256}, {256}}, ge::DT_FLOAT, ge::FORMAT_ND},
	143	+ {{{1}, {1}}, ge::DT_INT64, ge::FORMAT_ND, true, &seedValue},
	144	+ {{{1}, {1}}, ge::DT_INT64, ge::FORMAT_ND, true, &offsetValue},
	145	+ },
	146	+ {
	147	+ {{{256}, {256}}, ge::DT_FLOAT, ge::FORMAT_ND},
	148	+ },
	149	+ {
	150	+ {"lambd", Ops::Math::AnyValue::CreateFrom<float>(0.0f)},
	151	+ },
	152	+ &compileInfo);
	153	+
	154	+ std::vector<size_t> expectWorkspaces = {0};
	155	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, 0, expectWorkspaces);
	156	+}

Arandom/stateless_exponential/tests/ut/op_kernel/CMakeLists.txt+33-0

@@ -0,0 +1,33 @@
	1	+# ----------------------------------------------------------------------------
	2	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+# CANN Open Software License Agreement Version 2.0 (the "License").
	5	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+# See LICENSE in the root of the software repository for the full text of the License.
	9	+# ----------------------------------------------------------------------------
	10	+
	11	+if(UT_TEST_ALL OR OP_KERNEL_UT)
	12	+ set(KERNEL_STAGING_DIR ${CMAKE_CURRENT_BINARY_DIR}/kernel_dep_staging)
	13	+ file(MAKE_DIRECTORY ${KERNEL_STAGING_DIR}/stateless_exponential/arch35)
	14	+
	15	+ # stateless_exponential kernel reuses sim_thread_exponential_simt.h, included as
	16	+ # "../sim_thread_exponential/arch35/sim_thread_exponential_simt.h" (build/staging layout).
	17	+ execute_process(COMMAND ${CMAKE_COMMAND} -E create_symlink
	18	+ ${PROJECT_SOURCE_DIR}/random/sim_thread_exponential/op_kernel
	19	+ ${KERNEL_STAGING_DIR}/sim_thread_exponential)
	20	+ # sim_thread_exponential_simt.h internally includes "../../random_common/...".
	21	+ execute_process(COMMAND ${CMAKE_COMMAND} -E create_symlink
	22	+ ${PROJECT_SOURCE_DIR}/random/random_common/op_kernel
	23	+ ${KERNEL_STAGING_DIR}/random_common)
	24	+
	25	+ set(stateless_exponential_tiling_files
	26	+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/arch35/stateless_exponential_tiling_arch35.cpp
	27	+ ${PROJECT_SOURCE_DIR}/random/random_common/op_host/arch35/random_tiling_arch35.cpp)
	28	+ AddOpTestCase(
	29	+ stateless_exponential
	30	+ "ascend950"
	31	+ "-DTestUtDefaultTilingStruct=RandomUnifiedSimtTilingDataStruct -I${KERNEL_STAGING_DIR}/stateless_exponential -I${KERNEL_STAGING_DIR}/stateless_exponential/arch35"
	32	+ "${stateless_exponential_tiling_files}")
	33	+endif()

Arandom/stateless_exponential/tests/ut/op_kernel/stateless_exponential_tiling.h+57-0

@@ -0,0 +1,57 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+#ifndef STATELESS_EXPONENTIAL_TILING_H
	12	+#define STATELESS_EXPONENTIAL_TILING_H
	13	+
	14	+#include <cstdint>
	15	+#include <cstring>
	16	+
	17	+#include "../../../../random_common/op_kernel/arch35/random_unified_tiling_data_arch35.h"
	18	+#include "kernel_tiling/kernel_tiling.h"
	19	+
	20	+static inline unsigned long long __mul_i32toi64(unsigned int lhs, unsigned int rhs)
	21	+{
	22	+ return static_cast<unsigned long long>(lhs) * rhs;
	23	+}
	24	+
	25	+#define __aicore__
	26	+#ifdef __NPU_TILING__
	27	+inline[aicore] void InitTilingData(const __gm__ uint8_t* tiling, RandomUnifiedSimtTilingDataStruct* constData)
	28	+{
	29	+ const __gm__ uint32_t* src = reinterpret_cast<const __gm__ uint32_t*>(tiling);
	30	+ uint32_t* dst = reinterpret_cast<uint32_t*>(constData);
	31	+ for (size_t i = 0; i < sizeof(RandomUnifiedSimtTilingDataStruct) / sizeof(uint32_t); ++i) {
	32	+ (dst + i) = (src + i);
	33	+ }
	34	+}
	35	+#else
	36	+inline void InitTilingData(uint8_t* tiling, RandomUnifiedSimtTilingDataStruct* constData)
	37	+{
	38	+ std::memcpy(constData, tiling, sizeof(RandomUnifiedSimtTilingDataStruct));
	39	+}
	40	+#endif // __NPU_TILING__
	41	+
	42	+#define CONVERT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer) \
	43	+ __ubuf__ tilingStruct* tilingDataPointer = reinterpret_cast<__ubuf__ tilingStruct*>( \
	44	+ reinterpret_cast<__ubuf__ uint8_t*>(tilingPointer));
	45	+
	46	+#define INIT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer) \
	47	+ CONVERT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer);
	48	+
	49	+#define GET_TILING_DATA_WITH_STRUCT(tilingStruct, tilingData, tilingArg) \
	50	+ tilingStruct tilingData; \
	51	+ InitTilingData(tilingArg, &tilingData)
	52	+
	53	+#define GET_TILING_DATA(tilingData, tilingArg) \
	54	+ RandomUnifiedSimtTilingDataStruct tilingData; \
	55	+ InitTilingData(tilingArg, &tilingData)
	56	+
	57	+#endif // STATELESS_EXPONENTIAL_TILING_H

Arandom/stateless_exponential/tests/ut/op_kernel/test_stateless_exponential.cpp+180-0

@@ -0,0 +1,180 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file test_stateless_exponential.cpp
	13	+ * \brief StatelessExponential kernel UT (CPU simulation via ICPU_RUN_KF).
	14	+ */
	15	+
	16	+#include <cstdint>
	17	+#include <cmath>
	18	+#include <cstring>
	19	+#include "gtest/gtest.h"
	20	+#include "tikicpulib.h"
	21	+#include "../../../../random_common/op_kernel/arch35/random_unified_tiling_data_arch35.h"
	22	+
	23	+extern "C" __global__ __aicore__ void stateless_exponential(GM_ADDR self, GM_ADDR seed, GM_ADDR offset, GM_ADDR selfOut,
	24	+ GM_ADDR workspace, GM_ADDR tiling);
	25	+
	26	+namespace {
	27	+constexpr uint32_t kNumBlocks = 1;
	28	+constexpr uint64_t kTilingKeyFp32 = 3;
	29	+constexpr uint64_t kTilingKeyFp16 = 1;
	30	+constexpr uint64_t kTilingKeyBf16 = 2;
	31	+constexpr int64_t kElementCount = 256;
	32	+constexpr int64_t kSeed = 42;
	33	+constexpr int64_t kOffset = 0;
	34	+// SIMT_THREAD_GROUP_SIZE used by host-side execution policy; keep consistent so the
	35	+// CPU-simulated kernel maps all elements onto the launched threads.
	36	+constexpr uint32_t kSimtThreadGroupSize = 256;
	37	+
	38	+inline size_t Align32(size_t size) { return (size + 31U) / 32U * 32U; }
	39	+
	40	+// The kernel writes via ProcessWithSplitBlocks, which iterates splitBlocks[0..splitBlockCount).
	41	+// splitBlockCount must be >= 1, otherwise no element is written.
	42	+void FillTiling(RandomUnifiedSimtTilingDataStruct* tilingData, int64_t seed, int64_t offset)
	43	+{
	44	+ std::memset(tilingData, 0, sizeof(RandomUnifiedSimtTilingDataStruct));
	45	+ tilingData->usedCoreNum = kNumBlocks;
	46	+ tilingData->outputSize = kElementCount;
	47	+ tilingData->seed = seed;
	48	+ tilingData->offset = offset;
	49	+ tilingData->prob = 1.0f; // lambd = 1.0
	50	+ tilingData->splitBlockCount = 1;
	51	+ tilingData->splitBlocks[0].numel = kElementCount;
	52	+ tilingData->splitBlocks[0].gmOffset = 0;
	53	+ tilingData->splitBlocks[0].grid = 1;
	54	+ tilingData->splitBlocks[0].totalThreads = kSimtThreadGroupSize;
	55	+ tilingData->splitBlocks[0].kernelOffset = offset;
	56	+}
	57	+} // namespace
	58	+
	59	+class StatelessExponentialKernelTest : public testing::Test {};
	60	+
	61	+TEST_F(StatelessExponentialKernelTest, smoke_float32)
	62	+{
	63	+ auto* self = static_cast<uint8_t>(AscendC::GmAlloc(Align32(kElementCount sizeof(float))));
	64	+ auto* seed = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int64_t))));
	65	+ auto* offset = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int64_t))));
	66	+ auto* workspace = static_cast<uint8_t>(AscendC::GmAlloc(Align32(16 1024 * 1024)));
	67	+ auto* tiling = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(RandomUnifiedSimtTilingDataStruct))));
	68	+
	69	+ std::memset(self, 0, kElementCount * sizeof(float));
	70	+ reinterpret_cast<int64_t>(seed) = kSeed;
	71	+ reinterpret_cast<int64_t>(offset) = kOffset;
	72	+ FillTiling(reinterpret_cast<RandomUnifiedSimtTilingDataStruct*>(tiling), kSeed, kOffset);
	73	+
	74	+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
	75	+ ICPU_SET_TILING_KEY(kTilingKeyFp32);
	76	+ // In-place: self is both the input and output tensor.
	77	+ ICPU_RUN_KF(stateless_exponential, kNumBlocks, self, seed, offset, self, workspace, tiling);
	78	+
	79	+ auto* out = reinterpret_cast<float*>(self);
	80	+ for (int64_t i = 0; i < kElementCount; ++i) {
	81	+ EXPECT_TRUE(std::isfinite(out[i])) << "Element " << i << " not finite: " << out[i];
	82	+ EXPECT_GT(out[i], 0.0f) << "Element " << i << " <= 0: " << out[i];
	83	+ }
	84	+
	85	+ AscendC::GmFree(self);
	86	+ AscendC::GmFree(seed);
	87	+ AscendC::GmFree(offset);
	88	+ AscendC::GmFree(workspace);
	89	+ AscendC::GmFree(tiling);
	90	+}
	91	+
	92	+TEST_F(StatelessExponentialKernelTest, smoke_float16)
	93	+{
	94	+ auto* self = static_cast<uint8_t>(AscendC::GmAlloc(Align32(kElementCount sizeof(half))));
	95	+ auto* seed = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int64_t))));
	96	+ auto* offset = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int64_t))));
	97	+ auto* workspace = static_cast<uint8_t>(AscendC::GmAlloc(Align32(16 1024 * 1024)));
	98	+ auto* tiling = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(RandomUnifiedSimtTilingDataStruct))));
	99	+
	100	+ std::memset(self, 0, kElementCount * sizeof(half));
	101	+ reinterpret_cast<int64_t>(seed) = kSeed;
	102	+ reinterpret_cast<int64_t>(offset) = kOffset;
	103	+ FillTiling(reinterpret_cast<RandomUnifiedSimtTilingDataStruct*>(tiling), kSeed, kOffset);
	104	+
	105	+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
	106	+ ICPU_SET_TILING_KEY(kTilingKeyFp16);
	107	+ ICPU_RUN_KF(stateless_exponential, kNumBlocks, self, seed, offset, self, workspace, tiling);
	108	+
	109	+ auto* out = reinterpret_cast<half*>(self);
	110	+ for (int64_t i = 0; i < kElementCount; ++i) {
	111	+ float v = static_cast<float>(out[i]);
	112	+ EXPECT_TRUE(std::isfinite(v)) << "Element " << i << " not finite";
	113	+ EXPECT_GT(v, 0.0f) << "Element " << i << " <= 0";
	114	+ }
	115	+
	116	+ AscendC::GmFree(self);
	117	+ AscendC::GmFree(seed);
	118	+ AscendC::GmFree(offset);
	119	+ AscendC::GmFree(workspace);
	120	+ AscendC::GmFree(tiling);
	121	+}
	122	+
	123	+TEST_F(StatelessExponentialKernelTest, smoke_bfloat16)
	124	+{
	125	+ auto* self = static_cast<uint8_t>(AscendC::GmAlloc(Align32(kElementCount sizeof(bfloat16_t))));
	126	+ auto* seed = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int64_t))));
	127	+ auto* offset = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int64_t))));
	128	+ auto* workspace = static_cast<uint8_t>(AscendC::GmAlloc(Align32(16 1024 * 1024)));
	129	+ auto* tiling = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(RandomUnifiedSimtTilingDataStruct))));
	130	+
	131	+ std::memset(self, 0, kElementCount * sizeof(bfloat16_t));
	132	+ reinterpret_cast<int64_t>(seed) = kSeed;
	133	+ reinterpret_cast<int64_t>(offset) = kOffset;
	134	+ FillTiling(reinterpret_cast<RandomUnifiedSimtTilingDataStruct*>(tiling), kSeed, kOffset);
	135	+
	136	+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
	137	+ ICPU_SET_TILING_KEY(kTilingKeyBf16);
	138	+ ICPU_RUN_KF(stateless_exponential, kNumBlocks, self, seed, offset, self, workspace, tiling);
	139	+
	140	+ auto* out = reinterpret_cast<bfloat16_t*>(self);
	141	+ for (int64_t i = 0; i < kElementCount; ++i) {
	142	+ float v = static_cast<float>(out[i]);
	143	+ EXPECT_TRUE(std::isfinite(v)) << "Element " << i << " not finite";
	144	+ EXPECT_GT(v, 0.0f) << "Element " << i << " <= 0";
	145	+ }
	146	+
	147	+ AscendC::GmFree(self);
	148	+ AscendC::GmFree(seed);
	149	+ AscendC::GmFree(offset);
	150	+ AscendC::GmFree(workspace);
	151	+ AscendC::GmFree(tiling);
	152	+}
	153	+
	154	+TEST_F(StatelessExponentialKernelTest, determinism)
	155	+{
	156	+ auto* self1 = static_cast<uint8_t>(AscendC::GmAlloc(Align32(kElementCount sizeof(float))));
	157	+ auto* self2 = static_cast<uint8_t>(AscendC::GmAlloc(Align32(kElementCount sizeof(float))));
	158	+ auto* seed = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int64_t))));
	159	+ auto* offset = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int64_t))));
	160	+ auto* workspace = static_cast<uint8_t>(AscendC::GmAlloc(Align32(16 1024 * 1024)));
	161	+ auto* tiling = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(RandomUnifiedSimtTilingDataStruct))));
	162	+
	163	+ reinterpret_cast<int64_t>(seed) = 12345;
	164	+ reinterpret_cast<int64_t>(offset) = 4;
	165	+ FillTiling(reinterpret_cast<RandomUnifiedSimtTilingDataStruct*>(tiling), 12345, 4);
	166	+
	167	+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
	168	+ ICPU_SET_TILING_KEY(kTilingKeyFp32);
	169	+ ICPU_RUN_KF(stateless_exponential, kNumBlocks, self1, seed, offset, self1, workspace, tiling);
	170	+ ICPU_RUN_KF(stateless_exponential, kNumBlocks, self2, seed, offset, self2, workspace, tiling);
	171	+
	172	+ EXPECT_EQ(std::memcmp(self1, self2, kElementCount * sizeof(float)), 0);
	173	+
	174	+ AscendC::GmFree(self1);
	175	+ AscendC::GmFree(self2);
	176	+ AscendC::GmFree(seed);
	177	+ AscendC::GmFree(offset);
	178	+ AscendC::GmFree(workspace);
	179	+ AscendC::GmFree(tiling);
	180	+}

Mrandom/stateless_sample_multinomial/CMakeLists.txt+1-1

	@@ -14,5 +14,5 @@ add_all_modules_sources(OPTYPE stateless_sample_multinomial
14	ACLNNTYPE aclnn_exclude	14	ACLNNTYPE aclnn_exclude
15	COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT}	15	COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT}
16	TILING_DIR ${SUPPORT_TILING_DIR}	16	TILING_DIR ${SUPPORT_TILING_DIR}
17	- DEPENDENCIES random_common sim_thread_exponential	17	+ DEPENDENCIES random_common sim_thread_exponential stateless_exponential
18	DISABLE_IN_OPP TRUE)	18	DISABLE_IN_OPP TRUE)

Mrandom/stateless_sample_multinomial/op_api/aclnn_multinomial.cpp+32-213

	@@ -9,6 +9,7 @@
9	*/	9	*/
10	#include "aclnn_multinomial.h"	10	#include "aclnn_multinomial.h"
11	#include "stateless_sample_multinomial.h"	11	#include "stateless_sample_multinomial.h"
		12	+#include "random/stateless_exponential/op_api/stateless_exponential.h"
12	#include "multinomial_with_replacement.h"	13	#include "multinomial_with_replacement.h"
13	#include "math/reduce_sum/op_api/reduce_sum_op.h"	14	#include "math/reduce_sum/op_api/reduce_sum_op.h"
14	#include "math/abs/op_api/abs.h"	15	#include "math/abs/op_api/abs.h"
	@@ -315,69 +316,6 @@ static const aclTensor* Run950AicoreMultinomialWithReplacement(const aclTensor*
315	return multinomialOut;	316	return multinomialOut;
316	}	317	}
317		318
318	-static const aclTensor* RunDavidMultinomialReplaceMent(const aclTensor* selfContiguous, int64_t numsamples,
319	- const aclTensor* randomUniform, const aclTensor* out,
320	- aclOpExecutor* executor)
321	-{
322	- // trun weight into probability
323	- int64_t dimNum = static_cast<int64_t>(selfContiguous->GetViewShape().GetDimNum());
324	- int64_t lastDim = MakeWrapDim(-1, dimNum);
325	- const int64_t dim[] = {lastDim};
326	- auto dimArray = executor->AllocIntArray(dim, 1);
327	-
328	- auto dType = selfContiguous->GetDataType();
329	- auto computeDtype = (dType == DataType::DT_BF16 \|\| dType == DataType::DT_FLOAT16) ? DataType::DT_FLOAT : dType;
330	-
331	- const aclTensor* selfCompute = (dType != computeDtype) ? l0op::Cast(selfContiguous, computeDtype, executor) :
332	- selfContiguous;
333	-
334	- auto sumSelf = l0op::ReduceSumOp(selfCompute, dimArray, true, executor);
335	- CHECK_RET(sumSelf != nullptr, nullptr);
336	-
337	- auto divSumSelf = l0op::RealDiv(selfCompute, sumSelf, executor);
338	- CHECK_RET(divSumSelf != nullptr, nullptr);
339	-
340	- // trun probability into point on [0, 1], shape = original_shape + {1}
341	- const aclTensor* dimTensor = nullptr;
342	- if (lastDim == 0 \|\| lastDim > INT32_MAX) {
343	- dimTensor = executor->ConvertToTensor(&lastDim, 1, DataType::DT_INT64);
344	- } else {
345	- dimTensor = executor->ConvertToTensor(&lastDim, 1, DataType::DT_INT32);
346	- }
347	- CHECK_RET(dimTensor != nullptr, nullptr);
348	- auto accumSelf = l0op::Cumsum(divSumSelf, dimTensor, executor);
349	- CHECK_RET(accumSelf != nullptr, nullptr);
350	-
351	- auto unsqueezeAccum = l0op::UnsqueezeNd(accumSelf, -1, executor);
352	- CHECK_RET(unsqueezeAccum != nullptr, nullptr);
353	-
354	- auto randomCompute = (randomUniform->GetDataType() != computeDtype) ?
355	- l0op::Cast(randomUniform, computeDtype, executor) :
356	- randomUniform;
357	- CHECK_RET(randomCompute != nullptr, nullptr);
358	-
359	- int64_t newShape[dimNum + 1];
360	- newShape[0] = numsamples;
361	- for (int64_t i = 0; i < dimNum; i++) {
362	- newShape[i] = 1;
363	- }
364	- newShape[dimNum] = numsamples;
365	- auto newShapeArray = executor->AllocIntArray(newShape, dimNum + 1);
366	- auto reshapeRandom = l0op::Reshape(randomCompute, newShapeArray, executor);
367	- CHECK_RET(reshapeRandom != nullptr, nullptr);
368	-
369	- // caculate the point RandomUniform on interval unsqueezeAccum
370	- auto greaterEqual = l0op::GreaterEqual(reshapeRandom, unsqueezeAccum, executor);
371	- CHECK_RET(greaterEqual != nullptr, nullptr);
372	-
373	- auto castGreaterEqual = l0op::Cast(greaterEqual, out->GetDataType(), executor);
374	- CHECK_RET(castGreaterEqual != nullptr, nullptr);
375	-
376	- auto multinomialOut = l0op::ReduceSumOp(castGreaterEqual, dimArray, false, executor);
377	- CHECK_RET(multinomialOut != nullptr, nullptr);
378	- return multinomialOut;
379	-}
380	-
381	const aclTensor* GetRandomUniformNoReplaceMent(const aclTensor* selfContiguous, const int64_t seed,	319	const aclTensor* GetRandomUniformNoReplaceMent(const aclTensor* selfContiguous, const int64_t seed,
382	const int64_t offset, aclOpExecutor* executor)	320	const int64_t offset, aclOpExecutor* executor)
383	{	321	{
	@@ -439,85 +377,6 @@ static const aclTensor* Run950AicoreMultinomialWithoutReplacement(const aclTenso
439	return multinomialOut;	377	return multinomialOut;
440	}	378	}
441		379
442	-static const aclTensor* RunDavidMultinomialNoReplaceMent(const aclTensor* selfContiguous, int64_t numsamples,
443	- const aclTensor* randomUniform, const aclTensor* out,
444	- aclOpExecutor* executor)
445	-{
446	- auto dType = selfContiguous->GetDataType();
447	- auto computeDtype = (dType == DataType::DT_BF16 \|\| dType == DataType::DT_FLOAT16) ? DataType::DT_FLOAT : dType;
448	-
449	- const aclTensor* selfCompute = (dType != computeDtype) ? l0op::Cast(selfContiguous, computeDtype, executor) :
450	- selfContiguous;
451	- CHECK_RET(selfCompute != nullptr, nullptr);
452	-
453	- auto randomCompute = (randomUniform->GetDataType() != computeDtype) ?
454	- l0op::Cast(randomUniform, computeDtype, executor) :
455	- randomUniform;
456	- CHECK_RET(randomCompute != nullptr, nullptr);
457	-
458	- const aclTensor* oneTensor;
459	- const aclTensor* positiveMinTensor;
460	- if (computeDtype == DataType::DT_DOUBLE) {
461	- const double one = 1.0;
462	- oneTensor = executor->ConvertToTensor(&one, 1, DataType::DT_DOUBLE);
463	- const double positiveMin = 1e-7;
464	- positiveMinTensor = executor->ConvertToTensor(&positiveMin, 1, DataType::DT_DOUBLE);
465	- } else {
466	- const float one = 1.0f;
467	- oneTensor = executor->ConvertToTensor(&one, 1, DataType::DT_FLOAT);
468	- const float positiveMin = 1e-7f;
469	- positiveMinTensor = executor->ConvertToTensor(&positiveMin, 1, DataType::DT_FLOAT);
470	- }
471	-
472	- auto subRandom = l0op::Sub(oneTensor, randomCompute, executor);
473	- CHECK_RET(subRandom != nullptr, nullptr);
474	-
475	- const float logBase = -1.0f;
476	- const float logScale = 1.0f;
477	- const float logShift = 0.0f;
478	- auto logRandom = l0op::Log(subRandom, logBase, logScale, logShift, executor);
479	- CHECK_RET(logRandom != nullptr, nullptr);
480	-
481	- auto absLog = l0op::Abs(logRandom, executor);
482	- CHECK_RET(absLog != nullptr, nullptr);
483	-
484	- auto exponential = l0op::Add(absLog, positiveMinTensor, executor);
485	- CHECK_RET(exponential != nullptr, nullptr);
486	-
487	- auto divExponential = l0op::RealDiv(selfCompute, exponential, executor);
488	- CHECK_RET(divExponential != nullptr, nullptr);
489	-
490	- int64_t dimNum = static_cast<int64_t>(selfContiguous->GetViewShape().GetDimNum());
491	- int64_t lastDim = MakeWrapDim(-1, dimNum);
492	-
493	- // 获取索引并进行结果排序
494	- const aclTensor* multinomialOutInt32 = nullptr;
495	- if (numsamples == 1) {
496	- multinomialOutInt32 = l0op::ArgMaxV2(divExponential, lastDim, true, executor);
497	- CHECK_RET(multinomialOutInt32 != nullptr, nullptr);
498	- } else {
499	- auto topkOut = l0op::Topk(divExponential, numsamples, lastDim, true, true, op::DataType::DT_INT32, executor);
500	- auto indicesUnsorted = std::get<1>(topkOut);
501	- CHECK_RET(indicesUnsorted != nullptr, nullptr);
502	-
503	- // Sort Indices: Cast(int32->fp32) -> TopK(ascending) -> Cast(fp32->int32)
504	- auto indicesFloat = l0op::Cast(indicesUnsorted, DataType::DT_FLOAT, executor);
505	- CHECK_RET(indicesFloat != nullptr, nullptr);
506	-
507	- auto sortedTopk = l0op::Topk(indicesFloat, numsamples, lastDim, false, true, op::DataType::DT_INT32, executor);
508	- auto sortedValues = std::get<0>(sortedTopk);
509	- CHECK_RET(sortedValues != nullptr, nullptr);
510	-
511	- multinomialOutInt32 = l0op::Cast(sortedValues, DataType::DT_INT32, executor);
512	- CHECK_RET(multinomialOutInt32 != nullptr, nullptr);
513	- }
514	-
515	- // 转换到输出类型（INT64）
516	- auto multinomialOut = l0op::Cast(multinomialOutInt32, out->GetDataType(), executor);
517	- CHECK_RET(multinomialOut != nullptr, nullptr);
518	- return multinomialOut;
519	-}
520	-
521	const aclTensor* RunMultinomialNoReplaceMent(const aclTensor* selfContiguous, int64_t numsamples,	380	const aclTensor* RunMultinomialNoReplaceMent(const aclTensor* selfContiguous, int64_t numsamples,
522	const aclTensor* randomUniform, const aclTensor* out,	381	const aclTensor* randomUniform, const aclTensor* out,
523	aclOpExecutor* executor)	382	aclOpExecutor* executor)
	@@ -599,8 +458,21 @@ aclnnStatus aclnnMultinomialGetWorkspaceSize(const aclTensor* self, int64_t nums
599	uniqueExecutor.get());	458	uniqueExecutor.get());
600	} else if (!replacement \|\| numsamples == 1) {	459	} else if (!replacement \|\| numsamples == 1) {
601	if (IsRegBase()) {	460	if (IsRegBase()) {
602	- auto exp1Random = l0op::Run950AicoreExponentialWithoutReplacement(selfContiguous, seed, offset, 1.0f,	461	+ // StatelessExponential takes tensor seed/offset; convert the scalar ones here.
603	- uniqueExecutor.get());	462	+ aclIntArray* seedList = uniqueExecutor->AllocIntArray(&seed, 1);
		463	+ CHECK_RET(seedList != nullptr, ACLNN_ERR_INNER_NULLPTR);
		464	+ auto seedTensor = uniqueExecutor->ConvertToTensor(seedList, op::DataType::DT_INT64);
		465	+ CHECK_RET(seedTensor != nullptr, ACLNN_ERR_INNER_NULLPTR);
		466	+ aclIntArray* offsetList = uniqueExecutor->AllocIntArray(&offset, 1);
		467	+ CHECK_RET(offsetList != nullptr, ACLNN_ERR_INNER_NULLPTR);
		468	+ auto offsetTensor = uniqueExecutor->ConvertToTensor(offsetList, op::DataType::DT_INT64);
		atomgit-bot7月14日 🟡 Medium Priority 在 `aclnnMultinomialGetWorkspaceSize` 函数的 RegBase 无放回路径（第 443-446 行新增代码）中，`uniqueExecutor->AllocIntArray` 和 `uniqueExecutor->ConvertToTensor` 的返回值 `seedList`、`seedTensor`、`offsetList`、`offsetTensor` 均未做空指针检查。若发生 OOM 等分配失败场景，这些指针将为 null，随后被传入 `l0op::StatelessExponential()`（第 452 行），导致该函数内部解引用空指针而崩溃，而非返回 ACLNN 错误码让上层优雅处理。变更行 → 受影响的合约：`uniqueExecutor->AllocIntArray` / `ConvertToTensor` 可能返回 null → `StatelessExponential` 收到 null tensor → 空指针解引用崩溃。修复方向：为 `seedList`、`offsetList`、`seedTensor`、`offsetTensor` 各添加 `CHECK_RET(xxx != nullptr, ACLNN_ERR_INNER_NULLPTR);`。建议：为 seedList / offsetList / seedTensor / offsetTensor 分别添加空指针检查。参考同文件中其他 ConvertToTensor 的检查模式（如第 295 行 `CHECK_RET(dimTensor != nullptr, nullptr)`）。
		469	+ CHECK_RET(offsetTensor != nullptr, ACLNN_ERR_INNER_NULLPTR);
		470	+ // Pre-allocate a buffer filled in-place so the original weights (selfContiguous) are preserved.
		471	+ auto expInput = uniqueExecutor->AllocTensor(selfContiguous->GetViewShape(), selfContiguous->GetDataType(),
		472	+ selfContiguous->GetViewFormat());
		473	+ CHECK_RET(expInput != nullptr, ACLNN_ERR_INNER_NULLPTR);
		474	+ auto exp1Random = l0op::StatelessExponential(expInput, seedTensor, offsetTensor, 1.0f,
		475	+ uniqueExecutor.get());
604	CHECK_RET(exp1Random != nullptr, ACLNN_ERR_PARAM_NULLPTR);	476	CHECK_RET(exp1Random != nullptr, ACLNN_ERR_PARAM_NULLPTR);
605	multinomialOut = Run950AicoreMultinomialWithoutReplacement(selfContiguous, numsamples, exp1Random,	477	multinomialOut = Run950AicoreMultinomialWithoutReplacement(selfContiguous, numsamples, exp1Random,
606	uniqueExecutor.get());	478	uniqueExecutor.get());
	@@ -613,9 +485,13 @@ aclnnStatus aclnnMultinomialGetWorkspaceSize(const aclTensor* self, int64_t nums
613	} else {	485	} else {
614	if (IsRegBase()) {	486	if (IsRegBase()) {
615	aclIntArray* seedList = uniqueExecutor->AllocIntArray(&seed, 1);	487	aclIntArray* seedList = uniqueExecutor->AllocIntArray(&seed, 1);
		488	+ CHECK_RET(seedList != nullptr, ACLNN_ERR_INNER_NULLPTR);
616	auto seedTensor = uniqueExecutor->ConvertToTensor(seedList, op::DataType::DT_INT64);	489	auto seedTensor = uniqueExecutor->ConvertToTensor(seedList, op::DataType::DT_INT64);
		490	+ CHECK_RET(seedTensor != nullptr, ACLNN_ERR_INNER_NULLPTR);
617	aclIntArray* offsetList = uniqueExecutor->AllocIntArray(&offset, 1);	491	aclIntArray* offsetList = uniqueExecutor->AllocIntArray(&offset, 1);
		492	+ CHECK_RET(offsetList != nullptr, ACLNN_ERR_INNER_NULLPTR);
618	auto offsetTensor = uniqueExecutor->ConvertToTensor(offsetList, op::DataType::DT_INT64);	493	auto offsetTensor = uniqueExecutor->ConvertToTensor(offsetList, op::DataType::DT_INT64);
		atomgit-bot7月14日 🟡 Medium Priority 在 `aclnnMultinomialGetWorkspaceSize` 函数的 RegBase 有放回路径（第 464-467 行新增代码）中，`uniqueExecutor->AllocIntArray` 和 `uniqueExecutor->ConvertToTensor` 的返回值 `seedList`、`seedTensor`、`offsetList`、`offsetTensor` 均未做空指针检查。这些值随后传入 `Run950AicoreMultinomialWithReplacement()`（第 468 行），若指针为 null 将导致内部解引用崩溃而非返回错误码。变更行 → 受影响的合约：`AllocIntArray` / `ConvertToTensor` 可能返回 null → `Run950AicoreMultinomialWithReplacement` 收到 null tensor → 空指针解引用崩溃。修复方向：为 `seedList`、`offsetList`、`seedTensor`、`offsetTensor` 各添加 `CHECK_RET` 空指针检查。建议：为 seedList / offsetList / seedTensor / offsetTensor 分别添加空指针检查。参考同文件中已有的 CHECK_RET 模式。
		494	+ CHECK_RET(offsetTensor != nullptr, ACLNN_ERR_INNER_NULLPTR);
619	multinomialOut = Run950AicoreMultinomialWithReplacement(selfContiguous, numsamples, seedTensor,	495	multinomialOut = Run950AicoreMultinomialWithReplacement(selfContiguous, numsamples, seedTensor,
620	offsetTensor, uniqueExecutor.get());	496	offsetTensor, uniqueExecutor.get());
621	} else {	497	} else {
	@@ -642,64 +518,6 @@ aclnnStatus aclnnMultinomial(void* workspace, uint64_t workspaceSize, aclOpExecu
642	return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);	518	return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
643	}	519	}
644		520
645	-static std::pair<const aclTensor, const aclTensor> BuildDavidCounter(const aclTensor* seedTensor,
646	- const aclTensor* offsetTensor, int64_t offset,
647	- aclOpExecutor* executor)
648	-{
649	- auto seedU64 = l0op::Cast(seedTensor, op::DataType::DT_UINT64, executor);
650	- if (seedU64 == nullptr)
651	- return {nullptr, nullptr};
652	- auto offsetU64 = l0op::Cast(offsetTensor, op::DataType::DT_UINT64, executor);
653	- if (offsetU64 == nullptr)
654	- return {nullptr, nullptr};
655	-
656	- FVector<int64_t> offsetVector{0, static_cast<int64_t>(offset)};
657	- aclIntArray* offsetList = executor->AllocIntArray(offsetVector.data(), 2);
658	- auto tmpTensor = executor->ConvertToTensor(offsetList, op::DataType::DT_UINT64);
659	- auto resultAddOut = l0op::Add(offsetU64, tmpTensor, executor);
660	- if (resultAddOut == nullptr)
661	- return {nullptr, nullptr};
662	-
663	- return {seedU64, resultAddOut};
664	-}
665	-
666	-static const aclTensor* GetDavidRandomUniformReplaceMentTensor(const aclTensor* selfContiguous, int64_t numsamples,
667	- const aclTensor* seedTensor,
668	- const aclTensor* offsetTensor, int64_t offset,
669	- aclOpExecutor* executor)
670	-{
671	- auto [seedU64, resultAddOut] = BuildDavidCounter(seedTensor, offsetTensor, offset, executor);
672	- CHECK_RET(seedU64 != nullptr && resultAddOut != nullptr, nullptr);
673	-
674	- const int64_t randAShape[] = {numsamples};
675	- auto randAShapeArray = executor->AllocIntArray(randAShape, 1);
676	- op::Shape shape;
677	- op::ToShape(randAShapeArray->GetData(), randAShapeArray->Size(), shape);
678	- auto shapeTensor = executor->AllocTensor(shape, selfContiguous->GetDataType(), selfContiguous->GetViewFormat());
679	- CHECK_RET(shapeTensor != nullptr, nullptr);
680	-
681	- int32_t alg = 1;
682	- auto randomUniform = l0op::StatelessRandomUniformV2(shapeTensor, seedU64, resultAddOut, alg, executor);
683	- CHECK_RET(randomUniform != nullptr, nullptr);
684	- return randomUniform;
685	-}
686	-
687	-static const aclTensor* GetDavidRandomUniformNoReplaceMentTensor(const aclTensor* selfContiguous,
688	- const aclTensor* seedTensor,
689	- const aclTensor* offsetTensor, int64_t offset,
690	- aclOpExecutor* executor)
691	-{
692	- auto [seedU64, resultAddOut] = BuildDavidCounter(seedTensor, offsetTensor, offset, executor);
693	- CHECK_RET(seedU64 != nullptr && resultAddOut != nullptr, nullptr);
694	-
695	- int32_t alg = 1;
696	- auto statelessUniform = l0op::StatelessRandomUniformV2(selfContiguous, seedU64, resultAddOut, alg, executor);
697	- CHECK_RET(statelessUniform != nullptr, nullptr);
698	- auto randomUniform = l0op::Cast(statelessUniform, DataType::DT_FLOAT, executor);
699	- CHECK_RET(randomUniform != nullptr, nullptr);
700	- return randomUniform;
701	-}
702	-
703	static const aclTensor* AddOffsetTensor(const aclTensor* offsetTensor, int64_t offset, aclOpExecutor* executor)	521	static const aclTensor* AddOffsetTensor(const aclTensor* offsetTensor, int64_t offset, aclOpExecutor* executor)
704	{	522	{
705	FVector<int64_t> tmpVector = {static_cast<int64_t>(offset)};	523	FVector<int64_t> tmpVector = {static_cast<int64_t>(offset)};
	@@ -743,7 +561,7 @@ aclnnStatus aclnnMultinomialTensorGetWorkspaceSize(const aclTensor* self, int64_
743	CHECK_RET(offsetAddOut != nullptr, ACLNN_ERR_INNER_NULLPTR);	561	CHECK_RET(offsetAddOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
744		562
745	const aclTensor* multinomialOut;	563	const aclTensor* multinomialOut;
746	- if (!CheckSocVersionGe910B() \|\| selfSize <= CPU_NPU_BOUNDARY) {	564	+ if (UseAicpuPath(self, selfSize)) {
747	multinomialOut = l0op::MultinomialWithReplacementTensor(selfContiguous, numsamples, replacement, seedTensor,	565	multinomialOut = l0op::MultinomialWithReplacementTensor(selfContiguous, numsamples, replacement, seedTensor,
748	offsetAddOut, uniqueExecutor.get());	566	offsetAddOut, uniqueExecutor.get());
749	} else if (!replacement \|\| numsamples == 1) {	567	} else if (!replacement \|\| numsamples == 1) {
	@@ -755,11 +573,14 @@ aclnnStatus aclnnMultinomialTensorGetWorkspaceSize(const aclTensor* self, int64_
755	multinomialOut = RunMultinomialNoReplaceMent(selfContiguous, numsamples, randomUniform, out,	573	multinomialOut = RunMultinomialNoReplaceMent(selfContiguous, numsamples, randomUniform, out,
756	uniqueExecutor.get());	574	uniqueExecutor.get());
757	} else {	575	} else {
758	- randomUniform = GetDavidRandomUniformNoReplaceMentTensor(selfContiguous, seedTensor, offsetTensor, offset,	576	+ auto expInput = uniqueExecutor->AllocTensor(selfContiguous->GetViewShape(), selfContiguous->GetDataType(),
759	- uniqueExecutor.get());	577	+ selfContiguous->GetViewFormat());
760	- CHECK_RET(randomUniform != nullptr, ACLNN_ERR_INNER_NULLPTR);	578	+ CHECK_RET(expInput != nullptr, ACLNN_ERR_INNER_NULLPTR);
761	- multinomialOut = RunDavidMultinomialNoReplaceMent(selfContiguous, numsamples, randomUniform, out,	579	+ auto exp1Random = l0op::StatelessExponential(expInput, seedTensor, offsetAddOut, 1.0f,
762	- uniqueExecutor.get());	580	+ uniqueExecutor.get());
		581	+ CHECK_RET(exp1Random != nullptr, ACLNN_ERR_PARAM_NULLPTR);
		582	+ multinomialOut = Run950AicoreMultinomialWithoutReplacement(selfContiguous, numsamples, exp1Random,
		583	+ uniqueExecutor.get());
763	}	584	}
764	} else {	585	} else {
765	const aclTensor* randomUniform = nullptr;	586	const aclTensor* randomUniform = nullptr;
	@@ -770,11 +591,9 @@ aclnnStatus aclnnMultinomialTensorGetWorkspaceSize(const aclTensor* self, int64_
770	multinomialOut = RunMultinomialReplaceMent(selfContiguous, numsamples, randomUniform, out,	591	multinomialOut = RunMultinomialReplaceMent(selfContiguous, numsamples, randomUniform, out,
771	uniqueExecutor.get());	592	uniqueExecutor.get());
772	} else {	593	} else {
773	- randomUniform = GetDavidRandomUniformReplaceMentTensor(selfContiguous, numsamples, seedTensor, offsetTensor,	594	+ // RegBase: use the fused 950 AICore kernel directly (offsetAddOut already carries offsetTensor + offset).
774	- offset, uniqueExecutor.get());	595	+ multinomialOut = Run950AicoreMultinomialWithReplacement(selfContiguous, numsamples, seedTensor,
775	- CHECK_RET(randomUniform != nullptr, ACLNN_ERR_INNER_NULLPTR);	596	+ offsetAddOut, uniqueExecutor.get());
776	- multinomialOut = RunDavidMultinomialReplaceMent(selfContiguous, numsamples, randomUniform, out,
777	- uniqueExecutor.get());
778	}	597	}
779	}	598	}
780	CHECK_RET(multinomialOut != nullptr, ACLNN_ERR_PARAM_NULLPTR);	599	CHECK_RET(multinomialOut != nullptr, ACLNN_ERR_PARAM_NULLPTR);

Mrandom/stateless_sample_multinomial/op_api/stateless_sample_multinomial.cpp+5-35

	@@ -31,14 +31,9 @@ using namespace op;
31	namespace l0op {	31	namespace l0op {
32		32
33	OP_TYPE_REGISTER(StatelessSampleMultinomial);	33	OP_TYPE_REGISTER(StatelessSampleMultinomial);
34	-OP_TYPE_REGISTER(SimThreadExponential);
35		34
36	-const aclTensor* StatelessSampleMultinomial(	35	+const aclTensor* StatelessSampleMultinomial(const aclTensor* xTensor, const aclTensor* seedTensor,
37	- const aclTensor* xTensor,	36	+ const aclTensor* offsetTensor, int64_t numsamples, aclOpExecutor* executor)
38	- const aclTensor* seedTensor,
39	- const aclTensor* offsetTensor,
40	- int64_t numsamples,
41	- aclOpExecutor* executor)
42	{	37	{
43	L0_DFX(StatelessSampleMultinomial, xTensor, seedTensor, offsetTensor);	38	L0_DFX(StatelessSampleMultinomial, xTensor, seedTensor, offsetTensor);
44		39
	@@ -49,34 +44,9 @@ const aclTensor* StatelessSampleMultinomial(
49	aclTensor* out = executor->AllocTensor(outShape, DataType::DT_INT64, xTensor->GetViewFormat());	44	aclTensor* out = executor->AllocTensor(outShape, DataType::DT_INT64, xTensor->GetViewFormat());
50	CHECK_RET(out != nullptr, nullptr);	45	CHECK_RET(out != nullptr, nullptr);
51		46
52	- auto ret = ADD_TO_LAUNCHER_LIST_AICORE(	47	+ auto ret = ADD_TO_LAUNCHER_LIST_AICORE(StatelessSampleMultinomial, OP_ATTR_NAMES({"num_samples"}),
53	- StatelessSampleMultinomial,	48	+ OP_INPUT(xTensor, seedTensor, offsetTensor), OP_OUTPUT(out),
54	- OP_ATTR_NAMES({"num_samples"}),	49	+ OP_ATTR(numsamples));
55	- OP_INPUT(xTensor, seedTensor, offsetTensor),
56	- OP_OUTPUT(out),
57	- OP_ATTR(numsamples));
58	- CHECK_RET(ret == ACLNN_SUCCESS, nullptr);
59	-
60	- return out;
61	-}
62	-
63	-const aclTensor* Run950AicoreExponentialWithoutReplacement(
64	- const aclTensor* self, int64_t seed, int64_t offset, float lambd, aclOpExecutor* executor)
65	-{
66	- aclTensor* out = executor->AllocTensor(self->GetViewShape(), self->GetDataType(), self->GetViewFormat());
67	- CHECK_RET(out != nullptr, nullptr);
68	-
69	- auto shape = self->GetViewShape();
70	- int64_t count = 1;
71	- for (size_t i = 0; i < shape.GetDimNum(); i++) {
72	- count *= shape.GetDim(i);
73	- }
74	-
75	- auto ret = ADD_TO_LAUNCHER_LIST_AICORE(
76	- SimThreadExponential,
77	- OP_INPUT(out),
78	- OP_OUTPUT(out),
79	- OP_ATTR(count, lambd, seed, offset));
80	CHECK_RET(ret == ACLNN_SUCCESS, nullptr);	50	CHECK_RET(ret == ACLNN_SUCCESS, nullptr);
81		51
82	return out;	52	return out;

Mrandom/stateless_sample_multinomial/op_api/stateless_sample_multinomial.h+2-20

	@@ -32,26 +32,8 @@ namespace l0op {
32	* @param executor Op executor	32	* @param executor Op executor
33	* @return Output tensor with shape {numDist, numsamples}, dtype DT_INT64	33	* @return Output tensor with shape {numDist, numsamples}, dtype DT_INT64
34	*/	34	*/
35	-const aclTensor* StatelessSampleMultinomial(	35	+const aclTensor* StatelessSampleMultinomial(const aclTensor* xTensor, const aclTensor* seedTensor,
36	- const aclTensor* xTensor,	36	+ const aclTensor* offsetTensor, int64_t numsamples, aclOpExecutor* executor);
37	- const aclTensor* seedTensor,
38	- const aclTensor* offsetTensor,
39	- int64_t numsamples,
40	- aclOpExecutor* executor);
41	-
42	-/**
43	- * @brief Generate exponential random samples on 950 AICore for multinomial without replacement.
44	- * Uses SimThreadExponential kernel with Philox RNG (offset aligned to 4).
45	- *
46	- * @param self Input tensor, determines output shape and dtype
47	- * @param seed RNG seed (scalar)
48	- * @param offset RNG offset (scalar, will be aligned to multiple of 4)
49	- * @param lambd Lambda parameter for exponential distribution
50	- * @param executor Op executor
51	- * @return Output tensor with same shape and dtype as self
52	- */
53	-const aclTensor* Run950AicoreExponentialWithoutReplacement(
54	- const aclTensor* self, int64_t seed, int64_t offset, float lambd, aclOpExecutor* executor);
55		37
56	} // namespace l0op	38	} // namespace l0op
57		39

Mrandom/stateless_sample_multinomial/op_host/arch35/stateless_sample_multinomial_tiling.cpp+23-33

	@@ -33,40 +33,28 @@ OpTilingConfig StatelessSampleMultinomialTiling::BuildOpConfig()
33	{	33	{
34	OpTilingConfig config;	34	OpTilingConfig config;
35		35
36	- config.inputCheckRules = {	36	+ config.inputCheckRules = {{INPUT_IDX_X, {{ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16}, -1, {}, nullptr}},
37	- {INPUT_IDX_X, {{ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16}, -1, {}, nullptr}},	37	+ {INPUT_IDX_SEED, {{ge::DT_INT64}, 1, {}, nullptr}},
38	- {INPUT_IDX_SEED, {{ge::DT_INT64}, 1, {}, nullptr}},	38	+ {INPUT_IDX_OFFSET, {{ge::DT_INT64}, 1, {}, nullptr}}};
39	- {INPUT_IDX_OFFSET, {{ge::DT_INT64}, 1, {}, nullptr}}};	39	+ config.outputCheckRules = {{OUTPUT_IDX_Y, {{ge::DT_INT64}, -1, {}, nullptr}}};
40	- config.outputCheckRules = {
41	- {OUTPUT_IDX_Y, {{ge::DT_INT64}, -1, {}, nullptr}}};
42		40
43	config.getOutputSize = [](gert::TilingContext* ctx, int64_t& size) {	41	config.getOutputSize = [](gert::TilingContext* ctx, int64_t& size) {
44	auto xShape = ctx->GetInputShape(INPUT_IDX_X);	42	auto xShape = ctx->GetInputShape(INPUT_IDX_X);
45	- OP_CHECK_IF(xShape == nullptr,	43	+ OP_CHECK_IF(xShape == nullptr, OP_LOGE(ctx->GetNodeName(), "get x shape failed"), return ge::GRAPH_FAILED);
46	- OP_LOGE(ctx->GetNodeName(), "get x shape failed"), return ge::GRAPH_FAILED);
47	int64_t numDist = (xShape->GetStorageShape().GetDimNum() == 2) ? xShape->GetStorageShape().GetDim(0) : 1;	44	int64_t numDist = (xShape->GetStorageShape().GetDimNum() == 2) ? xShape->GetStorageShape().GetDim(0) : 1;
48	- auto numsamplesPtr = ctx->GetAttrs()->GetInt(0); // attr index 0: "num_samples"	45	+ auto numsamplesPtr = ctx->GetAttrs()->GetInt(0); // attr index 0: "num_samples"
49	- OP_CHECK_IF(numsamplesPtr == nullptr,	46	+ OP_CHECK_IF(numsamplesPtr == nullptr, OP_LOGE(ctx->GetNodeName(), "get num_samples attr failed"),
50	- OP_LOGE(ctx->GetNodeName(), "get num_samples attr failed"), return ge::GRAPH_FAILED);	47	+ return ge::GRAPH_FAILED);
51	int64_t realSize = numDist * (*numsamplesPtr);	48	int64_t realSize = numDist * (*numsamplesPtr);
52	size = (realSize + RANDOM_NUM_PER_COUNTER - 1) / RANDOM_NUM_PER_COUNTER;	49	size = (realSize + RANDOM_NUM_PER_COUNTER - 1) / RANDOM_NUM_PER_COUNTER;
53	return ge::GRAPH_SUCCESS;	50	return ge::GRAPH_SUCCESS;
54	};	51	};
55		52
56	- config.getSeedAndOffset = [](gert::TilingContext* ctx, int64_t& seed, int64_t& offset) {	53	+ // seed/offset are device-computed (offset is an l0op::Add output) with no host value at
57	- gert::Shape seedShape;	54	+ // tiling time; reading them faults. Tiling stores 0; the kernel reads real values from GM.
58	- auto ret = ExtractTensorValue(ctx, INPUT_IDX_SEED, seedShape);	55	+ config.getSeedAndOffset = [](gert::TilingContext* /ctx/, int64_t& seed, int64_t& offset) {
59	- OP_CHECK_IF(ret != ge::GRAPH_SUCCESS,	56	+ seed = 0;
60	- OP_LOGE(ctx->GetNodeName(), "get seed value failed"), return ge::GRAPH_FAILED);	57	+ offset = 0;
61	- seed = static_cast<int64_t>(seedShape.GetDim(0));
62	- gert::Shape offsetShape;
63	- ret = ExtractTensorValue(ctx, INPUT_IDX_OFFSET, offsetShape);
64	- OP_CHECK_IF(ret != ge::GRAPH_SUCCESS,
65	- OP_LOGE(ctx->GetNodeName(), "get offset value failed"), return ge::GRAPH_FAILED);
66	- offset = static_cast<int64_t>(offsetShape.GetDim(0));
67	- OP_CHECK_IF(offset % RANDOM_NUM_PER_COUNTER != 0,
68	- OP_LOGE(ctx->GetNodeName(), "The offset must be a multiple of 4, but got %ld", offset),
69	- return ge::GRAPH_FAILED);
70	return ge::GRAPH_SUCCESS;	58	return ge::GRAPH_SUCCESS;
71	};	59	};
72		60
	@@ -87,14 +75,16 @@ ge::graphStatus StatelessSampleMultinomialTiling::UniqueProcess()
87	int64_t numCategories = xShape->GetStorageShape().GetDim(xShape->GetStorageShape().GetDimNum() - 1);	75	int64_t numCategories = xShape->GetStorageShape().GetDim(xShape->GetStorageShape().GetDimNum() - 1);
88		76
89	auto numsamplesPtr = context_->GetAttrs()->GetInt(0);	77	auto numsamplesPtr = context_->GetAttrs()->GetInt(0);
90	- OP_CHECK_IF(numsamplesPtr == nullptr,	78	+ OP_CHECK_IF(numsamplesPtr == nullptr, OP_LOGE(context_->GetNodeName(), "get num_samples attr failed"),
91	- OP_LOGE(context_->GetNodeName(), "get num_samples attr failed"), return ge::GRAPH_FAILED);	79	+ return ge::GRAPH_FAILED);
92	- simtTilingData_.from = *numsamplesPtr; // num_samples	80	+ simtTilingData_.from = *numsamplesPtr; // num_samples
93	- simtTilingData_.extraInt64Param1 = numDist * (numsamplesPtr); // numDist numsamples	81	+ simtTilingData_.extraInt64Param1 = numDist * (numsamplesPtr); // numDist numsamples
94	- simtTilingData_.range = static_cast<uint64_t>(numCategories); // numCategories	82	+ simtTilingData_.range = static_cast<uint64_t>(numCategories); // numCategories
95	- simtTilingData_.splitBlockCount =	83	+ simtTilingData_.splitBlockCount = (static_cast<uint64_t>(numDist) * static_cast<uint64_t>(numCategories) >
96	- (static_cast<uint64_t>(numDist) * static_cast<uint64_t>(numCategories) > UINT32_MAX \|\|	84	+ UINT32_MAX \|\|
97	- static_cast<uint64_t>(simtTilingData_.extraInt64Param1) > UINT32_MAX) ? 1 : 0; // index type uint64 or uint32	85	+ static_cast<uint64_t>(simtTilingData_.extraInt64Param1) > UINT32_MAX) ?
		86	+ 1 :
		87	+ 0; // index type uint64 or uint32
98		88
99	return ge::GRAPH_SUCCESS;	89	return ge::GRAPH_SUCCESS;
100	}	90	}

Mrandom/stateless_sample_multinomial/op_kernel/arch35/stateless_sample_multinomial_impl.h+30-37

	@@ -26,12 +26,9 @@ constexpr static uint32_t UNROLL_FACTOR = 4;
26		26
27	template <typename XT, typename IndexT, uint16_t THREAD_LAUNCH_BOUND>	27	template <typename XT, typename IndexT, uint16_t THREAD_LAUNCH_BOUND>
28	__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_LAUNCH_BOUND) inline void SimtUniformRandomBinarySearch(	28	__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_LAUNCH_BOUND) inline void SimtUniformRandomBinarySearch(
29	- __gm__ volatile int64_t* outputGM,	29	+ __gm__ volatile int64_t* outputGM, __gm__ volatile XT* xGM, IndexT elementNum, int64_t seed, IndexT numsamples,
30	- __gm__ volatile XT* xGM,	30	+ uint32_t numCat, uint64_t nsMagic, uint64_t nsShift, IndexT samplesAligned, uint32_t baseOffsetLo,
31	- IndexT elementNum, int64_t seed,	31	+ uint32_t baseOffsetHi)
32	- IndexT numsamples, uint32_t numCat,
33	- uint64_t nsMagic, uint64_t nsShift, IndexT samplesAligned,
34	- uint32_t baseOffsetLo, uint32_t baseOffsetHi)
35	{	32	{
36	uint32_t key[ALG_KEY_SIZE] = {0, 0};	33	uint32_t key[ALG_KEY_SIZE] = {0, 0};
37	key[0] = static_cast<uint32_t>(seed);	34	key[0] = static_cast<uint32_t>(seed);
	@@ -42,11 +39,9 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_LAUNCH_BOUND) inline void SimtUniform
42		39
43	IndexT elementNumAligned = elementNum & ~static_cast<IndexT>(UNROLL_FACTOR - 1);	40	IndexT elementNumAligned = elementNum & ~static_cast<IndexT>(UNROLL_FACTOR - 1);
44		41
45	- for (IndexT baseIndex = idx * UNROLL_FACTOR; baseIndex < elementNum;	42	+ for (IndexT baseIndex = idx * UNROLL_FACTOR; baseIndex < elementNum; baseIndex += stride * UNROLL_FACTOR) {
46	- baseIndex += stride * UNROLL_FACTOR) {	43	+ uint32_t count = (baseIndex < elementNumAligned) ? UNROLL_FACTOR :
47	-	44	+ static_cast<uint32_t>(elementNum - baseIndex);
48	- uint32_t count = (baseIndex < elementNumAligned) ?
49	- UNROLL_FACTOR : static_cast<uint32_t>(elementNum - baseIndex);
50		45
51	IndexT d = static_cast<IndexT>(Simt::UintDiv(static_cast<uint64_t>(baseIndex), nsMagic, nsShift));	46	IndexT d = static_cast<IndexT>(Simt::UintDiv(static_cast<uint64_t>(baseIndex), nsMagic, nsShift));
52	IndexT s = baseIndex - d * numsamples;	47	IndexT s = baseIndex - d * numsamples;
	@@ -54,11 +49,8 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_LAUNCH_BOUND) inline void SimtUniform
54	uint64_t subsequence = static_cast<uint64_t>(d) * samplesAligned + s;	49	uint64_t subsequence = static_cast<uint64_t>(d) * samplesAligned + s;
55		50
56	for (uint32_t k = 0; k < count; k++) {	51	for (uint32_t k = 0; k < count; k++) {
57	- uint32_t counterTmp[ALG_COUNTER_SIZE] = {	52	+ uint32_t counterTmp[ALG_COUNTER_SIZE] = {baseOffsetLo, baseOffsetHi, static_cast<uint32_t>(subsequence),
58	- baseOffsetLo, baseOffsetHi,	53	+ static_cast<uint32_t>(subsequence >> 32)};
59	- static_cast<uint32_t>(subsequence),
60	- static_cast<uint32_t>(subsequence >> 32)
61	- };
62	PhiloxRandomSimt(key, counterTmp, counterTmp);	54	PhiloxRandomSimt(key, counterTmp, counterTmp);
63		55
64	XT u = static_cast<XT>(counterTmp[0] * RAND_2POW32_INV + RAND_2POW32_INV_HALF);	56	XT u = static_cast<XT>(counterTmp[0] * RAND_2POW32_INV + RAND_2POW32_INV_HALF);
	@@ -97,28 +89,30 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_LAUNCH_BOUND) inline void SimtUniform
97	template <typename XT>	89	template <typename XT>
98	class StatelessSampleMultinomialOp {	90	class StatelessSampleMultinomialOp {
99	public:	91	public:
100	- __aicore__ inline StatelessSampleMultinomialOp() {};	92	+ __aicore__ inline StatelessSampleMultinomialOp(){};
101	- __aicore__ inline void Init(GM_ADDR y, GM_ADDR x, GM_ADDR workspace,	93	+ __aicore__ inline void Init(GM_ADDR y, GM_ADDR x, GM_ADDR seed, GM_ADDR offset, GM_ADDR workspace,
102	- const RandomUnifiedSimtTilingDataStruct* __restrict tilingData,	94	+ const RandomUnifiedSimtTilingDataStruct* __restrict tilingData, TPipe* pipe);
103	- TPipe* pipe);
104	__aicore__ inline void Process();	95	__aicore__ inline void Process();
105		96
106	private:	97	private:
107	const RandomUnifiedSimtTilingDataStruct* tilingData_;	98	const RandomUnifiedSimtTilingDataStruct* tilingData_;
108	GlobalTensor<int64_t> outputGM_;	99	GlobalTensor<int64_t> outputGM_;
109	GM_ADDR xGM_;	100	GM_ADDR xGM_;
		101	+ GM_ADDR seedGM_;
		102	+ GM_ADDR offsetGM_;
110	uint32_t blockIdx_;	103	uint32_t blockIdx_;
111	};	104	};
112		105
113	template <typename XT>	106	template <typename XT>
114	__aicore__ inline void StatelessSampleMultinomialOp<XT>::Init(	107	__aicore__ inline void StatelessSampleMultinomialOp<XT>::Init(
115	- GM_ADDR y, GM_ADDR x, GM_ADDR workspace,	108	+ GM_ADDR y, GM_ADDR x, GM_ADDR seed, GM_ADDR offset, GM_ADDR workspace,
116	- const RandomUnifiedSimtTilingDataStruct* __restrict tilingData,	109	+ const RandomUnifiedSimtTilingDataStruct* __restrict tilingData, TPipe* pipe)
117	- TPipe* pipe)
118	{	110	{
119	tilingData_ = tilingData;	111	tilingData_ = tilingData;
120	outputGM_.SetGlobalBuffer((__gm__ int64_t*)y);	112	outputGM_.SetGlobalBuffer((__gm__ int64_t*)y);
121	xGM_ = x;	113	xGM_ = x;
		114	+ seedGM_ = seed;
		115	+ offsetGM_ = offset;
122	blockIdx_ = GetBlockIdx();	116	blockIdx_ = GetBlockIdx();
123	}	117	}
124		118
	@@ -134,29 +128,28 @@ __aicore__ inline void StatelessSampleMultinomialOp<XT>::Process()
134	uint64_t elementNum = static_cast<uint64_t>(tilingData_->extraInt64Param1);	128	uint64_t elementNum = static_cast<uint64_t>(tilingData_->extraInt64Param1);
135	bool useUint64Index = (tilingData_->splitBlockCount != 0);	129	bool useUint64Index = (tilingData_->splitBlockCount != 0);
136		130
		131	+ // Read real seed/offset from GM (tiling filled 0 placeholders because the offset is a
		132	+ // device-computed intermediate with no host value at tiling time; see the tiling impl).
		133	+ int64_t realSeed = (reinterpret_cast<__gm__ int64_t>(seedGM_));
		134	+ int64_t realOffset = (reinterpret_cast<__gm__ int64_t>(offsetGM_));
		135	+
137	uint64_t nsMagic, nsShift;	136	uint64_t nsMagic, nsShift;
138	GetUintDivMagicAndShift(nsMagic, nsShift, numsamples);	137	GetUintDivMagicAndShift(nsMagic, nsShift, numsamples);
139		138
140	- uint64_t baseOffset = (static_cast<uint64_t>(tilingData_->offset) + VEC_4 - 1) / VEC_4;	139	+ uint64_t baseOffset = (static_cast<uint64_t>(realOffset) + VEC_4 - 1) / VEC_4;
141	uint32_t baseOffsetLo = static_cast<uint32_t>(baseOffset);	140	uint32_t baseOffsetLo = static_cast<uint32_t>(baseOffset);
142	uint32_t baseOffsetHi = static_cast<uint32_t>(baseOffset >> 32);	141	uint32_t baseOffsetHi = static_cast<uint32_t>(baseOffset >> 32);
143		142
144	- uint64_t samplesAligned = ((numsamples + SAMPLES_ALIGNMENT - 1)	143	+ uint64_t samplesAligned = ((numsamples + SAMPLES_ALIGNMENT - 1) / SAMPLES_ALIGNMENT) * SAMPLES_ALIGNMENT;
145	- / SAMPLES_ALIGNMENT) * SAMPLES_ALIGNMENT;
146		144
147	if (useUint64Index) {	145	if (useUint64Index) {
148	- asc_vf_call<SimtUniformRandomBinarySearch<XT, uint64_t, CORE_THREAD_NUM_U64>>(dim3(CORE_THREAD_NUM_U64),	146	+ asc_vf_call<SimtUniformRandomBinarySearch<XT, uint64_t, CORE_THREAD_NUM_U64>>(
149	- (__gm__ volatile int64_t*)(outputGM_.GetPhyAddr()),	147	+ dim3(CORE_THREAD_NUM_U64), (__gm__ volatile int64_t)(outputGM_.GetPhyAddr()), (__gm__ volatile XT)(xGM_),
150	- (__gm__ volatile XT*)(xGM_),	148	+ elementNum, realSeed, numsamples, numCat, nsMagic, nsShift, samplesAligned, baseOffsetLo, baseOffsetHi);
151	- elementNum, tilingData_->seed,
152	- numsamples, numCat, nsMagic, nsShift, samplesAligned,
153	- baseOffsetLo, baseOffsetHi);
154	} else {	149	} else {
155	- asc_vf_call<SimtUniformRandomBinarySearch<XT, uint32_t, CORE_THREAD_NUM_U32>>(dim3(CORE_THREAD_NUM_U32),	150	+ asc_vf_call<SimtUniformRandomBinarySearch<XT, uint32_t, CORE_THREAD_NUM_U32>>(
156	- (__gm__ volatile int64_t*)(outputGM_.GetPhyAddr()),	151	+ dim3(CORE_THREAD_NUM_U32), (__gm__ volatile int64_t)(outputGM_.GetPhyAddr()), (__gm__ volatile XT)(xGM_),
157	- (__gm__ volatile XT*)(xGM_),	152	+ static_cast<uint32_t>(elementNum), realSeed, static_cast<uint32_t>(numsamples), numCat, nsMagic, nsShift,
158	- static_cast<uint32_t>(elementNum), tilingData_->seed,
159	- static_cast<uint32_t>(numsamples), numCat, nsMagic, nsShift,
160	static_cast<uint32_t>(samplesAligned), baseOffsetLo, baseOffsetHi);	153	static_cast<uint32_t>(samplesAligned), baseOffsetLo, baseOffsetHi);
161	}	154	}
162	}	155	}

Mrandom/stateless_sample_multinomial/op_kernel/stateless_sample_multinomial.cpp+3-3

	@@ -19,8 +19,8 @@ using namespace StatelessSampleMultinomial;
19		19
20	#define STATELESS_SAMPLE_MULTINOMIAL_DEFAULT_TILING_KEY 100	20	#define STATELESS_SAMPLE_MULTINOMIAL_DEFAULT_TILING_KEY 100
21		21
22	-__global__ __aicore__ void stateless_sample_multinomial(	22	+__global__ __aicore__ void stateless_sample_multinomial(GM_ADDR x, GM_ADDR seed, GM_ADDR offset, GM_ADDR y,
23	- GM_ADDR x, GM_ADDR seed, GM_ADDR offset, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)	23	+ GM_ADDR workspace, GM_ADDR tiling)
24	{	24	{
25	REGISTER_TILING_DEFAULT(RandomUnifiedSimtTilingDataStruct);	25	REGISTER_TILING_DEFAULT(RandomUnifiedSimtTilingDataStruct);
26	KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIV_1_0);	26	KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIV_1_0);
	@@ -29,7 +29,7 @@ __global__ __aicore__ void stateless_sample_multinomial(
29		29
30	if (TILING_KEY_IS(STATELESS_SAMPLE_MULTINOMIAL_DEFAULT_TILING_KEY)) {	30	if (TILING_KEY_IS(STATELESS_SAMPLE_MULTINOMIAL_DEFAULT_TILING_KEY)) {
31	StatelessSampleMultinomialOp<DTYPE_X> op;	31	StatelessSampleMultinomialOp<DTYPE_X> op;
32	- op.Init(y, x, workspace, &tilingData, &pipe);	32	+ op.Init(y, x, seed, offset, workspace, &tilingData, &pipe);
33	op.Process();	33	op.Process();
34	}	34	}
35	}	35	}