From 6f0fbeaab92c80d956cbd2a2591fedc79923bd82 Mon Sep 17 00:00:00 2001 From: machangjun Date: Sun, 25 Feb 2024 17:41:28 +0800 Subject: [PATCH] =?UTF-8?q?=E5=A2=9E=E5=8A=A0santizizer=20sample?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- sample/README.md | 17 ++++++++ sample/build/build.sh | 4 ++ sample/sanitizer_sample/Racecheck/Makefile | 28 +++++++++++++ sample/sanitizer_sample/Racecheck/main.cpp | 42 +++++++++++++++++++ .../Racecheck/raw_error_kernel.cpp | 28 +++++++++++++ 5 files changed, 119 insertions(+) create mode 100644 sample/sanitizer_sample/Racecheck/Makefile create mode 100644 sample/sanitizer_sample/Racecheck/main.cpp create mode 100644 sample/sanitizer_sample/Racecheck/raw_error_kernel.cpp diff --git a/sample/README.md b/sample/README.md index b46181b02..4f49fec5a 100644 --- a/sample/README.md +++ b/sample/README.md @@ -16,6 +16,23 @@ source set_env.sh 大概在2月份会补齐各类工具样例和文档 +### sanitizer_sample目录介绍 + +此目录下为sanitizer对应的样例库，包含竞争检测和内存检测相关的样例。 + +#### Racecheck目录介绍 + +Racecheck为竞争检测相关的样例。 + +raw_error_kernel.cpp文件为UB上先读后写竞争和GM上先写后读竞争问题的样例。 + + +运行阶段： + +``` +/usr/local/Ascend/ascend-toolkit/latest/tools/mssanitizer/bin/mssanitizer --tool=racecheck ./raw_error.fatbin +``` + ## 其他 现msprof使能仿真时，还需要额外调整，具体如下： 1. 编译阶段：在```sample\normal_sample\vec_only```相对路径下的```Makefile```文件中修改如下内容： diff --git a/sample/build/build.sh b/sample/build/build.sh index 80e17f8d5..ca6689ab7 100644 --- a/sample/build/build.sh +++ b/sample/build/build.sh @@ -16,4 +16,8 @@ cp -f libruntime_camodel.so.$(arch) libruntime_camodel.so # cd ${TOP_DIR}/normal_sample/vec_only make +mv *.fatbin ${TOP_DIR}/build + +cd ${TOP_DIR}/sanitizer_sample/Racecheck +make mv *.fatbin ${TOP_DIR}/build \ No newline at end of file diff --git a/sample/sanitizer_sample/Racecheck/Makefile b/sample/sanitizer_sample/Racecheck/Makefile new file mode 100644 index 000000000..ac230761b --- /dev/null +++ b/sample/sanitizer_sample/Racecheck/Makefile @@ -0,0 +1,28 @@ +ASCEND_HOME_PATH ?= /usr/local/Ascend/ascend-toolkit/latest + +COMPILER := ${ASCEND_HOME_PATH}/compiler/ccec_compiler/bin/ccec +HOST_COMPILER := g++ +COMPILER_FLAG := -xcce -O2 -std=c++17 --cce-enable-sanitizer -g +HOST_COMPILER_FLAG := -O2 -std=c++17 +LINK_FLAG := --cce-fatobj-link --cce-enable-sanitizer +DAV_FLAG := --cce-aicore-arch=dav-c220-vec +ASCENDC_INC_FLAG := -I${ASCEND_HOME_PATH}/compiler/tikcpp/tikcfw -I${ASCEND_HOME_PATH}/compiler/tikcpp/tikcfw/impl -I${ASCEND_HOME_PATH}/compiler/tikcpp/tikcfw/interface -I${ASCEND_HOME_PATH}/include +HOST_INC_FLAG := -I${ASCEND_HOME_PATH}/include +LINK_LIBS := -L${ASCEND_HOME_PATH}/lib64 -lruntime -lascendcl -lstdc++ + +all: build + +build: raw_error_kernel.o main.o raw_error.fatbin + +raw_error_kernel.o: raw_error_kernel.cpp + $(COMPILER) $(COMPILER_FLAG) $(DAV_FLAG) $(ASCENDC_INC_FLAG) -o $@ -c $^ + +main.o: main.cpp + $(HOST_COMPILER) $(HOST_COMPILER_FLAG) $(HOST_INC_FLAG) -o $@ -c $^ + +raw_error.fatbin: raw_error_kernel.o main.o + $(COMPILER) $(LINK_FLAG) $(DAV_FLAG) -o $@ $^ ${LINK_LIBS} + +.PHONY: clean +clean: + rm *.o raw_error.fatbin diff --git a/sample/sanitizer_sample/Racecheck/main.cpp b/sample/sanitizer_sample/Racecheck/main.cpp new file mode 100644 index 000000000..cd96d8e08 --- /dev/null +++ b/sample/sanitizer_sample/Racecheck/main.cpp @@ -0,0 +1,42 @@ +#include +#include "acl/acl.h" + +#define ACL_ERROR_NONE 0 + +#define CHECK_ACL(x) \ + do { \ + aclError __ret = x; \ + if (__ret != ACL_ERROR_NONE) { \ + std::cerr << __FILE__ << ":" << __LINE__ << " aclError:" << __ret << std::endl; \ + } \ + } while (0); + +extern "C" void raw_error_kernel_do(uint32_t blockDim, void *l2ctrl, void *stream, uint8_t *gmInput, uint8_t *gmOutput); + +int main(void) +{ + CHECK_ACL(aclInit(nullptr)); + aclrtContext context; + int32_t deviceId = 0; + CHECK_ACL(aclrtSetDevice(deviceId)); + CHECK_ACL(aclrtCreateContext(&context, deviceId)); + aclrtStream stream = nullptr; + CHECK_ACL(aclrtCreateStream(&stream)); + + uint8_t *gmInput = nullptr; + uint8_t *gmOutput = nullptr; + CHECK_ACL(aclrtMalloc((void**)&gmInput, 256, ACL_MEM_MALLOC_HUGE_FIRST)); + CHECK_ACL(aclrtMalloc((void**)&gmOutput, 256, ACL_MEM_MALLOC_HUGE_FIRST)); + + uint64_t blockDim = 1UL; + raw_error_kernel_do(blockDim, nullptr, stream, gmInput, gmOutput); + CHECK_ACL(aclrtSynchronizeStream(stream)); + + CHECK_ACL(aclrtFree(gmInput)); + CHECK_ACL(aclrtFree(gmOutput)); + CHECK_ACL(aclrtDestroyStream(stream)); + CHECK_ACL(aclrtDestroyContext(context)); + CHECK_ACL(aclrtResetDevice(deviceId)); + CHECK_ACL(aclFinalize()); + return 0; +} \ No newline at end of file diff --git a/sample/sanitizer_sample/Racecheck/raw_error_kernel.cpp b/sample/sanitizer_sample/Racecheck/raw_error_kernel.cpp new file mode 100644 index 000000000..b1db85609 --- /dev/null +++ b/sample/sanitizer_sample/Racecheck/raw_error_kernel.cpp @@ -0,0 +1,28 @@ +#include "kernel_operator.h" +#include "acl/acl.h" +using namespace AscendC; + +constexpr int32_t BYTESIZE_EXAMPLE = 256; +constexpr int32_t BUFFER_NUM = 1; +constexpr int32_t NUM_DATA = BYTESIZE_EXAMPLE / sizeof(half); + +extern "C" __global__ __aicore__ void raw_error_kernel(__gm__ uint8_t *gmInput, __gm__ uint8_t *gmOutput) { + TPipe pipe; + TQue xQue; + GlobalTensor xInGm, xOutGm; + pipe.InitBuffer(xQue, BUFFER_NUM, BYTESIZE_EXAMPLE); + LocalTensor xLocal = xQue.AllocTensor(); + xInGm.SetGlobalBuffer((__gm__ half*)gmInput, NUM_DATA); + xOutGm.SetGlobalBuffer((__gm__ half*)gmOutput, NUM_DATA); + DataCopy(xLocal, xInGm, NUM_DATA); + // 17行为对UB进行写入，22行为对UB进行读，由于中间没有阻塞，UB上存在先写后读的竞争。解决方法为借助Que，先入队然后出队 + // xQue.EnQue(xLocal); + // LocalTensor deQueLocal = xQue.DeQue(); + // DataCopy(xOutGm, deQueLocal, NUM_DATA); + DataCopy(xOutGm, xLocal, NUM_DATA); +} + +extern "C" void raw_error_kernel_do(uint32_t blockDim, void *l2ctrl, void *stream, uint8_t *gmInput, uint8_t *gmOutput) +{ + raw_error_kernel<<>>(gmInput, gmOutput); +} -- Gitee