Star1308
1425
代码
代码
Issues118
Pull Requests239
Actions
项目讨论
Wiki
分析
Star1308
1425
已合并
test: add UTs for random operators #2209
fenglin28创建于 4月13日
test: add UTs for random operators #2209
已合并
fenglin28创建于 4月13日
从fenglin28/ops-math_privatepatch_math_random_opapi_kernel_ut合入到cann/ops-mathmaster
共 30 个文件变更+1511-110
Mrandom/drop_out_do_mask/tests/ut/op_kernel/CMakeLists.txt+14-7
@@ -1,5 +1,5 @@
1# ----------------------------------------------------------------------------1# ----------------------------------------------------------------------------
2-# Copyright (c) 2025 Huawei Technologies Co., Ltd.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 of3# 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").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.5# Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -8,9 +8,16 @@
8# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
9# ----------------------------------------------------------------------------9# ----------------------------------------------------------------------------
10 10 
11-file(GLOB CURRENT_SOURCE_DIRS LIST_DIRECTORIES true ${CMAKE_CURRENT_SOURCE_DIR}/*)11+if (UT_TEST_ALL OR OP_KERNEL_UT)
12-foreach(SUB_DIR ${CURRENT_SOURCE_DIRS})12+ set(drop_out_do_mask_tiling_files
13- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")13+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/arch35/drop_out_do_mask_tiling.cpp
14- add_subdirectory(${SUB_DIR})14+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/arch35/drop_out_do_mask_tiling_arch35.cpp)
15- endif()15+ AddOpsTestCase(
16-endforeach()16+ OP_NAME drop_out_do_mask
17+ SOC_VERSION "ascend950"
18+ OTHER_COMPILE_OPTIONS "-DDTYPE_X=float"
19+ TILING_SRC_FILES ${drop_out_do_mask_tiling_files}
20+ UT_SRC_FILES
21+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_kernel/drop_out_do_mask_apt.cpp
22+ ${CMAKE_CURRENT_SOURCE_DIR}/test_drop_out_do_mask.cpp)
23+endif()
Arandom/drop_out_do_mask/tests/ut/op_kernel/test_drop_out_do_mask.cpp+90-0
@@ -0,0 +1,90 @@
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 <cstdint>
12+#include <cstring>
13+#include "gtest/gtest.h"
14+#include "tikicpulib.h"
15+#include "../../../op_host/arch35/drop_out_do_mask_tiling_arch35.h"
16+ 
17+extern "C" __global__ __aicore__ void drop_out_do_mask(
18+ GM_ADDR x, GM_ADDR mask, GM_ADDR prob, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling);
19+ 
20+namespace {
21+constexpr uint64_t kTilingKey = 100;
22+constexpr uint32_t kNumBlocks = 1;
23+constexpr int64_t kElementCount = 256;
24+constexpr size_t kMaskBytes = 32;
25+constexpr size_t kWorkspaceBytes = 32;
26+ 
27+size_t Align32(size_t size)
28+{
29+ return (size + 31U) / 32U * 32U;
30+}
31+ 
32+uint8_t *AllocGm(size_t size)
33+{
34+ return reinterpret_cast<uint8_t *>(AscendC::GmAlloc(Align32(size)));
35+}
36+ 
37+void InitInput(float *data)
38+{
39+ for (int64_t i = 0; i < kElementCount; ++i) {
40+ data[i] = static_cast<float>(i) + 0.5f;
41+ }
42+}
43+} // namespace
44+ 
45+class DropOutDoMaskKernelUT : public testing::Test {
46+};
47+ 
48+TEST_F(DropOutDoMaskKernelUT, keep_prob_one_copies_input_to_output)
49+{
50+ const size_t dataBytes = static_cast<size_t>(kElementCount) * sizeof(float);
51+ 
52+ optiling::DropOutDoMaskForAscendCTilingData tilingData;
53+ tilingData.set_usedCoreNum(1);
54+ tilingData.set_normBlockData(kElementCount);
55+ tilingData.set_tailBlockData(kElementCount);
56+ tilingData.set_ubFactor(kElementCount);
57+ tilingData.set_normBlockLoop(1);
58+ tilingData.set_normBlockTail(kElementCount);
59+ tilingData.set_tailBlockLoop(1);
60+ tilingData.set_tailBlockTail(kElementCount);
61+ tilingData.set_epsilon(1.0e-6f);
62+ const size_t tilingBytes = static_cast<size_t>(tilingData.GetDataSize());
63+ 
64+ uint8_t *x = AllocGm(dataBytes);
65+ uint8_t *mask = AllocGm(kMaskBytes);
66+ uint8_t *prob = AllocGm(sizeof(float));
67+ uint8_t *y = AllocGm(dataBytes);
68+ uint8_t *workspace = AllocGm(kWorkspaceBytes);
69+ uint8_t *tiling = AllocGm(tilingBytes);
70+ 
71+ InitInput(reinterpret_cast<float *>(x));
72+ std::memset(mask, 0xFF, Align32(kMaskBytes));
73+ *reinterpret_cast<float *>(prob) = 1.0f;
74+ std::memset(y, 0, Align32(dataBytes));
75+ std::memset(tiling, 0, Align32(tilingBytes));
76+ tilingData.SaveToBuffer(tiling, tilingBytes);
77+ 
78+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
79+ ICPU_SET_TILING_KEY(kTilingKey);
80+ ICPU_RUN_KF(drop_out_do_mask, kNumBlocks, x, mask, prob, y, workspace, tiling);
81+ 
82+ EXPECT_EQ(0, std::memcmp(y, x, dataBytes));
83+ 
84+ AscendC::GmFree(x);
85+ AscendC::GmFree(mask);
86+ AscendC::GmFree(prob);
87+ AscendC::GmFree(y);
88+ AscendC::GmFree(workspace);
89+ AscendC::GmFree(tiling);
90+}
Arandom/drop_out_do_mask_v3/tests/ut/op_kernel/CMakeLists.txt+24-0
@@ -0,0 +1,24 @@
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}/drop_out_do_mask_v3/arch35)
14+ execute_process(COMMAND ${CMAKE_COMMAND} -E create_symlink
15+ ${PROJECT_SOURCE_DIR}/random/random_common/op_kernel
16+ ${KERNEL_STAGING_DIR}/random_common)
17+ 
18+ set(drop_out_do_mask_v3_tiling_files
19+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/arch35/drop_out_do_mask_v3_tiling_arch35.cpp
20+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../../random_common/op_host/arch35/random_tiling_arch35.cpp)
21+ AddOpTestCase(drop_out_do_mask_v3 "ascend950"
22+ "-DDTYPE_X=float -DTestUtDefaultTilingStruct=RandomUnifiedTilingDataStruct -I${KERNEL_STAGING_DIR}/drop_out_do_mask_v3/arch35"
23+ "${drop_out_do_mask_v3_tiling_files}")
24+endif()
Arandom/drop_out_do_mask_v3/tests/ut/op_kernel/test_drop_out_do_mask_v3.cpp+86-0
@@ -0,0 +1,86 @@
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 <cstdint>
12+#include <cstring>
13+#include "gtest/gtest.h"
14+#include "tikicpulib.h"
15+#include "../../../op_host/arch35/drop_out_do_mask_v3_tiling_arch35.h"
16+ 
17+extern "C" __global__ __aicore__ void drop_out_do_mask_v3(
18+ GM_ADDR x, GM_ADDR mask, GM_ADDR prob, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling);
19+ 
20+namespace {
21+constexpr uint64_t kTilingKey = 100;
22+constexpr uint32_t kNumBlocks = 1;
23+constexpr int64_t kElementCount = 256;
24+constexpr size_t kMaskBytes = 32;
25+constexpr size_t kWorkspaceBytes = 32;
26+ 
27+size_t Align32(size_t size)
28+{
29+ return (size + 31U) / 32U * 32U;
30+}
31+ 
32+uint8_t *AllocGm(size_t size)
33+{
34+ return reinterpret_cast<uint8_t *>(AscendC::GmAlloc(Align32(size)));
35+}
36+ 
37+void InitInput(float *data)
38+{
39+ for (int64_t i = 0; i < kElementCount; ++i) {
40+ data[i] = static_cast<float>(i) + 1.0f;
41+ }
42+}
43+} // namespace
44+ 
45+class DropOutDoMaskV3KernelUT : public testing::Test {
46+};
47+ 
48+TEST_F(DropOutDoMaskV3KernelUT, keep_prob_one_copies_input_to_output)
49+{
50+ const size_t dataBytes = static_cast<size_t>(kElementCount) * sizeof(float);
51+ const size_t tilingBytes = sizeof(RandomUnifiedTilingDataStruct);
52+ 
53+ uint8_t *x = AllocGm(dataBytes);
54+ uint8_t *mask = AllocGm(kMaskBytes);
55+ uint8_t *prob = AllocGm(sizeof(float));
56+ uint8_t *y = AllocGm(dataBytes);
57+ uint8_t *workspace = AllocGm(kWorkspaceBytes);
58+ uint8_t *tiling = AllocGm(tilingBytes);
59+ 
60+ InitInput(reinterpret_cast<float *>(x));
61+ std::memset(mask, 0xFF, Align32(kMaskBytes));
62+ *reinterpret_cast<float *>(prob) = 1.0f;
63+ std::memset(y, 0, Align32(dataBytes));
64+ std::memset(tiling, 0, Align32(tilingBytes));
65+ 
66+ auto *tilingData = reinterpret_cast<RandomUnifiedTilingDataStruct *>(tiling);
67+ tilingData->usedCoreNum = 1;
68+ tilingData->normalCoreProNum = kElementCount;
69+ tilingData->tailCoreProNum = kElementCount;
70+ tilingData->singleBufferSize = kElementCount;
71+ tilingData->outputSize = kElementCount;
72+ tilingData->keepProb = 1.0f;
73+ 
74+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
75+ ICPU_SET_TILING_KEY(kTilingKey);
76+ ICPU_RUN_KF(drop_out_do_mask_v3, kNumBlocks, x, mask, prob, y, workspace, tiling);
77+ 
78+ EXPECT_EQ(0, std::memcmp(y, x, dataBytes));
79+ 
80+ AscendC::GmFree(x);
81+ AscendC::GmFree(mask);
82+ AscendC::GmFree(prob);
83+ AscendC::GmFree(y);
84+ AscendC::GmFree(workspace);
85+ AscendC::GmFree(tiling);
86+}
Arandom/drop_out_do_mask_v3_d/tests/ut/op_kernel/CMakeLists.txt+27-0
@@ -0,0 +1,27 @@
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}/drop_out_do_mask_v3_d/arch35)
14+ execute_process(COMMAND ${CMAKE_COMMAND} -E create_symlink
15+ ${PROJECT_SOURCE_DIR}/random/random_common/op_kernel
16+ ${KERNEL_STAGING_DIR}/random_common)
17+ execute_process(COMMAND ${CMAKE_COMMAND} -E create_symlink
18+ ${PROJECT_SOURCE_DIR}/random/drop_out_do_mask_v3/op_kernel
19+ ${KERNEL_STAGING_DIR}/drop_out_do_mask_v3)
20+ 
21+ set(drop_out_do_mask_v3_d_tiling_files
22+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/arch35/drop_out_do_mask_v3_d_tiling_arch35.cpp
23+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../../random_common/op_host/arch35/random_tiling_arch35.cpp)
24+ AddOpTestCase(drop_out_do_mask_v3_d "ascend950"
25+ "-DDTYPE_X=float -DTestUtDefaultTilingStruct=RandomUnifiedTilingDataStruct -I${KERNEL_STAGING_DIR}/drop_out_do_mask_v3_d/arch35 -I${KERNEL_STAGING_DIR}/drop_out_do_mask_v3_d"
26+ "${drop_out_do_mask_v3_d_tiling_files}")
27+endif()
Arandom/drop_out_do_mask_v3_d/tests/ut/op_kernel/test_drop_out_do_mask_v3_d.cpp+83-0
@@ -0,0 +1,83 @@
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 <cstdint>
12+#include <cstring>
13+#include "gtest/gtest.h"
14+#include "tikicpulib.h"
15+#include "../../../op_host/arch35/drop_out_do_mask_v3_d_tiling_arch35.h"
16+ 
17+extern "C" __global__ __aicore__ void drop_out_do_mask_v3_d(
18+ GM_ADDR x, GM_ADDR mask, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling);
19+ 
20+namespace {
21+constexpr uint64_t kTilingKey = 100;
22+constexpr uint32_t kNumBlocks = 1;
23+constexpr int64_t kElementCount = 256;
24+constexpr size_t kMaskBytes = 32;
25+constexpr size_t kWorkspaceBytes = 32;
26+ 
27+size_t Align32(size_t size)
28+{
29+ return (size + 31U) / 32U * 32U;
30+}
31+ 
32+uint8_t *AllocGm(size_t size)
33+{
34+ return reinterpret_cast<uint8_t *>(AscendC::GmAlloc(Align32(size)));
35+}
36+ 
37+void InitInput(float *data)
38+{
39+ for (int64_t i = 0; i < kElementCount; ++i) {
40+ data[i] = static_cast<float>(i) + 2.0f;
41+ }
42+}
43+} // namespace
44+ 
45+class DropOutDoMaskV3DKernelUT : public testing::Test {
46+};
47+ 
48+TEST_F(DropOutDoMaskV3DKernelUT, keep_prob_one_copies_input_to_output)
49+{
50+ const size_t dataBytes = static_cast<size_t>(kElementCount) * sizeof(float);
51+ const size_t tilingBytes = sizeof(RandomUnifiedTilingDataStruct);
52+ 
53+ uint8_t *x = AllocGm(dataBytes);
54+ uint8_t *mask = AllocGm(kMaskBytes);
55+ uint8_t *y = AllocGm(dataBytes);
56+ uint8_t *workspace = AllocGm(kWorkspaceBytes);
57+ uint8_t *tiling = AllocGm(tilingBytes);
58+ 
59+ InitInput(reinterpret_cast<float *>(x));
60+ std::memset(mask, 0xFF, Align32(kMaskBytes));
61+ std::memset(y, 0, Align32(dataBytes));
62+ std::memset(tiling, 0, Align32(tilingBytes));
63+ 
64+ auto *tilingData = reinterpret_cast<RandomUnifiedTilingDataStruct *>(tiling);
65+ tilingData->usedCoreNum = 1;
66+ tilingData->normalCoreProNum = kElementCount;
67+ tilingData->tailCoreProNum = kElementCount;
68+ tilingData->singleBufferSize = kElementCount;
69+ tilingData->outputSize = kElementCount;
70+ tilingData->keepProb = 1.0f;
71+ 
72+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
73+ ICPU_SET_TILING_KEY(kTilingKey);
74+ ICPU_RUN_KF(drop_out_do_mask_v3_d, kNumBlocks, x, mask, y, workspace, tiling);
75+ 
76+ EXPECT_EQ(0, std::memcmp(y, x, dataBytes));
77+ 
78+ AscendC::GmFree(x);
79+ AscendC::GmFree(mask);
80+ AscendC::GmFree(y);
81+ AscendC::GmFree(workspace);
82+ AscendC::GmFree(tiling);
83+}
Mrandom/drop_out_v3/tests/ut/op_kernel/CMakeLists.txt+14-7
@@ -1,5 +1,5 @@
1# ----------------------------------------------------------------------------1# ----------------------------------------------------------------------------
2-# Copyright (c) 2025 Huawei Technologies Co., Ltd.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 of3# 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").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.5# Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -8,9 +8,16 @@
8# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
9# ----------------------------------------------------------------------------9# ----------------------------------------------------------------------------
10 10 
11-file(GLOB CURRENT_SOURCE_DIRS LIST_DIRECTORIES true ${CMAKE_CURRENT_SOURCE_DIR}/*)11+if (UT_TEST_ALL OR OP_KERNEL_UT)
12-foreach(SUB_DIR ${CURRENT_SOURCE_DIRS})12+ set(KERNEL_STAGING_DIR ${CMAKE_CURRENT_BINARY_DIR}/kernel_dep_staging)
13- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")13+ file(MAKE_DIRECTORY ${KERNEL_STAGING_DIR}/drop_out_v3/arch35)
14- add_subdirectory(${SUB_DIR})14+ execute_process(COMMAND ${CMAKE_COMMAND} -E create_symlink
15- endif()15+ ${PROJECT_SOURCE_DIR}/random/random_common/op_kernel
16-endforeach()16+ ${KERNEL_STAGING_DIR}/random_common)
17+ 
18+ set(drop_out_v3_tiling_files
19+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/arch35/drop_out_v3_tiling_arch35.cpp)
20+ AddOpTestCase(drop_out_v3 "ascend950"
21+ "-DDTYPE_P=float -I${KERNEL_STAGING_DIR}/drop_out_v3/arch35"
22+ "${drop_out_v3_tiling_files}")
23+endif()
Mrandom/random_common/op_host/arch35/random_tiling_arch35.cpp+1-1
@@ -89,7 +89,7 @@ ge::graphStatus ExtractTensorValue(const gert::TilingContext* context, const int
89}89}
90 90 
91RandomTilingArch35::RandomTilingArch35(gert::TilingContext* context, const OpTilingConfig& config)91RandomTilingArch35::RandomTilingArch35(gert::TilingContext* context, const OpTilingConfig& config)
92- : context_(context), config_(config)92+ : context_(context), config_(config), tilingData_{}, simtTilingData_{}
93{}93{}
94 94 
95ge::graphStatus RandomTilingArch35::DoTiling()95ge::graphStatus RandomTilingArch35::DoTiling()
Rrandom/random_standard_normal_v2/tests/ut/op_host/test_random_standard_normal_v2_tiling.cpp→random/random_standard_normal_v2/tests/ut/op_host/arch35/test_random_standard_normal_v2_tiling.cpp+33-27
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.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 of3 * 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").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.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -8,49 +8,53 @@
8 * See LICENSE in the root of the software repository for the full text of the License.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
10 10 
11-#include <iostream>
12#include <gtest/gtest.h>11#include <gtest/gtest.h>
12+#include <iostream>
13+#include <vector>
14+#include "../../../../op_host/arch35/random_standard_normal_v2_tiling_arch35.h"
13#include "tiling_context_faker.h"15#include "tiling_context_faker.h"
14#include "tiling_case_executor.h"16#include "tiling_case_executor.h"
15-#include "../../../op_host/arch35/random_standard_normal_v2_tiling_arch35.h"
16 17 
17using namespace std;18using namespace std;
18using namespace ge;19using namespace ge;
19 20 
20class RandomStandardNormalV2Tiling : public testing::Test {21class RandomStandardNormalV2Tiling : public testing::Test {
21- protected:22+protected:
22- static void SetUpTestCase() {23+ static void SetUpTestCase()
23- std::cout << "RandomStandardNormalV2 SetUp" << std::endl;24+ {
24- }25+ std::cout << "RandomStandardNormalV2 SetUp" << std::endl;
26+ }
25 27 
26- static void TearDownTestCase() {28+ static void TearDownTestCase()
27- std::cout << "RandomStandardNormalV2 TearDown" << std::endl;29+ {
28- }30+ std::cout << "RandomStandardNormalV2 TearDown" << std::endl;
31+ }
29};32};
30 33 
31TEST_F(RandomStandardNormalV2Tiling, random_standard_normal_v2_tiling_950_float_001)34TEST_F(RandomStandardNormalV2Tiling, random_standard_normal_v2_tiling_950_float_001)
32{35{
33 optiling::RandomStandardNormalV2CompileInfo compileInfo = {64, 196608};36 optiling::RandomStandardNormalV2CompileInfo compileInfo = {64, 196608};
34- gert::StorageShape shape_shape = {{2}, {2}};37+ gert::StorageShape shapeShape = {{2}, {2}};
35- gert::StorageShape offset_shape = {{1}, {1}};38+ gert::StorageShape offsetShape = {{1}, {1}};
36- gert::StorageShape out_shape = {{32, 512}, {32, 512}};39+ gert::StorageShape outShape = {{32, 512}, {32, 512}};
37 auto seed = Ops::Math::AnyValue::CreateFrom<int64_t>(10);40 auto seed = Ops::Math::AnyValue::CreateFrom<int64_t>(10);
38 auto seed2 = Ops::Math::AnyValue::CreateFrom<int64_t>(5);41 auto seed2 = Ops::Math::AnyValue::CreateFrom<int64_t>(5);
39 auto dtype = Ops::Math::AnyValue::CreateFrom<int64_t>(0);42 auto dtype = Ops::Math::AnyValue::CreateFrom<int64_t>(0);
40 43 
41- vector<int32_t> shape_value = {32, 512};44+ vector<int32_t> shapeValue = {32, 512};
42- vector<int64_t> offset_value = {0};45+ vector<int64_t> offsetValue = {0};
43 46 
44 gert::TilingContextPara tilingContextPara(47 gert::TilingContextPara tilingContextPara(
45- "RandomStandardNormalV2", {{shape_shape, ge::DT_INT32, ge::FORMAT_ND, true, shape_value.data()}, {offset_shape, ge::DT_INT64, ge::FORMAT_ND, true, offset_value.data()}},48+ "RandomStandardNormalV2",
46- {{out_shape, ge::DT_FLOAT, ge::FORMAT_ND}, {offset_shape, ge::DT_INT64, ge::FORMAT_ND}},49+ {{shapeShape, ge::DT_INT32, ge::FORMAT_ND, true, shapeValue.data()},
50+ {offsetShape, ge::DT_INT64, ge::FORMAT_ND, true, offsetValue.data()}},
51+ {{outShape, ge::DT_FLOAT, ge::FORMAT_ND}, {offsetShape, ge::DT_INT64, ge::FORMAT_ND}},
47 {gert::TilingContextPara::OpAttr("dtype", dtype),52 {gert::TilingContextPara::OpAttr("dtype", dtype),
48 gert::TilingContextPara::OpAttr("seed", seed),53 gert::TilingContextPara::OpAttr("seed", seed),
49 gert::TilingContextPara::OpAttr("seed2", seed2)},54 gert::TilingContextPara::OpAttr("seed2", seed2)},
50 &compileInfo);55 &compileInfo);
51 uint64_t expectTilingKey = 100;56 uint64_t expectTilingKey = 100;
52- string expectTilingData =57+ string expectTilingData = "64 256 256 10912 10 0 5 16384 0 0 0 0 ";
53- "64 256 256 10912 10 0 5 16384 0 0 0 ";
54 std::vector<size_t> expectWorkspaces = {0};58 std::vector<size_t> expectWorkspaces = {0};
55 ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);59 ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
56}60}
@@ -58,19 +62,21 @@ TEST_F(RandomStandardNormalV2Tiling, random_standard_normal_v2_tiling_950_float_
58TEST_F(RandomStandardNormalV2Tiling, random_standard_normal_v2_tiling_950_float_002)62TEST_F(RandomStandardNormalV2Tiling, random_standard_normal_v2_tiling_950_float_002)
59{63{
60 optiling::RandomStandardNormalV2CompileInfo compileInfo = {64, 196608};64 optiling::RandomStandardNormalV2CompileInfo compileInfo = {64, 196608};
61- gert::StorageShape shape_shape = {{2}, {2}};65+ gert::StorageShape shapeShape = {{2}, {2}};
62- gert::StorageShape offset_shape = {{1}, {1}};66+ gert::StorageShape offsetShape = {{1}, {1}};
63- gert::StorageShape out_shape = {{32, 512}, {32, 512}};67+ gert::StorageShape outShape = {{32, 512}, {32, 512}};
64 auto seed = Ops::Math::AnyValue::CreateFrom<int64_t>(0);68 auto seed = Ops::Math::AnyValue::CreateFrom<int64_t>(0);
65 auto seed2 = Ops::Math::AnyValue::CreateFrom<int64_t>(0);69 auto seed2 = Ops::Math::AnyValue::CreateFrom<int64_t>(0);
66 auto dtype = Ops::Math::AnyValue::CreateFrom<int64_t>(0);70 auto dtype = Ops::Math::AnyValue::CreateFrom<int64_t>(0);
67 71 
68- vector<int32_t> shape_value = {32, 512};72+ vector<int32_t> shapeValue = {32, 512};
69- vector<int64_t> offset_value = {0};73+ vector<int64_t> offsetValue = {0};
70 74 
71 gert::TilingContextPara tilingContextPara(75 gert::TilingContextPara tilingContextPara(
72- "RandomStandardNormalV2", {{shape_shape, ge::DT_INT32, ge::FORMAT_ND, true, shape_value.data()}, {offset_shape, ge::DT_INT64, ge::FORMAT_ND, true, offset_value.data()}},76+ "RandomStandardNormalV2",
73- {{out_shape, ge::DT_FLOAT, ge::FORMAT_ND}, {offset_shape, ge::DT_INT64, ge::FORMAT_ND}},77+ {{shapeShape, ge::DT_INT32, ge::FORMAT_ND, true, shapeValue.data()},
78+ {offsetShape, ge::DT_INT64, ge::FORMAT_ND, true, offsetValue.data()}},
79+ {{outShape, ge::DT_FLOAT, ge::FORMAT_ND}, {offsetShape, ge::DT_INT64, ge::FORMAT_ND}},
74 {gert::TilingContextPara::OpAttr("dtype", dtype),80 {gert::TilingContextPara::OpAttr("dtype", dtype),
75 gert::TilingContextPara::OpAttr("seed", seed),81 gert::TilingContextPara::OpAttr("seed", seed),
76 gert::TilingContextPara::OpAttr("seed2", seed2)},82 gert::TilingContextPara::OpAttr("seed2", seed2)},
@@ -79,4 +85,4 @@ TEST_F(RandomStandardNormalV2Tiling, random_standard_normal_v2_tiling_950_float_
79 TilingInfo tilingInfo;85 TilingInfo tilingInfo;
80 auto tilingRet = ExecuteTiling(tilingContextPara, tilingInfo);86 auto tilingRet = ExecuteTiling(tilingContextPara, tilingInfo);
81 EXPECT_EQ(tilingRet, true);87 EXPECT_EQ(tilingRet, true);
82-}88+}
Arandom/random_standard_normal_v2/tests/ut/op_kernel/CMakeLists.txt+25-0
@@ -0,0 +1,25 @@
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}/random_standard_normal_v2/arch35)
14+ execute_process(COMMAND ${CMAKE_COMMAND} -E create_symlink
15+ ${PROJECT_SOURCE_DIR}/random/random_common/op_kernel
16+ ${KERNEL_STAGING_DIR}/random_common)
17+ 
18+ set(random_standard_normal_v2_tiling_files
19+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/arch35/random_standard_normal_v2_tiling_arch35.cpp)
20+ AddOpTestCase(
21+ random_standard_normal_v2
22+ "ascend950"
23+ "-DDTYPE_Y=float -DDTYPE_OFFSET=int64_t -DTestUtDefaultTilingStruct=RandomUnifiedTilingDataStruct -I${KERNEL_STAGING_DIR}/random_standard_normal_v2/arch35"
24+ "${random_standard_normal_v2_tiling_files}")
25+endif()
Arandom/random_standard_normal_v2/tests/ut/op_kernel/test_random_standard_normal_v2.cpp+87-0
@@ -0,0 +1,87 @@
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 <cstdint>
12+#include <cstring>
13+#include <iostream>
14+#include "gtest/gtest.h"
15+#include "tikicpulib.h"
16+#include "../../../../random_common/op_kernel/arch35/random_unified_tiling_data_arch35.h"
17+ 
18+extern "C" __global__ __aicore__ void random_standard_normal_v2(
19+ GM_ADDR shape, GM_ADDR offset, GM_ADDR y, GM_ADDR count, GM_ADDR workspace, GM_ADDR tiling);
20+ 
21+namespace {
22+constexpr uint32_t kNumBlocks = 1;
23+constexpr uint64_t kTilingKey = 100;
24+constexpr int64_t kElementCount = 256;
25+constexpr int64_t kExpectedOffset = kElementCount * 256;
26+ 
27+inline size_t Align32(size_t size)
28+{
29+ return (size + 31U) / 32U * 32U;
30+}
31+} // namespace
32+ 
33+class RandomStandardNormalV2KernelTest : public testing::Test {
34+protected:
35+ static void SetUpTestCase()
36+ {
37+ std::cout << "RandomStandardNormalV2KernelTest SetUp" << std::endl;
38+ }
39+ 
40+ static void TearDownTestCase()
41+ {
42+ std::cout << "RandomStandardNormalV2KernelTest TearDown" << std::endl;
43+ }
44+};
45+ 
46+TEST_F(RandomStandardNormalV2KernelTest, smoke_float)
47+{
48+ auto* shape = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(2 * sizeof(int32_t))));
49+ auto* offset = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int64_t))));
50+ auto* y = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(kElementCount * sizeof(float))));
51+ auto* count = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int64_t))));
52+ auto* workspace = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(1024 * 1024)));
53+ auto* tiling = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(RandomUnifiedTilingDataStruct))));
54+ 
55+ std::memset(y, 0, kElementCount * sizeof(float));
56+ std::memset(count, 0, sizeof(int64_t));
57+ std::memset(tiling, 0, sizeof(RandomUnifiedTilingDataStruct));
58+ reinterpret_cast<int32_t*>(shape)[0] = 16;
59+ reinterpret_cast<int32_t*>(shape)[1] = 16;
60+ reinterpret_cast<int64_t*>(offset)[0] = 0;
61+ 
62+ auto* tilingData = reinterpret_cast<RandomUnifiedTilingDataStruct*>(tiling);
63+ tilingData->usedCoreNum = kNumBlocks;
64+ tilingData->normalCoreProNum = kElementCount;
65+ tilingData->tailCoreProNum = kElementCount;
66+ tilingData->singleBufferSize = kElementCount;
67+ tilingData->key[0] = 10;
68+ tilingData->key[1] = 0;
69+ tilingData->counter[0] = 5;
70+ tilingData->counter[1] = 0;
71+ tilingData->counter[2] = 0;
72+ tilingData->counter[3] = 0;
73+ tilingData->outputSize = kElementCount;
74+ 
75+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
76+ ICPU_SET_TILING_KEY(kTilingKey);
77+ ICPU_RUN_KF(random_standard_normal_v2, kNumBlocks, shape, offset, y, count, workspace, tiling);
78+ 
79+ EXPECT_EQ(reinterpret_cast<int64_t*>(offset)[0], kExpectedOffset);
80+ 
81+ AscendC::GmFree(shape);
82+ AscendC::GmFree(offset);
83+ AscendC::GmFree(y);
84+ AscendC::GmFree(count);
85+ AscendC::GmFree(workspace);
86+ AscendC::GmFree(tiling);
87+}
Arandom/random_uniform_int_v2/tests/ut/op_kernel/CMakeLists.txt+19-0
@@ -0,0 +1,19 @@
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(random_uniform_int_v2_tiling_files
13+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/arch35/random_uniform_int_v2_tiling_arch35.cpp)
14+ AddOpTestCase(
15+ random_uniform_int_v2
16+ "ascend950"
17+ "-DDTYPE_Y=int32_t"
18+ "${random_uniform_int_v2_tiling_files}")
19+endif()
Arandom/random_uniform_int_v2/tests/ut/op_kernel/test_random_uniform_int_v2.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+#include <cstdint>
12+#include <cstring>
13+#include <iostream>
14+#include "gtest/gtest.h"
15+#include "tikicpulib.h"
16+#include "../../../op_kernel/random_uniform_int_v2.cpp"
17+ 
18+namespace {
19+constexpr uint32_t kNumBlocks = 1;
20+constexpr uint32_t kOpType = 0;
21+constexpr int64_t kElementCount = 256;
22+constexpr int64_t kExpectedOffset = kElementCount * 256;
23+ 
24+inline size_t Align32(size_t size)
25+{
26+ return (size + 31U) / 32U * 32U;
27+}
28+} // namespace
29+ 
30+class RandomUniformIntV2KernelTest : public testing::Test {
31+protected:
32+ static void SetUpTestCase()
33+ {
34+ std::cout << "RandomUniformIntV2KernelTest SetUp" << std::endl;
35+ }
36+ 
37+ static void TearDownTestCase()
38+ {
39+ std::cout << "RandomUniformIntV2KernelTest TearDown" << std::endl;
40+ }
41+};
42+ 
43+TEST_F(RandomUniformIntV2KernelTest, smoke_int32)
44+{
45+ auto* shape = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(2 * sizeof(int32_t))));
46+ auto* min = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int32_t))));
47+ auto* max = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int32_t))));
48+ auto* inOffset = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int64_t))));
49+ auto* y = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(kElementCount * sizeof(int32_t))));
50+ auto* outOffset = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int64_t))));
51+ auto* workspace = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(1024 * 1024)));
52+ auto* tiling = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(RandomUniformIntV2TilingData4RegBase))));
53+ 
54+ std::memset(y, 0, kElementCount * sizeof(int32_t));
55+ std::memset(tiling, 0, sizeof(RandomUniformIntV2TilingData4RegBase));
56+ reinterpret_cast<int32_t*>(shape)[0] = 16;
57+ reinterpret_cast<int32_t*>(shape)[1] = 16;
58+ reinterpret_cast<int32_t*>(min)[0] = 2;
59+ reinterpret_cast<int32_t*>(max)[0] = 5;
60+ reinterpret_cast<int64_t*>(inOffset)[0] = 0;
61+ reinterpret_cast<int64_t*>(outOffset)[0] = 0;
62+ 
63+ auto* tilingData = reinterpret_cast<RandomUniformIntV2TilingData4RegBase*>(tiling);
64+ tilingData->blockNum = kNumBlocks;
65+ tilingData->normalCoreProNum = kElementCount;
66+ tilingData->tailCoreProNum = kElementCount;
67+ tilingData->singleUbSize = kElementCount;
68+ tilingData->seed = 10;
69+ tilingData->seed2 = 5;
70+ tilingData->outputSize = kElementCount;
71+ tilingData->range = 3;
72+ tilingData->lo = 2;
73+ 
74+ auto func = random_uniform_int_v2<kOpType>;
75+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
76+ ICPU_RUN_KF(func, kNumBlocks, shape, min, max, inOffset, y, outOffset, workspace, tiling);
77+ 
78+ EXPECT_EQ(reinterpret_cast<int64_t*>(outOffset)[0], kExpectedOffset);
79+ 
80+ AscendC::GmFree(shape);
81+ AscendC::GmFree(min);
82+ AscendC::GmFree(max);
83+ AscendC::GmFree(inOffset);
84+ AscendC::GmFree(y);
85+ AscendC::GmFree(outOffset);
86+ AscendC::GmFree(workspace);
87+ AscendC::GmFree(tiling);
88+}
Arandom/random_uniform_v2/tests/ut/op_kernel/CMakeLists.txt+25-0
@@ -0,0 +1,25 @@
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}/random_uniform_v2/arch35)
14+ execute_process(COMMAND ${CMAKE_COMMAND} -E create_symlink
15+ ${PROJECT_SOURCE_DIR}/random/random_common/op_kernel
16+ ${KERNEL_STAGING_DIR}/random_common)
17+ 
18+ set(random_uniform_v2_tiling_files
19+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/arch35/random_uniform_v2_tiling_arch35.cpp)
20+ AddOpTestCase(
21+ random_uniform_v2
22+ "ascend950"
23+ "-DDTYPE_Y=float -DTestUtDefaultTilingStruct=RandomUnifiedTilingDataStruct -I${KERNEL_STAGING_DIR}/random_uniform_v2/arch35"
24+ "${random_uniform_v2_tiling_files}")
25+endif()
Arandom/random_uniform_v2/tests/ut/op_kernel/test_random_uniform_v2.cpp+87-0
@@ -0,0 +1,87 @@
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 <cstdint>
12+#include <cstring>
13+#include <iostream>
14+#include "gtest/gtest.h"
15+#include "tikicpulib.h"
16+#include "../../../../random_common/op_kernel/arch35/random_unified_tiling_data_arch35.h"
17+ 
18+extern "C" __global__ __aicore__ void random_uniform_v2(
19+ GM_ADDR shape, GM_ADDR inOffset, GM_ADDR y, GM_ADDR outOffset, GM_ADDR workspace, GM_ADDR tiling);
20+ 
21+namespace {
22+constexpr uint32_t kNumBlocks = 1;
23+constexpr uint64_t kTilingKey = 100;
24+constexpr int64_t kElementCount = 256;
25+constexpr int64_t kExpectedOffset = kElementCount * 256;
26+ 
27+inline size_t Align32(size_t size)
28+{
29+ return (size + 31U) / 32U * 32U;
30+}
31+} // namespace
32+ 
33+class RandomUniformV2KernelTest : public testing::Test {
34+protected:
35+ static void SetUpTestCase()
36+ {
37+ std::cout << "RandomUniformV2KernelTest SetUp" << std::endl;
38+ }
39+ 
40+ static void TearDownTestCase()
41+ {
42+ std::cout << "RandomUniformV2KernelTest TearDown" << std::endl;
43+ }
44+};
45+ 
46+TEST_F(RandomUniformV2KernelTest, smoke_float)
47+{
48+ auto* shape = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(2 * sizeof(int32_t))));
49+ auto* inOffset = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int64_t))));
50+ auto* y = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(kElementCount * sizeof(float))));
51+ auto* outOffset = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int64_t))));
52+ auto* workspace = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(1024 * 1024)));
53+ auto* tiling = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(RandomUnifiedTilingDataStruct))));
54+ 
55+ std::memset(y, 0, kElementCount * sizeof(float));
56+ std::memset(tiling, 0, sizeof(RandomUnifiedTilingDataStruct));
57+ reinterpret_cast<int32_t*>(shape)[0] = 16;
58+ reinterpret_cast<int32_t*>(shape)[1] = 16;
59+ reinterpret_cast<int64_t*>(inOffset)[0] = 0;
60+ reinterpret_cast<int64_t*>(outOffset)[0] = 0;
61+ 
62+ auto* tilingData = reinterpret_cast<RandomUnifiedTilingDataStruct*>(tiling);
63+ tilingData->usedCoreNum = kNumBlocks;
64+ tilingData->normalCoreProNum = kElementCount;
65+ tilingData->tailCoreProNum = kElementCount;
66+ tilingData->singleBufferSize = kElementCount;
67+ tilingData->key[0] = 10;
68+ tilingData->key[1] = 0;
69+ tilingData->counter[0] = 5;
70+ tilingData->counter[1] = 0;
71+ tilingData->counter[2] = 0;
72+ tilingData->counter[3] = 0;
73+ tilingData->outputSize = kElementCount;
74+ 
75+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
76+ ICPU_SET_TILING_KEY(kTilingKey);
77+ ICPU_RUN_KF(random_uniform_v2, kNumBlocks, shape, inOffset, y, outOffset, workspace, tiling);
78+ 
79+ EXPECT_EQ(reinterpret_cast<int64_t*>(outOffset)[0], kExpectedOffset);
80+ 
81+ AscendC::GmFree(shape);
82+ AscendC::GmFree(inOffset);
83+ AscendC::GmFree(y);
84+ AscendC::GmFree(outOffset);
85+ AscendC::GmFree(workspace);
86+ AscendC::GmFree(tiling);
87+}
Mrandom/stateless_bernoulli/tests/ut/op_kernel/CMakeLists.txt+16-7
@@ -1,5 +1,5 @@
1# ----------------------------------------------------------------------------1# ----------------------------------------------------------------------------
2-# Copyright (c) 2025 Huawei Technologies Co., Ltd.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 of3# 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").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.5# Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -8,9 +8,18 @@
8# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
9# ----------------------------------------------------------------------------9# ----------------------------------------------------------------------------
10 10 
11-file(GLOB CURRENT_SOURCE_DIRS LIST_DIRECTORIES true ${CMAKE_CURRENT_SOURCE_DIR}/*)11+if(UT_TEST_ALL OR OP_KERNEL_UT)
12-foreach(SUB_DIR ${CURRENT_SOURCE_DIRS})12+ set(KERNEL_STAGING_DIR ${CMAKE_CURRENT_BINARY_DIR}/kernel_dep_staging)
13- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")13+ file(MAKE_DIRECTORY ${KERNEL_STAGING_DIR}/stateless_bernoulli/arch35)
14- add_subdirectory(${SUB_DIR})14+ execute_process(COMMAND ${CMAKE_COMMAND} -E create_symlink
15- endif()15+ ${PROJECT_SOURCE_DIR}/random/random_common/op_kernel
16-endforeach()16+ ${KERNEL_STAGING_DIR}/random_common)
17+ 
18+ set(stateless_bernoulli_tiling_files
19+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/arch35/stateless_bernoulli_tiling_arch35.cpp)
20+ AddOpTestCase(
21+ stateless_bernoulli
22+ "ascend950"
23+ "-DDTYPE_Y=float -I${KERNEL_STAGING_DIR}/stateless_bernoulli/arch35"
24+ "${stateless_bernoulli_tiling_files}")
25+endif()
Arandom/stateless_bernoulli/tests/ut/op_kernel/test_stateless_bernoulli.cpp+84-0
@@ -0,0 +1,84 @@
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 <cstdint>
12+#include <cstring>
13+#include "gtest/gtest.h"
14+#include "tikicpulib.h"
15+ 
16+__global__ __aicore__ void stateless_bernoulli(
17+ GM_ADDR shape, GM_ADDR prob, GM_ADDR seed, GM_ADDR offset, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling);
18+ 
19+namespace {
20+constexpr uint32_t kNumBlocks = 1;
21+constexpr uint64_t kTilingKey = 1001;
22+constexpr int64_t kElementCount = 256;
23+ 
24+struct BernoulliTilingData {
25+ uint64_t blockNum;
26+ uint64_t probTensorSize;
27+ uint64_t outputSize;
28+ uint64_t isProbScalar;
29+ int64_t seed;
30+ int64_t philoxOffset;
31+};
32+ 
33+inline size_t Align32(size_t size)
34+{
35+ return (size + 31U) / 32U * 32U;
36+}
37+} // namespace
38+ 
39+class StatelessBernoulliKernelTest : public testing::Test {
40+};
41+ 
42+TEST_F(StatelessBernoulliKernelTest, smoke_float)
43+{
44+ auto* shape = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(2 * sizeof(int32_t))));
45+ auto* prob = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(float))));
46+ auto* seed = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int64_t))));
47+ auto* offset = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int64_t))));
48+ auto* y = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(kElementCount * sizeof(float))));
49+ auto* workspace = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(16 * 1024 * 1024)));
50+ auto* tiling = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(BernoulliTilingData))));
51+ 
52+ std::memset(y, 0, kElementCount * sizeof(float));
53+ std::memset(tiling, 0, sizeof(BernoulliTilingData));
54+ reinterpret_cast<int32_t*>(shape)[0] = 16;
55+ reinterpret_cast<int32_t*>(shape)[1] = 16;
56+ *reinterpret_cast<float*>(prob) = 0.5f;
57+ *reinterpret_cast<int64_t*>(seed) = 42;
58+ *reinterpret_cast<int64_t*>(offset) = 0;
59+ 
60+ auto* tilingData = reinterpret_cast<BernoulliTilingData*>(tiling);
61+ tilingData->blockNum = kNumBlocks;
62+ tilingData->probTensorSize = 1;
63+ tilingData->outputSize = kElementCount;
64+ tilingData->isProbScalar = 1;
65+ tilingData->seed = 42;
66+ tilingData->philoxOffset = 0;
67+ 
68+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
69+ ICPU_SET_TILING_KEY(kTilingKey);
70+ ICPU_RUN_KF(stateless_bernoulli, kNumBlocks, shape, prob, seed, offset, y, workspace, tiling);
71+ 
72+ auto* yData = reinterpret_cast<float*>(y);
73+ for (int64_t i = 0; i < kElementCount; ++i) {
74+ EXPECT_TRUE(yData[i] == 0.0f || yData[i] == 1.0f);
75+ }
76+ 
77+ AscendC::GmFree(shape);
78+ AscendC::GmFree(prob);
79+ AscendC::GmFree(seed);
80+ AscendC::GmFree(offset);
81+ AscendC::GmFree(y);
82+ AscendC::GmFree(workspace);
83+ AscendC::GmFree(tiling);
84+}
Arandom/stateless_drop_out_gen_mask/tests/ut/op_kernel/CMakeLists.txt+26-0
@@ -0,0 +1,26 @@
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_drop_out_gen_mask/arch35)
14+ execute_process(COMMAND ${CMAKE_COMMAND} -E create_symlink
15+ ${PROJECT_SOURCE_DIR}/random/random_common/op_kernel
16+ ${KERNEL_STAGING_DIR}/random_common)
17+ 
18+ set(stateless_drop_out_gen_mask_tiling_files
19+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/arch35/stateless_drop_out_gen_mask_tiling_arch35.cpp
20+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../../random_common/op_host/arch35/random_tiling_arch35.cpp)
21+ AddOpTestCase(
22+ stateless_drop_out_gen_mask
23+ "ascend950"
24+ "-DDTYPE_PROB=float -DTestUtDefaultTilingStruct=RandomUnifiedTilingDataStruct -I${KERNEL_STAGING_DIR}/stateless_drop_out_gen_mask/arch35"
25+ "${stateless_drop_out_gen_mask_tiling_files}")
26+endif()
Arandom/stateless_drop_out_gen_mask/tests/ut/op_kernel/test_stateless_drop_out_gen_mask.cpp+80-0
@@ -0,0 +1,80 @@
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 <cstdint>
12+#include <cstring>
13+#include "gtest/gtest.h"
14+#include "tikicpulib.h"
15+#include "../../../../random_common/op_kernel/arch35/random_unified_tiling_data_arch35.h"
16+ 
17+extern "C" __global__ __aicore__ void stateless_drop_out_gen_mask(
18+ GM_ADDR shape, GM_ADDR prob, GM_ADDR seed, GM_ADDR seed1, GM_ADDR offset, GM_ADDR y, GM_ADDR workspace,
19+ GM_ADDR tiling);
20+ 
21+namespace {
22+constexpr uint32_t kNumBlocks = 1;
23+constexpr uint64_t kTilingKey = 100;
24+constexpr int64_t kElementCount = 256;
25+ 
26+inline size_t Align32(size_t size)
27+{
28+ return (size + 31U) / 32U * 32U;
29+}
30+} // namespace
31+ 
32+class StatelessDropOutGenMaskKernelTest : public testing::Test {
33+};
34+ 
35+TEST_F(StatelessDropOutGenMaskKernelTest, smoke_float)
36+{
37+ auto* shape = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(2 * sizeof(int32_t))));
38+ auto* prob = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(float))));
39+ auto* seed = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int64_t))));
40+ auto* seed1 = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int64_t))));
41+ auto* offset = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int64_t))));
42+ auto* y = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(kElementCount * sizeof(uint8_t))));
43+ auto* workspace = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(1024 * 1024)));
44+ auto* tiling = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(RandomUnifiedTilingDataStruct))));
45+ 
46+ std::memset(y, 0, kElementCount * sizeof(uint8_t));
47+ std::memset(tiling, 0, sizeof(RandomUnifiedTilingDataStruct));
48+ reinterpret_cast<int32_t*>(shape)[0] = 16;
49+ reinterpret_cast<int32_t*>(shape)[1] = 16;
50+ *reinterpret_cast<float*>(prob) = 0.5f;
51+ *reinterpret_cast<int64_t*>(seed) = 42;
52+ *reinterpret_cast<int64_t*>(seed1) = 0;
53+ *reinterpret_cast<int64_t*>(offset) = 0;
54+ 
55+ auto* tilingData = reinterpret_cast<RandomUnifiedTilingDataStruct*>(tiling);
56+ tilingData->usedCoreNum = kNumBlocks;
57+ tilingData->normalCoreProNum = kElementCount;
58+ tilingData->tailCoreProNum = kElementCount;
59+ tilingData->singleBufferSize = kElementCount;
60+ tilingData->key[0] = 42;
61+ tilingData->key[1] = 0;
62+ tilingData->counter[0] = 0;
63+ tilingData->counter[1] = 0;
64+ tilingData->counter[2] = 0;
65+ tilingData->counter[3] = 0;
66+ tilingData->outputSize = kElementCount;
67+ 
68+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
69+ ICPU_SET_TILING_KEY(kTilingKey);
70+ ICPU_RUN_KF(stateless_drop_out_gen_mask, kNumBlocks, shape, prob, seed, seed1, offset, y, workspace, tiling);
71+ 
72+ AscendC::GmFree(shape);
73+ AscendC::GmFree(prob);
74+ AscendC::GmFree(seed);
75+ AscendC::GmFree(seed1);
76+ AscendC::GmFree(offset);
77+ AscendC::GmFree(y);
78+ AscendC::GmFree(workspace);
79+ AscendC::GmFree(tiling);
80+}
Rrandom/stateless_random_choice_with_mask/tests/ut/op_host/test_stateless_random_choice_with_mask_tiling.cpp→random/stateless_random_choice_with_mask/tests/ut/op_host/arch35/test_stateless_random_choice_with_mask_tiling.cpp+59-54
@@ -8,84 +8,89 @@
8 * See LICENSE in the root of the software repository for the full text of the License.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
10 10 
11-#include <iostream>
12#include <gtest/gtest.h>11#include <gtest/gtest.h>
12+#include <iostream>
13+#include <vector>
14+#include "../../../../op_host/arch35/stateless_random_choice_with_mask_simt_tiling.h"
13#include "tiling_context_faker.h"15#include "tiling_context_faker.h"
14#include "tiling_case_executor.h"16#include "tiling_case_executor.h"
15-#include "../../../op_host/arch35/stateless_random_choice_with_mask_simt_tiling.h"
16 17 
17using namespace std;18using namespace std;
18using namespace ge;19using namespace ge;
19 20 
20class StatelessRandomChoiceWithMaskTiling : public testing::Test {21class StatelessRandomChoiceWithMaskTiling : public testing::Test {
21- protected:22+protected:
22- static void SetUpTestCase() {23+ static void SetUpTestCase()
23- std::cout << "StatelessRandomChoiceWithMask SetUp" << std::endl;24+ {
24- }25+ std::cout << "StatelessRandomChoiceWithMask SetUp" << std::endl;
26+ }
25 27 
26- static void TearDownTestCase() {28+ static void TearDownTestCase()
27- std::cout << "StatelessRandomChoiceWithMask TearDown" << std::endl;29+ {
28- }30+ std::cout << "StatelessRandomChoiceWithMask TearDown" << std::endl;
31+ }
29};32};
30 33 
31TEST_F(StatelessRandomChoiceWithMaskTiling, stateless_random_choice_with_mask_tiling_001)34TEST_F(StatelessRandomChoiceWithMaskTiling, stateless_random_choice_with_mask_tiling_001)
32{35{
33 optiling::StatelessRandomChoiceWithMaskCompileInfo compileInfo = {64, 196608};36 optiling::StatelessRandomChoiceWithMaskCompileInfo compileInfo = {64, 196608};
34- gert::StorageShape x_shape = {{1}, {1}};37+ gert::StorageShape xShape = {{1}, {1}};
35- gert::StorageShape count_shape = {{1}, {1}};38+ gert::StorageShape countShape = {{1}, {1}};
36- gert::StorageShape seed_shape = {{1}, {1}};39+ gert::StorageShape seedShape = {{1}, {1}};
37- gert::StorageShape offset_shape = {{1}, {1}};40+ gert::StorageShape offsetShape = {{1}, {1}};
38- gert::StorageShape y_shape = {{1, 1}, {1, 1}};41+ gert::StorageShape yShape = {{1, 1}, {1, 1}};
39- gert::StorageShape mask_shape = {{1}, {1}};42+ gert::StorageShape maskShape = {{1}, {1}};
40 43 
41- vector<int32_t> count_value = {20};44+ vector<int32_t> countValue = {20};
42- vector<int64_t> seed_value = {12};45+ vector<int64_t> seedValue = {12};
43- vector<int64_t> offset_value = {22};46+ vector<int64_t> offsetValue = {22};
44 47 
45 gert::TilingContextPara tilingContextPara(48 gert::TilingContextPara tilingContextPara(
46- "StatelessRandomChoiceWithMask", 49+ "StatelessRandomChoiceWithMask",
47- {50+ {{xShape, ge::DT_BOOL, ge::FORMAT_ND},
48- {x_shape, ge::DT_BOOL, ge::FORMAT_ND},51+ {countShape, ge::DT_INT32, ge::FORMAT_ND, true, countValue.data()},
49- {count_shape, ge::DT_INT32, ge::FORMAT_ND, true, count_value.data()},52+ {seedShape, ge::DT_INT64, ge::FORMAT_ND, true, seedValue.data()},
50- {seed_shape, ge::DT_INT64, ge::FORMAT_ND, true, seed_value.data()},53+ {offsetShape, ge::DT_INT64, ge::FORMAT_ND, true, offsetValue.data()}},
51- {offset_shape, ge::DT_INT64, ge::FORMAT_ND, true, offset_value.data()}54+ {{yShape, ge::DT_INT32, ge::FORMAT_ND}, {maskShape, ge::DT_BOOL, ge::FORMAT_ND}},
52- },
53- {
54- {y_shape, ge::DT_INT32, ge::FORMAT_ND},
55- {mask_shape, ge::DT_BOOL, ge::FORMAT_ND}
56- },
57 &compileInfo);55 &compileInfo);
58- uint64_t expectTilingKey = 0;56+ 
59- std::vector<size_t> expectWorkspaces = {16787712};57+ TilingInfo tilingInfo;
58+ auto tilingRet = ExecuteTiling(tilingContextPara, tilingInfo);
59+ EXPECT_EQ(tilingRet, true);
60+ EXPECT_EQ(tilingInfo.tilingKey, 0);
61+ EXPECT_EQ(tilingInfo.blockNum, 1);
62+ ASSERT_EQ(tilingInfo.workspaceSizes.size(), 1);
63+ EXPECT_EQ(tilingInfo.workspaceSizes[0], 16787456);
60}64}
61 65 
62TEST_F(StatelessRandomChoiceWithMaskTiling, stateless_random_choice_with_mask_tiling_002)66TEST_F(StatelessRandomChoiceWithMaskTiling, stateless_random_choice_with_mask_tiling_002)
63{67{
64 optiling::StatelessRandomChoiceWithMaskCompileInfo compileInfo = {64, 196608};68 optiling::StatelessRandomChoiceWithMaskCompileInfo compileInfo = {64, 196608};
65- gert::StorageShape x_shape = {{66561}, {66561}};69+ gert::StorageShape xShape = {{66561}, {66561}};
66- gert::StorageShape count_shape = {{1}, {1}};70+ gert::StorageShape countShape = {{1}, {1}};
67- gert::StorageShape seed_shape = {{1}, {1}};71+ gert::StorageShape seedShape = {{1}, {1}};
68- gert::StorageShape offset_shape = {{1}, {1}};72+ gert::StorageShape offsetShape = {{1}, {1}};
69- gert::StorageShape y_shape = {{1, 1}, {1, 1}};73+ gert::StorageShape yShape = {{1, 1}, {1, 1}};
70- gert::StorageShape mask_shape = {{1}, {1}};74+ gert::StorageShape maskShape = {{1}, {1}};
71 75 
72- vector<int32_t> count_value = {20};76+ vector<int32_t> countValue = {20};
73- vector<int64_t> seed_value = {12};77+ vector<int64_t> seedValue = {12};
74- vector<int64_t> offset_value = {22};78+ vector<int64_t> offsetValue = {22};
75 79 
76 gert::TilingContextPara tilingContextPara(80 gert::TilingContextPara tilingContextPara(
77- "StatelessRandomChoiceWithMask", 81+ "StatelessRandomChoiceWithMask",
78- {82+ {{xShape, ge::DT_BOOL, ge::FORMAT_ND},
79- {x_shape, ge::DT_BOOL, ge::FORMAT_ND},83+ {countShape, ge::DT_INT32, ge::FORMAT_ND, true, countValue.data()},
80- {count_shape, ge::DT_INT32, ge::FORMAT_ND, true, count_value.data()},84+ {seedShape, ge::DT_INT64, ge::FORMAT_ND, true, seedValue.data()},
81- {seed_shape, ge::DT_INT64, ge::FORMAT_ND, true, seed_value.data()},85+ {offsetShape, ge::DT_INT64, ge::FORMAT_ND, true, offsetValue.data()}},
82- {offset_shape, ge::DT_INT64, ge::FORMAT_ND, true, offset_value.data()}86+ {{yShape, ge::DT_INT32, ge::FORMAT_ND}, {maskShape, ge::DT_BOOL, ge::FORMAT_ND}},
83- },
84- {
85- {y_shape, ge::DT_INT32, ge::FORMAT_ND},
86- {mask_shape, ge::DT_BOOL, ge::FORMAT_ND}
87- },
88 &compileInfo);87 &compileInfo);
89- uint64_t expectTilingKey = 0;88+ 
90- std::vector<size_t> expectWorkspaces = {17838080};89+ TilingInfo tilingInfo;
91-}90+ auto tilingRet = ExecuteTiling(tilingContextPara, tilingInfo);
91+ EXPECT_EQ(tilingRet, true);
92+ EXPECT_EQ(tilingInfo.tilingKey, 0);
93+ EXPECT_EQ(tilingInfo.blockNum, 64);
94+ ASSERT_EQ(tilingInfo.workspaceSizes.size(), 1);
95+ EXPECT_EQ(tilingInfo.workspaceSizes[0], 17582080);
96+}
Arandom/stateless_random_choice_with_mask/tests/ut/op_kernel/CMakeLists.txt+19-0
@@ -0,0 +1,19 @@
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(stateless_random_choice_with_mask_tiling_files
13+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/arch35/stateless_random_choice_with_mask_simt_tiling.cpp)
14+ AddOpTestCase(
15+ stateless_random_choice_with_mask
16+ "ascend950"
17+ ""
18+ "${stateless_random_choice_with_mask_tiling_files}")
19+endif()
Arandom/stateless_random_choice_with_mask/tests/ut/op_kernel/test_stateless_random_choice_with_mask.cpp+89-0
@@ -0,0 +1,89 @@
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 <cstdint>
12+#include <cstring>
13+#include "gtest/gtest.h"
14+#include "tikicpulib.h"
15+#include "../../../op_kernel/stateless_random_choice_with_mask.cpp"
16+ 
17+namespace {
18+constexpr uint32_t kNumBlocks = 1;
19+constexpr uint32_t kSchMode = 0;
20+constexpr int64_t kM = 4;
21+constexpr int64_t kN = 4;
22+constexpr int64_t kInputSize = kM * kN;
23+constexpr int32_t kCount = 2;
24+ 
25+inline size_t Align32(size_t size)
26+{
27+ return (size + 31U) / 32U * 32U;
28+}
29+} // namespace
30+ 
31+class StatelessRandomChoiceWithMaskKernelTest : public testing::Test {
32+};
33+ 
34+TEST_F(StatelessRandomChoiceWithMaskKernelTest, smoke_test)
35+{
36+ auto* x = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(kInputSize * sizeof(int32_t))));
37+ auto* count = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int32_t))));
38+ auto* seed = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int64_t))));
39+ auto* offset = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int64_t))));
40+ auto* y = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(kCount * kN * sizeof(int32_t))));
41+ auto* mask = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(kCount * sizeof(int32_t))));
42+ auto* shapeOut = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(2 * sizeof(int32_t))));
43+ auto* workspace = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(16 * 1024 * 1024)));
44+ auto* tiling = static_cast<uint8_t*>(
45+ AscendC::GmAlloc(Align32(sizeof(StatelessRandomChoiceWithMaskSimtTilingData))));
46+ 
47+ auto* xData = reinterpret_cast<int32_t*>(x);
48+ for (int64_t i = 0; i < kInputSize; ++i) {
49+ xData[i] = 1;
50+ }
51+ *reinterpret_cast<int32_t*>(count) = kCount;
52+ *reinterpret_cast<int64_t*>(seed) = 42;
53+ *reinterpret_cast<int64_t*>(offset) = 0;
54+ std::memset(y, 0, Align32(kCount * kN * sizeof(int32_t)));
55+ std::memset(mask, 0, Align32(kCount * sizeof(int32_t)));
56+ std::memset(workspace, 0, Align32(16 * 1024 * 1024));
57+ 
58+ auto* tilingData = reinterpret_cast<StatelessRandomChoiceWithMaskSimtTilingData*>(tiling);
59+ std::memset(tilingData, 0, sizeof(StatelessRandomChoiceWithMaskSimtTilingData));
60+ tilingData->blockNum = kNumBlocks;
61+ tilingData->normalCoreProNum = kInputSize;
62+ tilingData->m = kM;
63+ tilingData->n = kN;
64+ tilingData->seed = 42;
65+ tilingData->offset = 0;
66+ tilingData->inputSize = kInputSize;
67+ tilingData->noZeroCalcCount = kInputSize;
68+ tilingData->noZeroWorkspaceSize = 4096;
69+ tilingData->randomWorkspaceSize = 4096;
70+ tilingData->ubSize = 65536;
71+ tilingData->count = kCount;
72+ tilingData->inputDim = 2;
73+ tilingData->inputShape[0] = kM;
74+ tilingData->inputShape[1] = kN;
75+ 
76+ auto func = stateless_random_choice_with_mask<kSchMode>;
77+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
78+ ICPU_RUN_KF(func, kNumBlocks, x, count, seed, offset, y, mask, shapeOut, workspace, tiling);
79+ 
80+ AscendC::GmFree(x);
81+ AscendC::GmFree(count);
82+ AscendC::GmFree(seed);
83+ AscendC::GmFree(offset);
84+ AscendC::GmFree(y);
85+ AscendC::GmFree(mask);
86+ AscendC::GmFree(shapeOut);
87+ AscendC::GmFree(workspace);
88+ AscendC::GmFree(tiling);
89+}
Arandom/stateless_random_normal_v2/tests/ut/op_kernel/CMakeLists.txt+19-0
@@ -0,0 +1,19 @@
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(stateless_random_normal_v2_tiling_files
13+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/arch35/stateless_random_normal_v2_tiling_arch35.cpp)
14+ AddOpTestCase(
15+ stateless_random_normal_v2
16+ "ascend950"
17+ ""
18+ "${stateless_random_normal_v2_tiling_files}")
19+endif()
Arandom/stateless_random_normal_v2/tests/ut/op_kernel/test_stateless_random_normal_v2.cpp+89-0
@@ -0,0 +1,89 @@
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 <cstdint>
12+#include <cstring>
13+#include "gtest/gtest.h"
14+#include "tikicpulib.h"
15+ 
16+extern "C" __global__ __aicore__ void stateless_random_normal_v2(
17+ GM_ADDR shape, GM_ADDR key, GM_ADDR counter, GM_ADDR alg, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling);
18+ 
19+namespace {
20+constexpr uint32_t kNumBlocks = 1;
21+constexpr uint64_t kTilingKey = 101;
22+constexpr int64_t kElementCount = 256;
23+ 
24+struct StatelessRandomNormalV2TilingLayout {
25+ uint32_t blockNum;
26+ uint32_t blockTilingSize;
27+ uint32_t tailBlockTilingSize;
28+ uint32_t ubTilingSize;
29+ uint32_t alg;
30+ uint32_t key[2];
31+ uint32_t counter[4];
32+};
33+ 
34+inline size_t Align32(size_t size)
35+{
36+ return (size + 31U) / 32U * 32U;
37+}
38+} // namespace
39+ 
40+class StatelessRandomNormalV2KernelTest : public testing::Test {
41+};
42+ 
43+TEST_F(StatelessRandomNormalV2KernelTest, smoke_float)
44+{
45+ auto* shape = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(2 * sizeof(int32_t))));
46+ auto* key = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(2 * sizeof(uint32_t))));
47+ auto* counter = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(4 * sizeof(uint32_t))));
48+ auto* alg = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int32_t))));
49+ auto* y = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(kElementCount * sizeof(float))));
50+ auto* workspace = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(16 * 1024 * 1024)));
51+ auto* tiling = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(StatelessRandomNormalV2TilingLayout))));
52+ 
53+ std::memset(y, 0, kElementCount * sizeof(float));
54+ std::memset(tiling, 0, sizeof(StatelessRandomNormalV2TilingLayout));
55+ reinterpret_cast<int32_t*>(shape)[0] = 16;
56+ reinterpret_cast<int32_t*>(shape)[1] = 16;
57+ reinterpret_cast<uint32_t*>(key)[0] = 42;
58+ reinterpret_cast<uint32_t*>(key)[1] = 0;
59+ reinterpret_cast<uint32_t*>(counter)[0] = 0;
60+ reinterpret_cast<uint32_t*>(counter)[1] = 0;
61+ reinterpret_cast<uint32_t*>(counter)[2] = 0;
62+ reinterpret_cast<uint32_t*>(counter)[3] = 0;
63+ *reinterpret_cast<int32_t*>(alg) = 1;
64+ 
65+ auto* tilingData = reinterpret_cast<StatelessRandomNormalV2TilingLayout*>(tiling);
66+ tilingData->blockNum = kNumBlocks;
67+ tilingData->blockTilingSize = kElementCount;
68+ tilingData->tailBlockTilingSize = kElementCount;
69+ tilingData->ubTilingSize = kElementCount;
70+ tilingData->alg = 1;
71+ tilingData->key[0] = 42;
72+ tilingData->key[1] = 0;
73+ tilingData->counter[0] = 0;
74+ tilingData->counter[1] = 0;
75+ tilingData->counter[2] = 0;
76+ tilingData->counter[3] = 0;
77+ 
78+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
79+ ICPU_SET_TILING_KEY(kTilingKey);
80+ ICPU_RUN_KF(stateless_random_normal_v2, kNumBlocks, shape, key, counter, alg, y, workspace, tiling);
81+ 
82+ AscendC::GmFree(shape);
83+ AscendC::GmFree(key);
84+ AscendC::GmFree(counter);
85+ AscendC::GmFree(alg);
86+ AscendC::GmFree(y);
87+ AscendC::GmFree(workspace);
88+ AscendC::GmFree(tiling);
89+}
Arandom/stateless_random_uniform_v2/tests/ut/op_kernel/CMakeLists.txt+19-0
@@ -0,0 +1,19 @@
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(stateless_random_uniform_v2_tiling_files
13+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/arch35/stateless_random_uniform_v2_tiling_arch35.cpp)
14+ AddOpTestCase(
15+ stateless_random_uniform_v2
16+ "ascend950"
17+ ""
18+ "${stateless_random_uniform_v2_tiling_files}")
19+endif()
Arandom/stateless_random_uniform_v2/tests/ut/op_kernel/test_stateless_random_uniform_v2.cpp+89-0
@@ -0,0 +1,89 @@
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 <cstdint>
12+#include <cstring>
13+#include "gtest/gtest.h"
14+#include "tikicpulib.h"
15+ 
16+extern "C" __global__ __aicore__ void stateless_random_uniform_v2(
17+ GM_ADDR shape, GM_ADDR key, GM_ADDR counter, GM_ADDR alg, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling);
18+ 
19+namespace {
20+constexpr uint32_t kNumBlocks = 1;
21+constexpr uint64_t kTilingKey = 101;
22+constexpr int64_t kElementCount = 256;
23+ 
24+struct StatelessRandomUniformV2TilingLayout {
25+ uint32_t blockNum;
26+ uint32_t blockTilingSize;
27+ uint32_t tailBlockTilingSize;
28+ uint32_t ubTilingSize;
29+ uint32_t alg;
30+ uint32_t key[2];
31+ uint32_t counter[4];
32+};
33+ 
34+inline size_t Align32(size_t size)
35+{
36+ return (size + 31U) / 32U * 32U;
37+}
38+} // namespace
39+ 
40+class StatelessRandomUniformV2KernelTest : public testing::Test {
41+};
42+ 
43+TEST_F(StatelessRandomUniformV2KernelTest, smoke_float)
44+{
45+ auto* shape = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(2 * sizeof(int32_t))));
46+ auto* key = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(2 * sizeof(uint32_t))));
47+ auto* counter = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(4 * sizeof(uint32_t))));
48+ auto* alg = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int32_t))));
49+ auto* y = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(kElementCount * sizeof(float))));
50+ auto* workspace = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(16 * 1024 * 1024)));
51+ auto* tiling = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(StatelessRandomUniformV2TilingLayout))));
52+ 
53+ std::memset(y, 0, kElementCount * sizeof(float));
54+ std::memset(tiling, 0, sizeof(StatelessRandomUniformV2TilingLayout));
55+ reinterpret_cast<int32_t*>(shape)[0] = 16;
56+ reinterpret_cast<int32_t*>(shape)[1] = 16;
57+ reinterpret_cast<uint32_t*>(key)[0] = 42;
58+ reinterpret_cast<uint32_t*>(key)[1] = 0;
59+ reinterpret_cast<uint32_t*>(counter)[0] = 0;
60+ reinterpret_cast<uint32_t*>(counter)[1] = 0;
61+ reinterpret_cast<uint32_t*>(counter)[2] = 0;
62+ reinterpret_cast<uint32_t*>(counter)[3] = 0;
63+ *reinterpret_cast<int32_t*>(alg) = 1;
64+ 
65+ auto* tilingData = reinterpret_cast<StatelessRandomUniformV2TilingLayout*>(tiling);
66+ tilingData->blockNum = kNumBlocks;
67+ tilingData->blockTilingSize = kElementCount;
68+ tilingData->tailBlockTilingSize = kElementCount;
69+ tilingData->ubTilingSize = kElementCount;
70+ tilingData->alg = 1;
71+ tilingData->key[0] = 42;
72+ tilingData->key[1] = 0;
73+ tilingData->counter[0] = 0;
74+ tilingData->counter[1] = 0;
75+ tilingData->counter[2] = 0;
76+ tilingData->counter[3] = 0;
77+ 
78+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
79+ ICPU_SET_TILING_KEY(kTilingKey);
80+ ICPU_RUN_KF(stateless_random_uniform_v2, kNumBlocks, shape, key, counter, alg, y, workspace, tiling);
81+ 
82+ AscendC::GmFree(shape);
83+ AscendC::GmFree(key);
84+ AscendC::GmFree(counter);
85+ AscendC::GmFree(alg);
86+ AscendC::GmFree(y);
87+ AscendC::GmFree(workspace);
88+ AscendC::GmFree(tiling);
89+}
Mrandom/stateless_randperm/tests/ut/op_kernel/CMakeLists.txt+22-7
@@ -1,5 +1,5 @@
1# ----------------------------------------------------------------------------1# ----------------------------------------------------------------------------
2-# Copyright (c) 2025 Huawei Technologies Co., Ltd.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 of3# 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").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.5# Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -8,9 +8,24 @@
8# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
9# ----------------------------------------------------------------------------9# ----------------------------------------------------------------------------
10 10 
11-file(GLOB CURRENT_SOURCE_DIRS LIST_DIRECTORIES true ${CMAKE_CURRENT_SOURCE_DIR}/*)11+if(UT_TEST_ALL OR OP_KERNEL_UT)
12-foreach(SUB_DIR ${CURRENT_SOURCE_DIRS})12+ # Kernel code uses #include "../../sort/arch35/sort_tiling_data.h" etc., which is designed
13- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")13+ # for the flat staging directory produced by kernel_src_copy during production builds.
14- add_subdirectory(${SUB_DIR})14+ # In UT the kernel is compiled in-place (with the extra op_kernel/ level), so we create a
15- endif()15+ # mirrored staging layout in the build directory and pass it via -I so the compiler can
16-endforeach()16+ # resolve the relative include paths on fallback.
17+ set(KERNEL_STAGING_DIR ${CMAKE_CURRENT_BINARY_DIR}/kernel_dep_staging)
18+ file(MAKE_DIRECTORY ${KERNEL_STAGING_DIR}/stateless_randperm/arch35)
19+ execute_process(COMMAND ${CMAKE_COMMAND} -E create_symlink
20+ ${PROJECT_SOURCE_DIR}/math/sort/op_kernel
21+ ${KERNEL_STAGING_DIR}/sort)
22+ 
23+ set(stateless_randperm_tiling_files
24+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/stateless_randperm_tiling_arch35.cpp
25+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/stateless_randperm_tiling_for_sort.cpp)
26+ AddOpTestCase(
27+ stateless_randperm
28+ "ascend950"
29+ "-DDTYPE_Y=int32_t -DASCENDC_TPL_KERNEL -I${KERNEL_STAGING_DIR}/stateless_randperm/arch35"
30+ "${stateless_randperm_tiling_files}")
31+endif()
Arandom/stateless_randperm/tests/ut/op_kernel/test_stateless_randperm.cpp+83-0
@@ -0,0 +1,83 @@
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 <cstdint>
12+#include <cstring>
13+#include "gtest/gtest.h"
14+#include "tikicpulib.h"
15+#include "../../../op_kernel/stateless_randperm_apt.cpp"
16+ 
17+namespace {
18+constexpr uint32_t kNumBlocks = 1;
19+constexpr int64_t kN = 8;
20+ 
21+inline size_t Align32(size_t size)
22+{
23+ return (size + 31U) / 32U * 32U;
24+}
25+} // namespace
26+ 
27+class StatelessRandpermKernelTest : public testing::Test {
28+};
29+ 
30+TEST_F(StatelessRandpermKernelTest, smoke_int32)
31+{
32+ auto* n = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int32_t))));
33+ auto* seed = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int64_t))));
34+ auto* offset = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int64_t))));
35+ auto* y = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(kN * sizeof(int32_t))));
36+ auto* workspace = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(16 * 1024 * 1024)));
37+ auto* tiling = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(StatelessRandpermTilingData))));
38+ 
39+ *reinterpret_cast<int32_t*>(n) = kN;
40+ *reinterpret_cast<int64_t*>(seed) = 42;
41+ *reinterpret_cast<int64_t*>(offset) = 0;
42+ std::memset(y, 0, kN * sizeof(int32_t));
43+ std::memset(workspace, 0, Align32(16 * 1024 * 1024));
44+ 
45+ auto* tilingData = reinterpret_cast<StatelessRandpermTilingData*>(tiling);
46+ std::memset(tilingData, 0, sizeof(StatelessRandpermTilingData));
47+ tilingData->n = kN;
48+ tilingData->randomBits = 32;
49+ tilingData->islandFactor = 1;
50+ tilingData->islandFactorTail = 1;
51+ tilingData->castFactor = 1;
52+ tilingData->castFactorTail = 1;
53+ tilingData->realCoreNum = kNumBlocks;
54+ tilingData->randomWkSizeByte = 1024;
55+ tilingData->subNTileCount = 1;
56+ tilingData->subNTile[0] = kN;
57+ tilingData->philoxKey[0] = 42;
58+ tilingData->philoxKey[1] = 0;
59+ tilingData->philoxOffset = 0;
60+ 
61+ tilingData->sortTilingData.numTileDataSize = kN;
62+ tilingData->sortTilingData.unsortedDimParallel = 1;
63+ tilingData->sortTilingData.lastDimTileNum = 1;
64+ tilingData->sortTilingData.sortLoopTimes = 1;
65+ tilingData->sortTilingData.lastDimNeedCore = 1;
66+ tilingData->sortTilingData.keyParams0 = 1;
67+ tilingData->sortTilingData.keyParams1 = 256;
68+ tilingData->sortTilingData.tmpUbSize = 4096;
69+ tilingData->sortTilingData.lastAxisNum = kN;
70+ tilingData->sortTilingData.unsortedDimNum = 1;
71+ 
72+ // Template params: randomType=2(int32), nIsInt32=1, schId=0, isInt32=1, isDescend=0
73+ auto func = stateless_randperm<2, 1, 0, 1, 0>;
74+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
75+ ICPU_RUN_KF(func, kNumBlocks, n, seed, offset, y, workspace, tiling);
76+ 
77+ AscendC::GmFree(n);
78+ AscendC::GmFree(seed);
79+ AscendC::GmFree(offset);
80+ AscendC::GmFree(y);
81+ AscendC::GmFree(workspace);
82+ AscendC::GmFree(tiling);
83+}
Arandom/truncated_normal_v2/tests/ut/op_kernel/CMakeLists.txt+25-0
@@ -0,0 +1,25 @@
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}/truncated_normal_v2/arch35)
14+ execute_process(COMMAND ${CMAKE_COMMAND} -E create_symlink
15+ ${PROJECT_SOURCE_DIR}/random/random_common/op_kernel
16+ ${KERNEL_STAGING_DIR}/random_common)
17+ 
18+ set(truncated_normal_v2_tiling_files
19+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/arch35/truncated_normal_v2_tiling_arch35.cpp)
20+ AddOpTestCase(
21+ truncated_normal_v2
22+ "ascend950"
23+ "-DDTYPE_Y=float -DDTYPE_OFFSET=int64_t -DTestUtDefaultTilingStruct=RandomUnifiedSimtTilingDataStruct -I${KERNEL_STAGING_DIR}/truncated_normal_v2/arch35"
24+ "${truncated_normal_v2_tiling_files}")
25+endif()
Arandom/truncated_normal_v2/tests/ut/op_kernel/test_truncated_normal_v2.cpp+89-0
@@ -0,0 +1,89 @@
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 <cmath>
12+#include <cstdint>
13+#include <cstring>
14+#include <iostream>
15+#include "gtest/gtest.h"
16+#include "tikicpulib.h"
17+#include "../../../../random_common/op_kernel/arch35/random_unified_tiling_data_arch35.h"
18+ 
19+extern "C" __global__ __aicore__ void truncated_normal_v2(
20+ GM_ADDR shape, GM_ADDR offset, GM_ADDR y, GM_ADDR offset_ref, GM_ADDR workspace, GM_ADDR tiling);
21+ 
22+namespace {
23+constexpr uint32_t kNumBlocks = 1;
24+constexpr uint64_t kTilingKey = 100;
25+constexpr int64_t kElementCount = 256;
26+constexpr int64_t kExpectedOffset = kElementCount * 256;
27+constexpr float kRandomThreadR = 2.0f;
28+ 
29+inline size_t Align32(size_t size)
30+{
31+ return (size + 31U) / 32U * 32U;
32+}
33+} // namespace
34+ 
35+class TruncatedNormalV2KernelTest : public testing::Test {
36+protected:
37+ static void SetUpTestCase()
38+ {
39+ std::cout << "TruncatedNormalV2KernelTest SetUp" << std::endl;
40+ }
41+ 
42+ static void TearDownTestCase()
43+ {
44+ std::cout << "TruncatedNormalV2KernelTest TearDown" << std::endl;
45+ }
46+};
47+ 
48+TEST_F(TruncatedNormalV2KernelTest, smoke_float)
49+{
50+ auto* shape = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(2 * sizeof(int32_t))));
51+ auto* offset = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int64_t))));
52+ auto* y = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(kElementCount * sizeof(float))));
53+ auto* offsetRef = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int64_t))));
54+ auto* workspace = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(1024 * 1024)));
55+ auto* tiling = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(RandomUnifiedSimtTilingDataStruct))));
56+ 
57+ std::memset(y, 0, kElementCount * sizeof(float));
58+ std::memset(offsetRef, 0, sizeof(int64_t));
59+ std::memset(tiling, 0, sizeof(RandomUnifiedSimtTilingDataStruct));
60+ reinterpret_cast<int32_t*>(shape)[0] = 16;
61+ reinterpret_cast<int32_t*>(shape)[1] = 16;
62+ reinterpret_cast<int64_t*>(offset)[0] = 0;
63+ reinterpret_cast<int64_t*>(offsetRef)[0] = 0;
64+ 
65+ auto* tilingData = reinterpret_cast<RandomUnifiedSimtTilingDataStruct*>(tiling);
66+ tilingData->usedCoreNum = kNumBlocks;
67+ tilingData->outputSize = kElementCount;
68+ tilingData->seed = 10;
69+ tilingData->offset = 5;
70+ 
71+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
72+ ICPU_SET_TILING_KEY(kTilingKey);
73+ ICPU_RUN_KF(truncated_normal_v2, kNumBlocks, shape, offset, y, offsetRef, workspace, tiling);
74+ 
75+ EXPECT_EQ(reinterpret_cast<int64_t*>(offset)[0], kExpectedOffset);
76+ EXPECT_EQ(reinterpret_cast<int64_t*>(offsetRef)[0], 0);
77+ 
78+ auto* yData = reinterpret_cast<float*>(y);
79+ for (int64_t i = 0; i < kElementCount; ++i) {
80+ EXPECT_LT(std::abs(yData[i]), kRandomThreadR);
81+ }
82+ 
83+ AscendC::GmFree(shape);
84+ AscendC::GmFree(offset);
85+ AscendC::GmFree(y);
86+ AscendC::GmFree(offsetRef);
87+ AscendC::GmFree(workspace);
88+ AscendC::GmFree(tiling);
89+}