diff --git a/sample/README.md b/sample/README.md index b46181b022bc117661612d8c9c4446f2b6f7984f..4f49fec5acad887d8fbc395f3b13980f18878a56 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 80e17f8d518cf774745be22318a394d34a247559..ca6689ab7b21c5faccec4143b36d063e050f6968 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 0000000000000000000000000000000000000000..ac230761b61a28fa17decd73803949826b91aee0 --- /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 0000000000000000000000000000000000000000..cd96d8e0858a06d217105f48610890220ab59a4d --- /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 0000000000000000000000000000000000000000..b1db856092d3d552654c2c640facf32a8e42ba88 --- /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); +}