已合并
【文档修复】快速入门文档修复 #3587
YANXI_ZHAO创建于 6月25日
【文档修复】快速入门文档修复 #3587
已合并
共 2 个文件变更+31-25
| @@ -117,38 +117,40 @@ | |||
| 117 | 117 | ||
| 118 | - **Host端代码实现**: | 118 | - **Host端代码实现**: |
| 119 | 119 | ||
| 120 | - Host端通过`<<<>>>`语法糖调用Device端核函数,示例代码如下: | 120 | + Host端通过`<<<>>>`语法糖调用Device端核函数,示例代码片段如下: |
| 121 | ```cpp | 121 | ```cpp |
| 122 | int32_t main(int argc, char const *argv[]) | 122 | int32_t main(int argc, char const *argv[]) |
| 123 | { | 123 | { |
| 124 | ... | 124 | ... |
| 125 | - constexpr uint32_t num_blocks = 8; | 125 | + constexpr uint32_t numBlocks = 8; |
| 126 | ... | 126 | ... |
| 127 | + | ||
| 128 | + // Create runtime stream by aclrtCreateStream API | ||
| 129 | + aclrtStream stream = nullptr; | ||
| 130 | + aclrtCreateStream(&stream); | ||
| 127 | 131 | ||
| 128 | // Allocate host and device memory, and copy input data from host to device | 132 | // Allocate host and device memory, and copy input data from host to device |
| 129 | - aclrtMallocHost((void **)(&z_host), total_byte_size); | 133 | + aclrtMalloc((void**)&xDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST); |
| 130 | - aclrtMalloc((void **)&x_device, total_byte_size, ACL_MEM_MALLOC_HUGE_FIRST); | 134 | + aclrtMalloc((void**)&yDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST); |
| 131 | - aclrtMalloc((void **)&y_device, total_byte_size, ACL_MEM_MALLOC_HUGE_FIRST); | 135 | + aclrtMalloc((void**)&zDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST); |
| 132 | - aclrtMalloc((void **)&z_device, total_byte_size, ACL_MEM_MALLOC_HUGE_FIRST); | 136 | + aclrtMallocHost((void**)&zHost, totalByteSize); |
| 133 | - aclrtMemcpy((uint8_t*)x_device, total_byte_size, x_host, total_byte_size, ACL_MEMCPY_HOST_TO_DEVICE); | 137 | + |
| 134 | - aclrtMemcpy((uint8_t*)y_device, total_byte_size, y_host, total_byte_size, ACL_MEMCPY_HOST_TO_DEVICE); | 138 | + aclrtMemcpy(xDevice, totalByteSize, x.data(), totalByteSize, ACL_MEMCPY_HOST_TO_DEVICE); |
| 139 | + aclrtMemcpy(yDevice, totalByteSize, y.data(), totalByteSize, ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 135 | 140 | ||
| 136 | // Launch kernel <<<numBlocks, dynUBufSize, stream>>> | 141 | // Launch kernel <<<numBlocks, dynUBufSize, stream>>> |
| 137 | // numBlocks : Number of blocks. Default to 8 in this example. | 142 | // numBlocks : Number of blocks. Default to 8 in this example. |
| 138 | // dynUBufSize : Dynamic unified buffer size. Default to 0 in this example. | 143 | // dynUBufSize : Dynamic unified buffer size. Default to 0 in this example. |
| 139 | - // nullptr : Runtime stream. Uses default stream in this example. | 144 | + // stream : Runtime stream. |
| 140 | 145 | ||
| 141 | - // Example One:Call add kernel which is implemented by SIMD C API | 146 | + // Example:Call add kernel which is implemented by SIMD C++ Basic API |
| 142 | - add_custom<<<num_blocks, 0>>>(xDevice, yDevice, zDevice); | 147 | + add_custom<blockLength><<<numBlocks, 0, stream>>>(xDevice, yDevice, zDevice); |
| 143 | - | ||
| 144 | - // Example Two:Call add kernel which is implemented by SIMD C++ Basic API | ||
| 145 | - // add_custom<blockLength><<<numBlocks, 0, stream>>>(xDevice, yDevice, zDevice); | ||
| 146 | 148 | ||
| 147 | // Wait for the add_custom kernel to complete | 149 | // Wait for the add_custom kernel to complete |
| 148 | - aclrtSynchronizeDevice(); | 150 | + aclrtSynchronizeStream(stream); |
| 149 | 151 | ||
| 150 | // Copy the result from device memory to host memory | 152 | // Copy the result from device memory to host memory |
| 151 | - aclrtMemcpy(z_host, total_byte_size, (uint8_t*)z_device, total_byte_size, ACL_MEMCPY_DEVICE_TO_HOST); | 153 | + aclrtMemcpy(zHost, totalByteSize, zDevice, totalByteSize, ACL_MEMCPY_DEVICE_TO_HOST); |
| 152 | ... | 154 | ... |
| 153 | } | 155 | } |
| 154 | ``` | 156 | ``` |
| @@ -156,7 +158,7 @@ | |||
| 156 | - **算子编译与运行**: | 158 | - **算子编译与运行**: |
| 157 | 159 | ||
| 158 | CMake配置文件示例: | 160 | CMake配置文件示例: |
| 159 | - ``` | 161 | + ```cmake |
| 160 | cmake_minimum_required(VERSION 3.16) | 162 | cmake_minimum_required(VERSION 3.16) |
| 161 | 163 | ||
| 162 | find_package(ASC REQUIRED) | 164 | find_package(ASC REQUIRED) |
| @@ -178,7 +180,7 @@ | |||
| 178 | ) | 180 | ) |
| 179 | ``` | 181 | ``` |
| 180 | 编译与执行示例: | 182 | 编译与执行示例: |
| 181 | - ``` | 183 | + ```bash |
| 182 | mkdir -p build && cd build; # 创建并进入build目录 | 184 | mkdir -p build && cd build; # 创建并进入build目录 |
| 183 | cmake ..;make -j; # 编译工程 | 185 | cmake ..;make -j; # 编译工程 |
| 184 | ./c_api_add_example # 运行样例 | 186 | ./c_api_add_example # 运行样例 |
| @@ -22,24 +22,28 @@ | |||
| 22 | 22 | ||
| 23 | - **Host端代码实现**: | 23 | - **Host端代码实现**: |
| 24 | 24 | ||
| 25 | - Host端通过<<<>>>语法糖调用Device端代码。 | 25 | + Host端通过<<<>>>语法糖调用Device端代码片段。 |
| 26 | ```cpp | 26 | ```cpp |
| 27 | int main(int argc, char const* argv[]) | 27 | int main(int argc, char const* argv[]) |
| 28 | { | 28 | { |
| 29 | - aclrtSetDevice(0); | 29 | + ... |
| 30 | + constexpr uint32_t numBlocks = 8; | ||
| 31 | + ... | ||
| 32 | + | ||
| 33 | + aclrtStream stream = nullptr; | ||
| 34 | + aclrtCreateStream(&stream); | ||
| 30 | // Launch kernel <<<numBlocks, dynUBufSize, stream>>> | 35 | // Launch kernel <<<numBlocks, dynUBufSize, stream>>> |
| 31 | // numBlocks : Number of blocks. Default to 8 in this example. | 36 | // numBlocks : Number of blocks. Default to 8 in this example. |
| 32 | // dynUBufSize : Dynamic unified buffer size. Default to 0 in this example. | 37 | // dynUBufSize : Dynamic unified buffer size. Default to 0 in this example. |
| 33 | - // nullptr : Runtime stream. Uses default stream in this example. | 38 | + // stream : Runtime stream. Uses stream created by aclrtCreateStream API in this example. |
| 34 | - hello_world<<<8, 0, nullptr>>>(); | 39 | + hello_world<<<numBlocks, 0, stream>>>(); |
| 35 | - aclrtSynchronizeDevice(); | 40 | + ... |
| 36 | - return 0; | ||
| 37 | } | 41 | } |
| 38 | ``` | 42 | ``` |
| 39 | 43 | ||
| 40 | - **算子编译与运行**: | 44 | - **算子编译与运行**: |
| 41 | 45 | ||
| 42 | - ``` | 46 | + ```bash |
| 43 | bisheng hello_world.asc --npu-arch=dav-2201 -o demo | 47 | bisheng hello_world.asc --npu-arch=dav-2201 -o demo |
| 44 | ./demo | 48 | ./demo |
| 45 | ``` | 49 | ``` |
🔵 Low Priority
变更在第 128-130 行新增了
aclrtCreateStream(&stream)调用,但在整个代码片段中(至第 156 行的}结束)未出现对应的aclrtDestroyStream(stream)调用。该文档作为入门教程,用户会直接复制示例代码;缺少销毁操作会示范一个资源泄漏模式,在生产代码中持续创建 stream 而不销毁将导致运行时资源耗尽。触发条件:用户按照此示例代码在循环或多次调用场景中创建 stream。
失败模式:stream 句柄泄漏,最终可能导致 stream 资源耗尽。
建议:在示例代码末尾(
aclrtMemcpy结果回传之后)增加 stream 销毁调用:aclrtDestroyStream(stream);,或者在...省略处附近添加注释提示用户需要销毁 stream。