已合并
修改 helloworld 样例/资料 #6507
louyuxuan创建于 21 天前
修改 helloworld 样例/资料 #6507
已合并
共 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 parameters | 32 | + // 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 sample | 35 | + // 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 complete | 38 | // 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 | ```bash | 73 | ```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 | ./demo | 75 | ./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 | ```bash | 62 | ```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 | ```bash | 61 | ```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() | |||
| 28 | int32_t main(int argc, char const* argv[]) | 27 | int32_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 | } |