已合并
为 vector_add 增加 msprof 输入输出元信息上报 #355
qq_51867290创建于 24 天前
为 vector_add 增加 msprof 输入输出元信息上报 #355
已合并
共 3 个文件变更+47-0
| @@ -24,6 +24,7 @@ target_link_libraries(vector_add | |||
| 24 | dl | 24 | dl |
| 25 | platform | 25 | platform |
| 26 | tiling_api | 26 | tiling_api |
| 27 | + msprofiler | ||
| 27 | ) | 28 | ) |
| 28 | 29 | ||
| 29 | install(TARGETS vector_add | 30 | install(TARGETS vector_add |
| @@ -46,3 +46,7 @@ Vector add completed successfully! | |||
| 46 | ```shell | 46 | ```shell |
| 47 | Vector add failed! | 47 | Vector add failed! |
| 48 | ``` | 48 | ``` |
| 49 | + | ||
| 50 | +## References | ||
| 51 | + | ||
| 52 | +- [torch_library_report_tensor](https://gitcode.com/cann/asc-devkit/blob/master/examples/01_simd_cpp_api/01_utilities/04_profiling/torch_library_report_tensor/torch_library_report_tensor.asc): 性能分析元信息上报参考样例。 | ||
| @@ -20,6 +20,7 @@ | |||
| 20 | #include <memory> | 20 | #include <memory> |
| 21 | #include <vector> | 21 | #include <vector> |
| 22 | #include "acl/acl.h" | 22 | #include "acl/acl.h" |
| 23 | +#include "acl/acl_prof.h" | ||
| 23 | #include "kernel_operator.h" | 24 | #include "kernel_operator.h" |
| 24 | #include "platform/platform_ascendc.h" | 25 | #include "platform/platform_ascendc.h" |
| 25 | 26 | ||
| @@ -43,6 +44,24 @@ struct AclrtFreeDeleter { | |||
| 43 | } | 44 | } |
| 44 | }; | 45 | }; |
| 45 | 46 | ||
| 47 | +constexpr const char* PROF_OP_NAME = "vector_add"; | ||
| 48 | +constexpr const char* PROF_OP_TYPE = "VectorAdd"; | ||
| 49 | +constexpr uint32_t PROF_INPUT_TENSOR = 0; | ||
| 50 | +constexpr uint32_t PROF_OUTPUT_TENSOR = 1; | ||
| 51 | +constexpr uint32_t PROF_VECTOR_SHAPE_DIM = 1; | ||
| 52 | + | ||
| 53 | +// Tensor metadata profiling types are provided by acl/acl_prof.h since CANN 9.1.0. | ||
| 54 | +aclprofTensor MakeVectorTensor(uint32_t type, int64_t numElements) | ||
| 55 | +{ | ||
| 56 | + aclprofTensor tensor{}; | ||
| 57 | + tensor.type = type; | ||
| 58 | + tensor.format = static_cast<uint32_t>(ACL_FORMAT_ND); | ||
| 59 | + tensor.dataType = static_cast<uint32_t>(ACL_FLOAT); | ||
| 60 | + tensor.shapeDim = PROF_VECTOR_SHAPE_DIM; | ||
| 61 | + tensor.shape[0] = static_cast<uint32_t>(numElements); | ||
| 62 | + return tensor; | ||
| 63 | +} | ||
| 64 | + | ||
| 46 | std::tuple<int64_t, int64_t, int64_t> calc_tiling_params(int64_t totalLength) | 65 | std::tuple<int64_t, int64_t, int64_t> calc_tiling_params(int64_t totalLength) |
| 47 | { | 66 | { |
| 48 | constexpr static int64_t MIN_ELEMS_PER_CORE = 1024; | 67 | constexpr static int64_t MIN_ELEMS_PER_CORE = 1024; |
| @@ -182,7 +201,30 @@ int run_vector_add(aclrtStream stream, int64_t numElements) | |||
| 182 | int64_t numBlocks, blockLength, tileSize; | 201 | int64_t numBlocks, blockLength, tileSize; |
| 183 | std::tie(numBlocks, blockLength, tileSize) = calc_tiling_params(numElements); | 202 | std::tie(numBlocks, blockLength, tileSize) = calc_tiling_params(numElements); |
| 184 | CHECK_ACL(aclrtSynchronizeStream(stream)); | 203 | CHECK_ACL(aclrtSynchronizeStream(stream)); |
| 204 | + aclprofTensor tensors[] = { | ||
S | |||
| 205 | + MakeVectorTensor(PROF_INPUT_TENSOR, numElements), | ||
| 206 | + MakeVectorTensor(PROF_INPUT_TENSOR, numElements), | ||
| 207 | + MakeVectorTensor(PROF_OUTPUT_TENSOR, numElements), | ||
| 208 | + }; | ||
| 209 | + aclprofTensorInfo tensorInfo = { | ||
| 210 | + aclprofStr2Id(PROF_OP_NAME), | ||
| 211 | + aclprofStr2Id(PROF_OP_TYPE), | ||
| 212 | + 0, | ||
| 213 | + sizeof(tensors) / sizeof(aclprofTensor), | ||
| 214 | + 0, | ||
| 215 | + static_cast<uint32_t>(numBlocks), | ||
| 216 | + stream, | ||
| 217 | + tensors | ||
| 218 | + }; | ||
| 219 | + aclprofEventAttributes attrs = { | ||
| 220 | + ACL_PROF_EVENT_ATTR_VERSION, | ||
| 221 | + sizeof(aclprofEventAttributes::message), | ||
| 222 | + ACL_PROF_MESSAGE_TYPE_TENSOR_INFO, | ||
| 223 | + &tensorInfo | ||
| 224 | + }; | ||
| 225 | + aclprofRangePushEx(&attrs); | ||
| 185 | add_kernel<float><<<numBlocks, nullptr, stream>>>(d_A, d_B, d_C, numElements, blockLength, tileSize); | 226 | add_kernel<float><<<numBlocks, nullptr, stream>>>(d_A, d_B, d_C, numElements, blockLength, tileSize); |
| 227 | + aclprofRangePop(); | ||
| 186 | CHECK_ACL(aclrtSynchronizeStream(stream)); | 228 | CHECK_ACL(aclrtSynchronizeStream(stream)); |
| 187 | 229 | ||
| 188 | CHECK_ACL(aclrtMemcpy(h_C.data(), size, d_C, size, ACL_MEMCPY_DEVICE_TO_HOST)); | 230 | CHECK_ACL(aclrtMemcpy(h_C.data(), size, d_C, size, ACL_MEMCPY_DEVICE_TO_HOST)); |
加个注释和readme说明一下这个特性要求的cann包的版本