From a4e9dccdebe1c87f6903fad7e2a4ad69e7161dfa Mon Sep 17 00:00:00 2001 From: zhengweifeng Date: Tue, 27 Feb 2024 16:37:25 +0800 Subject: [PATCH 01/13] =?UTF-8?q?=E5=A2=9E=E5=8A=A0=E9=9D=9E=E6=B3=95?= =?UTF-8?q?=E8=AF=BB=E5=86=99=E5=92=8C=E5=A4=9A=E6=A0=B8=E8=B8=A9=E8=B8=8F?= =?UTF-8?q?=E6=A0=B7=E4=BE=8B?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- .../illegal_read_and_write/Makefile | 29 ++++++++++++++ .../illegal_read_and_write_kernel.cpp | 25 ++++++++++++ .../illegal_read_and_write/main.cpp | 39 +++++++++++++++++++ sample/sanitizer_sample/out_of_bound/Makefile | 29 ++++++++++++++ .../sanitizer_sample/out_of_bound/kernel.cpp | 30 ++++++++++++++ sample/sanitizer_sample/out_of_bound/main.cpp | 39 +++++++++++++++++++ .../out_of_bound/out_of_bound_kernel.cpp | 30 ++++++++++++++ 7 files changed, 221 insertions(+) create mode 100644 sample/sanitizer_sample/illegal_read_and_write/Makefile create mode 100644 sample/sanitizer_sample/illegal_read_and_write/illegal_read_and_write_kernel.cpp create mode 100644 sample/sanitizer_sample/illegal_read_and_write/main.cpp create mode 100644 sample/sanitizer_sample/out_of_bound/Makefile create mode 100644 sample/sanitizer_sample/out_of_bound/kernel.cpp create mode 100644 sample/sanitizer_sample/out_of_bound/main.cpp create mode 100644 sample/sanitizer_sample/out_of_bound/out_of_bound_kernel.cpp diff --git a/sample/sanitizer_sample/illegal_read_and_write/Makefile b/sample/sanitizer_sample/illegal_read_and_write/Makefile new file mode 100644 index 000000000..dd683bf26 --- /dev/null +++ b/sample/sanitizer_sample/illegal_read_and_write/Makefile @@ -0,0 +1,29 @@ +ASCEND_HOME_PATH ?= /usr/local/Ascend/ascend-toolkit/latest +TOP_DIR ?= / + +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${TOP_DIR}/third_party/inc +LINK_LIBS := -L${TOP_DIR}/third_party/lib -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/illegal_read_and_write/illegal_read_and_write_kernel.cpp b/sample/sanitizer_sample/illegal_read_and_write/illegal_read_and_write_kernel.cpp new file mode 100644 index 000000000..650ceb128 --- /dev/null +++ b/sample/sanitizer_sample/illegal_read_and_write/illegal_read_and_write_kernel.cpp @@ -0,0 +1,25 @@ +#include "kernel_operator.h" +#include "acl/acl.h" +using namespace AscendC; + +constexpr int32_t BYTESIZE = 256; +constexpr int32_t BYTESIZE_LARGE = 512; +constexpr int32_t NUM_DATA = BYTESIZE_EXAMPLE / sizeof(half); +constexpr int32_t NUM_DATA_LARGE = BYTESIZE_LARGE / sizeof(half); + +extern "C" __global__ __aicore__ void illegal_read_and_write_kernel(__gm__ uint8_t *gm) +{ + TPipe pipe; + TBuf xlm; + GlobalTensor xGm; + pipe.InitBuffer(xlm, BYTESIZE_LARGE); + LocalTensor xLm = xlm.Get(); + xGm.SetGlobalBuffer((__gm__ half *)gm, NUM_DATA); + DataCopy(xLm, xGm, NUM_DATA_LARGE); // 这里NUM_DATA_LARGE > NUM_DATA, 发生了对GM的越界非法读 + DataCopy(xGm, xLm, NUM_DATA_LARGE); // 这里NUM_DATA_LARGE > NUM_DATA, 发生了对GM的越界非法写 +} + +extern "C" void illegal_read_and_write_kernel_do(uint32_t blockDim, void *l2ctrl, void *stream, uint8_t *gm) +{ + illegal_read_and_write_kernel<<>>(gm); +} diff --git a/sample/sanitizer_sample/illegal_read_and_write/main.cpp b/sample/sanitizer_sample/illegal_read_and_write/main.cpp new file mode 100644 index 000000000..44138b3d9 --- /dev/null +++ b/sample/sanitizer_sample/illegal_read_and_write/main.cpp @@ -0,0 +1,39 @@ +#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 illegal_read_and_write_kernel_do(uint32_t blockDim, void *l2ctrl, void *stream, uint8_t *gm); + +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 *gm = nullptr; + CHECK_ACL(aclrtMalloc((void**)&gm, 256, ACL_MEM_MALLOC_HUGE_FIRST)); + + uint64_t blockDim = 1UL; + illegal_read_and_write_kernel_do(blockDim, nullptr, stream, gm); + CHECK_ACL(aclrtSynchronizeStream(stream)); + + CHECK_ACL(aclrtFree(gm)); + 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/out_of_bound/Makefile b/sample/sanitizer_sample/out_of_bound/Makefile new file mode 100644 index 000000000..dd683bf26 --- /dev/null +++ b/sample/sanitizer_sample/out_of_bound/Makefile @@ -0,0 +1,29 @@ +ASCEND_HOME_PATH ?= /usr/local/Ascend/ascend-toolkit/latest +TOP_DIR ?= / + +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${TOP_DIR}/third_party/inc +LINK_LIBS := -L${TOP_DIR}/third_party/lib -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/out_of_bound/kernel.cpp b/sample/sanitizer_sample/out_of_bound/kernel.cpp new file mode 100644 index 000000000..18bf12846 --- /dev/null +++ b/sample/sanitizer_sample/out_of_bound/kernel.cpp @@ -0,0 +1,30 @@ +#include "kernel_operator.h" +#include "acl/acl.h" +using namespace AscendC; + +constexpr int32_t BYTESIZE_EXAMPLE = 512; +constexpr int32_t NUM_DATA = 16; +constexpr int32_t CORE_OFFSET = 14; +constexpr int32_t LOOP_COUNT = 10; + +extern "C" __global__ __aicore__ void kernel(__gm__ uint8_t *gm) +{ + TPipe pipe; + TBuf xlm; + GlobalTensor xGm; + pipe.InitBuffer(xlm, BYTESIZE_EXAMPLE); + LocalTensor xLm = xlm.Get(); + for (int32_t i = 0; i < LOOP_COUNT; i++) + { + if (i == GetBlockIdx()) + { + xGm.SetGlobalBuffer((__gm__ half *)gm + i * CORE_OFFSET, NUM_DATA); + DataCopy(xGm, xLm, NUM_DATA); // 这里CORE_OFFSET < NUM_DATA, 多核写入GM出现内存踩踏 + } + } +} + +extern "C" void kernel_do(uint32_t blockDim, void *l2ctrl, void *stream, uint8_t *gm) +{ + raw_error_kernel<<>>(gm); +} diff --git a/sample/sanitizer_sample/out_of_bound/main.cpp b/sample/sanitizer_sample/out_of_bound/main.cpp new file mode 100644 index 000000000..53da5b54a --- /dev/null +++ b/sample/sanitizer_sample/out_of_bound/main.cpp @@ -0,0 +1,39 @@ +#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 out_of_bound_kernel_do(uint32_t blockDim, void *l2ctrl, void *stream, uint8_t *gm); + +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 *gm = nullptr; + CHECK_ACL(aclrtMalloc((void**)&gm, 512, ACL_MEM_MALLOC_HUGE_FIRST)); + + uint64_t blockDim = 10UL; + out_of_bound_kernel_do(blockDim, nullptr, stream, gm); + CHECK_ACL(aclrtSynchronizeStream(stream)); + + CHECK_ACL(aclrtFree(gm)); + 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/out_of_bound/out_of_bound_kernel.cpp b/sample/sanitizer_sample/out_of_bound/out_of_bound_kernel.cpp new file mode 100644 index 000000000..f0995f3f1 --- /dev/null +++ b/sample/sanitizer_sample/out_of_bound/out_of_bound_kernel.cpp @@ -0,0 +1,30 @@ +#include "kernel_operator.h" +#include "acl/acl.h" +using namespace AscendC; + +constexpr int32_t BYTESIZE_EXAMPLE = 512; +constexpr int32_t NUM_DATA = 16; +constexpr int32_t CORE_OFFSET = 14; +constexpr int32_t LOOP_COUNT = 10; + +extern "C" __global__ __aicore__ void out_of_bound_kernel(__gm__ uint8_t *gm) +{ + TPipe pipe; + TBuf xlm; + GlobalTensor xGm; + pipe.InitBuffer(xlm, BYTESIZE_EXAMPLE); + LocalTensor xLm = xlm.Get(); + for (int32_t i = 0; i < LOOP_COUNT; i++) + { + if (i == GetBlockIdx()) + { + xGm.SetGlobalBuffer((__gm__ half *)gm + i * CORE_OFFSET, NUM_DATA); + DataCopy(xGm, xLm, NUM_DATA); // 这里CORE_OFFSET < NUM_DATA, 多核写入GM出现内存踩踏 + } + } +} + +extern "C" void out_of_bound_kernel_do(uint32_t blockDim, void *l2ctrl, void *stream, uint8_t *gm) +{ + out_of_bound_kernel<<>>(gm); +} -- Gitee From 3954b920bcb7c53153a206357a7cd3e2aefd1fd1 Mon Sep 17 00:00:00 2001 From: zhengweifeng Date: Tue, 27 Feb 2024 08:50:27 +0000 Subject: [PATCH 02/13] =?UTF-8?q?=E5=88=A0=E9=99=A4=E6=96=87=E4=BB=B6=20sa?= =?UTF-8?q?mple/sanitizer=5Fsample/out=5Fof=5Fbound/kernel.cpp?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- .../sanitizer_sample/out_of_bound/kernel.cpp | 30 ------------------- 1 file changed, 30 deletions(-) delete mode 100644 sample/sanitizer_sample/out_of_bound/kernel.cpp diff --git a/sample/sanitizer_sample/out_of_bound/kernel.cpp b/sample/sanitizer_sample/out_of_bound/kernel.cpp deleted file mode 100644 index 18bf12846..000000000 --- a/sample/sanitizer_sample/out_of_bound/kernel.cpp +++ /dev/null @@ -1,30 +0,0 @@ -#include "kernel_operator.h" -#include "acl/acl.h" -using namespace AscendC; - -constexpr int32_t BYTESIZE_EXAMPLE = 512; -constexpr int32_t NUM_DATA = 16; -constexpr int32_t CORE_OFFSET = 14; -constexpr int32_t LOOP_COUNT = 10; - -extern "C" __global__ __aicore__ void kernel(__gm__ uint8_t *gm) -{ - TPipe pipe; - TBuf xlm; - GlobalTensor xGm; - pipe.InitBuffer(xlm, BYTESIZE_EXAMPLE); - LocalTensor xLm = xlm.Get(); - for (int32_t i = 0; i < LOOP_COUNT; i++) - { - if (i == GetBlockIdx()) - { - xGm.SetGlobalBuffer((__gm__ half *)gm + i * CORE_OFFSET, NUM_DATA); - DataCopy(xGm, xLm, NUM_DATA); // 这里CORE_OFFSET < NUM_DATA, 多核写入GM出现内存踩踏 - } - } -} - -extern "C" void kernel_do(uint32_t blockDim, void *l2ctrl, void *stream, uint8_t *gm) -{ - raw_error_kernel<<>>(gm); -} -- Gitee From e5a6458c9823e3eac296cea423555d231629b0f4 Mon Sep 17 00:00:00 2001 From: zhengweifeng Date: Tue, 27 Feb 2024 16:37:25 +0800 Subject: [PATCH 03/13] =?UTF-8?q?=E5=A2=9E=E5=8A=A0=E9=9D=9E=E6=B3=95?= =?UTF-8?q?=E8=AF=BB=E5=86=99=E5=92=8C=E5=A4=9A=E6=A0=B8=E8=B8=A9=E8=B8=8F?= =?UTF-8?q?=E6=A0=B7=E4=BE=8B?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- sample/build/build.sh | 10 ++++++++++ .../sanitizer_sample/illegal_read_and_write/Makefile | 4 ++++ .../illegal_read_and_write_kernel.cpp | 8 +++++--- sample/sanitizer_sample/out_of_bound/Makefile | 4 ++++ .../out_of_bound/out_of_bound_kernel.cpp | 8 +++++--- 5 files changed, 28 insertions(+), 6 deletions(-) diff --git a/sample/build/build.sh b/sample/build/build.sh index 299e04d9f..12175d487 100644 --- a/sample/build/build.sh +++ b/sample/build/build.sh @@ -16,4 +16,14 @@ mv *.fatbin ${TOP_DIR}/build # matmul_leakyrelu cd ${TOP_DIR}/normal_sample/mix make +mv *.fatbin ${TOP_DIR}/build + +# illegal_read_and_write +cd ${TOP_DIR}/sanitizer_sample/illegal_read_and_write +make +mv *.fatbin ${TOP_DIR}/build + +# out_of_bound +cd ${TOP_DIR}/sanitizer_sample/out_of_bound +make mv *.fatbin ${TOP_DIR}/build \ No newline at end of file diff --git a/sample/sanitizer_sample/illegal_read_and_write/Makefile b/sample/sanitizer_sample/illegal_read_and_write/Makefile index dd683bf26..1331e388a 100644 --- a/sample/sanitizer_sample/illegal_read_and_write/Makefile +++ b/sample/sanitizer_sample/illegal_read_and_write/Makefile @@ -9,7 +9,11 @@ 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${TOP_DIR}/third_party/inc +<<<<<<< HEAD LINK_LIBS := -L${TOP_DIR}/third_party/lib -lruntime -lascendcl -lstdc++ +======= +LINK_LIBS := -L${ASCEND_HOME_PATH}/lib64 -lruntime -lascendcl -lstdc++ +>>>>>>> 0ca8de6 (增加非法读写和多核踩踏样例) all: build diff --git a/sample/sanitizer_sample/illegal_read_and_write/illegal_read_and_write_kernel.cpp b/sample/sanitizer_sample/illegal_read_and_write/illegal_read_and_write_kernel.cpp index 650ceb128..6fd83c518 100644 --- a/sample/sanitizer_sample/illegal_read_and_write/illegal_read_and_write_kernel.cpp +++ b/sample/sanitizer_sample/illegal_read_and_write/illegal_read_and_write_kernel.cpp @@ -4,7 +4,7 @@ using namespace AscendC; constexpr int32_t BYTESIZE = 256; constexpr int32_t BYTESIZE_LARGE = 512; -constexpr int32_t NUM_DATA = BYTESIZE_EXAMPLE / sizeof(half); +constexpr int32_t NUM_DATA = BYTESIZE / sizeof(half); constexpr int32_t NUM_DATA_LARGE = BYTESIZE_LARGE / sizeof(half); extern "C" __global__ __aicore__ void illegal_read_and_write_kernel(__gm__ uint8_t *gm) @@ -14,9 +14,11 @@ extern "C" __global__ __aicore__ void illegal_read_and_write_kernel(__gm__ uint8 GlobalTensor xGm; pipe.InitBuffer(xlm, BYTESIZE_LARGE); LocalTensor xLm = xlm.Get(); + // 第18行给xGm分配了BYTESIZE字节的内存，但是第20、21行DataCopy搬运了BYTESIZE_LARGE字节的内存， + // BYTESIZE_LARGE > BYTESIZE,导致对xGm的越界非法读写，解决方法是将NUM_DATA替换成BYTESIZE_LARGE xGm.SetGlobalBuffer((__gm__ half *)gm, NUM_DATA); - DataCopy(xLm, xGm, NUM_DATA_LARGE); // 这里NUM_DATA_LARGE > NUM_DATA, 发生了对GM的越界非法读 - DataCopy(xGm, xLm, NUM_DATA_LARGE); // 这里NUM_DATA_LARGE > NUM_DATA, 发生了对GM的越界非法写 + DataCopy(xLm, xGm, NUM_DATA_LARGE); + DataCopy(xGm, xLm, NUM_DATA_LARGE); } extern "C" void illegal_read_and_write_kernel_do(uint32_t blockDim, void *l2ctrl, void *stream, uint8_t *gm) diff --git a/sample/sanitizer_sample/out_of_bound/Makefile b/sample/sanitizer_sample/out_of_bound/Makefile index dd683bf26..1331e388a 100644 --- a/sample/sanitizer_sample/out_of_bound/Makefile +++ b/sample/sanitizer_sample/out_of_bound/Makefile @@ -9,7 +9,11 @@ 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${TOP_DIR}/third_party/inc +<<<<<<< HEAD LINK_LIBS := -L${TOP_DIR}/third_party/lib -lruntime -lascendcl -lstdc++ +======= +LINK_LIBS := -L${ASCEND_HOME_PATH}/lib64 -lruntime -lascendcl -lstdc++ +>>>>>>> 0ca8de6 (增加非法读写和多核踩踏样例) all: build diff --git a/sample/sanitizer_sample/out_of_bound/out_of_bound_kernel.cpp b/sample/sanitizer_sample/out_of_bound/out_of_bound_kernel.cpp index f0995f3f1..540ebd0ea 100644 --- a/sample/sanitizer_sample/out_of_bound/out_of_bound_kernel.cpp +++ b/sample/sanitizer_sample/out_of_bound/out_of_bound_kernel.cpp @@ -2,7 +2,7 @@ #include "acl/acl.h" using namespace AscendC; -constexpr int32_t BYTESIZE_EXAMPLE = 512; +constexpr int32_t BYTESIZE = 512; constexpr int32_t NUM_DATA = 16; constexpr int32_t CORE_OFFSET = 14; constexpr int32_t LOOP_COUNT = 10; @@ -12,14 +12,16 @@ extern "C" __global__ __aicore__ void out_of_bound_kernel(__gm__ uint8_t *gm) TPipe pipe; TBuf xlm; GlobalTensor xGm; - pipe.InitBuffer(xlm, BYTESIZE_EXAMPLE); + pipe.InitBuffer(xlm, BYTESIZE); LocalTensor xLm = xlm.Get(); for (int32_t i = 0; i < LOOP_COUNT; i++) { if (i == GetBlockIdx()) { + // 这里第22行CORE_OFFSET < NUM_DATA, 第23行多核写入GM时，写入的size大于偏移，导致出现内存踩踏 + // 解决方法时将CORE_OFFSET替换成NUM_DATA xGm.SetGlobalBuffer((__gm__ half *)gm + i * CORE_OFFSET, NUM_DATA); - DataCopy(xGm, xLm, NUM_DATA); // 这里CORE_OFFSET < NUM_DATA, 多核写入GM出现内存踩踏 + DataCopy(xGm, xLm, NUM_DATA); } } } -- Gitee From 379746ef0dfade58871873801f80c0b24e218fea Mon Sep 17 00:00:00 2001 From: zhengweifeng Date: Tue, 27 Feb 2024 16:37:25 +0800 Subject: [PATCH 04/13] =?UTF-8?q?=E5=A2=9E=E5=8A=A0=E9=9D=9E=E6=B3=95?= =?UTF-8?q?=E8=AF=BB=E5=86=99=E5=92=8C=E5=A4=9A=E6=A0=B8=E8=B8=A9=E8=B8=8F?= =?UTF-8?q?=E6=A0=B7=E4=BE=8B?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- sample/sanitizer_sample/illegal_read_and_write/Makefile | 4 ---- sample/sanitizer_sample/out_of_bound/Makefile | 4 ---- 2 files changed, 8 deletions(-) diff --git a/sample/sanitizer_sample/illegal_read_and_write/Makefile b/sample/sanitizer_sample/illegal_read_and_write/Makefile index 1331e388a..42d437359 100644 --- a/sample/sanitizer_sample/illegal_read_and_write/Makefile +++ b/sample/sanitizer_sample/illegal_read_and_write/Makefile @@ -9,11 +9,7 @@ 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${TOP_DIR}/third_party/inc -<<<<<<< HEAD -LINK_LIBS := -L${TOP_DIR}/third_party/lib -lruntime -lascendcl -lstdc++ -======= LINK_LIBS := -L${ASCEND_HOME_PATH}/lib64 -lruntime -lascendcl -lstdc++ ->>>>>>> 0ca8de6 (增加非法读写和多核踩踏样例) all: build diff --git a/sample/sanitizer_sample/out_of_bound/Makefile b/sample/sanitizer_sample/out_of_bound/Makefile index 1331e388a..42d437359 100644 --- a/sample/sanitizer_sample/out_of_bound/Makefile +++ b/sample/sanitizer_sample/out_of_bound/Makefile @@ -9,11 +9,7 @@ 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${TOP_DIR}/third_party/inc -<<<<<<< HEAD -LINK_LIBS := -L${TOP_DIR}/third_party/lib -lruntime -lascendcl -lstdc++ -======= LINK_LIBS := -L${ASCEND_HOME_PATH}/lib64 -lruntime -lascendcl -lstdc++ ->>>>>>> 0ca8de6 (增加非法读写和多核踩踏样例) all: build -- Gitee From 885bd56d7e6728c033c923242a61c5b12cdb61e5 Mon Sep 17 00:00:00 2001 From: zhengweifeng Date: Tue, 27 Feb 2024 09:35:37 +0000 Subject: [PATCH 05/13] update sample/sanitizer_sample/out_of_bound/out_of_bound_kernel.cpp. Signed-off-by: zhengweifeng --- sample/sanitizer_sample/out_of_bound/out_of_bound_kernel.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sample/sanitizer_sample/out_of_bound/out_of_bound_kernel.cpp b/sample/sanitizer_sample/out_of_bound/out_of_bound_kernel.cpp index 540ebd0ea..1c248d303 100644 --- a/sample/sanitizer_sample/out_of_bound/out_of_bound_kernel.cpp +++ b/sample/sanitizer_sample/out_of_bound/out_of_bound_kernel.cpp @@ -19,7 +19,7 @@ extern "C" __global__ __aicore__ void out_of_bound_kernel(__gm__ uint8_t *gm) if (i == GetBlockIdx()) { // 这里第22行CORE_OFFSET < NUM_DATA, 第23行多核写入GM时，写入的size大于偏移，导致出现内存踩踏 - // 解决方法时将CORE_OFFSET替换成NUM_DATA + // 解决方法是将CORE_OFFSET替换成NUM_DATA xGm.SetGlobalBuffer((__gm__ half *)gm + i * CORE_OFFSET, NUM_DATA); DataCopy(xGm, xLm, NUM_DATA); } -- Gitee From 9af7eba4380a0d988987c64584d7dc6442eb674e Mon Sep 17 00:00:00 2001 From: zhengweifeng Date: Tue, 27 Feb 2024 09:35:37 +0000 Subject: [PATCH 06/13] update sample/sanitizer_sample/out_of_bound/out_of_bound_kernel.cpp. Signed-off-by: zhengweifeng --- sample/sanitizer_sample/illegal_read_and_write/Makefile | 8 ++++---- sample/sanitizer_sample/out_of_bound/Makefile | 8 ++++---- 2 files changed, 8 insertions(+), 8 deletions(-) diff --git a/sample/sanitizer_sample/illegal_read_and_write/Makefile b/sample/sanitizer_sample/illegal_read_and_write/Makefile index 42d437359..e50060c74 100644 --- a/sample/sanitizer_sample/illegal_read_and_write/Makefile +++ b/sample/sanitizer_sample/illegal_read_and_write/Makefile @@ -13,17 +13,17 @@ LINK_LIBS := -L${ASCEND_HOME_PATH}/lib64 -lruntime -lascendcl all: build -build: raw_error_kernel.o main.o raw_error.fatbin +build: illegal_read_and_write_kernel.o main.o illegal_read_and_write.fatbin -raw_error_kernel.o: raw_error_kernel.cpp +illegal_read_and_write_kernel.o: illegal_read_and_write_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 +illegal_read_and_write.fatbin: illegal_read_and_write_kernel.o main.o $(COMPILER) $(LINK_FLAG) $(DAV_FLAG) -o $@ $^ ${LINK_LIBS} .PHONY: clean clean: - rm *.o raw_error.fatbin + rm *.o illegal_read_and_write.fatbin diff --git a/sample/sanitizer_sample/out_of_bound/Makefile b/sample/sanitizer_sample/out_of_bound/Makefile index 42d437359..5cbdfe69e 100644 --- a/sample/sanitizer_sample/out_of_bound/Makefile +++ b/sample/sanitizer_sample/out_of_bound/Makefile @@ -13,17 +13,17 @@ LINK_LIBS := -L${ASCEND_HOME_PATH}/lib64 -lruntime -lascendcl all: build -build: raw_error_kernel.o main.o raw_error.fatbin +build: out_of_bound_kernel.o main.o out_of_bound.fatbin -raw_error_kernel.o: raw_error_kernel.cpp +out_of_bound_kernel.o: out_of_bound_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 +out_of_bound.fatbin: out_of_bound_kernel.o main.o $(COMPILER) $(LINK_FLAG) $(DAV_FLAG) -o $@ $^ ${LINK_LIBS} .PHONY: clean clean: - rm *.o raw_error.fatbin + rm *.o out_of_bound.fatbin -- Gitee From 45142c76d34630f0d9e0d190521ac828fc95611a Mon Sep 17 00:00:00 2001 From: zhengweifeng Date: Tue, 27 Feb 2024 11:10:14 +0000 Subject: [PATCH 07/13] update sample/sanitizer_sample/illegal_read_and_write/Makefile. Signed-off-by: zhengweifeng --- sample/sanitizer_sample/illegal_read_and_write/Makefile | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sample/sanitizer_sample/illegal_read_and_write/Makefile b/sample/sanitizer_sample/illegal_read_and_write/Makefile index e50060c74..a2cdd20ad 100644 --- a/sample/sanitizer_sample/illegal_read_and_write/Makefile +++ b/sample/sanitizer_sample/illegal_read_and_write/Makefile @@ -8,7 +8,7 @@ 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${TOP_DIR}/third_party/inc +HOST_INC_FLAG := -I${ASCEND_HOME_PATH}/include LINK_LIBS := -L${ASCEND_HOME_PATH}/lib64 -lruntime -lascendcl -lstdc++ all: build -- Gitee From bdb3311c45cf4514a971551a85113858b41a15a2 Mon Sep 17 00:00:00 2001 From: zhengweifeng Date: Tue, 27 Feb 2024 11:10:55 +0000 Subject: [PATCH 08/13] update sample/sanitizer_sample/out_of_bound/Makefile. Signed-off-by: zhengweifeng --- sample/sanitizer_sample/out_of_bound/Makefile | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sample/sanitizer_sample/out_of_bound/Makefile b/sample/sanitizer_sample/out_of_bound/Makefile index 5cbdfe69e..58c903341 100644 --- a/sample/sanitizer_sample/out_of_bound/Makefile +++ b/sample/sanitizer_sample/out_of_bound/Makefile @@ -8,7 +8,7 @@ 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${TOP_DIR}/third_party/inc +HOST_INC_FLAG := -I${ASCEND_HOME_PATH}/include LINK_LIBS := -L${ASCEND_HOME_PATH}/lib64 -lruntime -lascendcl -lstdc++ all: build -- Gitee From 8f0e6dfaf3839ccbb44fd177a4eb5246ba8c575f Mon Sep 17 00:00:00 2001 From: zhengweifeng Date: Tue, 27 Feb 2024 11:19:45 +0000 Subject: [PATCH 09/13] update sample/sanitizer_sample/out_of_bound/out_of_bound_kernel.cpp. Signed-off-by: zhengweifeng --- sample/sanitizer_sample/out_of_bound/out_of_bound_kernel.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sample/sanitizer_sample/out_of_bound/out_of_bound_kernel.cpp b/sample/sanitizer_sample/out_of_bound/out_of_bound_kernel.cpp index 1c248d303..36d964823 100644 --- a/sample/sanitizer_sample/out_of_bound/out_of_bound_kernel.cpp +++ b/sample/sanitizer_sample/out_of_bound/out_of_bound_kernel.cpp @@ -18,7 +18,7 @@ extern "C" __global__ __aicore__ void out_of_bound_kernel(__gm__ uint8_t *gm) { if (i == GetBlockIdx()) { - // 这里第22行CORE_OFFSET < NUM_DATA, 第23行多核写入GM时，写入的size大于偏移，导致出现内存踩踏 + // 这里第23行CORE_OFFSET < NUM_DATA, 第24行多核写入GM时，写入的size大于偏移，导致出现内存踩踏 // 解决方法是将CORE_OFFSET替换成NUM_DATA xGm.SetGlobalBuffer((__gm__ half *)gm + i * CORE_OFFSET, NUM_DATA); DataCopy(xGm, xLm, NUM_DATA); -- Gitee From 19da32fd35af3b293b7bafde31d11d4f8c69ddbd Mon Sep 17 00:00:00 2001 From: zhengweifeng Date: Tue, 27 Feb 2024 09:35:37 +0000 Subject: [PATCH 10/13] update sample/sanitizer_sample/out_of_bound/out_of_bound_kernel.cpp. Signed-off-by: zhengweifeng --- sample/sanitizer_sample/illegal_read_and_write/Makefile | 1 - sample/sanitizer_sample/out_of_bound/Makefile | 1 - 2 files changed, 2 deletions(-) diff --git a/sample/sanitizer_sample/illegal_read_and_write/Makefile b/sample/sanitizer_sample/illegal_read_and_write/Makefile index a2cdd20ad..4d810639e 100644 --- a/sample/sanitizer_sample/illegal_read_and_write/Makefile +++ b/sample/sanitizer_sample/illegal_read_and_write/Makefile @@ -1,5 +1,4 @@ ASCEND_HOME_PATH ?= /usr/local/Ascend/ascend-toolkit/latest -TOP_DIR ?= / COMPILER := ${ASCEND_HOME_PATH}/compiler/ccec_compiler/bin/ccec HOST_COMPILER := g++ diff --git a/sample/sanitizer_sample/out_of_bound/Makefile b/sample/sanitizer_sample/out_of_bound/Makefile index 58c903341..22e06f70f 100644 --- a/sample/sanitizer_sample/out_of_bound/Makefile +++ b/sample/sanitizer_sample/out_of_bound/Makefile @@ -1,5 +1,4 @@ ASCEND_HOME_PATH ?= /usr/local/Ascend/ascend-toolkit/latest -TOP_DIR ?= / COMPILER := ${ASCEND_HOME_PATH}/compiler/ccec_compiler/bin/ccec HOST_COMPILER := g++ -- Gitee From a7fd6e70bb9485842e5e3aff8389d9010536b65f Mon Sep 17 00:00:00 2001 From: zhengweifeng Date: Tue, 27 Feb 2024 09:35:37 +0000 Subject: [PATCH 11/13] update sample/sanitizer_sample/out_of_bound/out_of_bound_kernel.cpp. Signed-off-by: zhengweifeng --- .../illegal_read_and_write/illegal_read_and_write_kernel.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sample/sanitizer_sample/illegal_read_and_write/illegal_read_and_write_kernel.cpp b/sample/sanitizer_sample/illegal_read_and_write/illegal_read_and_write_kernel.cpp index 6fd83c518..f4094da92 100644 --- a/sample/sanitizer_sample/illegal_read_and_write/illegal_read_and_write_kernel.cpp +++ b/sample/sanitizer_sample/illegal_read_and_write/illegal_read_and_write_kernel.cpp @@ -15,7 +15,7 @@ extern "C" __global__ __aicore__ void illegal_read_and_write_kernel(__gm__ uint8 pipe.InitBuffer(xlm, BYTESIZE_LARGE); LocalTensor xLm = xlm.Get(); // 第18行给xGm分配了BYTESIZE字节的内存，但是第20、21行DataCopy搬运了BYTESIZE_LARGE字节的内存， - // BYTESIZE_LARGE > BYTESIZE,导致对xGm的越界非法读写，解决方法是将NUM_DATA替换成BYTESIZE_LARGE + // BYTESIZE_LARGE > BYTESIZE,导致对xGm的越界非法读写，解决方法是将第20、21行NUM_DATA_LARGE替换成NUM_DATA xGm.SetGlobalBuffer((__gm__ half *)gm, NUM_DATA); DataCopy(xLm, xGm, NUM_DATA_LARGE); DataCopy(xGm, xLm, NUM_DATA_LARGE); -- Gitee From b6f88f17f40b2eafd6d5a0c3a8b7e44bdc99f859 Mon Sep 17 00:00:00 2001 From: zhengweifeng Date: Tue, 27 Feb 2024 09:35:37 +0000 Subject: [PATCH 12/13] update sample/sanitizer_sample/out_of_bound/out_of_bound_kernel.cpp. Signed-off-by: zhengweifeng --- .../illegal_read_and_write_kernel.cpp | 6 ++++-- .../out_of_bound/out_of_bound_kernel.cpp | 15 +++++---------- 2 files changed, 9 insertions(+), 12 deletions(-) diff --git a/sample/sanitizer_sample/illegal_read_and_write/illegal_read_and_write_kernel.cpp b/sample/sanitizer_sample/illegal_read_and_write/illegal_read_and_write_kernel.cpp index f4094da92..8a85df5cc 100644 --- a/sample/sanitizer_sample/illegal_read_and_write/illegal_read_and_write_kernel.cpp +++ b/sample/sanitizer_sample/illegal_read_and_write/illegal_read_and_write_kernel.cpp @@ -14,11 +14,13 @@ extern "C" __global__ __aicore__ void illegal_read_and_write_kernel(__gm__ uint8 GlobalTensor xGm; pipe.InitBuffer(xlm, BYTESIZE_LARGE); LocalTensor xLm = xlm.Get(); - // 第18行给xGm分配了BYTESIZE字节的内存，但是第20、21行DataCopy搬运了BYTESIZE_LARGE字节的内存， - // BYTESIZE_LARGE > BYTESIZE,导致对xGm的越界非法读写，解决方法是将第20、21行NUM_DATA_LARGE替换成NUM_DATA xGm.SetGlobalBuffer((__gm__ half *)gm, NUM_DATA); DataCopy(xLm, xGm, NUM_DATA_LARGE); DataCopy(xGm, xLm, NUM_DATA_LARGE); + // 第17行给xGm分配了BYTESIZE字节的内存，但是第18、19行DataCopy搬运了BYTESIZE_LARGE字节的内存 + // BYTESIZE_LARGE > BYTESIZE，导致对xGm的越界非法读写，以下是正确写法 + // DataCopy(xLm, xGm, NUM_DATA); + // DataCopy(xGm, xLm, NUM_DATA); } extern "C" void illegal_read_and_write_kernel_do(uint32_t blockDim, void *l2ctrl, void *stream, uint8_t *gm) diff --git a/sample/sanitizer_sample/out_of_bound/out_of_bound_kernel.cpp b/sample/sanitizer_sample/out_of_bound/out_of_bound_kernel.cpp index 36d964823..cb54dc322 100644 --- a/sample/sanitizer_sample/out_of_bound/out_of_bound_kernel.cpp +++ b/sample/sanitizer_sample/out_of_bound/out_of_bound_kernel.cpp @@ -14,16 +14,11 @@ extern "C" __global__ __aicore__ void out_of_bound_kernel(__gm__ uint8_t *gm) GlobalTensor xGm; pipe.InitBuffer(xlm, BYTESIZE); LocalTensor xLm = xlm.Get(); - for (int32_t i = 0; i < LOOP_COUNT; i++) - { - if (i == GetBlockIdx()) - { - // 这里第23行CORE_OFFSET < NUM_DATA, 第24行多核写入GM时，写入的size大于偏移，导致出现内存踩踏 - // 解决方法是将CORE_OFFSET替换成NUM_DATA - xGm.SetGlobalBuffer((__gm__ half *)gm + i * CORE_OFFSET, NUM_DATA); - DataCopy(xGm, xLm, NUM_DATA); - } - } + xGm.SetGlobalBuffer((__gm__ half *)gm + GetBlockIdx() * CORE_OFFSET, NUM_DATA); + // 这里第17行CORE_OFFSET < NUM_DATA, 第21行多核写入GM时，写入的size大于偏移，导致出现内存踩踏 + // 以下是正确写法 + // xGm.SetGlobalBuffer((__gm__ half *)gm + GetBlockIdx() * NUM_DATA, NUM_DATA); + DataCopy(xGm, xLm, NUM_DATA); } extern "C" void out_of_bound_kernel_do(uint32_t blockDim, void *l2ctrl, void *stream, uint8_t *gm) -- Gitee From 5a6c592b8c0e305dd9d2dd58e75635073b3875e0 Mon Sep 17 00:00:00 2001 From: zhengweifeng Date: Tue, 27 Feb 2024 09:35:37 +0000 Subject: [PATCH 13/13] update sample/sanitizer_sample/out_of_bound/out_of_bound_kernel.cpp. Signed-off-by: zhengweifeng --- sample/build/build.sh | 4 ++-- .../{ => memcheck}/illegal_read_and_write/Makefile | 0 .../illegal_read_and_write/illegal_read_and_write_kernel.cpp | 0 .../{ => memcheck}/illegal_read_and_write/main.cpp | 0 sample/sanitizer_sample/{ => memcheck}/out_of_bound/Makefile | 0 sample/sanitizer_sample/{ => memcheck}/out_of_bound/main.cpp | 0 .../{ => memcheck}/out_of_bound/out_of_bound_kernel.cpp | 0 7 files changed, 2 insertions(+), 2 deletions(-) rename sample/sanitizer_sample/{ => memcheck}/illegal_read_and_write/Makefile (100%) rename sample/sanitizer_sample/{ => memcheck}/illegal_read_and_write/illegal_read_and_write_kernel.cpp (100%) rename sample/sanitizer_sample/{ => memcheck}/illegal_read_and_write/main.cpp (100%) rename sample/sanitizer_sample/{ => memcheck}/out_of_bound/Makefile (100%) rename sample/sanitizer_sample/{ => memcheck}/out_of_bound/main.cpp (100%) rename sample/sanitizer_sample/{ => memcheck}/out_of_bound/out_of_bound_kernel.cpp (100%) diff --git a/sample/build/build.sh b/sample/build/build.sh index 1680ac621..b393f205f 100644 --- a/sample/build/build.sh +++ b/sample/build/build.sh @@ -19,12 +19,12 @@ make mv *.fatbin ${TOP_DIR}/build # illegal_read_and_write -cd ${TOP_DIR}/sanitizer_sample/illegal_read_and_write +cd ${TOP_DIR}/sanitizer_sample/memcheck/illegal_read_and_write make mv *.fatbin ${TOP_DIR}/build # out_of_bound -cd ${TOP_DIR}/sanitizer_sample/out_of_bound +cd ${TOP_DIR}/sanitizer_sample/memcheck/out_of_bound make mv *.fatbin ${TOP_DIR}/build diff --git a/sample/sanitizer_sample/illegal_read_and_write/Makefile b/sample/sanitizer_sample/memcheck/illegal_read_and_write/Makefile similarity index 100% rename from sample/sanitizer_sample/illegal_read_and_write/Makefile rename to sample/sanitizer_sample/memcheck/illegal_read_and_write/Makefile diff --git a/sample/sanitizer_sample/illegal_read_and_write/illegal_read_and_write_kernel.cpp b/sample/sanitizer_sample/memcheck/illegal_read_and_write/illegal_read_and_write_kernel.cpp similarity index 100% rename from sample/sanitizer_sample/illegal_read_and_write/illegal_read_and_write_kernel.cpp rename to sample/sanitizer_sample/memcheck/illegal_read_and_write/illegal_read_and_write_kernel.cpp diff --git a/sample/sanitizer_sample/illegal_read_and_write/main.cpp b/sample/sanitizer_sample/memcheck/illegal_read_and_write/main.cpp similarity index 100% rename from sample/sanitizer_sample/illegal_read_and_write/main.cpp rename to sample/sanitizer_sample/memcheck/illegal_read_and_write/main.cpp diff --git a/sample/sanitizer_sample/out_of_bound/Makefile b/sample/sanitizer_sample/memcheck/out_of_bound/Makefile similarity index 100% rename from sample/sanitizer_sample/out_of_bound/Makefile rename to sample/sanitizer_sample/memcheck/out_of_bound/Makefile diff --git a/sample/sanitizer_sample/out_of_bound/main.cpp b/sample/sanitizer_sample/memcheck/out_of_bound/main.cpp similarity index 100% rename from sample/sanitizer_sample/out_of_bound/main.cpp rename to sample/sanitizer_sample/memcheck/out_of_bound/main.cpp diff --git a/sample/sanitizer_sample/out_of_bound/out_of_bound_kernel.cpp b/sample/sanitizer_sample/memcheck/out_of_bound/out_of_bound_kernel.cpp similarity index 100% rename from sample/sanitizer_sample/out_of_bound/out_of_bound_kernel.cpp rename to sample/sanitizer_sample/memcheck/out_of_bound/out_of_bound_kernel.cpp -- Gitee