写在最前面

本章节开始，将进行Ascend C进阶课程内容与CUDA的对比学习。

算子分析

算子分析有两种方法：


以上两种方法，只有参数上的区别。

Tiling

1.Tiling结构体定义：用于存储Tiling相关参数。
2.在host端使用__global__ aicore void add_custom函数进行算子自定义。
3.对于Tiling的数据在host端进行内存分配，然后将数据从host复制到device中，使用acldrtMallocHost和acldrtMemcpy函数进行内存操作。

对比学习CUDA

Tiling结构体定义:

在CUDA中，可以定义C++结构体来存储kernel的执行参数，比如每个tile的大小、迭代次数等。这些参数可以在启动kernel时传递给它。
自定义算子:

在CUDA中，自定义算子通常是通过编写CUDA kernel函数实现的。这些函数在GPU上并行执行，并通过执行配置（即线程块的大小和数量）来定义。
内存管理:

CUDA中的内存管理涉及到主机（CPU）和设备（GPU）之间的内存分配和数据传输。使用cudaMalloc, cudaFree来管理设备内存，使用cudaMemcpy来在主机和设备之间传输数据。
对于较大的数据结构，可能需要使用cudaMallocPitch和cudaMemcpy2D来优化内存访问模式并减少碎片化。
Host与Device间的数据传输:

在CUDA中，数据从主机内存传输到设备内存（或反向）是通过cudaMemcpy函数执行的，该函数有多个变体，如cudaMemcpyAsync，可以与CUDA流一起使用以实现异步数据传输。
Kernel函数执行:

在CUDA中，执行kernel函数需要定义线程格（grid）和线程块（block）的维度。每个线程块可以认为是一个“tile”，它在执行期间处理输入数据的一个子集。
这里给出一个在CUDA中如何使用的样例代码：

// CUDA样例代码

// 定义Tiling相关的配置结构体
struct TilingConfig {
    int tileWidth;
    int tileHeight;
    int iterations;
};

// CUDA Kernel函数，实现自定义的并行计算
__global__ void customKernel(float* input, float* output, TilingConfig config) {
    // 计算当前线程的全局索引
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int idy = blockIdx.y * blockDim.y + threadIdx.y;

    // 简单的示例操作：输入元素加倍
    if(idx < config.tileWidth && idy < config.tileHeight) {
        for(int iter = 0; iter < config.iterations; ++iter) {
            // 这里可以根据实际的Tiling逻辑进行数据处理
            output[idy * config.tileWidth + idx] = 2 * input[idy * config.tileWidth + idx];
        }
    }
}

// Host代码
int main() {
    TilingConfig config = {32, 32, 100}; // 假设每个tile是32x32，迭代100次
    size_t size = config.tileWidth * config.tileHeight * sizeof(float);
    float *input, *output;

    // 在主机和设备上分配内存
    cudaMallocManaged(&input, size);
    cudaMallocManaged(&output, size);

    // 初始化输入数据
    for(int i = 0; i < config.tileWidth * config.tileHeight; ++i) {
        input[i] = i; // 示例数据
    }

    // 配置执行参数
    dim3 blocks(1, 1); // 使用单个block
    dim3 threads(config.tileWidth, config.tileHeight); // 每个block的尺寸

    // 启动Kernel
    customKernel<<<blocks, threads>>>(input, output, config);

    // 等待GPU完成
    cudaDeviceSynchronize();

    // 释放内存
    cudaFree(input);
    cudaFree(output);

    return 0;
}

注：从下面章节开始，由于概念部分的内容与基础阶段高度类似，故部分删除，仅展示对比学习的代码部分。

动态shape场景的Tiling解析函数

Ascend C中，CPU模式和NPU模式涉及到指针转化，用宏函数 CONVERT_TILING_DATA 将_ubuf_ uint8_t* 转化为__ubuf_tilingstruct*，CPU模式和NPU模式之间逻辑的区别，用宏函数INIT_TILING DATA区分tiling data不同的初始化过程。
在CUDA中实现类似的功能，我们可以使用不同的函数来处理CPU和GPU模式下的数据转换和初始化过程：

// 定义Tiling结构体，用于存储Tiling操作的参数
struct TilingData {
    dim3 gridSize;
    dim3 blockSize;
    int tileWidth;
    int tileHeight;
    // 可以添加更多的Tiling参数
};

// 用于转换指针的函数
TilingData* convertToTilingData(uint8_t* ptr) {
    // 在CUDA中，你可以直接转换指针类型
    return reinterpret_cast<TilingData*>(ptr);
}

// 初始化TilingData的函数
void initTilingData(TilingData* tilingData, int totalWidth, int totalHeight, int maxTileWidth, int maxTileHeight) {
    // 根据输入和硬件限制计算Tiling参数
    // 这里的逻辑将与Ascend C中的宏 INIT_TILING_DATA 类似
    tilingData->tileWidth = min(maxTileWidth, totalWidth);
    tilingData->tileHeight = min(maxTileHeight, totalHeight);
    // ... 计算其他参数
}

// 示例函数，展示如何在CUDA中使用这些函数
void exampleUsage(uint8_t* ptr, int totalWidth, int totalHeight) {
    // 转换指针
    TilingData* tilingData = convertToTilingData(ptr);

    // 初始化TilingData
    initTilingData(tilingData, totalWidth, totalHeight, 16, 16); // 假设最大Tile大小为16x16

    // 根据tilingData来配置和启动CUDA内核
    // ...
}

CPU下算子调试

使用GDB进行调试:

在Ascend C环境中，可以使用GDB来调试CPU模式下的代码。调试前需要先设置环境变量，然后启动GDB，并设置断点、运行程序、查看变量值、继续执行到下一个断点等。
提供了常用的GDB命令，如run, break, continue, print, list, backtrace, display以及退出调试的quit。
使用打印信息进行调试:

在Ascend C的CPU模式下，推荐使用标准C/C++的打印函数，如printf, std::cout，因为NPU模式可能不支持这些标准输出函数。
与CUDA进行调试的对比：

使用GDB调试:

在CUDA中，可以使用类似的GDB调试方法，但专门用于CUDA的GDB版本称为cuda-gdb。cuda-gdb允许开发者调试运行在GPU上的CUDA代码。
CUDA的调试命令与传统的GDB相似，但还包括了一些特定于CUDA的命令，如检查GPU寄存器的状态或者在特定的CUDA线程上设置断点。
使用打印信息进行调试:

在CUDA中，可以在设备代码中使用printf来打印调试信息。从CUDA 5.0开始，printf可以在设备代码中使用，虽然有一些限制，如在CUDA内核中不支持宽字符和浮点数的格式化。

NPU下算子调试

调试输出文件:

在NPU模式下，调试产生的主要输出文件包括*.dump和*.vcd文件，这些文件记录了执行过程中的详细信息。
日志文件:

对于AI Core，调试信息被记录在*_core*_summary_log文件中，而波形数据则保存在*_core0_wave.vcd文件中。
性能分析:

调试工具可以帮助开发者理解NPU中的Model执行过程和性能瓶颈。通过分析日志和波形文件，可以识别出AI Core的利用率等关键性能指标。
波形分析工具:

使用专门的波形分析工具可以可视化和分析*.vcd文件，为开发者提供了执行细节的图形化表示，便于识别和解决问题。

与CUDA进行对比：

调试输出:

在CUDA中，调试输出通常不包括波形文件。CUDA调试更多地依赖于cuda-gdb或Nsight等工具来提供代码执行的洞察。
性能分析:

CUDA提供了Nsight Compute和Nsight Systems等工具来分析GPU上的kernel性能，这些工具提供了详细的性能指标，如内存访问模式、执行时间和流水线效率等。
波形分析:

CUDA通常不提供波形分析工具，因为GPU架构和波形分析在NPU硬件调试中更为常见。CUDA的性能分析侧重于通过软件工具来优化代码和内存访问模式。
日志文件:

CUDA的调试和性能分析信息通常通过控制台输出或可视化界面展示，而不是通过日志文件。

Ascend C算子交付件

Host侧


总结：
Header文件 (<KERNEL_NAME>_tiling.h):

这个头文件定义了与Tiling相关的数据结构，这些结构用于存储和处理Tiling操作所需的参数。
它包含了宏定义和函数声明，这些都是为了设置和管理Tiling过程所必需的。
源文件 (<KERNEL_NAME>.cpp):

源文件包含了实现Tiling逻辑的函数，例如如何分割数据以及如何处理不同的数据块。
它还包含了形状推导函数，这些函数基于算子的输入形状、逻辑和属性来确定输出形状。
除此之外，源文件还定义了算子的注册过程，包括输入、输出和属性的规格，以及与Tiling和形状推导相关的函数。
这些文件组成了算子的Host侧代码的基础，确保了在CPU上可以准确地设置和执行相应的NPU计算任务。

对比CUDA中的类似概念：

CUDA Kernel文件:

在CUDA中，算子（或称为kernel）的实现通常包含在.cu文件中，这与Ascend C中的.cpp文件相似。
与<KERNEL_NAME>_tiling.h类似，CUDA中可能会有一个头文件（通常是.cuh或.h），定义了kernel函数需要的所有数据结构和辅助函数。
CUDA Kernel注册和执行:

CUDA没有类似于Ascend C的算子注册机制，因为CUDA kernel的调用更直接，通过CUDA Runtime API或CUDA Driver API直接从Host代码中调用。
形状推导在CUDA中不是内置功能，它在深度学习框架中由框架本身处理，或者需要开发者自行在Host代码中实现。
Tiling和形状推导:

CUDA中的Tiling策略需要开发者在kernel函数中实现，并且可能需要在Host代码中进行复杂的计算来确定grid和block的尺寸。
CUDA也允许在kernel中通过内置变量获取当前线程的索引和维度，这些变量有助于实现Tiling逻辑。
Kernel侧

Kernel侧代码示例中展示了在Ascend平台上，如何使用Tiling技术来处理动态shape的数据。这里的关键步骤包括获取Tiling数据、初始化计算操作，并根据Tiling键（TILING_KEY_IS）来决定是否执行计算过程。
对比CUDA
在CUDA中，处理动态shape数据的过程类似，但没有直接的Tiling键或宏。您需要在kernel函数中手动管理数据的分块和处理。样例代码如下：

__global__ void myKernel(float* input, float* output, int totalLength, int tileNum) {
    extern __shared__ float tile[]; // 使用共享内存作为tile

    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int tileSize = totalLength / tileNum; // 假设totalLength可以整除tileNum
    int tileOffset = blockIdx.x * tileSize;

    // 判断当前线程是否在当前tile的范围内
    if (idx < totalLength) {
        // 根据当前block的tile进行处理
        
        tile[threadIdx.x] = input[tileOffset + threadIdx.x];
        __syncthreads(); // 同步当前block的线程

        // 进一步处理tile中的数据
        // 然后写回到全局内存
        output[tileOffset + threadIdx.x] = tile[threadIdx.x];
    }
}

// Host代码
void launchMyKernel(float* input, float* output, int totalLength, int tileNum) {
    int tileSize = totalLength / tileNum; // 计算每个tile的大小
    dim3 blockSize(tileSize); // 假设blockSize可以整除tileSize
    dim3 gridSize(tileNum);

    // 这里假设每个tile正好由一个block处理，因此共享内存的大小为一个tile的大小
    myKernel<<<gridSize, blockSize, tileSize * sizeof(float)>>>(input, output, totalLength, tileNum);
    cudaDeviceSynchronize(); // 等待kernel执行完成
}

在这个CUDA示例中，定义了一个kernel函数myKernel，它使用共享内存来存储每个tile的数据。通过计算tileSize和tileOffset来确定每个线程应该处理的数据。在Host函数launchMyKernel中，根据总数据长度和tile数量来计算grid和block的大小，并启动kernel。

UT/ST测试

当开发完成后，需要进行测试工作。详情如下：

UT（单元测试）流程：
单元测试是针对程序的最小可测试部分进行的测试，通常是对单个函数或方法的验证。
单元测试的关键目的是验证该代码片段的逻辑是否按预期工作。
UT测试需要达到尽可能高的代码覆盖率，理想情况下应达到100%，以确保所有代码路径都经过测试。
ST（系统测试）流程：
系统测试是在软件开发过程中较晚的阶段进行的，用于验证整个系统是否满足规定的需求。
ST测试的重点在于确保软件作为一个整体运行时的行为符合预期。
在ST测试中，通常使用.json格式的测试用例文件来指导测试流程，并利用专业的测试软件或框架来执行。
对比CUDA

在CUDA中进行单元测试（UT）和系统测试（ST）通常涉及到编写特定的测试代码，以验证算子的正确性和性能。单元测试通常针对单一功能或特定条件，而系统测试则更全面，考虑到整个系统的行为。
CUDA UT测试：
在CUDA中，单元测试可能使用如Google Test等框架，以确保单个CUDA kernel的正确性。
CUDA代码的测试可能需要模拟GPU上的数据处理过程，以及验证内存拷贝和同步操作是否正确执行。
CUDA ST测试：
对于CUDA应用程序的系统测试，可能涉及到整个应用程序的编译、链接、运行以及输出验证。
系统测试通常包括运行实际的数据集，验证是否达到了性能和准确性的要求。
例如，可以通过NVIDIA Nsight Systems进行性能分析，以及使用自动化脚本来测试不同的输入条件和工作负载。
如下给出一个Google Test来进行UT测试的代码部分。

// add.cu
__global__ void addKernel(int *c, const int *a, const int *b, size_t n) {
    size_t index = threadIdx.x + blockIdx.x * blockDim.x;
    if (index < n) {
        c[index] = a[index] + b[index];
    }
}

// add_test.cu
#include <gtest/gtest.h>
#include "add.cu"

class AddKernelTest : public ::testing::Test {
protected:
    void SetUp() override {
        // 初始化数据和分配内存
        cudaMalloc(&d_a, size * sizeof(int));
        cudaMalloc(&d_b, size * sizeof(int));
        cudaMalloc(&d_c, size * sizeof(int));
        
        // ... 初始化d_a和d_b的数据
    }

    void TearDown() override {
        // 释放内存
        cudaFree(d_a);
        cudaFree(d_b);
        cudaFree(d_c);
    }

    int *d_a = nullptr;
    int *d_b = nullptr;
    int *d_c = nullptr;
    size_t size = 1024;
};

TEST_F(AddKernelTest, CanAddVectors) {
    // 调用CUDA内核
    addKernel<<<(size + 255) / 256, 256>>>(d_c, d_a, d_b, size);
    
    // 等待CUDA操作完成
    cudaDeviceSynchronize();
    
    // 检查结果（这里简化了检查过程）
    int *h_c = new int[size];
    cudaMemcpy(h_c, d_c, size * sizeof(int), cudaMemcpyDeviceToHost);
    
    for (size_t i = 0; i < size; ++i) {
        ASSERT_EQ(h_c[i], /* 期望的值 */);
    }
    
    delete[] h_c;
}

CUDA系统测试示例
系统测试通常是在更高级别上进行的，您可能会有一个测试驱动程序，该程序执行完整的应用程序流程。这种类型的测试往往不在代码级别上，而是在构建或部署流程中执行：

# 系统测试脚本示例 (bash)
#!/bin/bash

# 编译CUDA程序
nvcc -o my_app main.cu

# 运行程序，可能会涉及多个不同的输入集合
./my_app input1.dat
./my_app input2.dat
# ...

# 检查输出文件是否符合预期
# ...

# 报告测试结果
echo "All system tests passed."

在CUDA中进行UT和ST测试需要注意，由于GPU操作通常是异步的，因此在测试中需要确保相关操作已完成。这可以通过调用cudaDeviceSynchronize()来实现，这会等待直到GPU完成所有前面的操作。