Star627
236
代码
代码
Issues34
Pull Requests81
Actions
项目讨论
Wiki
分析
Star627
236
已开启
【PR】: 简要描述 #1777
Achilles_d创建于 14 天前
【PR】: 简要描述 #1777
已开启
Achilles_d创建于 14 天前
从Achilles_d/graph-autofusion_locktest-thj合入到cann/graph-autofusionmaster
共 5 个文件变更+11-189
Msuper_kernel/CMakeLists.txt+1-1
@@ -72,7 +72,7 @@ target_link_options(ascendsk PRIVATE
72 $<$<CONFIG:Release>:-s>72 $<$<CONFIG:Release>:-s>
73)73)
74target_link_libraries(ascendsk PRIVATE c_sec_headers runtime_headers)74target_link_libraries(ascendsk PRIVATE c_sec_headers runtime_headers)
75-target_link_libraries(ascendsk PUBLIC sk_scope_dav_2201 sk_scope_dav_3510 acl_rt acl_rtc ascend_dump runtime error_manager c_sec unified_dlog ascendc_runtime profapi profimpl msprofiler json)75+target_link_libraries(ascendsk PUBLIC acl_rt acl_rtc ascend_dump runtime error_manager c_sec unified_dlog ascendc_runtime profapi profimpl msprofiler json)
76 76 
77if(ENABLE_CPP_UTEST OR ENABLE_RDV_UTEST)77if(ENABLE_CPP_UTEST OR ENABLE_RDV_UTEST)
78 add_subdirectory(tests/aot)78 add_subdirectory(tests/aot)
Msuper_kernel/kernel/CMakeLists.txt+0-27
@@ -54,30 +54,3 @@ add_custom_command(
54add_custom_target(generate_entry_stub54add_custom_target(generate_entry_stub
55 DEPENDS ${GENERATED_SK_ENTRY_STUB_CPP}55 DEPENDS ${GENERATED_SK_ENTRY_STUB_CPP}
56)56)
57- 
58-foreach(SK_KERNEL_ARCH IN LISTS SK_KERNEL_ARCHS)
59- string(REPLACE "-" "_" SK_KERNEL_ARCH_SUFFIX ${SK_KERNEL_ARCH})
60- set(SK_SCOPE_TARGET sk_scope_${SK_KERNEL_ARCH_SUFFIX})
61- 
62- add_library(${SK_SCOPE_TARGET} STATIC
63- sk_scope_kernel.asc
64- )
65- 
66- target_compile_options(${SK_SCOPE_TARGET} PRIVATE
67- $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${SK_KERNEL_ARCH}>
68- )
69- 
70- target_compile_definitions(${SK_SCOPE_TARGET} PRIVATE
71- $<$<COMPILE_LANGUAGE:ASC>:SK_SCOPE_SYMBOL_SUFFIX=_${SK_KERNEL_ARCH_SUFFIX}>
72- )
73- 
74- target_include_directories(${SK_SCOPE_TARGET} PRIVATE
75- ${CMAKE_CURRENT_SOURCE_DIR}/../src/aot/
76- )
77- target_link_libraries(${SK_SCOPE_TARGET} PRIVATE
78- $<BUILD_INTERFACE:runtime_headers>
79- $<BUILD_INTERFACE:asc_kernel_headers>
80- $<BUILD_INTERFACE:kernel_tiling_headers>
81- )
82- set_property(TARGET ${SK_SCOPE_TARGET} PROPERTY POSITION_INDEPENDENT_CODE ON)
83-endforeach()
Dsuper_kernel/kernel/sk_scope_kernel.asc+0-44
@@ -1,44 +0,0 @@
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-/* !
12- * \file sk_scope_kernel.asc
13- * \brief
14- */
15- 
16-#include "kernel_operator.h"
17-#include "sk_common.h"
18- 
19-#define SK_SCOPE_CONCAT_INNER(a, b) a##b
20-#define SK_SCOPE_CONCAT(a, b) SK_SCOPE_CONCAT_INNER(a, b)
21- 
22-#ifdef SK_SCOPE_SYMBOL_SUFFIX
23-#define SK_SCOPE_SYMBOL(name) SK_SCOPE_CONCAT(name, SK_SCOPE_SYMBOL_SUFFIX)
24-#else
25-#define SK_SCOPE_SYMBOL(name) name
26-#endif
27- 
28-extern "C" __global__ __vector__ void SK_SCOPE_SYMBOL(sk_scope_kernel_begin)(ScopeKernelArgs args) {}
29- 
30-extern "C" __global__ __vector__ void SK_SCOPE_SYMBOL(sk_scope_kernel_end)(ScopeKernelArgs args) {}
31- 
32-extern "C" __global__ __vector__ void SK_SCOPE_SYMBOL(sk_placeholder_kernel)(ScopeKernelArgs args) {}
33- 
34-extern "C" void SK_SCOPE_SYMBOL(sk_scope_kernel_begin_do)(void *stream, ScopeKernelArgs args) {
35- SK_SCOPE_SYMBOL(sk_scope_kernel_begin)<<<1, nullptr, stream>>>(args);
36-}
37- 
38-extern "C" void SK_SCOPE_SYMBOL(sk_scope_kernel_end_do)(void *stream, ScopeKernelArgs args) {
39- SK_SCOPE_SYMBOL(sk_scope_kernel_end)<<<1, nullptr, stream>>>(args);
40-}
41- 
42-extern "C" void SK_SCOPE_SYMBOL(sk_placeholder_kernel_do)(void *stream, ScopeKernelArgs args) {
43- SK_SCOPE_SYMBOL(sk_placeholder_kernel)<<<1, nullptr, stream>>>(args);
44-}
Msuper_kernel/src/aot/sk_scope_launch.cpp+10-87
@@ -21,101 +21,24 @@
21#include "sk_common.h"21#include "sk_common.h"
22#include "sk_log.h"22#include "sk_log.h"
23 23 
24-extern "C" void sk_scope_kernel_begin_do_dav_2201(void *stream, ScopeKernelArgs args);
25-extern "C" void sk_scope_kernel_end_do_dav_2201(void *stream, ScopeKernelArgs args);
26-extern "C" void sk_placeholder_kernel_do_dav_2201(void *stream, ScopeKernelArgs args);
27-extern "C" void sk_scope_kernel_begin_do_dav_3510(void *stream, ScopeKernelArgs args);
28-extern "C" void sk_scope_kernel_end_do_dav_3510(void *stream, ScopeKernelArgs args);
29-extern "C" void sk_placeholder_kernel_do_dav_3510(void *stream, ScopeKernelArgs args);
30- 
31-typedef void (*ScopeFuncImpl)(void *, ScopeKernelArgs);
32- 
33namespace {24namespace {
34 25 
35-struct ScopeKernelImpls {26+bool IsScopeNameValid(const char *scopeName) {
36- ScopeFuncImpl begin;
37- ScopeFuncImpl end;
38- ScopeFuncImpl placeholder;
39-};
40- 
41-ScopeKernelImpls GetScopeKernelImpls() {
42- const SkKernelArch arch = GetCurrentSkKernelArch();
43- SK_DLOGI("Scope kernel implementation selected: arch=%s, symbolSuffix=%s", to_string(arch),
44- GetSkKernelArchSymbolSuffix(arch));
45- switch (arch) {
46- case SkKernelArch::DAV_3510:
47- return {sk_scope_kernel_begin_do_dav_3510, sk_scope_kernel_end_do_dav_3510, sk_placeholder_kernel_do_dav_3510};
48- case SkKernelArch::DAV_2201:
49- default:
50- return {sk_scope_kernel_begin_do_dav_2201, sk_scope_kernel_end_do_dav_2201, sk_placeholder_kernel_do_dav_2201};
51- }
52-}
53- 
54-aclError LaunchScopeKernelImpl(const char *scopeName, aclrtStream stream, ScopeFuncImpl scopeKernelImpl,
55- const char *markerType) {
56- ScopeKernelArgs args;
57 if (scopeName == nullptr) {27 if (scopeName == nullptr) {
58- scopeName = DEFAULT_SK_SCOPE_NAME;28+ return true;
59 }29 }
60- size_t len = strlen(scopeName);30+ return strlen(scopeName) <= MAX_SCOPE_NAME_LEN - 1;
61- size_t maxLen = MAX_SCOPE_NAME_LEN - 1;
62- if (len > maxLen) {
63- SK_DLOGE("Failed to prepare scope kernel arguments: markerType=%s, scopeNameLength=%zu, maxLength=%zu", markerType,
64- len, maxLen);
65- return ACL_ERROR_INVALID_PARAM;
66- }
67- errno_t ret = memcpy_s(args.name, maxLen, scopeName, len);
68- if (ret != 0) {
69- SK_DLOGE("Failed to prepare scope kernel arguments: markerType=%s, memcpy_s ret=%d", markerType, ret);
70- return ACL_ERROR_INVALID_PARAM;
71- }
72- args.name[len] = '\0';
73- SK_DLOGI("Dispatch scope kernel: markerType=%s, scopeName=%s, stream=%p", markerType, scopeName, stream);
74- scopeKernelImpl(stream, args);
75- return ACL_SUCCESS;
76}31}
77 32 
78} // namespace33} // namespace
79 34 
80aclError LaunchScopeKernel(const char *scopeName, aclrtStream stream, bool isBegin) {35aclError LaunchScopeKernel(const char *scopeName, aclrtStream stream, bool isBegin) {
81- ScopeKernelImpls impls = GetScopeKernelImpls();36+ (void)stream;
82- if (isBegin) {37+ (void)isBegin;
83- SK_DLOGI("Scope begin marker kernel dispatch started");38+ if (!IsScopeNameValid(scopeName)) {
84- aclError ret = LaunchScopeKernelImpl(scopeName, stream, impls.begin, "begin");39+ SK_DLOGE("Scope name exceeds the maximum length: scopeNameLength=%zu, maxLength=%u", strlen(scopeName),
85- if (ret != ACL_SUCCESS) {40+ MAX_SCOPE_NAME_LEN - 1);
86- SK_DLOGE("Scope marker kernel processing failed: markerType=begin, ret=%d", ret);41+ return ACL_ERROR_INVALID_PARAM;
87- return ret;
88- }
89- ret = LaunchScopeKernelImpl(scopeName, stream, impls.placeholder, "placeholder");
90- if (ret != ACL_SUCCESS) {
91- SK_DLOGE("Scope marker kernel processing failed: markerType=placeholder, ret=%d", ret);
92- return ret;
93- }
94- ret = LaunchScopeKernelImpl(scopeName, stream, impls.placeholder, "placeholder");
95- if (ret != ACL_SUCCESS) {
96- SK_DLOGE("Scope marker kernel processing failed: markerType=placeholder, ret=%d", ret);
97- return ret;
98- }
99- SK_DLOGI("Scope begin marker kernels dispatched");
100- return ret;
101- } else {
102- SK_DLOGI("Scope end marker kernel dispatch started");
103- aclError ret = LaunchScopeKernelImpl(scopeName, stream, impls.placeholder, "placeholder");
104- if (ret != ACL_SUCCESS) {
105- SK_DLOGE("Scope marker kernel processing failed: markerType=placeholder, ret=%d", ret);
106- return ret;
107- }
108- ret = LaunchScopeKernelImpl(scopeName, stream, impls.placeholder, "placeholder");
109- if (ret != ACL_SUCCESS) {
110- SK_DLOGE("Scope marker kernel processing failed: markerType=placeholder, ret=%d", ret);
111- return ret;
112- }
113- ret = LaunchScopeKernelImpl(scopeName, stream, impls.end, "end");
114- if (ret != ACL_SUCCESS) {
115- SK_DLOGE("Scope marker kernel processing failed: markerType=end, ret=%d", ret);
116- return ret;
117- }
118- SK_DLOGI("Scope end marker kernels dispatched");
119- return ret;
120 }42 }
43+ return ACL_SUCCESS;
121}44}
Msuper_kernel/tests/aot/ut/stub/ut_common_stubs.cpp+0-30
@@ -398,36 +398,6 @@ extern "C" aclrtBinHandle AscendGetEntryBinHandle() {
398 return reinterpret_cast<aclrtBinHandle>(0x1234);398 return reinterpret_cast<aclrtBinHandle>(0x1234);
399}399}
400 400 
401-void sk_scope_kernel_begin_do_dav_2201(void *stream, ScopeKernelArgs args) {
402- (void)stream;
403- (void)args;
404-}
405- 
406-void sk_scope_kernel_end_do_dav_2201(void *stream, ScopeKernelArgs args) {
407- (void)stream;
408- (void)args;
409-}
410- 
411-void sk_placeholder_kernel_do_dav_2201(void *stream, ScopeKernelArgs args) {
412- (void)stream;
413- (void)args;
414-}
415- 
416-void sk_scope_kernel_begin_do_dav_3510(void *stream, ScopeKernelArgs args) {
417- (void)stream;
418- (void)args;
419-}
420- 
421-void sk_scope_kernel_end_do_dav_3510(void *stream, ScopeKernelArgs args) {
422- (void)stream;
423- (void)args;
424-}
425- 
426-void sk_placeholder_kernel_do_dav_3510(void *stream, ScopeKernelArgs args) {
427- (void)stream;
428- (void)args;
429-}
430- 
431void SkUtSetAclrtGetSocName(const char *socName) {401void SkUtSetAclrtGetSocName(const char *socName) {
432 g_aclrtGetSocName = socName;402 g_aclrtGetSocName = socName;
433}403}