已合并
docs: SIMT文档链接锚点修正与文案优化 #3599
qixingkai创建于 6月25日
docs: SIMT文档链接锚点修正与文案优化 #3599
已合并
共 17 个文件变更+27-26
| @@ -85,7 +85,7 @@ | |||
| 85 | 85 | ||
| 86 | - [性能调优](编程指南/调试调优/性能调优.md) | 86 | - [性能调优](编程指南/调试调优/性能调优.md) |
| 87 | 87 | ||
| 88 | - - [高级编程](./编程指南/高级编程/硬件实现/硬件实现.md) | 88 | + - [高级编程](./编程指南/高级编程/高级编程.md) |
| 89 | - [硬件实现](./编程指南/高级编程/硬件实现/硬件实现.md) | 89 | - [硬件实现](./编程指南/高级编程/硬件实现/硬件实现.md) |
| 90 | - [基本架构](./编程指南/高级编程/硬件实现/基本架构.md) | 90 | - [基本架构](./编程指南/高级编程/硬件实现/基本架构.md) |
| 91 | - [架构规格](./编程指南/高级编程/硬件实现/架构规格/架构规格.md) | 91 | - [架构规格](./编程指南/高级编程/硬件实现/架构规格/架构规格.md) |
| @@ -306,7 +306,7 @@ | |||
| 306 | 306 | ||
| 307 | - [执行控制](算子实践参考/SIMD与SIMT混合算子性能优化/执行控制/线程块数量配置优化.md) | 307 | - [执行控制](算子实践参考/SIMD与SIMT混合算子性能优化/执行控制/线程块数量配置优化.md) |
| 308 | - [线程块数量配置优化](算子实践参考/SIMD与SIMT混合算子性能优化/执行控制/线程块数量配置优化.md) | 308 | - [线程块数量配置优化](算子实践参考/SIMD与SIMT混合算子性能优化/执行控制/线程块数量配置优化.md) |
| 309 | - - [SIMT算子性能优化](算子实践参考/SIMT算子性能优化/内存访问/访存合并.md) | 309 | + - [SIMT算子性能优化](算子实践参考/SIMT算子性能优化/SIMT算子性能优化.md) |
| 310 | - [内存访问](算子实践参考/SIMT算子性能优化/内存访问/访存合并.md) | 310 | - [内存访问](算子实践参考/SIMT算子性能优化/内存访问/访存合并.md) |
| 311 | - [访存合并](算子实践参考/SIMT算子性能优化/内存访问/访存合并.md) | 311 | - [访存合并](算子实践参考/SIMT算子性能优化/内存访问/访存合并.md) |
| 312 | - [避免UB的Bank冲突](算子实践参考/SIMT算子性能优化/内存访问/避免UB的Bank冲突.md) | 312 | - [避免UB的Bank冲突](算子实践参考/SIMT算子性能优化/内存访问/避免UB的Bank冲突.md) |
| @@ -17,8 +17,8 @@ | |||
| 17 | - **算子设计**： | 17 | - **算子设计**： |
| 18 | 18 | ||
| 19 | - **Device端核函数编程接口** | 19 | - **Device端核函数编程接口** |
| 20 | - - 核函数定义：通过 [\_\_global\_\_](../../../编程指南/语言扩展层/SIMT-BuiltIn关键字.md) 修饰符声明。 | 20 | + - 核函数定义：通过 [\_\_global\_\_](../../../编程指南/语言扩展层/SIMT-BuiltIn关键字.md)修饰符声明。 |
| 21 | - - 数据划分：使用内置关键字 [threadIdx、blockIdx、blockDim](../../../编程指南/语言扩展层/SIMT-BuiltIn关键字.md) 计算线程索引，并为每个线程分配需要处理的数据元素。 | 21 | + - 数据划分：使用内置变量[threadIdx、blockIdx、blockDim](../../../编程指南/语言扩展层/SIMT-BuiltIn关键字.md)计算线程索引，并为每个线程分配需要处理的数据元素。 |
| 22 | - 数据搬入：无需额外接口，直接通过指针访问即可。 | 22 | - 数据搬入：无需额外接口，直接通过指针访问即可。 |
| 23 | - 数据计算：根据`index`中的索引值及操作符`[]`读取输入数据。 | 23 | - 数据计算：根据`index`中的索引值及操作符`[]`读取输入数据。 |
| 24 | - 数据搬出：无需额外接口，直接通过指针访问即可。 | 24 | - 数据搬出：无需额外接口，直接通过指针访问即可。 |
| @@ -69,6 +69,6 @@ struct alignas(4) Align4 { | |||
| 69 | | 16B | 32983.77 (alignof=16) | 38353.60 (alignof=4) | **1.16x** | | 69 | | 16B | 32983.77 (alignof=16) | 38353.60 (alignof=4) | **1.16x** | |
| 70 | | 32B | 60901.24 (alignof=32) | 77790.93 (alignof=4) | **1.28x** | | 70 | | 32B | 60901.24 (alignof=32) | 77790.93 (alignof=4) | **1.28x** | |
| 71 | 71 | ||
| 72 | -测试结果表明：**在数据总量不变的时候，相同大小的结构体，类型对齐可以提升搬运效率**。以4B结构体为例，随着结构体类型对齐位宽从1B对齐提升至4B，`非对齐Task Duration/对齐Task Duration`的耗时倍数逐步下降。此外，同大小的4B结构体，字段组合形式、有无填充位，都不会对搬运效率产生影响。 | 72 | +测试结果表明：**在数据总量不变的时候，相同大小的结构体，类型对齐可以提升搬运效率**。以4B结构体为例，随着结构体类型对齐位宽从1B对齐提升至4B，`非对齐Task Duration/对齐Task Duration`逐步下降，说明非对齐与对齐的耗时差距逐步缩小。此外，同大小的4B结构体，字段组合形式、有无填充位，都不会对搬运效率产生影响。 |
| 73 | 73 | ||
| 74 | 【总结】SIMT算子使用结构体搬运数据时，结构体类型对齐能让访存指令以更大位宽合并完成搬运，降低搬运耗时。Ascend 950PR/Ascend 950DT支持1B、2B、4B、8B位宽的访存指令，因此建议使用时在不超过结构体大小的情况下，尽量选择更高的位宽。 | 74 | 【总结】SIMT算子使用结构体搬运数据时，结构体类型对齐能让访存指令以更大位宽合并完成搬运，降低搬运耗时。Ascend 950PR/Ascend 950DT支持1B、2B、4B、8B位宽的访存指令，因此建议使用时在不超过结构体大小的情况下，尽量选择更高的位宽。 |
| @@ -105,7 +105,7 @@ gather_strided_kernel<<<64, MAX_THREAD_NUM, 0, stream>>>( | |||
| 105 | input_total_length, index_total_length); | 105 | input_total_length, index_total_length); |
| 106 | ``` | 106 | ``` |
| 107 | 107 | ||
| 108 | -上述实现对小数据量场景也使用64个物理核，每核只处理`16384 / 64 = 256`个元素，此时计算本身耗时很短，64核启动带来的调度固定开销在总耗时中占比升高。该实现在小数据量场景下的性能数据如下： | 108 | +上述实现对小数据量场景也使用64个物理核，每核只处理`16384 / 64 = 256`个元素，此时计算本身耗时很短，启动64个线程块带来的调度固定开销在总耗时中占比升高。该实现在小数据量场景下的性能数据如下： |
| 109 | 109 | ||
| 110 | | 线程块数量 | 每个线程块的线程数 | Task Duration（us） | | 110 | | 线程块数量 | 每个线程块的线程数 | Task Duration（us） | |
| 111 | | :-------: | :---------------: | :---------------: | | 111 | | :-------: | :---------------: | :---------------: | |
| @@ -43,7 +43,7 @@ kernel_name<<<blocks_per_grid, threads_per_block, dyn_ubuf_size, stream>>>(argum | |||
| 43 | 43 | ||
| 44 | - blocks\_per\_grid：int或[dim3](https://gitcode.com/cann/asc-devkit/blob/master/docs/guide/编程指南/语言扩展层/SIMT-BuiltIn关键字.md#内置结构体)类型，用于指定网格（grid）的维度与规模，blocks\_per\_grid.x \* blocks\_per\_grid.y \* blocks\_per\_grid.z等于启动的线程块总数，不得大于65535。 | 44 | - blocks\_per\_grid：int或[dim3](https://gitcode.com/cann/asc-devkit/blob/master/docs/guide/编程指南/语言扩展层/SIMT-BuiltIn关键字.md#内置结构体)类型，用于指定网格（grid）的维度与规模，blocks\_per\_grid.x \* blocks\_per\_grid.y \* blocks\_per\_grid.z等于启动的线程块总数，不得大于65535。 |
| 45 | - threads\_per\_block：int或[dim3](https://gitcode.com/cann/asc-devkit/blob/master/docs/guide/编程指南/语言扩展层/SIMT-BuiltIn关键字.md#内置结构体)类型，用于指定每个线程块（block）的维度与规模，threads\_per\_block.x \* threads\_per\_block.y \* threads\_per\_block.z等于每个线程块包含的线程数，需要小于等于\_\_launch\_bounds\_\_配置。 | 45 | - threads\_per\_block：int或[dim3](https://gitcode.com/cann/asc-devkit/blob/master/docs/guide/编程指南/语言扩展层/SIMT-BuiltIn关键字.md#内置结构体)类型，用于指定每个线程块（block）的维度与规模，threads\_per\_block.x \* threads\_per\_block.y \* threads\_per\_block.z等于每个线程块包含的线程数，需要小于等于\_\_launch\_bounds\_\_配置。 |
| 46 | -- dyn\_ubuf\_size：size_t类型，该参数指定除静态分配的内存外，本次调用为每个线程块动态分配的共享内存字节数，单位为bytes，默认设置为0。具体用法请参考[共享内存](./内存层级.md#共享内存unified-buffer)中的“动态申请”方式； | 46 | +- <a name="li12421238101815"></a>dyn\_ubuf\_size：size_t类型，该参数指定除静态分配的内存外，本次调用为每个线程块动态分配的共享内存字节数，单位为bytes，默认设置为0。具体用法请参考[共享内存](./内存层级.md#section66329146410)中的“动态申请”方式； |
| 47 | - stream：aclrtStream类型指针，指定关联的流，用于维护异步操作的执行顺序。 | 47 | - stream：aclrtStream类型指针，指定关联的流，用于维护异步操作的执行顺序。 |
| 48 | 48 | ||
| 49 | ## Grid与Thread索引内置变量 | 49 | ## Grid与Thread索引内置变量 |
| @@ -52,7 +52,7 @@ kernel_name<<<blocks_per_grid, threads_per_block, dyn_ubuf_size, stream>>>(argum | |||
| 52 | 52 | ||
| 53 | 常用的内置索引变量说明如下： | 53 | 常用的内置索引变量说明如下： |
| 54 | - threadIdx：获取当前线程在其所属线程块内的索引。threadIdx.x，threadIdx.y，threadIdx.z分别表示当前线程在3个维度的索引，threadIdx.x的范围为[0, blockDim.x)，threadIdx.y的范围为[0, blockDim.y)，threadIdx.z的范围为[0, blockDim.z)。 | 54 | - threadIdx：获取当前线程在其所属线程块内的索引。threadIdx.x，threadIdx.y，threadIdx.z分别表示当前线程在3个维度的索引，threadIdx.x的范围为[0, blockDim.x)，threadIdx.y的范围为[0, blockDim.y)，threadIdx.z的范围为[0, blockDim.z)。 |
| 55 | -- blockDim：获取线程块中配置的三维层次结构，即启动VF时配置的dim3结构体实例值。blockDim.x，blockDim.y，blockDim.z分别表示线程块中三个维度的线程数。在内核启动的执行配置中指定，对应核函数配置中的threads_per_block参数。 | 55 | +- blockDim：获取线程块中配置的三维层次结构，即核函数启动时配置的dim3结构体实例值。blockDim.x，blockDim.y，blockDim.z分别表示线程块中三个维度的线程数。在核函数启动的执行配置中指定，对应核函数配置中的threads_per_block参数。 |
| 56 | - blockIdx: 获取当前线程块在其所属网格中的索引，表示当前线程所在的线程块在整个网格中的位置坐标。 | 56 | - blockIdx: 获取当前线程块在其所属网格中的索引，表示当前线程所在的线程块在整个网格中的位置坐标。 |
| 57 | - gridDim： 表示整个计算任务在各个维度上分别由多少个线程块构成。在内核启动的执行配置中指定，对应核函数配置中的blocks_per_grid参数。各个维度上线程块关系需满足gridDim.x * gridDim.y * gridDim.z <= 65535。 | 57 | - gridDim： 表示整个计算任务在各个维度上分别由多少个线程块构成。在内核启动的执行配置中指定，对应核函数配置中的blocks_per_grid参数。各个维度上线程块关系需满足gridDim.x * gridDim.y * gridDim.z <= 65535。 |
| 58 | 58 | ||
| @@ -95,16 +95,16 @@ | |||
| 95 | - 对于二维线程块，其线程ID为二维结构，其计算公式为： | 95 | - 对于二维线程块，其线程ID为二维结构，其计算公式为： |
| 96 | 96 | ||
| 97 | ``` | 97 | ``` |
| 98 | - thread_id_x = blockIdx.x * blockDim.x + threadIdx.x； | 98 | + thread_id_x = blockIdx.x * blockDim.x + threadIdx.x; |
| 99 | - thread_id_y = blockIdx.y * blockDim.y + threadIdx.y； | 99 | + thread_id_y = blockIdx.y * blockDim.y + threadIdx.y; |
| 100 | ``` | 100 | ``` |
| 101 | 101 | ||
| 102 | - 对于三维线程块，其线程ID为三维结构，其计算公式为： | 102 | - 对于三维线程块，其线程ID为三维结构，其计算公式为： |
| 103 | 103 | ||
| 104 | ``` | 104 | ``` |
| 105 | - thread_id_x = blockIdx.x * blockDim.x + threadIdx.x； | 105 | + thread_id_x = blockIdx.x * blockDim.x + threadIdx.x; |
| 106 | - thread_id_y = blockIdx.y * blockDim.y + threadIdx.y； | 106 | + thread_id_y = blockIdx.y * blockDim.y + threadIdx.y; |
| 107 | - thread_id_z = blockIdx.z * blockDim.z + threadIdx.z； | 107 | + thread_id_z = blockIdx.z * blockDim.z + threadIdx.z; |
| 108 | ``` | 108 | ``` |
| 109 | 109 | ||
| 110 | 当前提供了以下仅在Device上可用的int类型的内置变量： | 110 | 当前提供了以下仅在Device上可用的int类型的内置变量： |
| @@ -10,4 +10,4 @@ SIMT编程基于AI Core的硬件能力实现，当前，SIMT语言扩展层支 | |||
| 10 | - [地址空间谓词函数](https://gitcode.com/cann/asc-devkit/blob/master/docs/api/SIMT-API/地址空间谓词函数/地址空间谓词函数.md)：提供判断输入指针所属地址空间的函数接口集合。 | 10 | - [地址空间谓词函数](https://gitcode.com/cann/asc-devkit/blob/master/docs/api/SIMT-API/地址空间谓词函数/地址空间谓词函数.md)：提供判断输入指针所属地址空间的函数接口集合。 |
| 11 | - [地址空间转换函数](https://gitcode.com/cann/asc-devkit/blob/master/docs/api/SIMT-API/地址空间转换函数/地址空间转换函数.md)：提供不同地址空间地址值与指针之间相互转换的函数接口集合。 | 11 | - [地址空间转换函数](https://gitcode.com/cann/asc-devkit/blob/master/docs/api/SIMT-API/地址空间转换函数/地址空间转换函数.md)：提供不同地址空间地址值与指针之间相互转换的函数接口集合。 |
| 12 | - [协作组](https://gitcode.com/cann/asc-devkit/blob/master/docs/api/SIMT-API/协作组/协作组.md)：提供更细粒度的线程间协作的API接口。 | 12 | - [协作组](https://gitcode.com/cann/asc-devkit/blob/master/docs/api/SIMT-API/协作组/协作组.md)：提供更细粒度的线程间协作的API接口。 |
| 13 | -- [调测接口](https://gitcode.com/cann/asc-devkit/blob/master/docs/api/Utils-API/调测接口/printf-290.md)：SIMT VF调试场景下使用的相关接口。 | 13 | +- [调测接口](https://gitcode.com/cann/asc-devkit/blob/master/docs/api/Utils-API/调测接口/调测接口.md)：SIMT调试场景下使用的相关接口。 |
| @@ -9,7 +9,7 @@ Ascend C支持SIMD与SIMT混合编程，实现向量级并行与线程级并行 | |||
| 9 | 9 | ||
| 10 | 如上图硬件架构所示，AIV核的计算单元包括SIMD和SIMT两种。针对这两种硬件计算单元，Ascend C抽象出了Vector Function的软件概念，用于表示在SIMT或SIMD硬件计算单元上执行的特定功能代码段。用户通过编写SIMD Vector Function或SIMT Vector Function来调用对应的执行单元完成计算任务，在算子核函数中调用不同类型的Vector Function，以实现SIMD/SIMT硬件单元的切换使用。 | 10 | 如上图硬件架构所示，AIV核的计算单元包括SIMD和SIMT两种。针对这两种硬件计算单元，Ascend C抽象出了Vector Function的软件概念，用于表示在SIMT或SIMD硬件计算单元上执行的特定功能代码段。用户通过编写SIMD Vector Function或SIMT Vector Function来调用对应的执行单元完成计算任务，在算子核函数中调用不同类型的Vector Function，以实现SIMD/SIMT硬件单元的切换使用。 |
| 11 | 11 | ||
| 12 | -如下图和示例代码所示，simt\_func表示在SIMT硬件单元上执行的代码段，simd\_func表示在SIMD硬件单元上执行的代码段，代码段属性通过[\_\_simt\_vf\_\_](核函数与VF函数.md#zh-cn_topic_0000002571578013_section1780955884616)、[\_\_simd\_vf\_\_](核函数与VF函数.md#zh-cn_topic_0000002571578013_section2067153312577)来标识。 | 12 | +如下图和示例代码所示，simt\_func表示在SIMT硬件单元上执行的代码段，simd\_func表示在SIMD硬件单元上执行的代码段，代码段属性通过[\_\_simt\_vf\_\_](核函数与VF函数.md#zh-cn_topic_0000002571578013_section1780955884616)、[\_\_simd\_vf\_\_](../../../语言扩展层/SIMD-BuiltIn关键字.md#section192521344610)来标识。 |
| 13 | 13 | ||
| 14 | **图 2** SIMD与SIMT混合编程VF示意图 | 14 | **图 2** SIMD与SIMT混合编程VF示意图 |
| 15 |  | 15 |  |
| @@ -20,7 +20,7 @@ SIMD与SIMT混合编程中涉及的函数类型如下表所示： | |||
| 20 | - simd vf函数：能被核函数或者`__aicore__`函数通过`asc_vf_call`接口调用。simd vf函数内只能调用`__simd_callee__`函数和`__callee__`函数。 | 20 | - simd vf函数：能被核函数或者`__aicore__`函数通过`asc_vf_call`接口调用。simd vf函数内只能调用`__simd_callee__`函数和`__callee__`函数。 |
| 21 | - simt vf函数：使用__simt_vf__标记，能被核函数或者`__aicore__`函数通过`asc_vf_call`接口调用。simt vf函数内只能调用`__simt_callee__`函数和`__callee__`函数。 | 21 | - simt vf函数：使用__simt_vf__标记，能被核函数或者`__aicore__`函数通过`asc_vf_call`接口调用。simt vf函数内只能调用`__simt_callee__`函数和`__callee__`函数。 |
| 22 | - `__simd_callee__`子函数：simd vf函数内部的子函数，子函数可能有返回值或者通过引用传参。__simd_callee__函数内只能调用`__simd_callee__`函数和`__callee__`函数。 | 22 | - `__simd_callee__`子函数：simd vf函数内部的子函数，子函数可能有返回值或者通过引用传参。__simd_callee__函数内只能调用`__simd_callee__`函数和`__callee__`函数。 |
| 23 | -- `__simt_callee__`子函数：simt vf函内部的子函数，子函数可能有返回值或者通过引用传参。`__simt_callee__`函数内只能调用`__simt_callee__`函数和`__callee__`函数。 | 23 | +- `__simt_callee__`子函数：simt vf函数内部的子函数，子函数可能有返回值或者通过引用传参。`__simt_callee__`函数内只能调用`__simt_callee__`函数和`__callee__`函数。 |
| 24 | - `__callee__`公共函数：可供VF内部调用的公共子函数。`__callee__`函数内只能调用`__callee__`函数。 | 24 | - `__callee__`公共函数：可供VF内部调用的公共子函数。`__callee__`函数内只能调用`__callee__`函数。 |
| 25 | 25 | ||
| 26 | **图 1** SIMD与SIMT混合编程函数调用层级 | 26 | **图 1** SIMD与SIMT混合编程函数调用层级 |
| @@ -0,0 +1 @@ | |||
| 1 | +# 高级AI-Core编程模型 | ||
| @@ -24,7 +24,7 @@ | |||
| 24 | 24 | ||
| 25 | - 算子功能: | 25 | - 算子功能: |
| 26 | 26 | ||
| 27 | - 本样例详细展示了在SIMT实现函数中使用`printf()`接口的实践方式，实现对算子执行过程中打印每个线程的变量信息。 | 27 | + 本样例详细展示了在SIMT核函数中使用`printf()`接口的实践方式，实现对算子执行过程中打印每个线程的变量信息。 |
| 28 | 28 | ||
| 29 | 29 | ||
| 30 | - 算子实现: | 30 | - 算子实现: |
| @@ -82,9 +82,9 @@ | |||
| 82 | - 配置日志环境变量（仅NPU模式） | 82 | - 配置日志环境变量（仅NPU模式） |
| 83 | 本样例支持通过环境变量控制日志输出行为（以下环境变量仅在NPU模式下生效）： | 83 | 本样例支持通过环境变量控制日志输出行为（以下环境变量仅在NPU模式下生效）： |
| 84 | ```bash | 84 | ```bash |
| 85 | - export ASCEND_PROCESS_LOG_PATH=./log // 日志落盘路径 | 85 | + export ASCEND_PROCESS_LOG_PATH=./log # 日志落盘路径 |
| 86 | - export ASCEND_SLOG_PRINT_TO_STDOUT=1 // 控制日志是否打印到屏幕（1：开启，0：关闭） | 86 | + export ASCEND_SLOG_PRINT_TO_STDOUT=1 # 控制日志是否打印到屏幕（1：开启，0：关闭） |
| 87 | - export ASCEND_GLOBAL_LOG_LEVEL=1 // 控制日志级别 | 87 | + export ASCEND_GLOBAL_LOG_LEVEL=1 # 控制日志级别 |
| 88 | ``` | 88 | ``` |
| 89 | 89 | ||
| 90 | > 注意： | 90 | > 注意： |
| @@ -80,7 +80,7 @@ struct Bucket { | |||
| 80 | </table> | 80 | </table> |
| 81 | 81 | ||
| 82 | - 算子实现： | 82 | - 算子实现： |
| 83 | - InsertHashTable算子的实现流程为，每个warp处理一个键值对，warp中线程0负责根据key的哈希值寻找可用桶，线程0-31负责将value存入桶中。线程0寻找可用桶时，通过开放寻址法解决哈希冲突问题，并使用asc_atomic_cas()接口解决多线程冲突问题，使用asc_threadfence()接口解决核间数据同步问题。 | 83 | + InsertHashTable算子的实现流程为，每个warp处理一个键值对，warp中线程0负责根据key的哈希值寻找可用桶，线程0-31负责将value存入桶中。线程0寻找可用桶时，通过开放寻址法解决哈希冲突问题，并使用`asc_atomic_cas()`接口解决多线程冲突问题，使用`asc_threadfence()`接口解决核间数据同步问题。 |
| 84 | 84 | ||
| 85 | - 调用实现： | 85 | - 调用实现： |
| 86 | 使用内核调用符<<<>>>调用核函数。 | 86 | 使用内核调用符<<<>>>调用核函数。 |
Mexamples/03_simt_api/03_best_practices/01_execution_conf_optimizations/grid_dim_config/README.md+1-1
| @@ -215,7 +215,7 @@ for (uint32_t i = global_idx; i < index_total_length; i += stride) { | |||
| 215 | 215 | ||
| 216 | 下图直观展示了上表中Task Duration随线程块数量的变化趋势： | 216 | 下图直观展示了上表中Task Duration随线程块数量的变化趋势： |
| 217 | 217 | ||
| 218 | -<img src="figures/fig2_case3_16384_curve.png" alt="Case 2 小数据量场景（16384） Task Duration 随线程块数量的变化" width="600" /> | 218 | +<img src="figures/fig2_case2_16384_curve.png" alt="Case 2 小数据量场景（16384） Task Duration 随线程块数量的变化" width="600" /> |
| 219 | 219 | ||
| 220 | 220 | ||
| 221 | **分析**： | 221 | **分析**： |
Rexamples/03_simt_api/03_best_practices/01_execution_conf_optimizations/grid_dim_config/figures/fig2_case3_16384_curve.png→examples/03_simt_api/03_best_practices/01_execution_conf_optimizations/grid_dim_config/figures/fig2_case2_16384_curve.png+0-0
| @@ -1,8 +1,8 @@ | |||
| 1 | # 样例列表 | 1 | # 样例列表 |
| 2 | | 目录名称 | 功能描述 | | 2 | | 目录名称 | 功能描述 | |
| 3 | | ------------------------------------- | ------------------------------------- | | 3 | | ------------------------------------- | ------------------------------------- | |
| 4 | -| [00_introduction](./00_introduction) | 基于Ascend C SIMT编程的简单的示例，通过Ascend C编程语言实现了自定义算子，分别给出对应的<<<>>>直调实现，适合初学者 | | 4 | +| [00_introduction](./00_introduction) | 基于Ascend C SIMT编程的简单的示例，通过Ascend C编程语言实现了自定义算子，分别给出对应的<<<>>>直调实现，适合初学者。 | |
| 5 | -| [01_utilities](./01_utilities) | 基于Ascend C SIMT的调试调优样例，通过printf、assert、clock等API介绍上板打印、打点、CPU孪生调试等系统工具使用方法，适用于调试阶段 | | 5 | +| [01_utilities](./01_utilities) | 基于Ascend C SIMT的调试调优样例，通过printf、assert、clock等API介绍上板打印、打点、CPU孪生调试等系统工具使用方法，适用于调试阶段。 | |
| 6 | | [02_features](./02_features) | 展示Ascend C SIMT编程的重点特性的样例，包含基础编程、内存资源管理、线程同步、原子操作等特性。 | | 6 | | [02_features](./02_features) | 展示Ascend C SIMT编程的重点特性的样例，包含基础编程、内存资源管理、线程同步、原子操作等特性。 | |
| 7 | | [03_best_practices](./03_best_practices) | 展示Ascend C SIMT编程的最佳实践。 | | 7 | | [03_best_practices](./03_best_practices) | 展示Ascend C SIMT编程的最佳实践。 | |
| 8 | 8 | ||
| @@ -1,8 +1,8 @@ | |||
| 1 | # Example List | 1 | # Example List |
| 2 | | Directory Name | Description | | 2 | | Directory Name | Description | |
| 3 | | ------------------------------------- | ------------------------------------- | | 3 | | ------------------------------------- | ------------------------------------- | |
| 4 | -| [00_introduction](./00_introduction) | Simple examples based on Ascend C SIMT programming, implementing custom operators through Ascend C programming language, providing corresponding <<<>>> direct invocation implementations, suitable for beginners | | 4 | +| [00_introduction](./00_introduction) | Simple examples based on Ascend C SIMT programming, implementing custom operators through Ascend C programming language, providing corresponding <<<>>> direct invocation implementations, suitable for beginners. | |
| 5 | -| [01_utilities](./01_utilities) | Debugging and tuning examples based on Ascend C SIMT, introducing system tool usage methods such as on-board printing, timing, and CPU twin debugging through printf, assert, clock and other APIs, suitable for debugging phase | | 5 | +| [01_utilities](./01_utilities) | Debugging and tuning examples based on Ascend C SIMT, introducing system tool usage methods such as on-board printing, timing, and CPU twin debugging through printf, assert, clock and other APIs, suitable for debugging phase. | |
| 6 | | [02_features](./02_features) | Examples demonstrating key features of Ascend C SIMT programming, including basic programming, memory resource management, thread synchronization, atomic operations, and other features. | | 6 | | [02_features](./02_features) | Examples demonstrating key features of Ascend C SIMT programming, including basic programming, memory resource management, thread synchronization, atomic operations, and other features. | |
| 7 | | [03_best_practices](./03_best_practices) | Demonstrating best practices for Ascend C SIMT programming. | | 7 | | [03_best_practices](./03_best_practices) | Demonstrating best practices for Ascend C SIMT programming. | |
| 8 | 8 | ||