Star767
274
代码
代码
Issues51
Pull Requests150
Actions
项目讨论
Wiki
分析
Star767
274
已合并
【feat】: 新增ascir:Rand/Randn #1749
liyuewei创建于 8月15日
【feat】: 新增ascir:Rand/Randn #1749
已合并
liyuewei创建于 8月15日
从liyuewei/graph-autofusion0_local合入到cann/graph-autofusiondevelop
共 26 个文件变更+881-2
Mautofuse/att/base/att_const_values.h+2-0
@@ -371,6 +371,8 @@ inline const std::string kRoundToInt = "RoundToInt";
371inline const std::string kTruncToInt = "TruncToInt";371inline const std::string kTruncToInt = "TruncToInt";
372inline const std::string kTrunc = "Trunc";372inline const std::string kTrunc = "Trunc";
373inline const std::string kVmAdd = "VmAdd";373inline const std::string kVmAdd = "VmAdd";
374+inline const std::string kRand = "Rand";
375+inline const std::string kRandn = "Randn";
374inline const std::string kPlaceholder = "Placeholder";376inline const std::string kPlaceholder = "Placeholder";
375#define JOIN(a, b) a##b377#define JOIN(a, b) a##b
376#define JOIN_A_B_C(a, b, c) a##b##c378#define JOIN_A_B_C(a, b, c) a##b##c
Mautofuse/common/ascgraph_info_complete.cpp+2-0
@@ -104,6 +104,8 @@ static const std::map<std::string, af::ComputeType> kOpTypeToComputeType = {
104 {Scalar::Type, af::ComputeType::kComputeInvalid},104 {Scalar::Type, af::ComputeType::kComputeInvalid},
105 {Output::Type, af::ComputeType::kComputeInvalid},105 {Output::Type, af::ComputeType::kComputeInvalid},
106 {IndexExpr::Type, af::ComputeType::kComputeInvalid},106 {IndexExpr::Type, af::ComputeType::kComputeInvalid},
107+ {"Rand", af::ComputeType::kComputeInvalid},
108+ {"Randn", af::ComputeType::kComputeInvalid},
107 109 
108 {Load::Type, af::ComputeType::kComputeLoad},110 {Load::Type, af::ComputeType::kComputeLoad},
109 {Store::Type, af::ComputeType::kComputeStore},111 {Store::Type, af::ComputeType::kComputeStore},
Mautofuse/compiler/py_module/pyascir.h+2-0
@@ -146,6 +146,8 @@ PyMODINIT_FUNC PyInit_pyautofuse(void);
146 OP(Axpy) \146 OP(Axpy) \
147 OP(CopySign) \147 OP(CopySign) \
148 OP(Ceil2Int) \148 OP(Ceil2Int) \
149+ OP(Rand) \
150+ OP(Randn) \
149 OP(RoundToInt) \151 OP(RoundToInt) \
150 OP(Square) \152 OP(Square) \
151 OP(Sinh) \153 OP(Sinh) \
Mautofuse/compiler/python/ascir_api.py+52-0
@@ -2272,3 +2272,55 @@ def I1e(
2272 return _common_in_1_out_1_normal_op(2272 return _common_in_1_out_1_normal_op(
2273 "I1e", owner_graph, x, axis=axis, size=size, stride=stride2273 "I1e", owner_graph, x, axis=axis, size=size, stride=stride
2274 )2274 )
2275+ 
2276+ 
2277+def Rand(
2278+ owner_graph: ascir.HintGraph,
2279+ *,
2280+ dtype: ascir.dtypes,
2281+ axis: List[ascir.Axis],
2282+ size: Optional[List[ascir.SizeExpr]] = None,
2283+ stride: Optional[List[ascir.SizeExpr]] = None,
2284+) -> ascir.OpsOperatorOutput:
2285+ """
2286+ Rand - generate uniform random numbers using PhiloxRandom algorithm.
2287+ Uses fixed philoxKey = {0, 0} and philoxCounter = {0, 0, 0, 0} for continuous mode.
2288+ Supports float data type.
2289+ """
2290+ meta = _get_metadata(owner_graph)
2291+ name = _generate_op_name(owner_graph, "rand")
2292+ rand_op = ascir.ops.Rand(name, owner_graph)
2293+ meta.ops.append(rand_op)
2294+ 
2295+ # Set dtype and view, then infer
2296+ rand_op.y.dtype = dtype
2297+ _infer_or_set_view(rand_op.y, axis, size, stride)
2298+ rand_op.infer_dtype()
2299+ 
2300+ return rand_op.y
2301+ 
2302+ 
2303+def Randn(
2304+ owner_graph: ascir.HintGraph,
2305+ *,
2306+ dtype: ascir.dtypes,
2307+ axis: List[ascir.Axis],
2308+ size: Optional[List[ascir.SizeExpr]] = None,
2309+ stride: Optional[List[ascir.SizeExpr]] = None,
2310+) -> ascir.OpsOperatorOutput:
2311+ """
2312+ Randn - generate normal random numbers using PhiloxRandom algorithm.
2313+ Uses fixed philoxKey = {0, 0} and philoxCounter = {0, 0, 0, 0} for continuous mode.
2314+ Supports uint32_t and int32_t data types.
2315+ """
2316+ meta = _get_metadata(owner_graph)
2317+ name = _generate_op_name(owner_graph, "randn")
2318+ randn_op = ascir.ops.Randn(name, owner_graph)
2319+ meta.ops.append(randn_op)
2320+ 
2321+ # Set dtype and view, then infer
2322+ randn_op.y.dtype = dtype
2323+ _infer_or_set_view(randn_op.y, axis, size, stride)
2324+ randn_op.infer_dtype()
2325+ 
2326+ return randn_op.y
Mautofuse/optimize/platform/common/base_alignment_strategy.cpp+3-1
@@ -371,7 +371,9 @@ af::Status BaseAlignmentStrategy::AlignVectorizedStrides(ascir::ImplGraph &impl_
371}371}
372 372 
373af::Status BaseAlignmentStrategy::InferAlignmentForOneNode(ascir::ImplGraph &, const af::AscNodePtr &node, bool &) {373af::Status BaseAlignmentStrategy::InferAlignmentForOneNode(ascir::ImplGraph &, const af::AscNodePtr &node, bool &) {
374- GE_ASSERT_TRUE(!node->inputs().empty(), "The inputs of %s(%s) is empty.", node->GetTypePtr(), node->GetNamePtr());374+ GE_ASSERT_TRUE(
375+ !node->inputs().empty() || af::ops::IsOps<af::ascir_op::Rand>(node) || af::ops::IsOps<af::ascir_op::Randn>(node),
376+ "The inputs of %s(%s) is empty.", node->GetTypePtr(), node->GetNamePtr());
375 GE_ASSERT_TRUE(!node->outputs().empty(), "The output of %s(%s) is empty.", node->GetTypePtr(), node->GetNamePtr());377 GE_ASSERT_TRUE(!node->outputs().empty(), "The output of %s(%s) is empty.", node->GetTypePtr(), node->GetNamePtr());
376 af::ComputeType compute_type = node->attr.api.compute_type;378 af::ComputeType compute_type = node->attr.api.compute_type;
377 auto it = compute_type_to_infer_func_.find(compute_type);379 auto it = compute_type_to_infer_func_.find(compute_type);
Mautofuse/tests/framework/share_graph/include/share_graph.h+2-0
@@ -127,6 +127,8 @@ struct ShareGraph {
127 static af::ComputeGraphPtr RemainderInt32StoreFusedGraph(size_t dims_size);127 static af::ComputeGraphPtr RemainderInt32StoreFusedGraph(size_t dims_size);
128 static af::ComputeGraphPtr FloorDivInt32StoreFusedGraph(size_t dims_size);128 static af::ComputeGraphPtr FloorDivInt32StoreFusedGraph(size_t dims_size);
129 static af::ComputeGraphPtr LoadTanhBf16FusedGraph(size_t dims_size);129 static af::ComputeGraphPtr LoadTanhBf16FusedGraph(size_t dims_size);
130+ static af::ComputeGraphPtr RandStoreFusedGraph(size_t dims_size);
131+ static af::ComputeGraphPtr RandnStoreFusedGraph(size_t dims_size);
130 static af::ComputeGraphPtr AddAbsScalarFusedGraph(size_t dims_size, af::DataType dtype);132 static af::ComputeGraphPtr AddAbsScalarFusedGraph(size_t dims_size, af::DataType dtype);
131 static af::ComputeGraphPtr AbsBrcAddFusedGraph(size_t dims_size);133 static af::ComputeGraphPtr AbsBrcAddFusedGraph(size_t dims_size);
132 static af::ComputeGraphPtr UbScalarBrcAbsAddFusedGraph(size_t dims_size);134 static af::ComputeGraphPtr UbScalarBrcAbsAddFusedGraph(size_t dims_size);
Mautofuse/tests/framework/share_graph/src/share_graph.cc+105-0
@@ -13335,4 +13335,109 @@ af::ComputeGraphPtr ShareGraph::ArgMaxFusedGraph(size_t dims_size) {
13335 return compute_graph;13335 return compute_graph;
13336}13336}
13337 13337 
13338+static void CreateRandStoreAscGraph(af::AscGraph &graph, size_t dims_size) {
13339+ auto s0 = graph.CreateSizeVar("s0");
13340+ auto z0 = graph.CreateAxis("z0", s0);
13341+ auto One = Symbol(1);
13342+ 
13343+ af::ascir_op::Rand rand_op("rand");
13344+ graph.AddNode(rand_op);
13345+ *rand_op.y.axis = {z0.id};
13346+ *rand_op.y.repeats = {s0};
13347+ *rand_op.y.strides = {One};
13348+ 
13349+ rand_op.attr.api.compute_type = af::ComputeType::kComputeElewise;
13350+ rand_op.attr.api.type = af::ApiType::kAPITypeCompute;
13351+ rand_op.attr.api.unit = af::ComputeUnit::kUnitVector;
13352+ rand_op.attr.sched.axis = {z0.id};
13353+ 
13354+ af::ascir_op::Store store("store");
13355+ store.x = rand_op.y;
13356+ store.attr.api.compute_type = ComputeType::kComputeStore;
13357+ store.attr.api.type = af::ApiType::kAPITypeCompute;
13358+ store.attr.api.unit = ComputeUnit::kUnitMTE3;
13359+ store.attr.sched.axis = {z0.id};
13360+ *store.y.axis = {z0.id};
13361+ *store.y.repeats = {s0};
13362+ *store.y.strides = {One};
13363+ store.y.dtype = af::DataType::DT_FLOAT;
13364+ 
13365+ af::ascir_op::Output y("output");
13366+ y.x = store.y;
13367+ y.ir_attr.SetIndex(0);
13368+}
13369+ 
13370+af::ComputeGraphPtr ShareGraph::RandStoreFusedGraph(size_t dims_size) {
13371+ auto builder = GraphBuilder("rand_store_test");
13372+ auto ascbc = builder.AddNode("ascbc", "AscGraph", 0, 1);
13373+ auto netoutput = builder.AddNode("netoutput", af::NETOUTPUT, 1, 0);
13374+ 
13375+ builder.AddDataEdge(ascbc, 0, netoutput, 0);
13376+ ComputeGraphPtr compute_graph = builder.GetGraph();
13377+ if (compute_graph == nullptr) {
13378+ return nullptr;
13379+ }
13380+ auto ascbc_node = compute_graph->FindNode("ascbc");
13381+ af::AscGraph sub_graph("rand_store");
13382+ CreateRandStoreAscGraph(sub_graph, dims_size);
13383+ 
13384+ std::string sub_graph_str;
13385+ af::AscGraphUtils::SerializeToReadable(sub_graph, sub_graph_str);
13386+ af::AttrUtils::SetStr(ascbc_node->GetOpDescBarePtr(), "ascgraph", sub_graph_str);
13387+ return compute_graph;
13388+}
13389+ 
13390+static void CreateRandnStoreAscGraph(af::AscGraph &graph, size_t dims_size, ge::DataType dtype) {
13391+ auto s0 = graph.CreateSizeVar("s0");
13392+ auto z0 = graph.CreateAxis("z0", s0);
13393+ auto One = Symbol(1);
13394+ 
13395+ af::ascir_op::Randn randn_op("randn");
13396+ graph.AddNode(randn_op);
13397+ *randn_op.y.axis = {z0.id};
13398+ *randn_op.y.repeats = {s0};
13399+ *randn_op.y.strides = {One};
13400+ randn_op.y.dtype = dtype;
13401+ 
13402+ randn_op.attr.api.compute_type = af::ComputeType::kComputeElewise;
13403+ randn_op.attr.api.type = af::ApiType::kAPITypeCompute;
13404+ randn_op.attr.api.unit = af::ComputeUnit::kUnitVector;
13405+ randn_op.attr.sched.axis = {z0.id};
13406+ 
13407+ af::ascir_op::Store store("store");
13408+ store.x = randn_op.y;
13409+ store.attr.api.compute_type = ComputeType::kComputeStore;
13410+ store.attr.api.type = af::ApiType::kAPITypeCompute;
13411+ store.attr.api.unit = ComputeUnit::kUnitMTE3;
13412+ store.attr.sched.axis = {z0.id};
13413+ *store.y.axis = {z0.id};
13414+ *store.y.repeats = {s0};
13415+ *store.y.strides = {One};
13416+ store.y.dtype = dtype;
13417+ 
13418+ af::ascir_op::Output y("output");
13419+ y.x = store.y;
13420+ y.ir_attr.SetIndex(0);
13421+}
13422+ 
13423+af::ComputeGraphPtr ShareGraph::RandnStoreFusedGraph(size_t dims_size) {
13424+ auto builder = GraphBuilder("randn_store_test");
13425+ auto ascbc = builder.AddNode("ascbc", "AscGraph", 0, 1);
13426+ auto netoutput = builder.AddNode("netoutput", af::NETOUTPUT, 1, 0);
13427+ 
13428+ builder.AddDataEdge(ascbc, 0, netoutput, 0);
13429+ ComputeGraphPtr compute_graph = builder.GetGraph();
13430+ if (compute_graph == nullptr) {
13431+ return nullptr;
13432+ }
13433+ auto ascbc_node = compute_graph->FindNode("ascbc");
13434+ af::AscGraph sub_graph("randn_store");
13435+ CreateRandnStoreAscGraph(sub_graph, dims_size, ge::DT_UINT32);
13436+ 
13437+ std::string sub_graph_str;
13438+ af::AscGraphUtils::SerializeToReadable(sub_graph, sub_graph_str);
13439+ af::AttrUtils::SetStr(ascbc_node->GetOpDescBarePtr(), "ascgraph", sub_graph_str);
13440+ return compute_graph;
13441+}
13442+ 
13338} // namespace ascir13443} // namespace ascir
Mautofuse/tests/v35/CMakeLists.txt+3-0
@@ -34,6 +34,9 @@ target_include_directories(test_codegen PRIVATE
34 ${CODE_ROOT_DIR}/v35/codegen/micro_api_call34 ${CODE_ROOT_DIR}/v35/codegen/micro_api_call
35 ${CODE_ROOT_DIR}/v35/codegen/reg_api_call35 ${CODE_ROOT_DIR}/v35/codegen/reg_api_call
36 ${CODE_ROOT_DIR}/v35/codegen/vec_func_call36 ${CODE_ROOT_DIR}/v35/codegen/vec_func_call
37+ ${CODE_ROOT_DIR}
38+ ${CODE_ROOT_DIR}/ascir
39+ ${CODE_ROOT_DIR}/v35/ascir
37)40)
38 41 
39target_sources(test_codegen PRIVATE42target_sources(test_codegen PRIVATE
Mautofuse/tests/v35/st/backend_e2e_v2/CMakeLists.txt+2-0
@@ -215,3 +215,5 @@ add_subdirectory(expm_test)
215add_subdirectory(round_to_int_float_to_int32_test)215add_subdirectory(round_to_int_float_to_int32_test)
216add_subdirectory(trunc_to_int_bf16_to_int32_test)216add_subdirectory(trunc_to_int_bf16_to_int32_test)
217# add_subdirectory(remainder_bf16_test)217# add_subdirectory(remainder_bf16_test)
218+add_subdirectory(rand_store_test)
219+add_subdirectory(randn_store_test)
Aautofuse/tests/v35/st/backend_e2e_v2/rand_store_test/CMakeLists.txt+8-0
@@ -0,0 +1,8 @@
1+backend_e2e_st_test(rand_store_test
2+ CODEGEN rand_store_backend_generator.cpp
3+ KERNEL_SRC
4+ rand_store_test_kernel.cpp
5+ rand_store_test_tiling.cpp
6+ autofuse_tiling_data.h
7+ TEST_SRC test_e2e_rand_store_kernel.cpp)
8+target_include_directories(rand_store_test_e2e_v2 PRIVATE ${CODE_ROOT_DIR}/v35/ascendc/api_regbase)
Aautofuse/tests/v35/st/backend_e2e_v2/rand_store_test/rand_store_backend_generator.cpp+79-0
@@ -0,0 +1,79 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under
4+ * the terms and conditions of CANN Open Software License Agreement Version 2.0
5+ * (the "License"). Please refer to the License for details. You may not use
6+ * this file except in compliance with the License. THIS SOFTWARE IS PROVIDED ON
7+ * AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS
9+ * FOR A PARTICULAR PURPOSE. See LICENSE in the root of the software repository
10+ * for the full text of the License.
11+ */
12+ 
13+#include <exception>
14+#include <fstream>
15+#include <gtest/gtest.h>
16+#include <iostream>
17+#include <map>
18+#include <string>
19+#include <vector>
20+ 
21+#include "backend_common.h"
22+#include "codegen.h"
23+#include "common/platform_context.h"
24+#include "optimize.h"
25+#include "runtime_stub.h"
26+#include "share_graph.h"
27+ 
28+class TestBackendRandStoreE2e : public testing::Test {
29+ protected:
30+ void SetUp() override {
31+ dlog_setlevel(ASCGEN_MODULE_NAME, DLOG_ERROR, 0);
32+ ge::PlatformContext::GetInstance().Reset();
33+ auto rand_stub_v2 = std::make_shared<af::RuntimeStubV2>();
34+ ge::RuntimeStub::SetInstance(rand_stub_v2);
35+ }
36+ void TearDown() override {
37+ dlog_setlevel(ASCGEN_MODULE_NAME, DLOG_ERROR, 0);
38+ ge::RuntimeStub::Reset();
39+ }
40+};
41+ 
42+TEST_F(TestBackendRandStoreE2e, RandStoreE2eCodegen) {
43+ bool gen_success = true;
44+ std::string rand_tiling_stub = R"(
45+#define REGISTER_TILING_DEFAULT(tiling)
46+#define GET_TILING_DATA(t, tiling) AutofuseTilingData t = *(AutofuseTilingData*)tiling;
47+)";
48+ 
49+ std::map<std::string, std::string> rand_shape_info({{"s0", "stub_s0"}});
50+ auto graph = ascir::ShareGraph::RandStoreFusedGraph(1);
51+ std::vector<std::string> rand_parts = splitString(KERNEL_SRC_LIST, ':');
52+ const std::string &kernel_src_file_name = rand_parts[0];
53+ const std::string &tiling_src_file_name = rand_parts[1];
54+ const std::string &tiling_data_src_file_name = rand_parts[2];
55+ 
56+ try {
57+ optimize::Optimizer rand_optimizer(optimize::OptimizerOptions{});
58+ codegen::Codegen rand_codegen(codegen::CodegenOptions{});
59+ 
60+ std::fstream rand_kernel_stream(kernel_src_file_name, std::ios::out);
61+ std::fstream rand_tiling_stream(tiling_src_file_name, std::ios::out);
62+ std::fstream rand_data_stream(tiling_data_src_file_name, std::ios::out);
63+ 
64+ std::vector<::ascir::ScheduledResult> rand_schedules;
65+ ascir::FusedScheduledResult rand_fused_result;
66+ rand_fused_result.node_idx_to_scheduled_results.push_back(rand_schedules);
67+ EXPECT_EQ(rand_optimizer.Optimize(graph, rand_fused_result), 0);
68+ codegen::CodegenResult rand_result;
69+ EXPECT_EQ(rand_codegen.Generate(rand_shape_info, rand_fused_result, rand_result), 0);
70+ EXPECT_NE(rand_result.kernel.find("Rand("), std::string::npos);
71+ rand_kernel_stream << rand_tiling_stub << RemoveSubDirInclude(rand_result.kernel);
72+ rand_tiling_stream << rand_result.tiling;
73+ rand_data_stream << rand_result.tiling_data;
74+ } catch (...) {
75+ gen_success = false;
76+ }
77+ 
78+ EXPECT_EQ(gen_success, true);
79+}
Aautofuse/tests/v35/st/backend_e2e_v2/rand_store_test/test_e2e_rand_store_kernel.cpp+60-0
@@ -0,0 +1,60 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under
4+ * the terms and conditions of CANN Open Software License Agreement Version 2.0
5+ * (the "License"). Please refer to the License for details. You may not use
6+ * this file except in compliance with the License. THIS SOFTWARE IS PROVIDED ON
7+ * AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS
9+ * FOR A PARTICULAR PURPOSE. See LICENSE in the root of the software repository
10+ * for the full text of the License.
11+ */
12+ 
13+#include <cmath>
14+#include <gtest/gtest.h>
15+ 
16+#include "autofuse_tiling_data.h"
17+#include "tikicpulib.h"
18+ 
19+extern "C" __global__ __aicore__ void rand_store_test(GM_ADDR y1, GM_ADDR workspace, GM_ADDR tiling);
20+extern "C" int64_t AutofuseTiling(uint32_t s0, AutofuseTilingData *tiling, uint32_t *workspaceSize, uint64_t *blockDim,
21+ uint32_t aiv_num, uint32_t ub_size);
22+ 
23+class E2EBackendRandStoreCode : public testing::Test, public testing::WithParamInterface<int> {};
24+ 
25+TEST_P(E2EBackendRandStoreCode, GenerateRandomNumbers) {
26+ auto test_size = GetParam();
27+ uint64_t block_dim = 48;
28+ 
29+ AutofuseTilingData tiling_data;
30+ float *y = static_cast<float *>(AscendC::GmAlloc(test_size * sizeof(float) + 32));
31+ 
32+ uint32_t ws_size = 0;
33+ AutofuseTiling(test_size, &tiling_data, &ws_size, &block_dim, 48, 192 * 1024);
34+ 
35+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
36+ ICPU_RUN_KF(rand_store_test, tiling_data.block_dim, reinterpret_cast<uint8_t *>(y), nullptr,
37+ reinterpret_cast<uint8_t *>(&tiling_data));
38+ 
39+ // Verify generated random numbers are in valid range [0, 1)
40+ uint32_t invalid_count = 0;
41+ uint32_t zero_count = 0;
42+ for (int i = 0; i < test_size; i++) {
43+ if (y[i] < 0.0F || y[i] >= 1.0F) {
44+ invalid_count++;
45+ }
46+ if (y[i] == 0.0F) {
47+ zero_count++;
48+ }
49+ }
50+ 
51+ // All values should be in [0, 1)
52+ EXPECT_EQ(invalid_count, 0U) << "Generated values outside [0, 1) range";
53+ 
54+ // Since it's random, we expect some non-zero values (extremely unlikely all are zero)
55+ EXPECT_LT(zero_count, test_size) << "All generated values are zero (unlikely for random generator)";
56+ 
57+ AscendC::GmFree(y);
58+}
59+ 
60+INSTANTIATE_TEST_SUITE_P(GenerateWithDifferentSize, E2EBackendRandStoreCode, ::testing::Values(32, 64, 128, 256, 512));
Aautofuse/tests/v35/st/backend_e2e_v2/randn_store_test/CMakeLists.txt+8-0
@@ -0,0 +1,8 @@
1+backend_e2e_st_test(randn_store_test
2+ CODEGEN randn_store_backend_generator.cpp
3+ KERNEL_SRC
4+ randn_store_test_kernel.cpp
5+ randn_store_test_tiling.cpp
6+ autofuse_tiling_data.h
7+ TEST_SRC test_e2e_randn_store_kernel.cpp)
8+target_include_directories(randn_store_test_e2e_v2 PRIVATE ${CODE_ROOT_DIR}/v35/ascendc/api_regbase)
Aautofuse/tests/v35/st/backend_e2e_v2/randn_store_test/randn_store_backend_generator.cpp+79-0
@@ -0,0 +1,79 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under
4+ * the terms and conditions of CANN Open Software License Agreement Version 2.0
5+ * (the "License"). Please refer to the License for details. You may not use
6+ * this file except in compliance with the License. THIS SOFTWARE IS PROVIDED ON
7+ * AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS
9+ * FOR A PARTICULAR PURPOSE. See LICENSE in the root of the software repository
10+ * for the full text of the License.
11+ */
12+ 
13+#include <exception>
14+#include <fstream>
15+#include <gtest/gtest.h>
16+#include <iostream>
17+#include <map>
18+#include <string>
19+#include <vector>
20+ 
21+#include "backend_common.h"
22+#include "codegen.h"
23+#include "common/platform_context.h"
24+#include "optimize.h"
25+#include "runtime_stub.h"
26+#include "share_graph.h"
27+ 
28+class TestBackendRandnStoreE2e : public testing::Test {
29+ protected:
30+ void SetUp() override {
31+ dlog_setlevel(ASCGEN_MODULE_NAME, DLOG_ERROR, 0);
32+ ge::PlatformContext::GetInstance().Reset();
33+ auto randn_stub_v2 = std::make_shared<af::RuntimeStubV2>();
34+ ge::RuntimeStub::SetInstance(randn_stub_v2);
35+ }
36+ void TearDown() override {
37+ dlog_setlevel(ASCGEN_MODULE_NAME, DLOG_ERROR, 0);
38+ ge::RuntimeStub::Reset();
39+ }
40+};
41+ 
42+TEST_F(TestBackendRandnStoreE2e, RandnStoreE2eCodegen) {
43+ bool gen_success = true;
44+ std::string randn_tiling_stub = R"(
45+#define REGISTER_TILING_DEFAULT(tiling)
46+#define GET_TILING_DATA(t, tiling) AutofuseTilingData t = *(AutofuseTilingData*)tiling;
47+)";
48+ 
49+ std::map<std::string, std::string> randn_shape_info({{"s0", "stub_s0"}});
50+ auto graph = ascir::ShareGraph::RandnStoreFusedGraph(1);
51+ std::vector<std::string> randn_parts = splitString(KERNEL_SRC_LIST, ':');
52+ const std::string &kernel_src_file_name = randn_parts[0];
53+ const std::string &tiling_src_file_name = randn_parts[1];
54+ const std::string &tiling_data_src_file_name = randn_parts[2];
55+ 
56+ try {
57+ optimize::Optimizer randn_optimizer(optimize::OptimizerOptions{});
58+ codegen::Codegen randn_codegen(codegen::CodegenOptions{});
59+ 
60+ std::fstream randn_kernel_stream(kernel_src_file_name, std::ios::out);
61+ std::fstream randn_tiling_stream(tiling_src_file_name, std::ios::out);
62+ std::fstream randn_data_stream(tiling_data_src_file_name, std::ios::out);
63+ 
64+ std::vector<::ascir::ScheduledResult> randn_schedules;
65+ ascir::FusedScheduledResult randn_fused_result;
66+ randn_fused_result.node_idx_to_scheduled_results.push_back(randn_schedules);
67+ EXPECT_EQ(randn_optimizer.Optimize(graph, randn_fused_result), 0);
68+ codegen::CodegenResult randn_result;
69+ EXPECT_EQ(randn_codegen.Generate(randn_shape_info, randn_fused_result, randn_result), 0);
70+ EXPECT_NE(randn_result.kernel.find("Rand("), std::string::npos);
71+ randn_kernel_stream << randn_tiling_stub << RemoveSubDirInclude(randn_result.kernel);
72+ randn_tiling_stream << randn_result.tiling;
73+ randn_data_stream << randn_result.tiling_data;
74+ } catch (...) {
75+ gen_success = false;
76+ }
77+ 
78+ EXPECT_EQ(gen_success, true);
79+}
Aautofuse/tests/v35/st/backend_e2e_v2/randn_store_test/test_e2e_randn_store_kernel.cpp+54-0
@@ -0,0 +1,54 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under
4+ * the terms and conditions of CANN Open Software License Agreement Version 2.0
5+ * (the "License"). Please refer to the License for details. You may not use
6+ * this file except in compliance with the License. THIS SOFTWARE IS PROVIDED ON
7+ * AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS
9+ * FOR A PARTICULAR PURPOSE. See LICENSE in the root of the software repository
10+ * for the full text of the License.
11+ */
12+ 
13+#include <gtest/gtest.h>
14+#include <limits>
15+ 
16+#include "autofuse_tiling_data.h"
17+#include "tikicpulib.h"
18+ 
19+extern "C" __global__ __aicore__ void randn_store_test(GM_ADDR y1, GM_ADDR workspace, GM_ADDR tiling);
20+extern "C" int64_t AutofuseTiling(uint32_t s0, AutofuseTilingData *tiling, uint32_t *workspaceSize, uint64_t *blockDim,
21+ uint32_t aiv_num, uint32_t ub_size);
22+ 
23+class E2EBackendRandnStoreCode : public testing::Test, public testing::WithParamInterface<int> {};
24+ 
25+TEST_P(E2EBackendRandnStoreCode, GenerateRandomIntegers) {
26+ auto test_size = GetParam();
27+ uint64_t block_dim = 48;
28+ 
29+ AutofuseTilingData tiling_data;
30+ uint32_t *y = static_cast<uint32_t *>(AscendC::GmAlloc(test_size * sizeof(uint32_t) + 32));
31+ 
32+ uint32_t ws_size = 0;
33+ AutofuseTiling(test_size, &tiling_data, &ws_size, &block_dim, 48, 192 * 1024);
34+ 
35+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
36+ ICPU_RUN_KF(randn_store_test, tiling_data.block_dim, reinterpret_cast<uint8_t *>(y), nullptr,
37+ reinterpret_cast<uint8_t *>(&tiling_data));
38+ 
39+ // Verify generated random integers are in valid uint32_t range
40+ uint32_t zero_count = 0;
41+ for (int i = 0; i < test_size; i++) {
42+ // All values should be valid uint32_t (automatically true by type)
43+ if (y[i] == 0U) {
44+ zero_count++;
45+ }
46+ }
47+ 
48+ // Since it's random, we expect some non-zero values (extremely unlikely all are zero)
49+ EXPECT_LT(zero_count, test_size) << "All generated values are zero (unlikely for random generator)";
50+ 
51+ AscendC::GmFree(y);
52+}
53+ 
54+INSTANTIATE_TEST_SUITE_P(GenerateWithDifferentSize, E2EBackendRandnStoreCode, ::testing::Values(32, 64, 128, 256, 512));
Aautofuse/tests/v35/ut/codegen/reg_api_call/test_codegen_rand_reg_api_call.cpp+249-0
@@ -0,0 +1,249 @@
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+#include "gtest/gtest.h"
12+ 
13+#include "ascendc_ir.h"
14+#include "ascir_ops.h"
15+#include "ascir_ops_utils.h"
16+#include "ascir/generator/v2_ascir_codegen_impl.h"
17+#include "codegen_kernel.h"
18+#include "common_utils.h"
19+#include "utils/api_call_factory.h"
20+#include "unary_output_api_call.h"
21+ 
22+using namespace ge;
23+using namespace af::ops;
24+using namespace af::ascir_op;
25+ 
26+namespace codegen {
27+namespace {
28+ 
29+// 创建 Rand 算子的测试图
30+void BuildRandGraph(af::AscGraph &graph, const af::Expression &s0, const af::Axis &z0, ge::DataType dtype) {
31+ Rand rand_op("Rand");
32+ graph.AddNode(rand_op);
33+ 
34+ *rand_op.y.axis = {z0.id};
35+ *rand_op.y.repeats = {s0};
36+ *rand_op.y.strides = {One};
37+}
38+ 
39+void InitRandAttrs(af::AscGraph &graph, ge::DataType dtype, const af::Axis &z0) {
40+ auto rand = graph.FindNode("Rand");
41+ rand->attr.api.compute_type = af::ComputeType::kComputeElewise;
42+ rand->attr.api.type = af::ApiType::kAPITypeCompute;
43+ rand->attr.api.unit = af::ComputeUnit::kUnitVector;
44+ rand->attr.sched.loop_axis = z0.id;
45+ 
46+ auto &rand_attr = rand->outputs[0].attr;
47+ rand_attr.vectorized_axis = {z0.id};
48+ rand_attr.vectorized_strides = {One};
49+ rand_attr.dtype = dtype;
50+ rand_attr.mem.position = af::Position::kPositionVecOut;
51+ rand_attr.mem.tensor_id = 0;
52+ rand_attr.mem.alloc_type = af::AllocType::kAllocTypeQueue;
53+ rand_attr.que.id = 1;
54+ rand_attr.opt.merge_scope = af::kIdNone;
55+}
56+ 
57+// 创建 Randn 算子的测试图
58+void BuildRandnGraph(af::AscGraph &graph, const af::Expression &s0, const af::Axis &z0, ge::DataType dtype) {
59+ Randn randn_op("Randn");
60+ graph.AddNode(randn_op);
61+ 
62+ *randn_op.y.axis = {z0.id};
63+ *randn_op.y.repeats = {s0};
64+ *randn_op.y.strides = {One};
65+}
66+ 
67+void InitRandnAttrs(af::AscGraph &graph, ge::DataType dtype, const af::Axis &z0) {
68+ auto randn = graph.FindNode("Randn");
69+ randn->attr.api.compute_type = af::ComputeType::kComputeElewise;
70+ randn->attr.api.type = af::ApiType::kAPITypeCompute;
71+ randn->attr.api.unit = af::ComputeUnit::kUnitVector;
72+ randn->attr.sched.loop_axis = z0.id;
73+ 
74+ auto &randn_attr = randn->outputs[0].attr;
75+ randn_attr.vectorized_axis = {z0.id};
76+ randn_attr.vectorized_strides = {One};
77+ randn_attr.dtype = dtype;
78+ randn_attr.mem.position = af::Position::kPositionVecOut;
79+ randn_attr.mem.tensor_id = 0;
80+ randn_attr.mem.alloc_type = af::AllocType::kAllocTypeQueue;
81+ randn_attr.que.id = 1;
82+ randn_attr.opt.merge_scope = af::kIdNone;
83+}
84+ 
85+} // namespace
86+ 
87+// 测试 Rand Codegen 基本信息
88+TEST(CodegenKernel, TestRandCodegenBasicInfo) {
89+ af::AscGraph graph("test_rand_graph");
90+ 
91+ auto s0 = graph.CreateSizeVar("s0");
92+ auto z0 = graph.CreateAxis("z0", s0);
93+ BuildRandGraph(graph, s0, z0, DT_FLOAT);
94+ InitRandAttrs(graph, DT_FLOAT, z0);
95+ 
96+ auto rand = graph.FindNode("Rand");
97+ 
98+ // 创建 Codegen 实例并测试基本信息
99+ af::ascir::RandAscIrCodegenImplV2 codegen;
100+ 
101+ // 测试 GetApiName - Rand 返回 "Rand"
102+ std::string api_name = codegen.GetApiName();
103+ EXPECT_EQ(api_name, "Rand");
104+ 
105+ // 测试 GetApiCallName
106+ std::string api_call_name = codegen.GetApiCallName();
107+ EXPECT_EQ(api_call_name, "UnaryOutputApiCall");
108+ 
109+ // 测试 LoadApiHeaderFiles
110+ auto headers = codegen.LoadApiHeaderFiles(false);
111+ EXPECT_EQ(headers.size(), 1);
112+ EXPECT_EQ(headers[0], "random_reg_base.h");
113+ 
114+ // 测试 IncludeApiHeaderFiles（应该包含 AscendC Philox 头文件）
115+ auto include_headers = codegen.IncludeApiHeaderFiles();
116+ EXPECT_EQ(include_headers.size(), 1);
117+ EXPECT_EQ(include_headers[0], "adv_api/math/philox.h");
118+}
119+ 
120+// 测试 Randn Codegen 基本信息
121+TEST(CodegenKernel, TestRandnCodegenBasicInfo) {
122+ af::AscGraph graph("test_randn_graph");
123+ 
124+ auto s0 = graph.CreateSizeVar("s0");
125+ auto z0 = graph.CreateAxis("z0", s0);
126+ BuildRandnGraph(graph, s0, z0, DT_UINT32);
127+ InitRandnAttrs(graph, DT_UINT32, z0);
128+ 
129+ auto randn = graph.FindNode("Randn");
130+ 
131+ // 创建 Codegen 实例并测试基本信息
132+ af::ascir::RandnAscIrCodegenImplV2 codegen;
133+ 
134+ // 测试 GetApiName - Randn 也返回 "Rand"
135+ std::string api_name = codegen.GetApiName();
136+ EXPECT_EQ(api_name, "Rand");
137+ 
138+ // 测试 GetApiCallName
139+ std::string api_call_name = codegen.GetApiCallName();
140+ EXPECT_EQ(api_call_name, "UnaryOutputApiCall");
141+ 
142+ // 测试 LoadApiHeaderFiles
143+ auto headers = codegen.LoadApiHeaderFiles(false);
144+ EXPECT_EQ(headers.size(), 1);
145+ EXPECT_EQ(headers[0], "random_reg_base.h");
146+ 
147+ // 测试 IncludeApiHeaderFiles（应该包含 AscendC Philox 头文件）
148+ auto include_headers = codegen.IncludeApiHeaderFiles();
149+ EXPECT_EQ(include_headers.size(), 1);
150+ EXPECT_EQ(include_headers[0], "adv_api/math/philox.h");
151+}
152+ 
153+// 测试 Rand 不同 dtype 支持
154+TEST(CodegenKernel, TestRandDifferentDtypes) {
155+ std::vector<ge::DataType> supported_dtypes = {DT_FLOAT};
156+ 
157+ for (auto dtype : supported_dtypes) {
158+ af::AscGraph graph("test_rand_dtype_graph");
159+ 
160+ auto s0 = graph.CreateSizeVar("s0");
161+ auto z0 = graph.CreateAxis("z0", s0);
162+ BuildRandGraph(graph, s0, z0, dtype);
163+ InitRandAttrs(graph, dtype, z0);
164+ 
165+ auto rand = graph.FindNode("Rand");
166+ EXPECT_EQ(rand->outputs[0].attr.dtype, dtype);
167+ }
168+}
169+ 
170+// 测试 Randn 不同 dtype 支持
171+TEST(CodegenKernel, TestRandnDifferentDtypes) {
172+ std::vector<ge::DataType> supported_dtypes = {DT_UINT32, DT_INT32};
173+ 
174+ for (auto dtype : supported_dtypes) {
175+ af::AscGraph graph("test_randn_dtype_graph");
176+ 
177+ auto s0 = graph.CreateSizeVar("s0");
178+ auto z0 = graph.CreateAxis("z0", s0);
179+ BuildRandnGraph(graph, s0, z0, dtype);
180+ InitRandnAttrs(graph, dtype, z0);
181+ 
182+ auto randn = graph.FindNode("Randn");
183+ EXPECT_EQ(randn->outputs[0].attr.dtype, dtype);
184+ }
185+}
186+ 
187+// 测试 Rand 生成的 API 调用字符串（简化版）
188+TEST(CodegenKernel, TestRandApiCallStringGeneration) {
189+ af::AscGraph graph("test_rand_api_call_graph");
190+ 
191+ auto s0 = graph.CreateSizeVar("s0");
192+ auto z0 = graph.CreateAxis("z0", s0);
193+ BuildRandGraph(graph, s0, z0, DT_FLOAT);
194+ InitRandAttrs(graph, DT_FLOAT, z0);
195+ 
196+ auto rand = graph.FindNode("Rand");
197+ 
198+ codegen::Tiler tiler;
199+ codegen::TPipe tpipe("tpipe", tiler);
200+ tpipe.AddTensor(rand->outputs[0]);
201+ 
202+ tiler.AddAxis(z0);
203+ tiler.AddSizeVar(af::SizeVar(s0));
204+ 
205+ // 使用 UnaryOutputApiCall 来生成调用字符串
206+ codegen::UnaryOutputApiCall call("Rand");
207+ EXPECT_EQ(call.Init(rand), 0);
208+ 
209+ std::string result;
210+ call.Generate(tpipe, {}, result);
211+ 
212+ // 验证生成的字符串包含关键信息
213+ EXPECT_NE(result.find("Rand"), std::string::npos);
214+ EXPECT_NE(result.find("local_0"), std::string::npos);
215+ EXPECT_NE(result.find("local_0_actual_size"), std::string::npos);
216+}
217+ 
218+// 测试 Randn 生成的 API 调用字符串
219+TEST(CodegenKernel, TestRandnApiCallStringGeneration) {
220+ af::AscGraph graph("test_randn_api_call_graph");
221+ 
222+ auto s0 = graph.CreateSizeVar("s0");
223+ auto z0 = graph.CreateAxis("z0", s0);
224+ BuildRandnGraph(graph, s0, z0, DT_UINT32);
225+ InitRandnAttrs(graph, DT_UINT32, z0);
226+ 
227+ auto randn = graph.FindNode("Randn");
228+ 
229+ codegen::Tiler tiler;
230+ codegen::TPipe tpipe("tpipe", tiler);
231+ tpipe.AddTensor(randn->outputs[0]);
232+ 
233+ tiler.AddAxis(z0);
234+ tiler.AddSizeVar(af::SizeVar(s0));
235+ 
236+ // 使用 UnaryOutputApiCall 来生成调用字符串
237+ codegen::UnaryOutputApiCall call("Rand");
238+ EXPECT_EQ(call.Init(randn), 0);
239+ 
240+ std::string result;
241+ call.Generate(tpipe, {}, result);
242+ 
243+ // 验证生成的字符串包含关键信息
244+ EXPECT_NE(result.find("Rand"), std::string::npos); // Randn 也调用 Rand
245+ EXPECT_NE(result.find("local_0"), std::string::npos);
246+ EXPECT_NE(result.find("local_0_actual_size"), std::string::npos);
247+}
248+ 
249+} // namespace codegen
Mautofuse/v35/ascendc/api_regbase/CMakeLists.txt+1-0
@@ -15,6 +15,7 @@ set(ascendc_api_regbase_extend_src
15 floor_div.h15 floor_div.h
16 pow.h16 pow.h
17 polygamma.h17 polygamma.h
18+ random.h
18 split.h19 split.h
19 where.h20 where.h
20 erf.h21 erf.h
Aautofuse/v35/ascendc/api_regbase/random.h+31-0
@@ -0,0 +1,31 @@
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+#ifndef __ASCENDC_API_REGBASE_RANDOM_H__
11+#define __ASCENDC_API_REGBASE_RANDOM_H__
12+ 
13+/**
14+ * @brief Rand - unified random number generation using AscendC PhiloxRandom API
15+ * Uses fixed philoxKey = {0, 0} and philoxCounter = {0, 0, 0, 0} for continuous mode
16+ * Supports float, float16, uint32_t, int32_t data types
17+ *
18+ * @tparam T data type, supports float, float16, uint32_t, int32_t
19+ * @param dst output tensor
20+ * @param size number of elements to generate
21+ */
22+template <typename T>
23+inline __aicore__ void Rand(const AscendC::LocalTensor<T> &dst, const uint32_t size) {
24+ // PhiloxKey and PhiloxCounter are plain C arrays: uint32_t[2] and uint32_t[4]
25+ AscendC::PhiloxKey philoxKey = {0, 0};
26+ AscendC::PhiloxCounter philoxCounter = {0, 0, 0, 0};
27+ 
28+ AscendC::PhiloxRandom<10>(dst, philoxKey, philoxCounter, size);
29+}
30+ 
31+#endif // __ASCENDC_API_REGBASE_RANDOM_H__
Mautofuse/v35/ascir/generator/ascir_builtin_ops_v2.cpp+10-0
@@ -501,6 +501,16 @@ REG_ASC_IR(Output).Impl(v2_soc_versions,
501 {{"T", TensorType{DT_INT8, DT_UINT8, DT_INT16, DT_UINT16, DT_INT32, DT_UINT32, DT_INT64,501 {{"T", TensorType{DT_INT8, DT_UINT8, DT_INT16, DT_UINT16, DT_INT32, DT_UINT32, DT_INT64,
502 DT_UINT64, DT_FLOAT16, DT_FLOAT, DT_BF16, DT_BOOL}}}});502 DT_UINT64, DT_FLOAT16, DT_FLOAT, DT_BF16, DT_BOOL}}}});
503 503 
504+REG_ASC_IR(Rand).Inputs({}).Output("y", "T").StartNode().Impl(
505+ v2_soc_versions, {af::ascir::AscIrImplCreator<af::ascir::RandAscIrAttImplV2>(),
506+ af::ascir::AscIrImplCreator<af::ascir::RandAscIrCodegenImplV2>(),
507+ {{"T", TensorType{DT_FLOAT}}}});
508+ 
509+REG_ASC_IR(Randn).Inputs({}).Output("y", "T").StartNode().Impl(
510+ v2_soc_versions, {af::ascir::AscIrImplCreator<af::ascir::RandnAscIrAttImplV2>(),
511+ af::ascir::AscIrImplCreator<af::ascir::RandnAscIrCodegenImplV2>(),
512+ {{"T", TensorType{DT_UINT32, DT_INT32}}}});
513+ 
504REG_ASC_IR(Workspace).Impl(v2_soc_versions,514REG_ASC_IR(Workspace).Impl(v2_soc_versions,
505 {af::ascir::AscIrImplCreator<af::ascir::WorkspaceAscIrAttImplV2>(),515 {af::ascir::AscIrImplCreator<af::ascir::WorkspaceAscIrAttImplV2>(),
506 af::ascir::AscIrImplCreator<af::ascir::WorkspaceAscIrCodegenImplV2>(),516 af::ascir::AscIrImplCreator<af::ascir::WorkspaceAscIrCodegenImplV2>(),
Mautofuse/v35/ascir/generator/v2_ascir_att_impl.h+2-0
@@ -199,6 +199,8 @@ REG_ASC_IR_ATT_V2_CLASS_DEFINE(LegendrePolynomialP);
199REG_ASC_IR_ATT_V2_CLASS_DEFINE(AiryAi);199REG_ASC_IR_ATT_V2_CLASS_DEFINE(AiryAi);
200REG_ASC_IR_ATT_V2_CLASS_DEFINE(Erfinv);200REG_ASC_IR_ATT_V2_CLASS_DEFINE(Erfinv);
201REG_ASC_IR_ATT_V2_CLASS_DEFINE(Softmax);201REG_ASC_IR_ATT_V2_CLASS_DEFINE(Softmax);
202+REG_ASC_IR_ATT_V2_CLASS_DEFINE(Rand);
203+REG_ASC_IR_ATT_V2_CLASS_DEFINE(Randn);
202} // namespace ascir204} // namespace ascir
203} // namespace af205} // namespace af
204 206 
Mautofuse/v35/ascir/generator/v2_ascir_codegen_impl.h+34-0
@@ -5300,6 +5300,40 @@ class SoftmaxAscIrCodegenImplV2 : public AscIrCodegenV2 {
5300 }5300 }
5301};5301};
5302 5302 
5303+/*********************************************************************************/
5304+class RandAscIrCodegenImplV2 : public AscIrCodegenV2 {
5305+ public:
5306+ [[nodiscard]] std::string GetApiCallName() const override {
5307+ return "UnaryOutputApiCall";
5308+ }
5309+ [[nodiscard]] std::string GetApiName() const override {
5310+ return "Rand";
5311+ }
5312+ [[nodiscard]] std::vector<std::string> LoadApiHeaderFiles([[maybe_unused]] bool is_dynamic) const override {
5313+ return {"random_reg_base.h"};
5314+ }
5315+ [[nodiscard]] std::vector<std::string> IncludeApiHeaderFiles() const override {
5316+ return {"adv_api/math/philox.h"};
5317+ }
5318+};
5319+ 
5320+/*********************************************************************************/
5321+class RandnAscIrCodegenImplV2 : public AscIrCodegenV2 {
5322+ public:
5323+ [[nodiscard]] std::string GetApiCallName() const override {
5324+ return "UnaryOutputApiCall";
5325+ }
5326+ [[nodiscard]] std::string GetApiName() const override {
5327+ return "Rand";
5328+ }
5329+ [[nodiscard]] std::vector<std::string> LoadApiHeaderFiles([[maybe_unused]] bool is_dynamic) const override {
5330+ return {"random_reg_base.h"};
5331+ }
5332+ [[nodiscard]] std::vector<std::string> IncludeApiHeaderFiles() const override {
5333+ return {"adv_api/math/philox.h"};
5334+ }
5335+};
5336+ 
5303/*********************************************************************************/5337/*********************************************************************************/
5304class UnsupportedAscIrCodegenImplV2 : public AscIrCodegenV2 {5338class UnsupportedAscIrCodegenImplV2 : public AscIrCodegenV2 {
5305 public:5339 public:
Mautofuse/v35/att/api_perf_register/ascir_api_perf_v2.cpp+4-0
@@ -914,6 +914,10 @@ ApiPerfRegister<ApiPerf> airy_ai_api_perf_v2(ApiPerfRegisterV2(kAiryAi, GetPerfF
914 &perf_param_table_v2, &tiling_schedule_config_table_v2));914 &perf_param_table_v2, &tiling_schedule_config_table_v2));
915ApiPerfRegister<ApiPerf> erfinv_api_perf_v2(ApiPerfRegisterV2(kErfinv, GetPerfFunc(kUnitVector), nullptr,915ApiPerfRegister<ApiPerf> erfinv_api_perf_v2(ApiPerfRegisterV2(kErfinv, GetPerfFunc(kUnitVector), nullptr,
916 &perf_param_table_v2, &tiling_schedule_config_table_v2));916 &perf_param_table_v2, &tiling_schedule_config_table_v2));
917+ApiPerfRegister<ApiPerf> rand_api_perf_v2(ApiPerfRegisterV2(kRand, GetPerfFunc(kUnitVector), nullptr,
918+ &perf_param_table_v2, &tiling_schedule_config_table_v2));
919+ApiPerfRegister<ApiPerf> randn_api_perf_v2(ApiPerfRegisterV2(kRandn, GetPerfFunc(kUnitVector), nullptr,
920+ &perf_param_table_v2, &tiling_schedule_config_table_v2));
917ApiPerfRegister<ApiPerf> square_api_perf_v2(ApiPerfRegisterV2(kSquare, GetPerfFunc(kSquare + "V2"), nullptr,921ApiPerfRegister<ApiPerf> square_api_perf_v2(ApiPerfRegisterV2(kSquare, GetPerfFunc(kSquare + "V2"), nullptr,
918 &perf_param_table_v2, &tiling_schedule_config_table_v2));922 &perf_param_table_v2, &tiling_schedule_config_table_v2));
919 923 
Mautofuse/v35/codegen/ascendc_reg_base_api_register.cpp+4-0
@@ -239,6 +239,9 @@ Register::Register() {
239 };239 };
240 const std::string kAscendcPolyGammaRegBaseStr = {240 const std::string kAscendcPolyGammaRegBaseStr = {
241#include "polygamma_reg_base.h"241#include "polygamma_reg_base.h"
242+ };
243+ const std::string kAscendcRandomRegBaseStr = {
244+#include "random_reg_base.h"
242 };245 };
243 const std::string kAscendcSignBitRegBaseStr = {246 const std::string kAscendcSignBitRegBaseStr = {
244#include "signbit_reg_base.h"247#include "signbit_reg_base.h"
@@ -363,6 +366,7 @@ Register::Register() {
363 {"log_ndtr_reg_base.h", kAscendcLogNdtrRegBaseStr},366 {"log_ndtr_reg_base.h", kAscendcLogNdtrRegBaseStr},
364 {"next_after_reg_base.h", kAscendcNextAfterRegBaseStr},367 {"next_after_reg_base.h", kAscendcNextAfterRegBaseStr},
365 {"polygamma_reg_base.h", kAscendcPolyGammaRegBaseStr},368 {"polygamma_reg_base.h", kAscendcPolyGammaRegBaseStr},
369+ {"random_reg_base.h", kAscendcRandomRegBaseStr},
366 {"signbit_reg_base.h", kAscendcSignBitRegBaseStr},370 {"signbit_reg_base.h", kAscendcSignBitRegBaseStr},
367 {"frexp_reg_base.h", kAscendcFrexpRegBaseStr},371 {"frexp_reg_base.h", kAscendcFrexpRegBaseStr},
368 {"shifted_chebyshev_polynomial_utils_reg_base.h", kAscendcShiftedChebyshevPolynomialUtilsRegBaseStr},372 {"shifted_chebyshev_polynomial_utils_reg_base.h", kAscendcShiftedChebyshevPolynomialUtilsRegBaseStr},
Aautofuse/v35/codegen/reg_api_call/unary_output_api_call.cpp+48-0
@@ -0,0 +1,48 @@
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+#include "unary_output_api_call.h"
12+#include "attr_utils.h"
13+#include "ascir_ops.h"
14+#include "common_utils.h"
15+#include "common/ge_common/debug/log.h"
16+#include "graph/ascendc_ir/utils/asc_tensor_utils.h"
17+#include "common/checker.h"
18+#include "api_call/utils/api_call_factory.h"
19+#include "codegen/expression_convert_struct.h"
20+ 
21+namespace codegen {
22+using namespace std;
23+using namespace af::ops;
24+using namespace af::ascir_op;
25+ 
26+Status UnaryOutputApiCall::Generate(const TPipe &tpipe, const std::vector<ascir::AxisId> &current_axis,
27+ const std::vector<std::reference_wrapper<const Tensor>> &inputs,
28+ const std::vector<std::reference_wrapper<const Tensor>> &outputs,
29+ std::string &result) const {
30+ // 验证：无输入，单输出
31+ GE_CHK_BOOL_RET_STATUS(inputs.empty(), af::FAILED, "UnaryOutputApiCall expects no inputs, but got %zu inputs",
32+ inputs.size());
33+ GE_CHK_BOOL_RET_STATUS(outputs.size() == 1, af::FAILED,
34+ "UnaryOutputApiCall expects exactly 1 output, but got %zu outputs", outputs.size());
35+ 
36+ auto y = outputs[0].get();
37+ 
38+ stringstream ss;
39+ // 生成 API 调用：Rand(local_0[0], local_0_actual_size);
40+ ss << this->api_name_ << "(" << y << "[" << tpipe.tiler.TensorVectorizedOffset(current_axis, y) << "], "
41+ << y.actual_size << ");" << std::endl;
42+ 
43+ result = ss.str();
44+ return af::SUCCESS;
45+}
46+ 
47+static ApiCallRegister<UnaryOutputApiCall> register_unary_output_api_call("UnaryOutputApiCall");
48+} // namespace codegen
Aautofuse/v35/codegen/reg_api_call/unary_output_api_call.h+34-0
@@ -0,0 +1,34 @@
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+#ifndef __AUTOFUSE_UNARY_OUTPUT_API_CALL_H__
11+#define __AUTOFUSE_UNARY_OUTPUT_API_CALL_H__
12+ 
13+#include "codegen_kernel.h"
14+ 
15+namespace codegen {
16+ 
17+/**
18+ * @brief UnaryOutputApiCall - API call generator for operators with no inputs and single output
19+ * Specifically designed for operators like Rand that generate data without input tensors
20+ */
21+class UnaryOutputApiCall final : public ApiCall {
22+ public:
23+ using ApiCall::Generate;
24+ explicit UnaryOutputApiCall(const std::string &api_name) : ApiCall(api_name) {}
25+ ~UnaryOutputApiCall() override = default;
26+ 
27+ Status Generate(const TPipe &tpipe, const std::vector<ascir::AxisId> &current_axis,
28+ const std::vector<std::reference_wrapper<const Tensor>> &inputs,
29+ const std::vector<std::reference_wrapper<const Tensor>> &outputs, std::string &result) const override;
30+};
31+ 
32+} // namespace codegen
33+ 
34+#endif // __AUTOFUSE_UNARY_OUTPUT_API_CALL_H__
Mscripts/test/run_autofuse_test.sh+3-1
@@ -898,7 +898,9 @@ build_backend() {
898 chebyshev_polynomial_v_store_test_e2e_v2 \898 chebyshev_polynomial_v_store_test_e2e_v2 \
899 chebyshev_polynomial_w_store_test_e2e_v2 \899 chebyshev_polynomial_w_store_test_e2e_v2 \
900 hermite_polynomial_h_store_test_e2e_v2 \900 hermite_polynomial_h_store_test_e2e_v2 \
901- hermite_polynomial_he_store_test_e2e_v2"901+ hermite_polynomial_he_store_test_e2e_v2 \
902+ rand_store_test_e2e_v2 \
903+ randn_store_test_e2e_v2"
902 fi904 fi
903 MAKE_TARGET_LIST_CODEGEN=$(echo "${MAKE_TARGET_LIST}" | sed 's/e2e/codegen/g')905 MAKE_TARGET_LIST_CODEGEN=$(echo "${MAKE_TARGET_LIST}" | sed 's/e2e/codegen/g')
904 CTEST_BACKEND_TEST1_REGEX=$(build_backend_test_regex ${MAKE_TARGET_LIST_CODEGEN})906 CTEST_BACKEND_TEST1_REGEX=$(build_backend_test_regex ${MAKE_TARGET_LIST_CODEGEN})