从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并发。

除了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应用。

---

二、软件栈对比：DTK到底兼容到什么程度

DTK（DCU Toolkit）是DCU硬件平台的开发工具包。官方说法是同时支持HIP和CUDA两种编程模型，能快速将CUDA和ROCm生态中的应用部署到DCU上。说白了就是：你写HIP能跑，写CUDA也能跑（通过GPUfusion兼容层）。

2.1 软件栈全景

DTK各组件对ROCm生态原生兼容，与CUDA生态高度兼容。

核心模块对照表：

模块	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__ 标记，这就是核函数。

// 设备端核函数
__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<<<gridDim, blockDim, sharedMemBytes, stream>>>(N, d_a);

核函数的规矩：

• 返回值必须是 void
• 参数可以是设备端指针、立即数、结构体或引用
• 以线程为基本单位执行，线程之间总是"并发"的
• 每个线程有独立编号

HIP核函数启动语法有三种写法（本质上等价）：

方式	语法
HIP宏	hipLaunchKernelGGL(kernel, numBlocks, dimBlocks, sharedMemBytes, stream, args...)
HIP Kernel语法	kernel<<<Dg, Db, Ns, S>>>(args...)
CUDA Kernel语法	kernel<<<Dg, Db, Ns, S>>>(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执行。块内线程可直接通信和同步

核函数中通过内置变量获取线程在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线程网格——计算全局索引：

int idx = blockDim.x * blockIdx.x + threadIdx.x;

2D线程网格——需要同时计算X和Y：

int idx_x = blockDim.x * blockIdx.x + threadIdx.x;
int idx_y = blockDim.y * blockIdx.y + threadIdx.y;

搞懂了1D和2D，3D加个Z轴同理。并行算法的核心就是把计算任务映射到这个3D-Grid上——这个映射关系选对了，性能提升往往是数量级的。

3.3 存储结构：三层内存模型

HIP存储结构是层次化的，直接对应硬件存储层次。搞懂了这个，一大半的性能问题都能自己排查了。

__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;
}

存储层次速查表：

修饰符	存储器	作用域	生命周期
（无修饰符）标量	寄存器	线程	线程
（无修饰符）数组	本地内存	线程	线程
__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 阻塞与非阻塞

// 阻塞拷贝——完成前主机端暂停
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）都通过队列执行。同一队列顺序执行，不同队列共享硬件但互不影响。

hipStream_t stream;
hipStreamCreate(&stream);
// ... 提交任务到stream ...
hipStreamDestroy(stream);

注意：使用NULL或0表示默认流。对默认流的操作会导致强制同步，影响其他流中的任务。

4.3 Pinned Memory

主机端默认分配的是Pageable内存。Pinned Memory（Page-Lock Memory）让DMA直接访问，大幅提升Host↔Device传输带宽。

double *h_a = NULL;
hipHostMalloc(&h_a, Nbytes);  // 分配pinned内存
// ... 使用 ...
hipHostFree(h_a);             // 释放

频繁传输的数据优先用Pinned Memory。带宽差距非常显著，这个优化是性价比最高的。

4.4 同步

hipDeviceSynchronize();           // 重型同步：阻塞主机直到所有设备流完成
hipStreamSynchronize(stream);     // 轻量同步：仅等待指定流

4.5 Event

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 多设备编程

int canAccessPeer;
hipDeviceCanAccessPeer(&canAccessPeer, device0, device1);
if (canAccessPeer) {
    hipDeviceEnablePeerAccess(device1, 0);
    // ... P2P内存访问 ...
    hipDeviceDisablePeerAccess(device1);
}

每个设备最多支持8个对等连接。P2P访问可以通过 hipMemcpyPeerAsync 实现设备间直接拷贝。

---

五、Runtime日志与异常调试

5.1 Runtime日志

出问题时开启Runtime日志做初步分析：

# 方式一：环境变量
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内存中转。

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：

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上写出真正跑得快的代码。