想象你在一座叫“N 城”（NVIDIA）的城市生活了很多年，早已熟练用当地“方言”写菜谱——这套方言就是 CUDA。突然有一天，公司告诉你：新产品要同时开进另一座叫“A 城”（AMD）的城市。你担心要重新学一门语言，翻开 A 城的菜谱却发现：这里的方言和 N 城惊人地相似。“西红柿炒蛋”在这里写作“番茄炒蛋”，做法几乎一模一样。更妙的是，A 城官方还提供了一位叫 hipify 的“翻译官”，能一次性把你的老菜谱批量翻成新方言。当然，翻译官也有翻不动的俚语——纹理内存、多卡通信这类“城市特产”仍要你亲自动手。本文就带你走一遍这次“跨城搬家”：认识 HIP 方言、摸清 ROCm 市政设施、学会呼叫翻译官、搞懂重点词汇对照表，最后在新厨房里调出合适的火候。

岗位关联

对 AI 从业者而言，掌握 HIP 迁移意味着你的训练与推理代码不再被单一 GPU 厂商锁死，在云资源选型与硬件采购上多一份主动权，是 AI 工程化落地的加分技能。

目录
  1. 可视化演示
  2. HIP概念认知：给 CUDA 写的“方言版”
  3. 硬件与生态：红城的市政工程
  4. 迁移工具链：请来方言翻译官 hipify
  5. API映射差异：方言对照小册子
  6. 不兼容特性：翻译不了的俚语
  7. 迁移实战与性能调优：新厨房试火候
  8. AI 应用场景
  9. Python 演示
  10. 常见误区
  11. 面试题

可视化演示

演示CUDA到HIP的代码迁移流程：API映射对照、线程层次模型、hipify自动转换管道，三种模式交互切换。

本节配有交互动效（进入视口后自动加载，约 3KB）。

API映射对照 cudaMalloc → hipMalloc cudaMemcpy → hipMemcpy cudaLaunch → hipLaunch 6组核心API一一映射，迁移成本极低

HIP概念认知：给 CUDA 写的“方言版”

你在一座叫 N 城（NVIDIA）的城市住了很多年，写菜谱用的是当地方言 CUDA。现在公司让你把产品搬到隔壁的 A 城（AMD），你原以为要重新学一门语言，翻开 A 城的菜谱却发现——“西红柿炒蛋”变成了“番茄炒蛋”，做法几乎一模一样。HIP（Heterogeneous-Computing Interface for Portability，异构计算可移植接口）就是 AMD 为 CUDA 开发者准备的那本“番茄版菜谱”：一门与 CUDA 高度同构的 GPU 编程语言。

说它同构并非夸张。CUDA 里最核心的语法要素在 HIP 中基本原样保留：__global__ 修饰的内核函数（可以理解成发给整个厨房的广播口令）、threadIdx/blockIdx 内建索引（灶台编号与灶台组号）、<<<grid, block>>> 内核启动语法（宣布“这一桌菜按 4 行 8 列同时开炒”）。差异主要集中在 API 前缀和少量头文件：cudaMalloc 变成了 hipMalloc，cuda_runtime.h 变成了 hip_runtime.h。

HIP 的跨平台能力来自“双后端”编译。用 hipcc 编译源码时，编译器根据目标平台选择后端：在 AMD GPU 上生成 ROCm 后端代码，直接调用 AMD 底层运行时；在 NVIDIA GPU 上则走 HIP-CUDA 后端，把 HIP 代码翻译回 CUDA 再交给 NVIDIA 工具链。这意味着你维护一份 HIP 源码，就能在两座城市的“厨房”里编译运行。对团队而言，迁移成本的核心在于“有多少代码必须人手工改”，公式如下——

就像搬家时，普通家具搬过去就能用，但那些 N 城独有的大件旧物才最费时。公式里的自动转换率越高，你需要动手的部分越少，这次跨城搬家也就越轻松。

MigrationCost=Nlines×(1−rauto)\displaystyle \text{MigrationCost} = N_{\text{lines}} \times (1 - r_{\text{auto}})
迁移成本 = 总代码行数 ×（1 − 自动转换率）。自动转换率越高，需要人工改动的行数越少。
  1. 总代码行数 \(N_{\text{lines}}\)待迁移的 CUDA 源文件总行数，是工程量估算的基数。
    NlinesN_{\text{lines}}
  2. 自动转换率 \(r_{\text{auto}}\)hipify 能自动改写的行数占总行数的比例，经验值通常在 80%～95% 之间。
    rauto=NautoNlinesr_{\text{auto}} = \frac{N_{\text{auto}}}{N_{\text{lines}}}
  3. 不可自动迁移比例 \((1 - r_{\text{auto}})\)剩下需要人工介入的部分，是迁移风险与工作量的主要来源。
    1−rauto1 - r_{\text{auto}}
  4. 迁移总成本总行数与不可自动迁移比例相乘，得到预估的人工维护成本，数值越小迁移越轻松。
    MigrationCost=Nlines×(1−rauto)\text{MigrationCost} = N_{\text{lines}} \times (1 - r_{\text{auto}})

硬件与生态：红城的市政工程

搬进 A 城，光会方言还不够。你得知道水电煤气在哪办、市政管网怎么走。ROCm（Radeon Open Compute）就是 A 城的整套市政系统——HIP 只是你厨房里那把新锅，锅要正常工作，靠的是从水管到燃气的全链路支撑。

ROCm 是 AMD 开源的 GPU 计算平台，和“CUDA Toolkit 加 NVIDIA 驱动”的关系类似。它自上而下分为几层：最上面是领域库（如 rocBLAS 线性代数库、rocFFT 傅里叶库），中间是 HIP 运行时与编译工具链（hipcc 基于 LLVM/Clang），下面是 AMDGPU 内核驱动，最底层是 AMD 的 CDNA（数据中心）与 RDNA（消费级）GPU 硬件。

这套分层决定了你“能不能搬”。ROCm 对 CDNA 架构的数据中心 GPU（如 MI300X、MI250）支持最完善，对部分 RDNA 工作站卡也有官方支持，但大量消费级游戏卡不在支持列表内。迁移前务必查阅官方 GPU 支持矩阵，否则装好驱动才发现“户口”办不下来，返工成本很高。另外，一次 GPU 任务的端到端时间是数据搬运与内核计算的叠加，ROCm 生态里同样要关心 PCIe 带宽与内核效率：

市政没接通，锅再新也做不了饭。迁移 HIP 的第一步不是写代码，而是确认你的硬件在 ROCm 的“服务区”内。

Ttotal=Ttransfer+Tkernel\displaystyle T_{\text{total}} = T_{\text{transfer}} + T_{\text{kernel}}
端到端任务时间 = 数据搬运时间 + 内核计算时间。任何一个成为瓶颈，都会拖慢整体性能。
  1. 端到端总耗时 \(T_{\text{total}}\)一次 GPU 任务从发起拷贝到内核执行完毕的总时间。
    TtotalT_{\text{total}}
  2. 数据搬运时间 \(T_{\text{transfer}}\)主机与设备间通过 PCIe 或显存拷贝数据的时间，受总线带宽限制。
    Ttransfer=DbytesBbusT_{\text{transfer}} = \frac{D_{\text{bytes}}}{B_{\text{bus}}}
  3. 内核计算时间 \(T_{\text{kernel}}\)GPU 上执行 kernel 的耗时，与占用率、访存效率和指令吞吐相关。
    TkernelT_{\text{kernel}}
  4. 两部分叠加总时间是搬运与计算的叠加，调优要先判断谁占大头，再对症下药。
    Ttotal=Ttransfer+TkernelT_{\text{total}} = T_{\text{transfer}} + T_{\text{kernel}}

迁移工具链：请来方言翻译官 hipify

确认水电没问题后，你请来了 A 城官方推荐的“方言翻译官”——hipify。这是一套自动迁移工具，能把 N 城方言的菜谱批量翻成 A 城方言。翻译官其实有两位：一位急性子，一位慢性子。

急性子叫 hipify-perl：它是一段 Perl 脚本，靠正则表达式做快速的文本级替换，优点是轻量、无需编译环境，适合第一轮粗翻。慢性子叫 hipify-clang：它基于 Clang 编译器对代码做语法树（AST）级分析，能识别变量作用域、宏展开、注释边界，替换更精准、误报更少，是官方推荐的正式迁移工具。

两者的底层逻辑都维护着一张“映射规则表”：cudaMalloc→hipMalloc 这类一对一机械替换直接执行；cudaMemcpyPeerToPeer 这类需要人工确认的“灰色条目”则不会被擅自翻译，而是给出行号与警告。hipify 运行后会输出转换后的 .hip.cpp 源文件、一份详细报告和一份 TODO 清单，自动转换率如下：

翻译官的工作效率可以用自动转换率衡量。hipify 的自动转换率通常超过 90%，但剩下的 10% 才是决定你周末加不加班的关键。

ConvRate=NautoNtotal\displaystyle \text{ConvRate} = \frac{N_{\text{auto}}}{N_{\text{total}}}
自动转换率 = 成功自动映射的 API 调用次数 ÷ 全部待映射 API 调用次数。它是衡量 hipify 工作效率的核心指标。
  1. 待映射 API 总数 \(N_{\text{total}}\)源文件中所有需要替换的 CUDA API 调用（如 cudaMalloc、cudaMemcpy 等）。
    NtotalN_{\text{total}}
  2. 成功自动映射数 \(N_{\text{auto}}\)hipify 依据映射规则表成功替换的 API 数量。
    NautoN_{\text{auto}}
  3. 转换率计算两者相除得到自动转换率，通常超过 90%，剩余部分会出现在 TODO 清单中。
    ConvRate=NautoNtotal\text{ConvRate} = \frac{N_{\text{auto}}}{N_{\text{total}}}

API映射差异：方言对照小册子

方言对照表是搬家时最实用的小册子。A 城与 N 城 90% 的词汇一一对应，只是叫法不同。运行时 API 的映射最直观：cudaMalloc→hipMalloc、cudaMemcpy→hipMemcpy（方向枚举跟着改为 hipMemcpyHostToDevice、hipMemcpyDeviceToHost）、cudaFree→hipFree、cudaStreamCreate→hipStreamCreate、cudaEventRecord→hipEventRecord、错误检查 cudaGetLastError→hipGetLastError。

内存管理方面还有 cudaMallocHost（锁页内存）→hipMallocHost、cudaMemset→hipMemset。流（Stream）与事件（Event）作为异步并发控制机制，在 HIP 中一一对应；流里排队的内核和事件之间的依赖关系完全沿用 CUDA 的语义。最省心的是内建变量：threadIdx、blockIdx、blockDim、gridDim 的语义与取值规则和 CUDA 完全一致，kernel 启动的 <<<grid, block>>> 语法也不变。也就是说，迁移时你不需要改任何“索引算位置的逻辑”，只改“申请和搬运数据的函数”。

衡量内存迁移效率时，带宽公式是核心：

这份词汇对照表就像“番茄”和“西红柿”：发音变了，菜的做法一点没变。换词不换语法，正是你能在两座城市间自由穿梭的最大底气。

BW=DbytesTcopy\displaystyle \text{BW} = \frac{D_{\text{bytes}}}{T_{\text{copy}}}
有效带宽 = 搬运数据量 ÷ 拷贝耗时。带宽越接近硬件峰值，说明内存搬运越高效。
  1. 搬运数据量 \(D_{\text{bytes}}\)一次 hipMemcpy 调用传输的字节总数，例如 1024×4 字节的浮点数组。
    Dbytes=1024×4D_{\text{bytes}} = 1024 \times 4
  2. 拷贝耗时 \(T_{\text{copy}}\)从调用 hipMemcpy 到返回的时间，可借助事件（hipEvent）精确计时。
    TcopyT_{\text{copy}}
  3. 带宽计算两者相除得到有效带宽，用于判断内存迁移是否高效。
    BW=DbytesTcopy\text{BW} = \frac{D_{\text{bytes}}}{T_{\text{copy}}}

不兼容特性：翻译不了的俚语

再全的词典也收不进所有俚语。A 城没有“豆汁儿”这个词，翻译官也束手无策。HIP 迁移中最费时的工作，就集中在这几类“城市特产”上。

第一类是纹理内存（Texture Memory）。CUDA 中通过 texture reference/texture object 绑定数据、再用 tex1D/tex2D 采样的写法，在 HIP 中只有部分对应，且 API 差异较大。常见策略是改用 __ldg 只读缓存加载，或直接退化为全局内存访问，牺牲一部分缓存加速来换取可移植性。第二类是多卡通信：NVIDIA 的 NCCL 对应 AMD 的 RCCL，虽然二者都实现 all-reduce 等集合通信算法，但初始化上下文、调用头文件不同，通信代码需要重写。

第三类是原子操作与复杂特性。atomicAdd 对 double 的支持、CUDA Graph 捕获机制、动态并行（Dynamic Parallelism）等特性在部分 AMD 硬件或 ROCm 版本上缺失或表现不同，需要查目标硬件的支持文档逐条核对。经验做法是：迁移前先估算“手工占比”，把风险前置：

遇到翻译不了的俚语，聪明的办法是“改说普通话”：用两座城市都懂的通用写法重写这块逻辑，而不是硬抠某个方言词。手工占比越低，迁移之路就越平坦。

ManualRatio=NmanualNtotal\displaystyle \text{ManualRatio} = \frac{N_{\text{manual}}}{N_{\text{total}}}
手工处理占比 = 需人工修改的代码行数 ÷ 总代码行数。占比越高，迁移风险越大。
  1. 总代码行数 \(N_{\text{total}}\)待迁移源文件的总行数，是统计分母。
    NtotalN_{\text{total}}
  2. 需手工处理行数 \(N_{\text{manual}}\)包含纹理内存、多卡通信、内联汇编等 hipify 无法自动转换的行。
    NmanualN_{\text{manual}}
  3. 手工占比计算两者相除得到占比，迁移前估算它有助于制定排期与风险预案。
    ManualRatio=NmanualNtotal\text{ManualRatio} = \frac{N_{\text{manual}}}{N_{\text{total}}}

迁移实战与性能调优：新厨房试火候

翻译完菜谱、搬进新厨房，最后一步是开火试菜。你会发现 A 城的灶台火力和 N 城不一样——同样的“三分钟爆炒”，在 N 城刚刚好，在 A 城可能糊了。这就是迁移后的性能调优。

完整迁移流程分五步：①环境准备——安装 ROCm 驱动、HIP SDK，确认 GPU 型号在支持矩阵内；②自动转换——先用 hipify-perl 快速粗翻，再用 hipify-clang 精转；③手工修正——处理报告中的纹理、多卡、原子操作等 TODO；④编译验证——改用 hipcc 编译，修复头文件、链接库与宏定义错误；⑤功能测试与调优——跑通正确性后再谈性能。

调优阶段有两大法宝：rocprof（HIP 程序性能分析器，类似 NVIDIA 的 Nsight Compute）与 rocm-smi（硬件状态监控）。常见优化手段包括：调整 hipcc 编译选项（如 -O3、--offload-arch 指定目标架构）、保证全局内存合并访问（coalescing）、调整 block 大小以提升 SM 占用率、用 hipMallocManaged 托管内存简化搬运、减少 H2D/D2H 拷贝次数。其中占用率是判断内核是否充分利用硬件的最直观指标：

搬家的终点不是钥匙到手，而是新厨房里第一桌菜端上桌。迁移同理：编译通过只是开始，用 rocprof 重新“试火候”，让每个内核都在 AMD 硬件上物尽其用，这次跨城搬家才算真正圆满。

Occupancy=AactiveAmax\displaystyle \text{Occupancy} = \frac{A_{\text{active}}}{A_{\text{max}}}
占用率 = 每个 SM 实际激活的线程束数 ÷ 硬件允许的最大线程束数。占用率越高，越能隐藏访存延迟。
  1. 实际激活线程束数 \(A_{\text{active}}\)kernel 运行时每个流多处理器（SM）上真正活跃的 warp 数量。
    AactiveA_{\text{active}}
  2. 最大线程束数 \(A_{\text{max}}\)由硬件规格与线程块资源（寄存器、共享内存）共同决定的上限。
    AmaxA_{\text{max}}
  3. 占用率计算两者相除得到占用率，一般建议达到 50% 以上，否则延迟隐藏不充分。
    Occupancy=AactiveAmax\text{Occupancy} = \frac{A_{\text{active}}}{A_{\text{max}}}

AI 应用场景

跨厂商 AI 训练集群统一算子库

阶段：训练基础设施搭建

解决的问题：很多深度学习框架依赖手写 CUDA 算子（如 FlashAttention 的自定义 kernel），当公司同时在 NVIDIA 与 AMD 集群上训练大模型时，同一套算子要维护两份代码。用 HIP 迁移后，算子库一份源码两处编译，开发与维护成本近乎减半。

具体案例：某自动驾驶公司模型团队拥有 200 张 A100 与 100 张 MI300X，原先只在 NVIDIA 集群训练。通过 hipify 迁移自研算子和自定义 loss kernel，配合 rocBLAS 替换 cuBLAS 调用，两周内完成算子库统一，双集群任务调度打通。

云推理服务多硬件适配

阶段：模型推理部署与优化

解决的问题：云厂商推理网关需要根据成本与库存动态调度 T4 与 MI250 实例，但推理 kernel 是纯 CUDA 实现，无法在 AMD 实例上运行。HIP 迁移让推理服务无侵入地适配两种硬件，提升资源池利用率。

具体案例：某云厂商的 Llama-3-70B 推理服务，使用 hipify 转换量化 kernel 与 KV-cache 管理逻辑，再把 NCCL 通信替换为 RCCL，最终在同一网关下实现 NVIDIA/AMD 实例无缝混跑。

科研超算平台软件栈移植

阶段：高性能计算平台运营

解决的问题：高校超算中心新采购 AMD 分区，但历史科研代码全部为 CUDA 编写。用 HIP 统一中间层，让用户透明地在两个分区提交作业，避免重复开发。

具体案例：某高校超算平台将材料模拟软件中的 CUDA kernel 批量迁移到 HIP，并用 CMake 实现 CUDA/HIP 双后端的自动编译判断，用户在 PBS 脚本中只需切换队列名即可调度到不同架构节点。

Python 演示

python80 行
#!/usr/bin/env python# -*- coding: utf-8 -*-"""mini_hipify: 用映射表模拟 CUDA 到 HIP 的自动迁移流程""" import sys # CUDA -> HIP 自动映射规则(部分)API_MAP = {    'cudaMalloc': 'hipMalloc',    'cudaFree': 'hipFree',    'cudaMemcpy': 'hipMemcpy',    'cudaMemcpyAsync': 'hipMemcpyAsync',    'cudaMemset': 'hipMemset',    'cudaStreamCreate': 'hipStreamCreate',    'cudaStreamDestroy': 'hipStreamDestroy',    'cudaEventCreate': 'hipEventCreate',    'cudaEventRecord': 'hipEventRecord',    'cudaEventSynchronize': 'hipEventSynchronize',    'cudaDeviceSynchronize': 'hipDeviceSynchronize',    'cudaGetLastError': 'hipGetLastError',} # 这些特性 hipify 无法自动处理, 需要人工介入(仅作演示)MANUAL_WARN = ('texture<', '__texture__', 'cudaMemcpyPeerToPeer')  def hipify_source(src):    total = auto = manual = 0    out_lines = []    for line in src.splitlines():        if 'cuda_runtime.h' in line:            line = line.replace('cuda_runtime.h', 'hip/hip_runtime.h')        hits = [k for k in API_MAP if k in line]        total += len(hits)        for k in hits:            line = line.replace(k, API_MAP[k])        auto += len(hits)        if any(w in line for w in MANUAL_WARN):            manual += 1            out_lines.append('// TODO: manual fix | ' + line)        else:            out_lines.append(line)    return '\n'.join(out_lines), total, auto, manual  def main():    sample = r'''#include <cuda_runtime.h> __global__ void add_one(float* d_a) {    int i = threadIdx.x + blockIdx.x * blockDim.x;    d_a[i] += 1.0f;} void run(float* h_a) {    float* d_a;    cudaMalloc(&d_a, 1024 * sizeof(float));    cudaMemcpy(d_a, h_a, 1024 * sizeof(float), cudaMemcpyHostToDevice);    add_one<<<4, 256>>>(d_a);    cudaDeviceSynchronize();    cudaMemcpy(h_a, d_a, 1024 * sizeof(float), cudaMemcpyDeviceToHost);    cudaFree(d_a);}'''    src = sample    if len(sys.argv) > 1:        with open(sys.argv[1], 'r', encoding='utf-8') as f:            src = f.read()     converted, total, auto, manual = hipify_source(src)    print(converted)    print('--- migration stats ---')    print('total API calls: %d' % total)    print('auto converted : %d' % auto)    print('auto rate      : %.1f%%' % (100.0 * auto / total if total else 0.0))    print('manual lines   : %d' % manual)  if __name__ == '__main__':    main()
目的：用一张映射表模拟 hipify 的核心逻辑——把 CUDA 源码逐行改写成 HIP 源码，并统计自动转换率与需人工处理的行数。
  1. API_MAP = {'cudaMalloc': 'hipMalloc', ...}模拟 hipify 内置的“方言映射规则表”，是自动转换的依据。
  2. if 'cuda_runtime.h' in line: line.replace(...)额外处理头文件替换，对应真实迁移中 #include <hip/hip_runtime.h>。
  3. hits = [k for k in API_MAP if k in line]逐行扫描，找出当前行命中的 CUDA API。
  4. line.replace(k, API_MAP[k])等价于翻译官查词典后做的“换词不换语法”。
  5. if any(w in line for w in MANUAL_WARN)命中纹理、多卡等特例时，自动标记 // TODO: manual fix，提醒人工介入。
  6. print('auto rate : %.1f%%' % ...)输出自动转换率，对应正文公式中的 ConvRate。

常见误区

误区我以为把代码里所有 cuda 开头的函数名替换成 hip 开头，就完成了迁移，hipify 干的就是这活儿。

正解前缀替换只解决了约 90% 的机械性 API 变化。真正的迁移还包括头文件（#include <hip/hip_runtime.h>）、宏定义、编译工具链（nvcc→hipcc），以及 cudaMemcpyPeerToPeer、纹理内存等没有一一对应 API 的特例。翻译官翻完只是草稿，你必须编译、跑测试、看报告，绝不能只做字符串替换。

为何常见：因为 90% 的 API 确实只是前缀差异，容易让人误以为迁移等于批量替换。

误区hipify 都帮我改好了，直接用 hipcc 编译就能跑，不用再看代码。

正解hipify 的输出是一份“草稿”。它无法处理头文件依赖差异、第三方库（NCCL/cuBLAS）替换、纹理内存绑定、内联 PTX 汇编等。官方流程要求转换后必须人工审查 hipify 报告中的 TODO 标记，并完成编译验证与功能测试。就像翻译官翻完的菜谱还要大厨试菜，确认口味没跑偏。

为何常见：工具自动化的“爽感”掩盖了它只是一个辅助工具的事实。

误区代码都迁移过来了，N 城跑多快，A 城就应该跑多快。

正解HIP 编译到不同后端，硬件架构（SM 与 CU 的调度方式、缓存大小）、编译器优化都不同。同样的 kernel 在 AMD 上可能需要调整 block 大小、访存模式、占用率。必须用 rocprof 重新分析，逐个 kernel 调优，性能才会达标。“语法兼容”不等于“性能等价”，就像同样的菜谱换了个灶台，火候必须重新试。

为何常见：人们把“语法兼容”误解为“性能等价”，忽略了硬件后端的巨大差异。

误区既然 ROCm 是 AMD 官方的计算平台，任何 AMD 显卡装上驱动就能跑 HIP 程序。

正解ROCm 目前主要支持 CDNA 数据中心 GPU 与部分 RDNA 工作站卡，很多消费级 RDNA 游戏卡不在官方支持列表，或在旧版本 ROCm 上存在功能缺失。迁移前要先查目标 GPU 的官方支持矩阵，否则就像买了新房才发现小区还没通燃气。

为何常见：AMD 官方宣传强调“开源、广泛支持”，让用户高估了支持范围。

误区HIP 不过是把 cuda 换成 hip，底层还是 CUDA 那一套。

正解HIP 有独立的运行时、编译器和生态。它在 AMD 上走 ROCm 后端，在 NVIDIA 上才转成 CUDA 后端。API 数量、错误码、部分语义都有差异；NCCL→RCCL、cuBLAS→rocBLAS 等生态组件也不互通。两座城市方言虽像，市政系统却各自独立。

为何常见：因为语法相似度太高，初学者容易忽略底层实现与生态的不同。

面试题

Q1. 在 HIP 编程模型中，与 CUDA 的 cudaMalloc 对应的内存分配函数是？

单选题 · easy

  • A.hipCalloc()
  • B.hipMalloc()← 正确答案
  • C.hipMemNew()
  • D.hipAllocNew()

答案：B。hipMalloc 就是 cudaMalloc 在 A 城方言里的说法，功能一致，都是给 GPU 设备内存分配空间。hipCalloc 是分配并清零的变体，hipMemNew 与 hipAllocNew 并不存在。

Q2. 下列哪个工具用于把 CUDA 代码自动转换为 HIP 代码？

单选题 · easy

  • A.nvcc 迁移工具
  • B.clang 迁移工具
  • C.hipify 迁移工具← 正确答案
  • D.rocprof 迁移工具

答案：C。hipify 是 AMD 官方的自动迁移工具，像方言翻译官一样批量改写 CUDA API。nvcc 是 NVIDIA 的编译器，clang 是通用编译器，rocprof 是性能分析器，都不是迁移工具。

Q3. 关于 HIP 中的线程索引变量，正确的说法是？

单选题 · mid

  • A.含义与取值规则一致← 正确答案
  • B.名称相同但语义不同
  • C.只在 AMD 上有效
  • D.两者毫无相同之处

答案：A。threadIdx、blockIdx 等在 HIP 与 CUDA 中含义和取值规则完全一致，就像“灶台编号”在两座城市说的是同一个数。区别只在于底层平台后端不同。

Q4. 下列哪类 API 在 hipify 迁移时最可能需要人工处理？

单选题 · mid

  • A.全局内存分配类
  • B.纹理内存相关类← 正确答案
  • C.内核启动语法类
  • D.错误检查函数类

答案：B。纹理内存的绑定与采样写法在两套 API 中差异最大，hipify 通常无法自动处理，需要手工改写。全局内存、内核启动、错误检查都有直接的对应关系。

Q5. HIP 直接沿用了 CUDA 的哪些语法要素？

多选题 · mid

  • A.全局内核声明语法← 正确答案
  • B.线程索引内建变量← 正确答案
  • C.专有内存分配调用
  • D.内核尖括号启动语法← 正确答案

答案：ABD。__global__ 内核声明、threadIdx 等内建变量、<<<grid, block>>> 启动语法在 HIP 中原样保留。而内存分配 API 是改名重灾区：cudaMalloc 变成了 hipMalloc，不属于“原样沿用”。

Q6. cudaFree 在 HIP 中对应的内存释放函数是？

单选题 · mid

  • A.hipRelease调用
  • B.hipFree函数← 正确答案
  • C.hipDelete函数
  • D.hipMalloc函数

答案：B。hipFree 与 cudaFree 一一对应，都是释放设备内存。hipRelease、hipDelete 在 HIP 中不存在；hipMalloc 是分配函数而非释放函数。

Q7. 下列哪些属于 ROCm 生态的核心组件？

多选题 · hard

  • A.hipcc 编译工具← 正确答案
  • B.rocprof 工具← 正确答案
  • C.nvcc 编译工具
  • D.ROCm 运行时库← 正确答案

答案：ABD。ROCm 生态包含 hipcc 编译器、HIP 运行时、rocprof 分析器等组件，它们是 A 城市政系统的水电煤。nvcc 是 NVIDIA 的编译器，属于 CUDA 生态，不是 ROCm 的一部分。

Q8. 关于 HIP 的跨平台理念，下列说法正确的是？

单选题 · easy

  • A.让代码只在本机运行
  • B.让代码只在单卡运行
  • C.让代码只在一家厂商运行
  • D.让同一份代码跨多家硬件← 正确答案

答案：D。HIP 的核心卖点就是“一份代码，两处编译”：在 AMD 平台走 ROCm 后端，在 NVIDIA 平台走 HIP-CUDA 后端。让同一份源码跨厂商运行，是 HIP 存在的根本意义。

Q9. 与 CUDA 生态中 NCCL 对应的 HIP 多卡通信库是？

单选题 · hard

  • A.RCCL 多卡库← 正确答案
  • B.cuBLAS 数值库
  • C.ROCm 算子编译库
  • D.NCCL 跨平台版

答案：A。RCCL 是 AMD 的集合通信库，对应 NVIDIA 的 NCCL，用于多卡 all-reduce 等通信。cuBLAS 是数值计算库，ROCm 的算子编译库是另一回事，NCCL 并没有跨平台版本。

Q10. 迁移后的 HIP 程序性能调优通常包括哪些手段？

多选题 · hard

  • A.调整编译优化参数← 正确答案
  • B.优化全局内存访问← 正确答案
  • C.提升核心占用率← 正确答案
  • D.增加错误处理分支

答案：ABC。调优三板斧：调整 hipcc 编译选项、优化内存合并访问、提升 SM 占用率。增加错误处理分支属于健壮性改进，与性能调优无关。