fix bug： stateless normal导致host重启-ops-math-AtomGit

Mrandom/stateless_normal/op_kernel/arch35/stateless_normal.h+24-25

	@@ -37,15 +37,13 @@ static constexpr uint16_t NUM_TWO = 2;
37	// For bf16/fp16: three-step rounding to match GPU normal_()→mul_(std)→add_(mean)	37	// For bf16/fp16: three-step rounding to match GPU normal_()→mul_(std)→add_(mean)
38	template <typename T, typename M_T, typename S_T>	38	template <typename T, typename M_T, typename S_T>
39	struct NormalTransform {	39	struct NormalTransform {
40	- __gm__ volatile M_T* meanGM_;	40	+ M_T meanVal_;
41	- __gm__ volatile S_T* stdevGM_;	41	+ S_T stdVal_;
42		42
43	- __aicore__ NormalTransform(__gm__ volatile M_T* meanGM, __gm__ volatile S_T* stdevGM)	43	+ __aicore__ NormalTransform(M_T meanVal, S_T stdVal) : meanVal_(meanVal), stdVal_(stdVal) {}
44	- : meanGM_(meanGM), stdevGM_(stdevGM) {}
45		44
46	- __simt_callee__ __aicore__ inline void operator()(	45	+ __simt_callee__ __aicore__ inline void operator()(__gm__ volatile T* outputGm, uint64_t li, const uint32_t* results,
47	- __gm__ volatile T* outputGm, uint64_t li, const uint32_t* results, uint32_t iStep,	46	+ uint32_t iStep, [[maybe_unused]] uint32_t unroll = 1)
48	- [[maybe_unused]] uint32_t unroll = 1)
49	{	47	{
50	uint32_t pairBase = (iStep / NUM_TWO) * NUM_TWO;	48	uint32_t pairBase = (iStep / NUM_TWO) * NUM_TWO;
51	float u1 = results[pairBase] * RAND_2POW32_INV + RAND_2POW32_INV_HALF;	49	float u1 = results[pairBase] * RAND_2POW32_INV + RAND_2POW32_INV_HALF;
	@@ -55,7 +53,7 @@ struct NormalTransform {
55	BoxMullerFloat(u1, u2, &z0, &z1);	53	BoxMullerFloat(u1, u2, &z0, &z1);
56	float z = (iStep % NUM_TWO == 0) ? z0 : z1;	54	float z = (iStep % NUM_TWO == 0) ? z0 : z1;
57		55
58	- outputGm[li] = static_cast<T>(z * static_cast<float>(stdevGM_[0]) + static_cast<float>(meanGM_[0]));	56	+ outputGm[li] = static_cast<T>(z * static_cast<float>(stdVal_) + static_cast<float>(meanVal_));
59	}	57	}
60	};	58	};
61		59
	@@ -68,37 +66,38 @@ struct NormalLauncher {
68	GM_ADDR yAddr_;	66	GM_ADDR yAddr_;
69	GM_ADDR meanAddr_;	67	GM_ADDR meanAddr_;
70	GM_ADDR stdevAddr_;	68	GM_ADDR stdevAddr_;
		69	+ AscendC::GlobalTensor<M_T> meanGlobal_;
		70	+ AscendC::GlobalTensor<S_T> stdGlobal_;
71		71
72	__aicore__ NormalLauncher(int64_t seed, int64_t realOffset, GM_ADDR y, GM_ADDR mean, GM_ADDR stdev)	72	__aicore__ NormalLauncher(int64_t seed, int64_t realOffset, GM_ADDR y, GM_ADDR mean, GM_ADDR stdev)
73	- : seed_(seed), realOffset_(realOffset), yAddr_(y), meanAddr_(mean), stdevAddr_(stdev) {}	73	+ : seed_(seed), realOffset_(realOffset), yAddr_(y), meanAddr_(mean), stdevAddr_(stdev)
		74	+ {}
74		75
75	- __aicore__ inline void operator()(	76	+ __aicore__ inline void operator()(const ExecutionPolicyKernel& policy, int64_t gmOffset, int64_t kernelOffset,
76	- const ExecutionPolicyKernel& policy,	77	+ int64_t numel, [[maybe_unused]] int64_t grid, int64_t totalThreads)
77	- int64_t gmOffset,
78	- int64_t kernelOffset,
79	- int64_t numel,
80	- [[maybe_unused]] int64_t grid,
81	- int64_t totalThreads)
82	{	78	{
83	__gm__ volatile T* gmPtr = reinterpret_cast<__gm__ volatile T*>(yAddr_) + gmOffset;	79	__gm__ volatile T* gmPtr = reinterpret_cast<__gm__ volatile T*>(yAddr_) + gmOffset;
84	// mean/stdev 始终按 float* 读取（L2 传入 DT_FLOAT tensor）	80	// mean/stdev 始终按 float* 读取（L2 传入 DT_FLOAT tensor）
85	- __gm__ volatile M_T* meanPtr = reinterpret_cast<__gm__ volatile M_T*>(meanAddr_);	81	+ __gm__ M_T* meanPtr = reinterpret_cast<__gm__ M_T*>(meanAddr_);
86	- __gm__ volatile S_T* stdevPtr = reinterpret_cast<__gm__ volatile S_T*>(stdevAddr_);	82	+ __gm__ S_T* stdevPtr = reinterpret_cast<__gm__ S_T*>(stdevAddr_);
87		83
88	- NormalTransform<T, M_T, S_T> transform(meanPtr, stdevPtr);	84	+ meanGlobal_.SetGlobalBuffer(meanPtr);
		85	+ stdGlobal_.SetGlobalBuffer(stdevPtr);
		86	+
		87	+ M_T meanVal = meanGlobal_(0);
		88	+ S_T stdVal = stdGlobal_(0);
		89	+
		90	+ NormalTransform<T, M_T, S_T> transform(meanVal, stdVal);
89	Simt::VF_CALL<PhiloxSimtKernelDiscontinuous<T, NormalTransform<T, M_T, S_T>>>(	91	Simt::VF_CALL<PhiloxSimtKernelDiscontinuous<T, NormalTransform<T, M_T, S_T>>>(
90	- Simt::Dim3(DEFAULT_SIMT_THREAD_NUM),	92	+ Simt::Dim3(DEFAULT_SIMT_THREAD_NUM), gmPtr, realOffset_ + kernelOffset, seed_, static_cast<uint64_t>(numel),
91	- gmPtr, realOffset_ + kernelOffset, seed_, static_cast<uint64_t>(numel),
92	policy.magic, policy.shift, static_cast<uint64_t>(totalThreads), transform);	93	policy.magic, policy.shift, static_cast<uint64_t>(totalThreads), transform);
93	}	94	}
94	};	95	};
95		96
96	// Entry point: called from stateless_normal.cpp	97	// Entry point: called from stateless_normal.cpp
97	template <typename T, typename M_T, typename S_T>	98	template <typename T, typename M_T, typename S_T>
98	-__aicore__ inline void Process(	99	+__aicore__ inline void Process(GM_ADDR seed, GM_ADDR offset, GM_ADDR y, GM_ADDR mean, GM_ADDR stdev,
99	- GM_ADDR seed, GM_ADDR offset,	100	+ const RandomUnifiedSimtTilingDataStruct* __restrict tilingData)
100	- GM_ADDR y, GM_ADDR mean, GM_ADDR stdev,
101	- const RandomUnifiedSimtTilingDataStruct* __restrict tilingData)
102	{	101	{
103	if (GetBlockIdx() >= static_cast<uint32_t>(tilingData->usedCoreNum)) {	102	if (GetBlockIdx() >= static_cast<uint32_t>(tilingData->usedCoreNum)) {
104	return;	103	return;

Arandom/stateless_normal/tests/ut/op_kernel/stateless_normal_tiling.h+57-0

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

Mrandom/stateless_normal/tests/ut/op_kernel/test_stateless_normal.cpp+8-13

	@@ -14,10 +14,10 @@
14	#include "gtest/gtest.h"	14	#include "gtest/gtest.h"
15	#include "tikicpulib.h"	15	#include "tikicpulib.h"
16	#include "../../../../random_common/op_kernel/arch35/random_unified_tiling_data_arch35.h"	16	#include "../../../../random_common/op_kernel/arch35/random_unified_tiling_data_arch35.h"
		17	+#include "./stateless_normal_tiling.h"
17		18
18	-extern "C" __global__ __aicore__ void stateless_normal(	19	+extern "C" __global__ __aicore__ void stateless_normal(GM_ADDR shape, GM_ADDR seed, GM_ADDR offset, GM_ADDR mean,
19	- GM_ADDR shape, GM_ADDR seed, GM_ADDR offset, GM_ADDR mean, GM_ADDR stdev,	20	+ GM_ADDR stdev, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling);
20	- GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling);
21		21
22	namespace {	22	namespace {
23		23
	@@ -39,10 +39,7 @@ constexpr int64_t MAX_THREADS_PER_SM = 2048;
39	constexpr int64_t BLOCKS_PER_SM = MAX_THREADS_PER_SM / GPU_BLOCK_SIZE;	39	constexpr int64_t BLOCKS_PER_SM = MAX_THREADS_PER_SM / GPU_BLOCK_SIZE;
40	constexpr int64_t MAX_GENERATOR_OFFSETS = 4;	40	constexpr int64_t MAX_GENERATOR_OFFSETS = 4;
41		41
42	-inline size_t Align32(size_t size)	42	+inline size_t Align32(size_t size) { return (size + 31U) / 32U * 32U; }
43	-{
44	- return (size + 31U) / 32U * 32U;
45	-}
46		43
47	// Compute execution policy for a split block (mirrors CalcExecutionPoliciesForBlocks)	44	// Compute execution policy for a split block (mirrors CalcExecutionPoliciesForBlocks)
48	inline void CalcBlockPolicy(int64_t numel, int64_t& grid, int64_t& totalThreads, int64_t& counterOffset)	45	inline void CalcBlockPolicy(int64_t numel, int64_t& grid, int64_t& totalThreads, int64_t& counterOffset)
	@@ -60,15 +57,14 @@ inline void CalcBlockPolicy(int64_t numel, int64_t& grid, int64_t& totalThreads,
60	}	57	}
61		58
62	// Fill a single split block with execution policy	59	// Fill a single split block with execution policy
63	-inline void FillSingleBlock(RandomUnifiedSimtTilingDataStruct* tilingData, int64_t numel, int64_t seedVal, int64_t offsetVal)	60	+inline void FillSingleBlock(RandomUnifiedSimtTilingDataStruct* tilingData, int64_t numel, int64_t seedVal,
		61	+ int64_t offsetVal)
64	{	62	{
65	tilingData->splitBlockCount = 1;	63	tilingData->splitBlockCount = 1;
66	tilingData->splitBlocks[0].numel = numel;	64	tilingData->splitBlocks[0].numel = numel;
67	tilingData->splitBlocks[0].gmOffset = 0;	65	tilingData->splitBlocks[0].gmOffset = 0;
68	int64_t counterOffset = 0;	66	int64_t counterOffset = 0;
69	- CalcBlockPolicy(numel, tilingData->splitBlocks[0].grid,	67	+ CalcBlockPolicy(numel, tilingData->splitBlocks[0].grid, tilingData->splitBlocks[0].totalThreads, counterOffset);
70	- tilingData->splitBlocks[0].totalThreads,
71	- counterOffset);
72	tilingData->splitBlocks[0].kernelOffset = offsetVal;	68	tilingData->splitBlocks[0].kernelOffset = offsetVal;
73	}	69	}
74		70
	@@ -78,8 +74,7 @@ inline float ReadAsFloat(const uint8_t* buf, int64_t idx)
78	}	74	}
79	} // namespace	75	} // namespace
80		76
81	-class StatelessNormalKernelTest : public testing::Test {	77	+class StatelessNormalKernelTest : public testing::Test {};
82	-};
83		78
84	// Test 1: standard normal N(0,1), mean/stdev broadcast to full tensor	79	// Test 1: standard normal N(0,1), mean/stdev broadcast to full tensor
85	TEST_F(StatelessNormalKernelTest, smoke_standard_normal)	80	TEST_F(StatelessNormalKernelTest, smoke_standard_normal)