已合并
【文档修复】快速入门文档修复 #3587
YANXI_ZHAO创建于 6月25日
【文档修复】快速入门文档修复 #3587
已合并
YANXI_ZHAO创建于 6月25日
2 个文件变更+31-25
@@ -117,38 +117,40 @@
117 117 
118 - **Host端代码实现**118 - **Host端代码实现**
119 119 
120- Host端通过`<<<>>>`语法糖调用Device端核函数,示例代码如下:120+ Host端通过`<<<>>>`语法糖调用Device端核函数,示例代码片段如下:
121 ```cpp121 ```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);
atomgit-bot
atomgit-botatomgit-bot6月25日

🔵 Low Priority

变更在第 128-130 行新增了 aclrtCreateStream(&stream) 调用,但在整个代码片段中(至第 156 行的 } 结束)未出现对应的 aclrtDestroyStream(stream) 调用。该文档作为入门教程,用户会直接复制示例代码;缺少销毁操作会示范一个资源泄漏模式,在生产代码中持续创建 stream 而不销毁将导致运行时资源耗尽。

触发条件:用户按照此示例代码在循环或多次调用场景中创建 stream。

失败模式:stream 句柄泄漏,最终可能导致 stream 资源耗尽。

建议:在示例代码末尾(aclrtMemcpy 结果回传之后)增加 stream 销毁调用:aclrtDestroyStream(stream);,或者在 ... 省略处附近添加注释提示用户需要销毁 stream。

likedislike
127 131
128 // Allocate host and device memory, and copy input data from host to device132 // 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 API146+ // 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 complete149 // Wait for the add_custom kernel to complete
148- aclrtSynchronizeDevice();150+ aclrtSynchronizeStream(stream);
149 151 
150 // Copy the result from device memory to host memory152 // 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 ```cpp26 ```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 demo47 bisheng hello_world.asc --npu-arch=dav-2201 -o demo
44 ./demo48 ./demo
45 ```49 ```