refactor(example): rename NPU_ARCH to NPU_SOC_VERSION with backward compat-ops-cv-AtomGit

已合并

refactor(example): rename NPU_ARCH to NPU_SOC_VERSION with backward compat #1079

liu-wei创建于 7月2日从liu-wei/ops-cv： 9.1.0_fast_kernel 合入到 cann/ops-cv： 9.1.0

refactor(example): rename NPU_ARCH to NPU_SOC_VERSION with backward compat #1079

已合并

liu-wei创建于 7月2日

从liu-wei/ops-cv 9.1.0_fast_kernel合入到cann/ops-cv 9.1.0

共 4 个文件变更+28-14

Mexamples/fast_kernel_launch_example/CMakeLists.txt+10-2

	@@ -23,8 +23,16 @@ include(cmake/torch.cmake)
23	include(cmake/torch_npu.cmake)	23	include(cmake/torch_npu.cmake)
24	include(cmake/func.cmake)	24	include(cmake/func.cmake)
25		25
26	-# set NPU architecture, default to ascend910b	26	+# set NPU SOC version, default to ascend910b
27	-set(NPU_ARCH "ascend910b" CACHE STRING "NPU architecture")	27	+if(NOT DEFINED NPU_SOC_VERSION)
		28	+ if(DEFINED NPU_ARCH)
		29	+ set(NPU_SOC_VERSION "${NPU_ARCH}" CACHE STRING "NPU SOC version")
		30	+ else()
		31	+ set(NPU_SOC_VERSION "ascend910b" CACHE STRING "NPU SOC version")
		32	+ endif()
		33	+else()
		34	+ set(NPU_SOC_VERSION "${NPU_SOC_VERSION}" CACHE STRING "NPU SOC version")
		35	+endif()
28	# set EXTENSION_MODULE_NAME to `ascend_ops`	36	# set EXTENSION_MODULE_NAME to `ascend_ops`
29	set(EXTENSION_MODULE_NAME "ascend_ops" CACHE STRING "Extension module name")	37	set(EXTENSION_MODULE_NAME "ascend_ops" CACHE STRING "Extension module name")
30		38

Mexamples/fast_kernel_launch_example/README.md+8-6

	@@ -31,6 +31,8 @@
31	3. 构建Wheel包 \| Build the Wheel：	31	3. 构建Wheel包 \| Build the Wheel：
32		32
33	```sh	33	```sh
		34	+ # NPU_SOC_VERSION设置编译款型, Atlas A2系列产品使用"ascend910b"（默认），Atlas A3系列产品使用"ascend910_93"，Ascend 950PR/Ascend 950DT产品使用"ascend950"
		35	+ export NPU_SOC_VERSION=ascend910b
34	# -n: non-isolated build (uses existing environment)	36	# -n: non-isolated build (uses existing environment)
35	python3 -m build --wheel -n	37	python3 -m build --wheel -n
36	```	38	```
	@@ -151,7 +153,7 @@ print("Verification successful!")
151		153
152	/**	154	/**
153	* 实现算子调用接口	155	* 实现算子调用接口
154	- * 在这个接口中, 需要完成NPU Kernel的调用	156	+ * 在这个接口中，需要完成NPU Kernel的调用
155	* 1. 计算出输出的Tensor的个数/Shape/Dtype(可以调用Meta函数实现，也可以直接实现)	157	* 1. 计算出输出的Tensor的个数/Shape/Dtype(可以调用Meta函数实现，也可以直接实现)
156	* 2. 计算Tiling：根据Shape得到如何分块计算	158	* 2. 计算Tiling：根据Shape得到如何分块计算
157	* 3. 调用NPU Kernel	159	* 3. 调用NPU Kernel
	@@ -164,9 +166,9 @@ print("Verification successful!")
164	const c10::OptionalDeviceGuard guard(x.device());	166	const c10::OptionalDeviceGuard guard(x.device());
165	auto z = add_meta(x, y);	167	auto z = add_meta(x, y);
166	auto stream = c10_npu::getCurrentNPUStream().stream(false);	168	auto stream = c10_npu::getCurrentNPUStream().stream(false);
167	- int64_t totalLength, blockDim, blockLength, tileSize;	169	+ int64_t totalLength, numBlocks, blockLength, tileSize;
168	totalLength = x.numel();	170	totalLength = x.numel();
169	- std::tie(blockDim, blockLength, tileSize) = calc_tiling_params(totalLength);	171	+ std::tie(numBlocks, blockLength, tileSize) = calc_tiling_params(totalLength);
170	auto x_ptr = (GM_ADDR)x.data_ptr();	172	auto x_ptr = (GM_ADDR)x.data_ptr();
171	auto y_ptr = (GM_ADDR)y.data_ptr();	173	auto y_ptr = (GM_ADDR)y.data_ptr();
172	auto z_ptr = (GM_ADDR)z.data_ptr();	174	auto z_ptr = (GM_ADDR)z.data_ptr();
	@@ -176,15 +178,15 @@ print("Verification successful!")
176	// 根据不同的数据类型，调用不同的NPU Kernel	178	// 根据不同的数据类型，调用不同的NPU Kernel
177	AT_DISPATCH_CASE(torch::kFloat32, [&] {	179	AT_DISPATCH_CASE(torch::kFloat32, [&] {
178	using scalar_t = float;	180	using scalar_t = float;
179	- add_kernel<scalar_t><<<blockDim, nullptr, stream>>>(x_ptr, y_ptr, z_ptr, totalLength, blockLength, tileSize);	181	+ add_kernel<scalar_t><<<numBlocks, nullptr, stream>>>(x_ptr, y_ptr, z_ptr, totalLength, blockLength, tileSize);
180	})	182	})
181	AT_DISPATCH_CASE(torch::kFloat16, [&] {	183	AT_DISPATCH_CASE(torch::kFloat16, [&] {
182	using scalar_t = half;	184	using scalar_t = half;
183	- add_kernel<scalar_t><<<blockDim, nullptr, stream>>>(x_ptr, y_ptr, z_ptr, totalLength, blockLength, tileSize);	185	+ add_kernel<scalar_t><<<numBlocks, nullptr, stream>>>(x_ptr, y_ptr, z_ptr, totalLength, blockLength, tileSize);
184	})	186	})
185	AT_DISPATCH_CASE(torch::kInt32, [&] {	187	AT_DISPATCH_CASE(torch::kInt32, [&] {
186	using scalar_t = int32_t;	188	using scalar_t = int32_t;
187	- add_kernel<scalar_t><<<blockDim, nullptr, stream>>>(x_ptr, y_ptr, z_ptr, totalLength, blockLength, tileSize);	189	+ add_kernel<scalar_t><<<numBlocks, nullptr, stream>>>(x_ptr, y_ptr, z_ptr, totalLength, blockLength, tileSize);
188	})	190	})
189	);	191	);
190	return 0;	192	return 0;

Mexamples/fast_kernel_launch_example/cmake/func.cmake+2-2

	@@ -14,8 +14,8 @@
14	macro(recursive_add_subdirectory)	14	macro(recursive_add_subdirectory)
15	file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)	15	file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
16	foreach(SUB_DIR ${CURRENT_DIRS})	16	foreach(SUB_DIR ${CURRENT_DIRS})
17	- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/${NPU_ARCH}/CMakeLists.txt")	17	+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/${NPU_SOC_VERSION}/CMakeLists.txt")
18	- add_subdirectory(${SUB_DIR}/${NPU_ARCH})	18	+ add_subdirectory(${SUB_DIR}/${NPU_SOC_VERSION})
19	endif()	19	endif()
20	endforeach()	20	endforeach()
21	endmacro()	21	endmacro()

Mexamples/fast_kernel_launch_example/setup.py+8-4

	@@ -106,9 +106,13 @@ class CMakeBuildCommand(Command):
106	TORCH_NPU_PATH = os.path.dirname(torch_npu.__file__)	106	TORCH_NPU_PATH = os.path.dirname(torch_npu.__file__)
107	logging.info(f"Using Torch NPU path: {TORCH_NPU_PATH}")	107	logging.info(f"Using Torch NPU path: {TORCH_NPU_PATH}")
108		108
109	- # Get NPU_ARCH from environment variable or set default	109	+ # Get NPU_SOC_VERSION from environment variable or set default
110	- NPU_ARCH = os.environ.get('NPU_ARCH', 'ascend910b')	110	+ npu_soc_version = os.environ.get('NPU_SOC_VERSION')
111	- logging.info(f"Using NPU_ARCH: {NPU_ARCH}")	111	+ if npu_soc_version is None:
		112	+ npu_soc_version = os.environ.get('NPU_ARCH', 'ascend910b')
		113	+ if 'NPU_ARCH' in os.environ:
		114	+ logging.info("NPU_ARCH is deprecated, please use NPU_SOC_VERSION.")
		115	+ logging.info(f"Using NPU_SOC_VERSION: {npu_soc_version}")
112		116
113	# Build the CMake project	117	# Build the CMake project
114	build_temp = os.path.join(os.getcwd(), 'build')	118	build_temp = os.path.join(os.getcwd(), 'build')
	@@ -116,7 +120,7 @@ class CMakeBuildCommand(Command):
116	'-DCMAKE_BUILD_TYPE=Release',	120	'-DCMAKE_BUILD_TYPE=Release',
117	f'-DTorch_DIR={Torch_DIR}',	121	f'-DTorch_DIR={Torch_DIR}',
118	f'-DTORCH_NPU_PATH={TORCH_NPU_PATH}',	122	f'-DTORCH_NPU_PATH={TORCH_NPU_PATH}',
119	- f'-DNPU_ARCH={NPU_ARCH}'	123	+ f'-DNPU_SOC_VERSION={npu_soc_version}'
120	]	124	]
121	subprocess.check_call(cmake_config_command, cwd=os.getcwd())	125	subprocess.check_call(cmake_config_command, cwd=os.getcwd())
122	subprocess.check_call(['cmake', '--build', build_temp, '--config', 'Release', '--parallel', num_jobs], cwd=os.getcwd())	126	subprocess.check_call(['cmake', '--build', build_temp, '--config', 'Release', '--parallel', num_jobs], cwd=os.getcwd())