当前位置： 首页 > news >正文

使用HIP编写GPU 算子向量加法

news 2026/5/16 6:24:57

HIP (Heterogeneous-compute Interface for Portability)来编写一个 GPU 算子（operator）。HIP 是 AMD 推出的 GPU 编程接口，类似 CUDA，但可在 AMD 和 NVIDIA GPU 上运行。下面我给你一个完整示例，演示如何写一个简单算子，比如向量加法（Vector Add），并用 HIP 编译运行。

实践

1️⃣ 基本 HIP 算子示例：向量加法

#include<hip/hip_runtime.h>#include<iostream>#defineN1024// 向量长度#defineTHREADS_PER_BLOCK256// HIP kernel: 向量加法__global__voidvector_add(constfloat*A,constfloat*B,float*C,intn){intidx=blockIdx.x*blockDim.x+threadIdx.x;if(idx<n){C[idx]=A[idx]+B[idx];}}intmain(){// 主机向量float*h_A=newfloat[N];float*h_B=newfloat[N];float*h_C=newfloat[N];// 初始化向量for(inti=0;i<N;i++){h_A[i]=i*1.0f;h_B[i]=(N-i)*1.0f;}// 设备向量float*d_A,*d_B,*d_C;hipMalloc(&d_A,N*sizeof(float));hipMalloc(&d_B,N*sizeof(float));hipMalloc(&d_C,N*sizeof(float));// 拷贝数据到设备hipMemcpy(d_A,h_A,N*sizeof(float),hipMemcpyHostToDevice);hipMemcpy(d_B,h_B,N*sizeof(float),hipMemcpyHostToDevice);// 计算线程块和网格大小intblocks=(N+THREADS_PER_BLOCK-1)/THREADS_PER_BLOCK;// 启动 kernelhipLaunchKernelGGL(vector_add,dim3(blocks),dim3(THREADS_PER_BLOCK),0,0,d_A,d_B,d_C,N);// 拷贝结果回主机hipMemcpy(h_C,d_C,N*sizeof(float),hipMemcpyDeviceToHost);// 验证结果for(inti=0;i<10;i++){// 打印前 10 个元素std::cout<<h_A[i]<<" + "<<h_B[i]<<" = "<<h_C[i]<<std::endl;}// 释放内存delete[]h_A;delete[]h_B;delete[]h_C;hipFree(d_A);hipFree(d_B);hipFree(d_C);return0;}

2️⃣ 编译 & 运行

# 编译hipcc vector_add.cpp-ovector_add# 运行./vector_add

输出前 10 个元素类似：

0 + 1024 = 1024 1 + 1023 = 1024 2 + 1022 = 1024 ...

3️⃣ 说明

  1. __global__修饰函数表示这是一个 GPU kernel。

  2. hipMalloc/hipFree对应 CUDA 的cudaMalloc/cudaFree。

  3. hipMemcpy负责主机和设备间数据拷贝。

  4. hipLaunchKernelGGL用来启动 kernel：

    hipLaunchKernelGGL(kernel_name,dim3(blocks),dim3(threads_per_block),shared_mem_bytes,stream,args...)

  5. HIP 的算子逻辑和 CUDA 几乎一致，只要把cudaXXX改成hipXXX就能跑在 AMD GPU。

运行时报错

/workspace/models/liuysh# hipcc vector_add.cpp -o vector_add vector_add.cpp:29:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 29 | hipMalloc(&d_A, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:30:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 30 | hipMalloc(&d_B, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:31:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 31 | hipMalloc(&d_C, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:34:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 34 | hipMemcpy(d_A, h_A, N * sizeof(float), hipMemcpyHostToDevice); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:35:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 35 | hipMemcpy(d_B, h_B, N * sizeof(float), hipMemcpyHostToDevice); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:44:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 44 | hipMemcpy(h_C, d_C, N * sizeof(float), hipMemcpyDeviceToHost); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:53:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ vector_add.cpp:53:19: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ vector_add.cpp:53:33: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ 9 warnings generated when compiling for gfx906. vector_add.cpp:29:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 29 | hipMalloc(&d_A, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:30:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 30 | hipMalloc(&d_B, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:31:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 31 | hipMalloc(&d_C, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:34:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 34 | hipMemcpy(d_A, h_A, N * sizeof(float), hipMemcpyHostToDevice); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:35:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 35 | hipMemcpy(d_B, h_B, N * sizeof(float), hipMemcpyHostToDevice); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:44:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 44 | hipMemcpy(h_C, d_C, N * sizeof(float), hipMemcpyDeviceToHost); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:53:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ vector_add.cpp:53:19: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ vector_add.cpp:53:33: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ 9 warnings generated when compiling for gfx926. vector_add.cpp:29:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 29 | hipMalloc(&d_A, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:30:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 30 | hipMalloc(&d_B, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:31:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 31 | hipMalloc(&d_C, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:34:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 34 | hipMemcpy(d_A, h_A, N * sizeof(float), hipMemcpyHostToDevice); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:35:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 35 | hipMemcpy(d_B, h_B, N * sizeof(float), hipMemcpyHostToDevice); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:44:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 44 | hipMemcpy(h_C, d_C, N * sizeof(float), hipMemcpyDeviceToHost); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:53:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ vector_add.cpp:53:19: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ vector_add.cpp:53:33: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ 9 warnings generated when compiling for gfx928. vector_add.cpp:29:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 29 | hipMalloc(&d_A, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:30:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 30 | hipMalloc(&d_B, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:31:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 31 | hipMalloc(&d_C, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:34:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 34 | hipMemcpy(d_A, h_A, N * sizeof(float), hipMemcpyHostToDevice); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:35:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 35 | hipMemcpy(d_B, h_B, N * sizeof(float), hipMemcpyHostToDevice); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:44:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 44 | hipMemcpy(h_C, d_C, N * sizeof(float), hipMemcpyDeviceToHost); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:53:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ vector_add.cpp:53:19: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ vector_add.cpp:53:33: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ 9 warnings generated when compiling for gfx936. vector_add.cpp:29:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 29 | hipMalloc(&d_A, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:30:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 30 | hipMalloc(&d_B, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:31:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 31 | hipMalloc(&d_C, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:34:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 34 | hipMemcpy(d_A, h_A, N * sizeof(float), hipMemcpyHostToDevice); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:35:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 35 | hipMemcpy(d_B, h_B, N * sizeof(float), hipMemcpyHostToDevice); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:44:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 44 | hipMemcpy(h_C, d_C, N * sizeof(float), hipMemcpyDeviceToHost); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:53:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ vector_add.cpp:53:19: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ vector_add.cpp:53:33: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ 9 warnings generated when compiling for gfx938. vector_add.cpp:29:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 29 | hipMalloc(&d_A, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:30:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 30 | hipMalloc(&d_B, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:31:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 31 | hipMalloc(&d_C, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:34:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 34 | hipMemcpy(d_A, h_A, N * sizeof(float), hipMemcpyHostToDevice); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:35:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 35 | hipMemcpy(d_B, h_B, N * sizeof(float), hipMemcpyHostToDevice); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:44:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 44 | hipMemcpy(h_C, d_C, N * sizeof(float), hipMemcpyDeviceToHost); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:53:5: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ vector_add.cpp:53:19: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ vector_add.cpp:53:33: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ 9 warnings generated when compiling for host. root@f10r1n06:/workspace/models/liuysh# ./vector_add 0 + 1024 = 1024 1 + 1023 = 1024 2 + 1022 = 1024 3 + 1021 = 1024 4 + 1020 = 1024 5 + 1019 = 1024 6 + 1018 = 1024 7 + 1017 = 1024 8 + 1016 = 1024 9 + 1015 = 1024 Warning: Resource leak detected by SharedSignalPool, 68 Signals leaked.

根据编译和运行日志来看，遇到的“报错”主要是编译时的警告和运行时的资源泄漏提示：

1️⃣ 编译时警告

日志里反复出现：

warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result]
  • 出现在hipMalloc、hipMemcpy、hipFree这几行。
  • 原因：从 HIP 5.x 版本开始，这些函数都加了[[nodiscard]]属性，意思是返回值不能忽略，因为返回值中可能包含错误码。
  • 你现在的代码写法直接调用了函数，没有检查返回值，所以会产生警告。

示例：

hipError_t err=hipMalloc(&d_A,N*sizeof(float));if(err!=hipSuccess){std::cerr<<"hipMalloc failed: "<<hipGetErrorString(err)<<std::endl;}

✅ 建议：编译通过，不会阻止程序运行，但在生产环境中最好检查返回值。

2️⃣ 运行时提示

Warning: Resource leak detected by SharedSignalPool, 68 Signals leaked.

  • 这是 HIP / ROCm 在运行 kernel 后检测到 GPU 资源没有完全释放。

  • 可能原因：

    1. kernel 启动后没有同步 (hipDeviceSynchronize())。
    2. 某些 GPU 信号或流被泄漏（在你的简单示例里可能是内部 runtime 的管理问题）。

  • 解决方法：

    • 在 kernel 调用后加同步，确保所有操作完成：
hipLaunchKernelGGL(vector_add,dim3(blocks),dim3(THREADS_PER_BLOCK),0,0,d_A,d_B,d_C,N);hipDeviceSynchronize();// 等待 kernel 执行完成
  • 对于资源释放问题，保证hipFree被调用即可。你原来的代码已经调用了hipFree(d_A)、hipFree(d_B)、hipFree(d_C)，所以主要是 runtime 内部 signal 还未完全回收。加同步通常可以消除警告。

✅ 总结

  1. 编译警告：nodiscard警告，可以通过检查返回值消除。
  2. 运行时警告：SharedSignalPool 资源泄漏，一般加hipDeviceSynchronize()后消失。
  3. 程序逻辑：你的向量加法结果是正确的，所以算子逻辑没有问题。
http://www.jsqmd.com/news/826521/

相关文章：

  • Anolis OS Linux Dirty Frag 漏洞安全声明
  • 终极炉石传说游戏优化插件：HsMod完整配置与使用指南
  • oracle的26版本及以下 Null的判断及空串判定
  • PNP、NPN、源型、漏型，一次全搞懂
  • BallonsTranslator：3分钟搞定漫画翻译的终极AI辅助工具
  • 从计数器到计时器：使用Spectator构建可观测性系统的实践指南
  • [GESP202512 C++ 三级] 判断题第 9 题
  • ±0.03mm的精度怎么保证？翌东塑胶用AI赋能质量管控升级
  • Minecraft服务器技能数据自动化管理：mcpskills-cli命令行工具实战指南
  • 02 Transformer 基础：Self-Attention 原理详解
  • 思源宋体TTF完全指南：7种字重免费解决中文排版难题
  • AI 智能体 “寒武纪”——OpenClaw 狂飙迭代，引领开源 Agent 商业化落地浪潮
  • 2026年山东大学软件学院创新项目实训博客（五）
  • 62-260515 AI 科技日报 (Qwen3.6 模型推理速度再提升，MTP加速至1.8倍)
  • 开源智能体框架xbrain：模块化设计与工程实践指南
  • 基于DB-GPT-Hub的文本到SQL微调实战：从原理到企业级部署
  • AI Agent Harness Engineering 的安全攻防：你的智能体如何被欺骗、劫持与利用
  • 指纹浏览器开发教程五：浏览器内存信息相关能力该怎么改
  • 基于MCP协议构建AI技能服务器：从原理到实战开发指南
  • SLO-Warden：基于错误预算的云原生服务稳定性自动化管理实践
  • 智能苔藓花园：用CircuitPython与NeoPixel打造会呼吸的天气可视化装置
  • Kaggle竞赛实战技能库：从数据预处理到模型集成的完整工程化实践
  • FAST开发方法在系统分析中四个阶段
  • Windows Cleaner：3步解决C盘爆红，快速释放系统空间的实用指南
  • 2026年公考软件大盘点：技术架构与用户体验深度评测
  • 04 AI 时代的岗位分工与协作机制
  • SpleeterGui终极指南：3步实现AI音乐人声分离的免费神器
  • 我们团队的技术债已经堆成山，我用这四步说服老板给时间重构
  • Swift集成飞书API：原生SDK实现iOS/macOS应用无缝协同
  • 使用git filter-repo删除已提交到git中的敏感信息，api key，配置文件等