diff --git a/sample/build/build.sh b/sample/build/build.sh index 40890330b626577d569b213ecb996bd192c62473..b393f205f41a74d7e2fbc2afe1086b515e6797ee 100644 --- a/sample/build/build.sh +++ b/sample/build/build.sh @@ -18,6 +18,16 @@ cd ${TOP_DIR}/normal_sample/mix make mv *.fatbin ${TOP_DIR}/build +# 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/memcheck/out_of_bound +make +mv *.fatbin ${TOP_DIR}/build + # illegal align sample for sanitizer cd ${TOP_DIR}/sanitizer_sample/memcheck/illegal_align make diff --git a/sample/sanitizer_sample/memcheck/illegal_read_and_write/Makefile b/sample/sanitizer_sample/memcheck/illegal_read_and_write/Makefile new file mode 100644 index 0000000000000000000000000000000000000000..4d810639efc60a737236f2a318e4b13cd02c00eb --- /dev/null +++ b/sample/sanitizer_sample/memcheck/illegal_read_and_write/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: illegal_read_and_write_kernel.o main.o illegal_read_and_write.fatbin + +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 $^ + +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 illegal_read_and_write.fatbin diff --git a/sample/sanitizer_sample/memcheck/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 new file mode 100644 index 0000000000000000000000000000000000000000..8a85df5cc049419d7238c8f01d070746f198bd2e --- /dev/null +++ b/sample/sanitizer_sample/memcheck/illegal_read_and_write/illegal_read_and_write_kernel.cpp @@ -0,0 +1,29 @@ +#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 / 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); + 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) +{ + illegal_read_and_write_kernel<<>>(gm); +} diff --git a/sample/sanitizer_sample/memcheck/illegal_read_and_write/main.cpp b/sample/sanitizer_sample/memcheck/illegal_read_and_write/main.cpp new file mode 100644 index 0000000000000000000000000000000000000000..44138b3d91aa7639b302cae9968560539e204dd5 --- /dev/null +++ b/sample/sanitizer_sample/memcheck/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/memcheck/out_of_bound/Makefile b/sample/sanitizer_sample/memcheck/out_of_bound/Makefile new file mode 100644 index 0000000000000000000000000000000000000000..22e06f70ff6828ddced8e5a0f9ce80145b1ea787 --- /dev/null +++ b/sample/sanitizer_sample/memcheck/out_of_bound/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: out_of_bound_kernel.o main.o out_of_bound.fatbin + +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 $^ + +out_of_bound.fatbin: out_of_bound_kernel.o main.o + $(COMPILER) $(LINK_FLAG) $(DAV_FLAG) -o $@ $^ ${LINK_LIBS} + +.PHONY: clean +clean: + rm *.o out_of_bound.fatbin diff --git a/sample/sanitizer_sample/memcheck/out_of_bound/main.cpp b/sample/sanitizer_sample/memcheck/out_of_bound/main.cpp new file mode 100644 index 0000000000000000000000000000000000000000..53da5b54a415ffc7eaa6a76cc3413d6d98aa1201 --- /dev/null +++ b/sample/sanitizer_sample/memcheck/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/memcheck/out_of_bound/out_of_bound_kernel.cpp b/sample/sanitizer_sample/memcheck/out_of_bound/out_of_bound_kernel.cpp new file mode 100644 index 0000000000000000000000000000000000000000..cb54dc3221598680d7322b4202fcb80cd5cd01af --- /dev/null +++ b/sample/sanitizer_sample/memcheck/out_of_bound/out_of_bound_kernel.cpp @@ -0,0 +1,27 @@ +#include "kernel_operator.h" +#include "acl/acl.h" +using namespace AscendC; + +constexpr int32_t BYTESIZE = 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); + LocalTensor xLm = xlm.Get(); + 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) +{ + out_of_bound_kernel<<>>(gm); +}