From 6f1f89cc3bc413d8f2bbb5a19e707ada720f6069 Mon Sep 17 00:00:00 2001 From: zhengweifeng Date: Tue, 27 Feb 2024 15:28:50 +0800 Subject: [PATCH 1/2] =?UTF-8?q?=E5=A2=9E=E5=8A=A0=E9=9D=9E=E6=B3=95?= =?UTF-8?q?=E8=AF=BB=E5=86=99=E4=B8=8E=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 --- .../illegl_read_and_write/Makefile | 29 ++++++++++++++ .../illegl_read_and_write/kernel.cpp | 25 ++++++++++++ .../illegl_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 +++++++++++++++++++ 6 files changed, 191 insertions(+) create mode 100644 sample/sanitizer_sample/illegl_read_and_write/Makefile create mode 100644 sample/sanitizer_sample/illegl_read_and_write/kernel.cpp create mode 100644 sample/sanitizer_sample/illegl_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 diff --git a/sample/sanitizer_sample/illegl_read_and_write/Makefile b/sample/sanitizer_sample/illegl_read_and_write/Makefile new file mode 100644 index 0000000000..dd683bf26d --- /dev/null +++ b/sample/sanitizer_sample/illegl_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/illegl_read_and_write/kernel.cpp b/sample/sanitizer_sample/illegl_read_and_write/kernel.cpp new file mode 100644 index 0000000000..b1f36a3caa --- /dev/null +++ b/sample/sanitizer_sample/illegl_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 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 kernel_do(uint32_t blockDim, void *l2ctrl, void *stream, uint8_t *gm) +{ + kernel<<>>(gm); +} diff --git a/sample/sanitizer_sample/illegl_read_and_write/main.cpp b/sample/sanitizer_sample/illegl_read_and_write/main.cpp new file mode 100644 index 0000000000..7e947b41b8 --- /dev/null +++ b/sample/sanitizer_sample/illegl_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 raw_error_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; + raw_error_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 0000000000..dd683bf26d --- /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 0000000000..18bf128469 --- /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 0000000000..c78b6b0f8a --- /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 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; + raw_error_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 -- Gitee From dcda333533e235ecd2a92a464084bd69c4d56639 Mon Sep 17 00:00:00 2001 From: zhengweifeng Date: Tue, 27 Feb 2024 07:53:41 +0000 Subject: [PATCH 2/2] =?UTF-8?q?=E5=9B=9E=E9=80=80=20'Pull=20Request=20!1?= =?UTF-8?q?=20:=20=E5=A2=9E=E5=8A=A0=E9=9D=9E=E6=B3=95=E8=AF=BB=E5=86=99?= =?UTF-8?q?=E3=80=81=E5=A4=9A=E6=A0=B8=E8=B8=A9=E8=B8=8F=E6=A0=B7=E4=BE=8B?= =?UTF-8?q?'?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- .../illegl_read_and_write/Makefile | 29 -------------- .../illegl_read_and_write/kernel.cpp | 25 ------------ .../illegl_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 ------------------- 6 files changed, 191 deletions(-) delete mode 100644 sample/sanitizer_sample/illegl_read_and_write/Makefile delete mode 100644 sample/sanitizer_sample/illegl_read_and_write/kernel.cpp delete mode 100644 sample/sanitizer_sample/illegl_read_and_write/main.cpp delete mode 100644 sample/sanitizer_sample/out_of_bound/Makefile delete mode 100644 sample/sanitizer_sample/out_of_bound/kernel.cpp delete mode 100644 sample/sanitizer_sample/out_of_bound/main.cpp diff --git a/sample/sanitizer_sample/illegl_read_and_write/Makefile b/sample/sanitizer_sample/illegl_read_and_write/Makefile deleted file mode 100644 index dd683bf26d..0000000000 --- a/sample/sanitizer_sample/illegl_read_and_write/Makefile +++ /dev/null @@ -1,29 +0,0 @@ -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/illegl_read_and_write/kernel.cpp b/sample/sanitizer_sample/illegl_read_and_write/kernel.cpp deleted file mode 100644 index b1f36a3caa..0000000000 --- a/sample/sanitizer_sample/illegl_read_and_write/kernel.cpp +++ /dev/null @@ -1,25 +0,0 @@ -#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 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 kernel_do(uint32_t blockDim, void *l2ctrl, void *stream, uint8_t *gm) -{ - kernel<<>>(gm); -} diff --git a/sample/sanitizer_sample/illegl_read_and_write/main.cpp b/sample/sanitizer_sample/illegl_read_and_write/main.cpp deleted file mode 100644 index 7e947b41b8..0000000000 --- a/sample/sanitizer_sample/illegl_read_and_write/main.cpp +++ /dev/null @@ -1,39 +0,0 @@ -#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 *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; - raw_error_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 deleted file mode 100644 index dd683bf26d..0000000000 --- a/sample/sanitizer_sample/out_of_bound/Makefile +++ /dev/null @@ -1,29 +0,0 @@ -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 deleted file mode 100644 index 18bf128469..0000000000 --- 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); -} diff --git a/sample/sanitizer_sample/out_of_bound/main.cpp b/sample/sanitizer_sample/out_of_bound/main.cpp deleted file mode 100644 index c78b6b0f8a..0000000000 --- a/sample/sanitizer_sample/out_of_bound/main.cpp +++ /dev/null @@ -1,39 +0,0 @@ -#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 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; - raw_error_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 -- Gitee