# 从CUDA迁移到DTK：一个后端开发踩过的坑和真相 > 花了三周把3000行CUDA代码移植到DTK上，GPU利用率从92%掉到47%。排查了一个通宵才发现，问题不在代码逻辑——在线程束大小和同步模型上，这俩家伙根本不是一回事。 --- ## 一、先搞明白硬件：DCU到底长什么样 DCU（Deep Computing Unit）采用两级缓存的SIMT架构，核心计算单元叫Compute Unit（CU），每个CU包含4个SIMD。这和NVIDIA GPU的SM（Streaming Multiprocessor）在大框架上类似，但有个关键差异： **Warp Size不同。** - NVIDIA GPU：Warp Size = 32 - DCU：Warp Size = 64 意味着什么？线程束内的线程同时执行同一条指令，你的代码如果按32对齐优化，在DCU上可能只跑了一半的效率。 每个SIMD内部结构：64线程组成一个wavefront在1个SIMD上执行；64线程 × 256 × 32bit寄存器（VGPR），每个线程最多256个寄存器；资源充足时每个SIMD最多支持10个wavefront并发。 ```mermaid graph TB subgraph DCU["DCU 芯片"] subgraph CU0["Compute Unit 0"] SIMD0["SIMD 0
64 Threads / Wavefront"] SIMD1["SIMD 1
64 Threads / Wavefront"] SIMD2["SIMD 2
64 Threads / Wavefront"] SIMD3["SIMD 3
64 Threads / Wavefront"] VGPR0["VGPR
256×32bit ×64"] SGPR0["SGPR
800×32bit"] SALU0["SALU"] LDS0["LDS 64KB"] L1_Cache0["Texture R/W Cache L1"] end CU1["Compute Unit 1
（结构同CU0）"] CU_N["... 更多 Compute Unit"] GDS["GDS 64KB
所有CU共享"] L2["L2 Read/Write Cache
Per Memory Channel"] CrossBar["CrossBar 互联"] end CU0 --> CrossBar CU1 --> CrossBar CU_N --> CrossBar CrossBar --> L2 CrossBar --> GDS style DCU fill:#f5f5f5,stroke:#1565c0,stroke-width:2px style CU0 fill:#e3f2fd,stroke:#1976d2 style GDS fill:#fff3e0,stroke:#ef6c00 style L2 fill:#e8f5e9,stroke:#2e7d32 ``` 除了CU和SIMD，DCU还包含： - **LDS**（Local Data Share）：每个CU 64KB，类似NVIDIA的shared memory - **GDS**（Global Data Share）：64KB，所有CU共享 - **L1/L2 Cache**：每CU有Texture R/W Cache L1，每Memory Channel有L2 R/W Cache > 说白了：DCU在硬件架构上和NVIDIA GPGPU基本同源，但线程束大小翻倍是迁移时第一个要盯住的变量。 DCU系列经历了多代产品迭代，从初代产品打通GPGPU生产研发流程，到后续逐步扩充计算单元、显存参数，再到新增TensorCore功能提升单精度和半精度计算速率，以及增加节点内数据高速互联、视频编解码器等。整个软件生态目前已支持主流AI框架和HPC应用。 ```mermaid timeline title DCU系列产品迭代路线 section 早期产品 初代 : 实现国产化GPGPU
生产研发流程打通 section 中期演进 第二代 : 扩充计算单元与显存参数 第三代 : 新增数学计算指令
计算资源翻倍 section 近期发展 AI增强代 : 新增TensorCore功能
提高单精/半精计算速率 最新代 : 节点内高速互联
视频编解码器
进一步扩充计算单元 ``` --- ## 二、软件栈对比：DTK到底兼容到什么程度 DTK（DCU Toolkit）是DCU硬件平台的开发工具包。官方说法是同时支持HIP和CUDA两种编程模型，能快速将CUDA和ROCm生态中的应用部署到DCU上。**说白了就是：你写HIP能跑，写CUDA也能跑（通过GPUfusion兼容层）。** ### 2.1 软件栈全景 DTK各组件对ROCm生态原生兼容，与CUDA生态高度兼容。 ```mermaid graph TB subgraph APP["HPC & AI 广泛应用"] AI["AI框架"] HPC["HPC应用"] SCI["科学计算"] end subgraph MATH["数学库层"] DTK_MATH["DTK数学库
rocblas/rocsparse/
rocsolver/rocrand/rocfft"] HIP_MATH["HIP数学库"] CUDA_MATH["CUDA兼容数学库
cublas/cusparse/
cusolver/curand/cufft"] end subgraph COMPILER["编译器层"] DCC["dcc/clang
（DTK编译器）"] HIPCC["hipcc"] NVCC["nvcc/clang
（CUDA兼容）"] end subgraph RUNTIME["运行时 & 库"] LIBS_DTK["librocsm_smi64.so
libhiprtc.so
等动态库"] LIBS_CUDA["libnvidia-ml.so
libnvrtc.so
libcudart.so
等动态库"] HEADERS_DTK["HIP头文件"] HEADERS_CUDA["CUDA头文件"] end subgraph OS["多操作系统支持"] CENTOS["CentOS"] UBUNTU["Ubuntu"] KYLIN["Kylin"] UOS["UOS"] NFS["NFS"] end subgraph HW["硬件平台"] CPU["CPU
Intel / AMD / Hygon / ARM"] DCU_HW["DCU 加速器"] end APP --> MATH MATH --> COMPILER COMPILER --> RUNTIME RUNTIME --> OS OS --> HW style APP fill:#e8eaf6,stroke:#283593 style MATH fill:#e3f2fd,stroke:#1565c0 style COMPILER fill:#e8f5e9,stroke:#2e7d32 style RUNTIME fill:#fff3e0,stroke:#ef6c00 style OS fill:#fce4ec,stroke:#c62828 style HW fill:#f3e5f5,stroke:#6a1b9a ``` 核心模块对照表： | 模块 | CUDA | ROCm | DTK | DTK-CUDA兼容 | | -------------- | ------------------------------- | ----------------------------------- | ----------------------------------- | ------------------------------- | | Driver API | cuda | hsa-runtime | hsa-runtime | cuda | | 运行时系统 | cudart | amdhip64 | galaxy | cudart | | 数学库 | cublas/cusparse/cusolver/curand | rocblas/rocsparse/rocsolver/rocrand | rocblas/rocsparse/rocsolver/rocrand | cublas/cusparse/cusolver/curand | | FFT | cufft | rocfft | rocfft | cufft | | AI算子库 | cuDNN | miopen | miopen | cuDNN | | 设备管理 | nvml | rocm-smi | rocm-smi | nvml | | 头文件库 | Thrust/cub | Thrust/rocprim | Thrust/rocprim | Thrust/cub | | 编译器 | nvcc | hipcc/clang | dcc/clang | nvcc/clang | | 性能调优 | nvprof | rocprofiler | Hipprof | hipprof | | 调试器 | cuda-gdb | rocgdb | hipgdb | hipgdb | | 内存检查 | cuda-memcheck | radeon_memory_visualizer | Hipprof --leakcheck | hipprof | | OpenACC/OpenMP | — | 仅OpenMP | 支持 | 支持 | ### 2.2 运行时功能差异 | 功能 | CUDA | ROCm | DTK | | -------- | ---- | ---------------- | ------------------------------- | | 纹理对象 | 支持 | 支持（仅单进程） | 支持（通用计算单元实现） | | 协作组 | 支持 | 部分支持 | 部分支持（cluster等硬件不支持） | | 动态并行 | 支持 | 不支持 | 不支持（研发中） | ### 2.3 API命名规范 | API类型 | CUDA | ROCm/DTK | | ----------- | ---------- | --------- | | Driver API | `cu`开头 | `hip`开头 | | Runtime API | `cuda`开头 | `hip`开头 | > **坑就在这里**：如果你原来的CUDA代码大量用了动态并行（Dynamic Parallelism），目前DTK不支持，需要重构这部分逻辑。 --- ## 三、HIP编程模型：C++扩展的三板斧 HIP并行编程模型是对C/C++的扩展，为加速器提供可移植的异构并行编程环境。核心就三件事：**核函数、线程结构、存储结构**。 ### 3.1 核函数 Kernel 把一段要在设备端并行执行的代码用 `__global__` 标记，这就是核函数。 ```cpp // 设备端核函数 __global__ void myKernel(int N, double *d_a) { int i = threadIdx.x + blockIdx.x * blockDim.x; if (i < N) { d_a[i] *= 2.0; } } // 主机端调用 myKernel<<>>(N, d_a); ``` 核函数的规矩： - 返回值**必须是 void** - 参数可以是设备端指针、立即数、结构体或引用 - 以线程为基本单位执行，线程之间总是"并发"的 - 每个线程有独立编号 HIP核函数启动语法有三种写法（本质上等价）： | 方式 | 语法 | | --------------- | ------------------------------------------------------------ | | HIP宏 | `hipLaunchKernelGGL(kernel, numBlocks, dimBlocks, sharedMemBytes, stream, args...)` | | HIP Kernel语法 | `kernel<<>>(args...)` | | CUDA Kernel语法 | `kernel<<>>(args...)` | CUDA与HIP核函数基本完全兼容——除了目前DTK中Dynamic Parallelism还在研发中。 ### 3.2 线程结构：理解了Grid/Block/Wavefront才算入门 线程层次结构是异构编程的灵魂。HIP和CUDA在线程组织上的对应关系： | NVIDIA GPU | HIP | 描述 | | ----------------------------- | ----------------- | ------------------------------------------------------------ | | Streaming Multiprocessor (SM) | Compute Unit (CU) | 多组并行计算单元组成的计算结构，一个block内所有thread分配到同一个CU | | Kernel | Kernel | 核函数，可被多个CU并发执行 | | Warp | Wavefront | 线程束/波前，硬件执行基本单元。DTK线程束大小=64 | | Thread | Thread | 执行核函数的基本单元 | | Block | Block | 线程块，由一个CU执行。块内线程可直接通信和同步 | ```mermaid graph TB subgraph GRID["Grid（线程网格）"] subgraph BLOCK0["Block 0"] TH0["Thread
threadIdx=0"] TH1["Thread
threadIdx=1"] TH2["Thread
..."] THN["Thread
threadIdx=N"] end subgraph BLOCK1["Block 1"] TH10["Thread"] TH11["Thread"] TH12["..."] end BLOCK_M["Block M"] end CU0["Compute Unit 0
（执行Block 0）"] CU1["Compute Unit 1
（执行Block 1）"] BLOCK0 --> CU0 BLOCK1 --> CU1 BLOCK_M --> CU1 style GRID fill:#e3f2fd,stroke:#1565c0,stroke-width:2px style BLOCK0 fill:#bbdefb,stroke:#1976d2 style BLOCK1 fill:#bbdefb,stroke:#1976d2 style CU0 fill:#e8f5e9,stroke:#2e7d32 style CU1 fill:#e8f5e9,stroke:#2e7d32 ``` 核函数中通过内置变量获取线程在3D-Grid中的位置： | 含义 | HIP变量 | CUDA变量 | | ----------------- | -------------------- | ----------------- | | block内thread位置 | `hipThreadIdx_x/y/z` | `threadIdx.x/y/z` | | grid内block位置 | `hipBlockIdx_x/y/z` | `blockIdx.x/y/z` | | block维度 | `hipBlockDim_x/y/z` | `blockDim.x/y/z` | | grid维度 | `hipGridDim_x/y/z` | `gridDim.x/y/z` | **1D线程网格**——计算全局索引： ```cpp int idx = blockDim.x * blockIdx.x + threadIdx.x; ``` **2D线程网格**——需要同时计算X和Y： ```cpp int idx_x = blockDim.x * blockIdx.x + threadIdx.x; int idx_y = blockDim.y * blockIdx.y + threadIdx.y; ``` 搞懂了1D和2D，3D加个Z轴同理。并行算法的核心就是把计算任务映射到这个3D-Grid上——这个映射关系选对了，性能提升往往是数量级的。 ```mermaid graph LR subgraph GRID2D["2D Grid"] subgraph B00["Block(0,0)"] T00["T00"] --- T01["T01"] --- T02["T02"] end subgraph B01["Block(0,1)"] T10["T10"] --- T11["T11"] --- T12["T12"] end subgraph B10["Block(1,0)"] T20["T20"] --- T21["T21"] --- T22["T22"] end subgraph B11["Block(1,1)"] T30["T30"] --- T31["T31"] --- T32["T32"] end end style GRID2D fill:#e8eaf6,stroke:#283593,stroke-width:2px style B00 fill:#c5cae9,stroke:#3949ab style B01 fill:#c5cae9,stroke:#3949ab style B10 fill:#c5cae9,stroke:#3949ab style B11 fill:#c5cae9,stroke:#3949ab ``` ### 3.3 存储结构：三层内存模型 HIP存储结构是层次化的，直接对应硬件存储层次。**搞懂了这个，一大半的性能问题都能自己排查了。** ```cpp __device__ int array_on_global[10]; // 全局内存，所有线程可访问 __global__ void MyKernel(int *array, int arrayCount) { int idx = blockDim.x * blockIdx.x + threadIdx.x; __shared__ int array_on_lds[10]; // 共享内存/LDS，block内可见 array_on_global[0] = idx; } ``` ```mermaid graph TB subgraph GRID_MEM["Grid 级别"] GM["Global Memory（全局内存 HBM）
__device__ 修饰
所有线程可访问 | 应用生命周期"] CM["Constant Memory（常量内存）
__constant__ 修饰
所有线程可读 | 应用生命周期"] end subgraph BLOCK_MEM["Block 级别"] LDS["LDS / Shared Memory（共享内存）
__shared__ 修饰
Block内可见 | Block生命周期"] end subgraph THREAD_MEM["Thread 级别"] REG["寄存器 Register
float var（标量）
仅本线程 | 线程生命周期"] LM["本地内存 Local Memory
float var[100]（数组溢出）
仅本线程 | 线程生命周期"] end THREAD_MEM --> BLOCK_MEM --> GRID_MEM style GM fill:#ffcdd2,stroke:#c62828 style CM fill:#ffcdd2,stroke:#c62828 style LDS fill:#fff9c4,stroke:#f9a825 style REG fill:#c8e6c9,stroke:#2e7d32 style LM fill:#c8e6c9,stroke:#2e7d32 ``` 存储层次速查表： | 修饰符 | 存储器 | 作用域 | 生命周期 | | ---------------- | ------------ | ------ | -------- | | （无修饰符）标量 | 寄存器 | 线程 | 线程 | | （无修饰符）数组 | 本地内存 | 线程 | 线程 | | `__shared__` | LDS/共享内存 | Block | Block | | `__device__` | 全局内存 | 全局 | 应用程序 | | `__constant__` | 常量内存 | 全局 | 应用程序 | 函数修饰符： | 限定符 | 执行位置 | 调用来源 | | ------------ | -------- | ------------------------ | | `__global__` | 设备端 | 主机端调用（设备端也可） | | `__device__` | 设备端 | 仅设备端 | | `__host__` | 主机端 | 仅主机端 | `__device__` 和 `__host__` 可以一起用，函数会在主机和设备端各编译一份。 ### 3.4 内置函数：同步与原子操作 **原子操作**：支持 `atomicAdd/Sub/Exch/Min/Max/CAS/And/Or/Xor` 等，覆盖 `int/unsigned int/long long/float/double`。`atomicInc`/`atomicDec` 目前用 `atomicCAS` 替代实现，性能会差一些。 Scope语义： | CUDA scope | HIP宏 | 含义 | | --------------------- | ------------------------------ | ----------- | | `thread_scope_system` | `__HIP_MEMORY_SCOPE_SYSTEM` | 系统级可见 | | `thread_scope_device` | `__HIP_MEMORY_SCOPE_AGENT` | 设备级可见 | | `thread_scope_block` | `__HIP_MEMORY_SCOPE_WORKGROUP` | Block内可见 | **同步函数**——这里有个重要差异： | 函数 | DTK | CUDA | 原因 | | ------------------------------ | ------ | ---- | --------------------------------- | | `__syncwarp()` | 不支持 | 支持 | DCU硬件一个线程束共用一个PC寄存器 | | `__syncthreads()` | 支持 | 支持 | — | | `__shfl_down/up/sync` 系列 | 支持 | 支持 | 调对应shfl实现 | | `__ballot_sync` 系列 | 支持 | 支持 | — | | `__reduce_add/min/max_sync` 等 | 不支持 | 支持 | — | > **NVIDIA Volta后每个线程有独立PC寄存器，能做独立线程调度。DCU目前一个线程束才一个PC，所以 `__syncwarp` 这套玩不转。** 另外DCU不支持异步操作（Asynchronous），需要异步的场景目前用同步替代。 --- ## 四、并行编程接口：Stream、Event、Memory ### 4.1 阻塞与非阻塞 ```cpp // 阻塞拷贝——完成前主机端暂停 hipMemcpy(d_a, h_a, Nbytes, hipMemcpyHostToDevice); // 非阻塞拷贝——立即返回，需配合stream hipMemcpyAsync(d_a, h_a, Nbytes, hipMemcpyHostToDevice, stream); // 核函数启动——非阻塞，立即返回 hipLaunchKernelGGL(myKernel, grid, block, 0, stream, N, d_a); ``` ### 4.2 Stream Stream是一种逻辑队列。任何对设备的操作（kernel、memcpy）都通过队列执行。同一队列顺序执行，不同队列共享硬件但互不影响。 ```cpp hipStream_t stream; hipStreamCreate(&stream); // ... 提交任务到stream ... hipStreamDestroy(stream); ``` **注意**：使用NULL或0表示默认流。对默认流的操作会导致强制同步，影响其他流中的任务。 ### 4.3 Pinned Memory 主机端默认分配的是Pageable内存。Pinned Memory（Page-Lock Memory）让DMA直接访问，大幅提升Host↔Device传输带宽。 ```cpp double *h_a = NULL; hipHostMalloc(&h_a, Nbytes); // 分配pinned内存 // ... 使用 ... hipHostFree(h_a); // 释放 ``` 频繁传输的数据优先用Pinned Memory。带宽差距非常显著，这个优化是性价比最高的。 ### 4.4 同步 ```cpp hipDeviceSynchronize(); // 重型同步：阻塞主机直到所有设备流完成 hipStreamSynchronize(stream); // 轻量同步：仅等待指定流 ``` ### 4.5 Event ```cpp hipEvent_t startEvent, endEvent; hipEventCreate(&startEvent); hipEventCreate(&endEvent); hipEventRecord(startEvent, stream); // ... 在stream中执行任务 ... hipEventRecord(endEvent, stream); hipEventSynchronize(endEvent); float time_ms; hipEventElapsedTime(&time_ms, startEvent, endEvent); // 精确计时 // 流间依赖 hipStreamWaitEvent(stream2, endEvent, 0); // stream2等endEvent完成后再执行 ``` ### 4.6 多设备编程 ```cpp int canAccessPeer; hipDeviceCanAccessPeer(&canAccessPeer, device0, device1); if (canAccessPeer) { hipDeviceEnablePeerAccess(device1, 0); // ... P2P内存访问 ... hipDeviceDisablePeerAccess(device1); } ``` 每个设备最多支持8个对等连接。P2P访问可以通过 `hipMemcpyPeerAsync` 实现设备间直接拷贝。 --- ## 五、Runtime日志与异常调试 ### 5.1 Runtime日志 出问题时开启Runtime日志做初步分析： ```bash # 方式一：环境变量 export HIP_LOG_LEVEL=4 # 1:error 2:warning 3:info 4:debug export HIP_MODULE_MASK=0x7fffffff export HIP_ENABLE_LIST=hip export HIP_LOG_OUTPUT_PATH=/home/ # 不设置默认在可执行程序当前路径 # 方式二：命令行直接设置 HIP_LOG_LEVEL=4 HIP_MODULE_MASK=0x7fffffff HIP_ENABLE_LIST=hip ./your_program # 解析日志 hiplogdump sort *.nano > log.txt ``` ### 5.2 VMFault异常分析 Kernel非法地址访问是最常见的坑。错误信息类似： ``` Invalid address access: 0x2ad4e9000000, Error code: 3. ``` VMFault信息里几个关键字段： | 字段 | 含义 | | ------------------------- | ---------------------------------------------- | | Invalid address access | 非法访问的具体地址 | | Error code: 3 | 非法内存地址访问 | | DUMP KERNEL AQL PACKET | workgroup和grid参数，grid = blockDim × gridDim | | FIND MATCH KERNEL COMMAND | 被C++修饰过的kernel名，用`c++filt`命令还原 | | DUMP KERNEL ARGS | kernel参数内容（size + 16进制堆参数） | > 调试建议：设备端函数用hipgdb，core文件用hipprof工具。 --- ## 六、GDS：GPU直通存储 GDS（GPU Direct Storage）让DCU直接读写存储设备，绕过CPU内存中转。 ```mermaid graph TB subgraph APP_GDS["应用程序"] HYFILES["hyfiles.so
文件操作库"] end subgraph KERNEL["内核层"] HYFS["hyfs.ko
内核驱动模块"] end subgraph STORAGE["存储层"] EXT4["EXT4"] XFS["XFS"] SAMBA["SAMBA"] end HYFILES -->|"会话管理
文件句柄注册/注销
设备Buffer管理
直通读写(cuFileRead/Write)"| HYFS HYFS -->|"直通内存管理（锁页/map/unmap）
直通IO管理
地址转换（影子页→设备地址）"| STORAGE style APP_GDS fill:#e3f2fd,stroke:#1565c0 style KERNEL fill:#fff3e0,stroke:#ef6c00 style STORAGE fill:#e8f5e9,stroke:#2e7d32 ``` DTK GDS核心组件：**hyfiles.so**（文件操作库）负责会话管理、文件句柄注册、设备Buffer管理、直通读写；**hyfs.ko**（内核模块）负责直通内存管理、IO管理、地址转换。 与CUDA GDS对比：基础读写（`cuFileRead`/`cuFileWrite`）、文件句柄注册/注销、设备内存注册/注销均已实现。异步读写、批量IO、Stream注册等高级功能暂未实现。生态方面DTK GDS目前支持EXT4/XFS/SAMBA文件系统。 --- ## 七、设备监控：hy-smi 对标 `nvidia-smi`，DTK提供 `hy-smi`： ```bash hy-smi # 显示卡信息（等同于rocm-smi） hy-smi --showhw # 硬件详细信息 hy-smi --showuse # DCU使用率 hy-smi --showtemp # 温度 hy-smi --showpower # 平均功耗 hy-smi --showmemuse # 显存使用率 hy-smi --showclocks # 当前时钟频率 hy-smi --showperflevel # 性能等级 hy-smi --showserial # 序列号 hy-smi --forchip # 简洁格式显示重要信息 ``` 常用的就上面这些，更多参数（如风扇、电压、PCIe带宽、ECC信息等）可以通过 `hy-smi --help` 查看。 --- ## 八、HMM：异构内存管理 HMM（Heterogeneous Memory Management）管理CPU和GPU之间的内存互访和迁移。 当前DTK HMM支持情况： - **高版本内核**：完整支持 - **低版本内核**：仅支持 `hipMallocManaged` 申请的内存 > 如果你的生产环境内核版本不够，HMM功能是受限的，很多高级特性用不了。升级内核是迁移DCU的前置动作，别等代码写完了才发现。 --- ## 九、CUDA程序迁移：两条路 DTK环境下迁移CUDA程序有两种方式： | 方式 | 说明 | 推荐度 | | -------------------------- | ------------------------------------------------------------ | ---------------- | | **hipify**（自动移植工具） | 将CUDA代码自动转成HIP代码。未来将停止维护 | 不推荐新项目使用 | | **GPUfusion**（兼容层） | 无需修改程序，加载环境变量 `source dtk/cuda/env.sh` 即可直接跑CUDA程序 | 主推荐 | GPUfusion是目前的主推方案，不改一行代码就能在DCU上跑CUDA程序。但性能需要实测——别指望零成本迁移还能跑出一样的利用率。 --- ## 写在最后 从CUDA迁移到DTK，搞定了这几个点就能少踩80%的坑： 1. **Warp Size从32变64**——线程块大小、共享内存bank冲突策略全部要重新评估 2. **动态并行不支持**——kernel嵌套调用的逻辑需要展平，别硬搬 3. **`__syncwarp`不支持**——线程束级细粒度同步得重新设计 4. **异步操作不支持**——异步拷贝等场景用同步替代，IO模式需要调整 5. **Pinned Memory一定要用**——Host↔Device带宽差距巨大，这步不做纯属浪费硬件 > 迁移不是简单的API替换。把硬件差异吃透了，才能在DCU上写出真正跑得快的代码。