已开启
【代码侦探Challenge01-DivCustom】新增DivCustom代码提交PR样例 #1052
guest_shao创建于 7月16日
【代码侦探Challenge01-DivCustom】新增DivCustom代码提交PR样例 #1052
已开启
共 3 个文件变更+114-0
| @@ -0,0 +1,10 @@ | |||
| 1 | +cmake_minimum_required(VERSION 3.16) | ||
| 2 | +find_package(ASC REQUIRED) | ||
| 3 | +project(kernel_samples LANGUAGES ASC CXX) | ||
| 4 | + | ||
| 5 | +add_executable(div_test | ||
| 6 | + div_custom.asc | ||
| 7 | +) | ||
| 8 | +target_compile_options(div_test PRIVATE | ||
| 9 | + $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-2201> | ||
| 10 | +) | ||
| @@ -0,0 +1,95 @@ | |||
| 1 | +// div_custom.asc | ||
| 2 | + | ||
| 3 | +#include <cstdint> | ||
| 4 | +#include <iostream> | ||
| 5 | +#include <vector> | ||
| 6 | +#include <algorithm> | ||
| 7 | +#include <iterator> | ||
| 8 | +#include "acl/acl.h" | ||
| 9 | +#include "kernel_operator.h" | ||
| 10 | + | ||
| 11 | +constexpr uint32_t BUFFER_NUM = 2; // tensor num for each queue | ||
| 12 | + | ||
| 13 | +struct DivCustomTilingData | ||
| 14 | +{ | ||
| 15 | + uint32_t totalLength; | ||
| 16 | + uint32_t tileNum; | ||
| 17 | +}; | ||
| 18 | + | ||
| 19 | +class KernelDiv { | ||
| 20 | +public: | ||
| 21 | + __aicore__ inline KernelDiv(){} | ||
| 22 | + __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLength, uint32_t tileNum) | ||
| 23 | + { | ||
| 24 | + // 请补充…… | ||
| 25 | + // 添加注释 | ||
| 26 | + } | ||
| 27 | + __aicore__ inline void Process() | ||
| 28 | + { | ||
| 29 | + // 请补充…… | ||
| 30 | + } | ||
| 31 | + | ||
| 32 | +private: | ||
| 33 | + __aicore__ inline void CopyIn(int32_t progress) | ||
| 34 | + { | ||
| 35 | + // 请补充…… | ||
| 36 | + } | ||
| 37 | + __aicore__ inline void Compute(int32_t progress) | ||
| 38 | + { | ||
| 39 | + // 请补充…… | ||
| 40 | + } | ||
| 41 | + __aicore__ inline void CopyOut(int32_t progress) | ||
| 42 | + { | ||
| 43 | + // 请补充…… | ||
| 44 | + } | ||
| 45 | + | ||
| 46 | +private: | ||
| 47 | + // 请补充…… | ||
| 48 | +}; | ||
| 49 | + | ||
| 50 | +__global__ __aicore__ void div_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, DivCustomTilingData tiling) | ||
| 51 | +{ | ||
| 52 | + // 请补充…… | ||
| 53 | +} | ||
| 54 | + | ||
| 55 | +std::vector<float> kernel_div(std::vector<float> &x, std::vector<float> &y) | ||
| 56 | +{ | ||
| 57 | + // 请补充…… | ||
| 58 | +} | ||
| 59 | + | ||
| 60 | +uint32_t VerifyResult(std::vector<float> &output, std::vector<float> &golden) | ||
| 61 | +{ | ||
| 62 | + auto printTensor = [](std::vector<float> &tensor, const char *name) { | ||
| 63 | + constexpr size_t maxPrintSize = 20; | ||
| 64 | + std::cout << name << ": "; | ||
| 65 | + std::copy(tensor.begin(), tensor.begin() + std::min(tensor.size(), maxPrintSize), | ||
| 66 | + std::ostream_iterator<float>(std::cout, " ")); | ||
| 67 | + if (tensor.size() > maxPrintSize) { | ||
| 68 | + std::cout << "..."; | ||
| 69 | + } | ||
| 70 | + std::cout << std::endl; | ||
| 71 | + }; | ||
| 72 | + printTensor(output, "Output"); | ||
| 73 | + printTensor(golden, "Golden"); | ||
| 74 | + if (std::equal(golden.begin(), golden.end(), output.begin())) { | ||
| 75 | + std::cout << "[Success] Case accuracy is verification passed." << std::endl; | ||
| 76 | + return 0; | ||
| 77 | + } else { | ||
| 78 | + std::cout << "[Failed] Case accuracy is verification failed!" << std::endl; | ||
| 79 | + return 1; | ||
| 80 | + } | ||
| 81 | +} | ||
| 82 | + | ||
| 83 | +int32_t main(int32_t argc, char *argv[]) | ||
| 84 | +{ | ||
| 85 | + constexpr uint32_t totalLength = 8 * 2048; | ||
| 86 | + constexpr float valueX = 1.2f; | ||
| 87 | + constexpr float valueY = 2.3f; | ||
| 88 | + std::vector<float> x(totalLength, valueX); | ||
| 89 | + std::vector<float> y(totalLength, valueY); | ||
| 90 | + | ||
| 91 | + // 请补充…… | ||
| 92 | + | ||
| 93 | + std::vector<float> golden(totalLength, valueX / valueY); | ||
| 94 | + return VerifyResult(output, golden); | ||
| 95 | +} | ||
| @@ -0,0 +1,9 @@ | |||||||||
| 1 | +#!/bin/bash | ||||||||
🟡 Medium Priority 变更行: 脚本缺少
建议:在 改动建议
![]() ![]() | |||||||||
| 2 | +# 激活cann环境,可根据实际情况修改,在线环境一般不用修改, | ||||||||
🟡 Medium Priority run.sh 脚本缺少
触发条件:运行环境中 建议在 建议:在 改动建议
![]() ![]() | |||||||||
| 3 | +source "$ASCEND_TOOLKIT_HOME/set_env.sh" | ||||||||
| 4 | +mkdir -p build | ||||||||
| 5 | +export ASC_DIR="$ASCEND_HOME_PATH/aarch64-linux/tikcpp/ascendc_kernel_cmake/" | ||||||||
| 6 | +cd build/ && \ | ||||||||
| 7 | +cmake .. && \ | ||||||||
| 8 | +make && \ | ||||||||
| 9 | +./div_test | ||||||||


🟡 Medium Priority
变更行:
div_custom.asc第 74 行std::equal(golden.begin(), golden.end(), output.begin())。std::equal(first1, last1, first2)仅比较[first1, last1)范围内的元素是否与从first2开始的元素一一相等,不检查两个区间的长度是否相同。因此:output.size() < golden.size(),output.begin()迭代器前进超过golden.size()步时会超出output的合法范围,触发未定义行为。output.size() > golden.size()且前golden.size()个元素恰好匹配,std::equal返回true,output尾部多余的(可能错误)元素被完全忽略,精度验证形同虚设。VerifyResult不在 README 列出的"你需要完成的内容"中,是挑战提供方应保证正确的函数,此缺陷会直接影响所有参与者的验证结果可信度。