已合并
修改 helloworld 样例/资料 #6507
louyuxuan创建于 21 天前
修改 helloworld 样例/资料 #6507
已合并
louyuxuan创建于 21 天前
共 4 个文件变更+48-35
@@ -29,13 +29,12 @@
29 int32_t main(int argc, char const *argv[])29 int32_t main(int argc, char const *argv[])
30 {30 {
31 ...31 ...
32- // 4. Configure kernel launch parameters32+ // Launch kernel <<<grid_dim, block_dim, dynamic_memory_size, stream>>>
33- constexpr uint32_t blocks_per_grid = 2; // Number of thread blocks (Grid size)33+ // grid_dim : Number of thread blocks (Grid size)
34- constexpr uint32_t threads_per_block = 32; // Number of threads per block (Block size)34+ // block_dim : Number of threads per block (Block size)
35- constexpr uint32_t dyn_ubuf_size = 0; // No dynamic memory required in this sample35+ // dynamic_memory_size : Dynamic unified buffer size. Default to 0 in this example
36- 36+ // stream : Runtime stream, pass nullptr to use default stream
37- // 5. Launch kernel <<<grid_dim, block_dim, dynamic_memory_size, stream>>>37+ hello_world<<<2, 32, 0, nullptr>>>();
38- hello_world<<<blocks_per_grid, threads_per_block, dyn_ubuf_size, stream>>>();
39 // Wait for kernel execution to complete38 // Wait for kernel execution to complete
40 ...39 ...
41 }40 }
@@ -43,8 +42,36 @@
43 42 
44- **算子编译与运行**:43- **算子编译与运行**:
45 44 
45+ 本样例完整代码(hello_world.asc)如下:
46+ 
47+ ```cpp
48+ #include "acl/acl.h"
49+ #include "simt_api/asc_simt.h"
50+ #include "utils/debug/asc_printf.h"
51+ 
52+ __global__ void hello_world()
53+ {
54+ if (threadIdx.x < 3) {
55+ printf("[blockIdx (%u/%u)][threadIdx (%u/%u)]: Hello World!\n", blockIdx.x, gridDim.x, threadIdx.x, blockDim.x);
56+ }
57+ }
58+ 
59+ int32_t main(int argc, char const* argv[])
60+ {
61+ aclInit(nullptr);
62+ aclrtSetDevice(0);
63+ hello_world<<<2, 32, 0, nullptr>>>();
64+ aclrtSynchronizeDevice();
65+ aclrtResetDevice(0);
66+ aclFinalize();
67+ return 0;
68+ }
69+ ```
70+ 
71+ 运行如下命令进行编译运行:
72+ 
46 ```bash73 ```bash
47- bisheng hello_world.asc --npu-arch=dav-3510 -o demo --enable-simt 74+ bisheng hello_world.asc --npu-arch=dav-3510 -o demo --enable-simt
48 ./demo75 ./demo
49 ```76 ```
50 77 
@@ -56,15 +56,16 @@
56 | `CMAKE_ASC_ARCHITECTURES` | `dav-3510` | NPU 架构:本样例仅支持 dav-3510(Ascend 950PR/Ascend 950DT) |56 | `CMAKE_ASC_ARCHITECTURES` | `dav-3510` | NPU 架构:本样例仅支持 dav-3510(Ascend 950PR/Ascend 950DT) |
57 57 
58- 执行结果58- 执行结果
59+ 
59 执行结果如下,说明执行成功。60 执行结果如下,说明执行成功。
60 61 
61 ```bash62 ```bash
62- [blockIdx (0/2)][threadIdx (2/32)]: Hello World!
63- [blockIdx (0/2)][threadIdx (1/32)]: Hello World!
64 [blockIdx (0/2)][threadIdx (0/32)]: Hello World!63 [blockIdx (0/2)][threadIdx (0/32)]: Hello World!
65- [blockIdx (1/2)][threadIdx (2/32)]: Hello World!64+ [blockIdx (0/2)][threadIdx (1/32)]: Hello World!
66- [blockIdx (1/2)][threadIdx (1/32)]: Hello World!65+ [blockIdx (0/2)][threadIdx (2/32)]: Hello World!
67 [blockIdx (1/2)][threadIdx (0/32)]: Hello World!66 [blockIdx (1/2)][threadIdx (0/32)]: Hello World!
67+ [blockIdx (1/2)][threadIdx (1/32)]: Hello World!
68+ [blockIdx (1/2)][threadIdx (2/32)]: Hello World!
68 ```69 ```
69 70 
70 输出共6行,每行对应一个线程的打印结果。`[blockIdx (X/2)]` 表示当前为第X个线程块(共2个线程块),`[threadIdx (Y/32)]` 表示线程块内第Y个线程,说明核函数已成功在AIV Core上执行。71 输出共6行,每行对应一个线程的打印结果。`[blockIdx (X/2)]` 表示当前为第X个线程块(共2个线程块),`[threadIdx (Y/32)]` 表示线程块内第Y个线程,说明核函数已成功在AIV Core上执行。
@@ -59,12 +59,12 @@ Execute the following steps in the root directory of this example to compile and
59 The execution result is as follows, indicating successful execution.59 The execution result is as follows, indicating successful execution.
60 60 
61 ```bash61 ```bash
62- [blockIdx (0/2)][threadIdx (2/32)]: Hello World!
63- [blockIdx (0/2)][threadIdx (1/32)]: Hello World!
64 [blockIdx (0/2)][threadIdx (0/32)]: Hello World!62 [blockIdx (0/2)][threadIdx (0/32)]: Hello World!
65- [blockIdx (1/2)][threadIdx (2/32)]: Hello World!63+ [blockIdx (0/2)][threadIdx (1/32)]: Hello World!
66- [blockIdx (1/2)][threadIdx (1/32)]: Hello World!64+ [blockIdx (0/2)][threadIdx (2/32)]: Hello World!
67 [blockIdx (1/2)][threadIdx (0/32)]: Hello World!65 [blockIdx (1/2)][threadIdx (0/32)]: Hello World!
66+ [blockIdx (1/2)][threadIdx (1/32)]: Hello World!
67+ [blockIdx (1/2)][threadIdx (2/32)]: Hello World!
68 ```68 ```
69 69 
70 The output consists of 6 lines, each corresponding to the print result of one thread. `[blockIdx (X/2)]` indicates the X-th thread block (out of 2 thread blocks in total), and `[threadIdx (Y/32)]` indicates the Y-th thread within the thread block. This shows that the kernel function has executed successfully on the AIV Core.70 The output consists of 6 lines, each corresponding to the print result of one thread. `[blockIdx (X/2)]` indicates the X-th thread block (out of 2 thread blocks in total), and `[threadIdx (Y/32)]` indicates the Y-th thread within the thread block. This shows that the kernel function has executed successfully on the AIV Core.
@@ -13,7 +13,6 @@
13 * \brief Hello World SIMT sample.13 * \brief Hello World SIMT sample.
14 */14 */
15 15 
16-#include <cstdio>
17#include "acl/acl.h"16#include "acl/acl.h"
18#include "simt_api/asc_simt.h"17#include "simt_api/asc_simt.h"
19#include "utils/debug/asc_printf.h"18#include "utils/debug/asc_printf.h"
@@ -28,24 +27,10 @@ __global__ void hello_world()
28int32_t main(int argc, char const* argv[])27int32_t main(int argc, char const* argv[])
29{28{
30 aclInit(nullptr);29 aclInit(nullptr);
31- int32_t deviceId = 0;30+ aclrtSetDevice(0);
32- aclrtSetDevice(deviceId);31+ hello_world<<<2, 32, 0, nullptr>>>();
33- aclrtStream stream = nullptr;32+ aclrtSynchronizeDevice();
34- aclrtCreateStream(&stream);33+ aclrtResetDevice(0);
35- 
36- constexpr uint32_t blocks_per_grid = 2;
37- constexpr uint32_t threads_per_block = 32;
38- constexpr uint32_t dyn_ubuf_size = 0;
39- hello_world<<<blocks_per_grid, threads_per_block, dyn_ubuf_size, stream>>>();
40- aclrtSynchronizeStream(stream);
41- 
42- const char* err = aclGetRecentErrMsg();
43- if (err != nullptr) {
44- fprintf(stderr, "%s\n", err);
45- }
46- 
47- aclrtDestroyStream(stream);
48- aclrtResetDevice(deviceId);
49 aclFinalize();34 aclFinalize();
50 return 0;35 return 0;
51}36}