[自动驾驶 SoC]-2 软硬协同

以GPU为例细说下AI芯片的软硬协同工作机制：

先说软件部分内核分派

1 编写和编译内核函数

1）编写内核函数

内核函数是运行在GPU上的函数，通常用CUDA或OpenCL编写

2）编译内核函数

使用专用编译器（如NVCC）将内核函数编译为GPU可执行的二进制代码。

2 准备运行时代码

编译后的二进制代码需要在运行时加载到主机程序中。

例如：

extern "C" {

    extern __global__ void myKernel(float *d_A);

}

int main() {

    float *d_A;

    size_t size = 256 * sizeof(float);

    // 分配设备内存

    cudaMalloc((void**)&d_A, size);

    // 内核参数设置

    dim3 grid(1);

    dim3 block(256);

    // 内存复制操作（如有需要）

    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);

    // 启动内核

    myKernel<<<grid, block>>>(d_A);

    // 设备到主机的内存复制操作

    cudaMemcpy(h_A, d_A, size, cudaMemcpyDeviceToHost);

    // 释放设备内存

    cudaFree(d_A);

    return 0;

}

3 形成命令队列

在CUDA中，启动内核时，运行时会形成内核启动命令并将其加入到CUDA流中。主机端运行时将内核启动命令等插入用于存放AQL(Architected Queuing Language)数据包的AQL队列中，AQL数据包含有与内核分派有关信息，如网格、工作组大小、内核函数信息等。

cudaStream_t stream;

cudaStreamCreate(&stream);

myKernel<<<grid, block, 0, stream>>>(d_A);

cudaStreamSynchronize(stream);

cudaStreamDestroy(stream);

4 调用驱动程序

CUDA编程接口（API）封装了对驱动程序的调用。这些API调用底层驱动程序接口，将内核启动命令和参数传递给GPU驱动程序。

1）内核启动命令封装

CUDA API 将内核启动命令封装成可传递给驱动程序的结构。

// Pseudo-code for CUDA kernel launch

cudaError_t cudaLaunchKernel(const void *func, dim3 gridDim, dim3 blockDim, void **args, size_t sharedMem, cudaStream_t stream) {

    // Translate API parameters to driver-level parameters

    CUfunction kernel = (CUfunction)func;

    void *config[] = {CU_LAUNCH_PARAM_BUFFER_POINTER, args, CU_LAUNCH_PARAM_END };

    // Call driver API

    return cuLaunchKernel(kernel, gridDim.x, gridDim.y, gridDim.z, blockDim.x, blockDim.y, blockDim.z, sharedMem, stream, NULL, config);

}

2）调用驱动程序接口

驱动程序接口负责将内核启动命令传递到GPU。

CUresult cuLaunchKernel(CUfunction f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, Unsigned int sharedMemBytes, CUstream hStream, void **kernelParams, void **extra) {

    // Implementation-specific details

    // Pass parameters to GPU hardware

}

5 内核分派到GPU硬件执行

接下来细说硬件分派过程

当驱动程序接收内核启动命令，生成命令缓冲区并填充相关信息，包括内核函数指针、参数、网格和块的配置等。

// Pseudo-code for driver-level kernel launch

CUresult cuLaunchKernel(CUfunction f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, Unsigned int sharedMemBytes, CUstream hStream, void **kernelParams, void **extra) {

    // Create command buffer and fill with kernel launch information

    CommandBuffer cmdBuffer;

    fillCommandBuffer(&cmdBuffer, f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes, hStream, kernelParams);

    // Submit command buffer to ACE

    submitCommandBufferToACE(&cmdBuffer);

    return CUDA_SUCCESS;

}

1 ACE接收并解析命令

Asynchronous Compute Engine (ACE) 负责接收命令缓冲区并解析内核启动命令。ACE将这些命令转化为GPU硬件可以理解的形式。

void submitCommandBufferToACE(CommandBuffer *cmdBuffer) {

    ACE ace = getAvailableACE();

    ace.processCommandBuffer(cmdBuffer);

}

void ACE::processCommandBuffer(CommandBuffer *cmdBuffer) {

    // Parse command buffer

    KernelLaunchCommand cmd = parseCommandBuffer(cmdBuffer);

    // Queue command for execution

    commandQueue.push(cmd);

}

2 调度与资源分配

ACE调度器从命令队列中提取内核启动命令，并根据资源可用性和调度策略将任务分配给合适的着色器引擎（Shader Engine，SE）。GPU将计算资源划分为若干个SE，根据深度神经网络算子工作负载不同，这些SE及其CU (Compute Unit)可能全部或部分被承载工作负载的工作组占用。

void ACE::scheduleCommands() {

    while (!commandQueue.empty()) {

        KernelLaunchCommand cmd = commandQueue.front();

        commandQueue.pop();

        ShaderEngine se = getAvailableShaderEngine();

        se.dispatchKernel(cmd);

    }

}

3 着色器引擎执行

Shader Engine接收内核启动命令，配置计算单元（如SM，Streaming Multiprocessor），并启动内核执行。SM流式多处理器是GPU架构中的一个更高级的组成单元，它由多个CU组成，并提供对这些CU的调度和管理；一个SM可以同时执行多个线程块（block），并在这些线程块之间进行上下文切换，以提高GPU的并行度和资源利用率。CU称为计算单元，是GPGPU架构中的基本计算单元,每个CU包含多个ALU（算术逻辑单元）、寄存器文件、共享内存和其他辅助单元，能够并行执行多个线程或工作项。

void ShaderEngine::dispatchKernel(KernelLaunchCommand &cmd) {

    // Configure shader engine with kernel launch parameters

    configureShaderEngine(cmd);

    // Execute kernel

    executeKernel(cmd);

}

void ShaderEngine::configureShaderEngine(KernelLaunchCommand &cmd) {

    // Set up grid and block dimensions

    setGridDim(cmd.gridDim);

    setBlockDim(cmd.blockDim);

    // Load kernel arguments

    loadKernelArguments(cmd.kernelParams);

}

void ShaderEngine::executeKernel(KernelLaunchCommand &cmd) {

    // Dispatch work items (threads) to execution units (CUDA cores)

    for (int blockIdx = 0; blockIdx < cmd.gridDim.x; ++blockIdx) {

        for (int threadIdx = 0; threadIdx < cmd.blockDim.x; ++threadIdx) {

            int smId = allocateSM();

            int coreId = allocateCore(smId);

            launchThread(smId, coreId, cmd.kernelFunction, cmd.kernelParams);

        }

    }

}

在Shader Engine中，内核函数被分发到不同的计算单元CU，以下是取值（Fetch）、译码（Decode）、执行（Execute）和写回（Writeback）步骤的详细实现：

1）取值（Fetch）

每个计算单元（如CUDA核心）从指令缓存中取出内核指令。

Instruction fetchInstruction(CUfunction kernelFunction) {

    // Fetch the next instruction from the kernel function's instruction cache

    return instructionCache[kernelFunction][programCounter++];

}

2）译码（Decode）

将取出的指令译码为可执行的操作。

DecodedInstruction decodeInstruction(Instruction instruction) {

    // Decode the instruction into an executable format

    return decodeTable[instruction.opcode](instruction);

}

3）执行（Execute）

执行译码后的指令，完成计算操作。

void executeInstruction(DecodedInstruction decodedInstruction, void **kernelParams) {

    // Execute the decoded instruction using the provided kernel parameters

    executionUnit[decodedInstruction.opcode](decodedInstruction, kernelParams);

}

4）写回（Writeback）

将执行结果写回到适当的存储位置（如寄存器文件或全局内存）

void writebackResult(DecodedInstruction decodedInstruction, ExecutionResult result) {

    // Write the result of the execution back to the appropriate location

    if (decodedInstruction.destination.isRegister) {

        registerFile[decodedInstruction.destination] = result;

    } else {

        globalMemory[decodedInstruction.destination] = result;

    }

}

4 硬件执行与事件处理

执行过程中和执行完毕后，硬件会生成事件。这些事件用于同步和错误处理。

void handleEvent(EventType event) {

    switch (event) {

        case KERNEL_LAUNCH_COMPLETE:

            // Handle kernel launch completion

            break;

        case MEMORY_COPY_COMPLETE:

            // Handle memory copy completion

            break;

        case ERROR_OCCURRED:

            // Handle errors

            break;

    }

}